GPU-accelerated UI toolkit (Vulkan)
git clone https://git.lucas.co/cce-ui.git
src/draw/rt_common.wgsl (11K)
1 // rt_common.wgsl — the path tracer's shared core (RT-renderer phases 2+4).
2 //
3 // Everything except the trace call: params, scene/material buffers,
4 // accumulation, RNG, sky, sampling, and cs_main. Binding 1 and
5 // `intersect_scene` come from whichever tier file is concatenated after
6 // this one at pipeline creation:
7 // - rt_bvh.wgsl — tier 1: a CPU-built BVH traversed in compute; runs
8 // on any device, no VK_KHR_ray_* required.
9 // - rt_query.wgsl — tier 2: hardware ray queries against a driver-built
10 // TLAS (VK_KHR_ray_query), engaging RT cores.
11 // One dispatch adds `spp` samples per pixel into the accumulation buffer
12 // (progressive refinement); the running mean is tone-mapped (clamped
13 // linear) into `out_img`, which the stage blits into the backdrop pane.
14
15 struct Params {
16 // Inverse of the raster path's proj*view*model: unprojects wgpu-style NDC
17 // (y up, z in [0,1]) into mesh space, so rays live in the same space as
18 // the triangles fed to `set_rt_scene`.
19 inv_mvp: mat4x4<f32>,
20 width: u32,
21 height: u32,
22 sample_index: u32,
23 max_bounces: u32,
24 // Samples per dispatch: 1 for the interactive viewport (one refinement
25 // step per frame), higher for offscreen/thumbnail rendering so a whole
26 // image needs only a few submits.
27 spp: u32,
28 _pad0: u32,
29 _pad1: u32,
30 _pad2: u32,
31 // The scene's image, a quad of two triangles whose material is marked
32 // textured (albedo.w). xyz = its top-left corner; w = its opacity, 0
33 // when there is no image to sample (a textured hit then lets the ray
34 // through).
35 img_origin: vec4<f32>,
36 // xyz = the top edge, corner to corner; w = the texture's width in
37 // texels.
38 img_u: vec4<f32>,
39 // xyz = the left edge, top to bottom; w = the texture's height.
40 img_v: vec4<f32>,
41 // What a camera ray that meets nothing shows: rgb linear, w = 1 when
42 // set. w = 0 shows the sky, as every miss did before. Only the camera
43 // ray — a bounce that leaves the scene still meets the sky, which is
44 // the light, so the backdrop changes what is seen behind the scene and
45 // not how the scene is lit.
46 background: vec4<f32>,
47 // The environment (`RtEnvironment`), xyz each: toward the sun (unit),
48 // the sun's radiance, the sky overhead, the sky below. Linear.
49 sun_dir: vec4<f32>,
50 sun_color: vec4<f32>,
51 sky_zenith: vec4<f32>,
52 sky_nadir: vec4<f32>,
53 }
54
55 @group(0) @binding(0) var<uniform> params: Params;
56
57 // Binding 1 belongs to the tier file: the BVH node buffer (tier 1) or the
58 // acceleration structure (tier 2).
59
60 // Positions in xyz; p0.w carries the material index (bitcast).
61 struct Tri {
62 p0: vec4<f32>,
63 p1: vec4<f32>,
64 p2: vec4<f32>,
65 }
66 @group(0) @binding(2) var<storage, read> tris: array<Tri>;
67
68 struct Material {
69 albedo: vec4<f32>,
70 emission: vec4<f32>,
71 }
72 @group(0) @binding(3) var<storage, read> materials: array<Material>;
73
74 // One vec4 per pixel: rgb = radiance sum, a = sample count.
75 @group(0) @binding(4) var<storage, read_write> accum: array<vec4<f32>>;
76
77 @group(0) @binding(5) var out_img: texture_storage_2d<rgba8unorm, write>;
78
79 // Primary-hit features for the denoiser (rt_denoise.wgsl), two vec4s per
80 // pixel: [2i] = (shading normal, hit t — 1e30 for sky), [2i+1] = (albedo, 0).
81 @group(0) @binding(6) var<storage, read_write> features: array<vec4<f32>>;
82
83 // The scene's image and its sampler. Always bound: to a 1x1 stand-in while
84 // the scene has no image.
85 @group(0) @binding(7) var img: texture_2d<f32>;
86 @group(0) @binding(8) var img_sampler: sampler;
87
88 // PCG (O'Neill) — one u32 of state per path, advanced per draw.
89 fn rand(state: ptr<function, u32>) -> f32 {
90 var s = *state * 747796405u + 2891336453u;
91 *state = s;
92 let word = ((s >> ((s >> 28u) + 4u)) ^ s) * 277803737u;
93 return f32((word >> 22u) ^ word) * (1.0 / 4294967295.0);
94 }
95
96 // The tier boundary: whichever tier file follows provides
97 // fn intersect_scene(ro: vec3<f32>, rd: vec3<f32>) -> HitInfo
98 struct HitInfo {
99 t: f32,
100 tri: u32,
101 }
102
103 // The environment (`RtEnvironment`): a vertical gradient plus one sun lobe.
104 // The scene's only light unless a material emits. Its default is the soft
105 // studio sky this was before it was a parameter.
106 fn sky(rd: vec3<f32>) -> vec3<f32> {
107 let t = clamp(rd.y * 0.5 + 0.5, 0.0, 1.0);
108 var s = mix(params.sky_nadir.rgb, params.sky_zenith.rgb, t);
109 s = s + params.sun_color.rgb * pow(max(dot(rd, params.sun_dir.xyz), 0.0), 48.0);
110 return s;
111 }
112
113 fn cosine_dir(n: vec3<f32>, r1: f32, r2: f32) -> vec3<f32> {
114 let a = 6.28318530718 * r1;
115 let r = sqrt(r2);
116 var up = vec3<f32>(1.0, 0.0, 0.0);
117 if abs(n.x) > 0.5 {
118 up = vec3<f32>(0.0, 1.0, 0.0);
119 }
120 let tangent = normalize(cross(n, up));
121 let bitangent = cross(n, tangent);
122 return normalize(
123 tangent * (r * cos(a)) + bitangent * (r * sin(a)) + n * sqrt(max(0.0, 1.0 - r2)),
124 );
125 }
126
127 @compute @workgroup_size(8, 8)
128 fn cs_main(@builtin(global_invocation_id) gid: vec3<u32>) {
129 if gid.x >= params.width || gid.y >= params.height {
130 return;
131 }
132 let idx = gid.y * params.width + gid.x;
133
134 var total = vec3<f32>(0.0);
135 for (var s: u32 = 0u; s < params.spp; s = s + 1u) {
136 var rng: u32 = (idx * 9781u) ^ ((params.sample_index + s) * 26699u) ^ 0x9e3779b9u;
137
138 // Jittered primary ray, unprojected through inv_mvp (NDC y up, z 0..1).
139 let jx = rand(&rng);
140 let jy = rand(&rng);
141 let ndc_x = (f32(gid.x) + jx) / f32(params.width) * 2.0 - 1.0;
142 let ndc_y = 1.0 - (f32(gid.y) + jy) / f32(params.height) * 2.0;
143 let p_near = params.inv_mvp * vec4<f32>(ndc_x, ndc_y, 0.0, 1.0);
144 let p_far = params.inv_mvp * vec4<f32>(ndc_x, ndc_y, 1.0, 1.0);
145 var ro = p_near.xyz / p_near.w;
146 var rd = normalize(p_far.xyz / p_far.w - ro);
147 // The angle one pixel subtends, for the image's mip level: the
148 // ray through the next pixel along, against this one.
149 let p_next = params.inv_mvp
150 * vec4<f32>(ndc_x + 2.0 / f32(params.width), ndc_y, 1.0, 1.0);
151 let pixel_angle = length(normalize(p_next.xyz / p_next.w - ro) - rd);
152 let eye = ro;
153
154 var radiance = vec3<f32>(0.0);
155 var throughput = vec3<f32>(1.0);
156 // Until the path first lands on something or leaves for the sky:
157 // what it lands on is the pixel's feature for the denoiser. Not
158 // "the first bounce" — a ray let through the image's clear texels
159 // has used one and landed on nothing.
160 var primary = s == 0u;
161 // Until the path first scatters it is the camera's own ray, and a
162 // pixel's footprint on what it hits is known.
163 var straight = true;
164 for (var bounce: u32 = 0u; bounce < params.max_bounces; bounce = bounce + 1u) {
165 let hit = intersect_scene(ro, rd);
166 if hit.t >= 1e30 {
167 if primary {
168 features[2u * idx] = vec4<f32>(0.0, 0.0, 0.0, 1e30);
169 features[2u * idx + 1u] = vec4<f32>(1.0, 1.0, 1.0, 0.0);
170 }
171 // `straight`, not `primary`: every sample's camera ray,
172 // where `primary` is the one sample that writes features.
173 if straight && params.background.w > 0.5 {
174 radiance = radiance + throughput * params.background.rgb;
175 } else {
176 radiance = radiance + throughput * sky(rd);
177 }
178 break;
179 }
180 let tri = tris[hit.tri];
181 let mat = materials[bitcast<u32>(tri.p0.w)];
182 let at = ro + rd * hit.t;
183 let albedo = mat.albedo.rgb;
184 if mat.albedo.w > 0.5 {
185 // The image: a picture carries its own light. What a ray
186 // finds there is the image's colour as it is, neither lit
187 // by the sky nor shadowed by the scene — the colour the
188 // raster pass draws — and the path ends on it. Where the
189 // image is clear the ray goes on as if nothing were there:
190 // by chance, in proportion, which over the samples is the
191 // image's own alpha.
192 let rel = at - params.img_origin.xyz;
193 let u = params.img_u.xyz;
194 let v = params.img_v.xyz;
195 let uv = vec2<f32>(dot(rel, u) / dot(u, u), dot(rel, v) / dot(v, v));
196 // The level whose texel is a pixel's footprint wide. By the
197 // footprint's SHORT axis: seen at a slant the long one is
198 // averaged by the samples, where a level chosen for it
199 // would blur both. A scattered ray has no footprint and
200 // takes a coarse level.
201 var lod = 3.0;
202 if straight {
203 let footprint = length(at - eye) * pixel_angle;
204 let texel = length(u) / max(params.img_u.w, 1.0);
205 lod = max(log2(footprint / max(texel, 1e-12)), 0.0);
206 }
207 let texel = textureSampleLevel(img, img_sampler, uv, lod);
208 if rand(&rng) >= texel.a * params.img_origin.w {
209 ro = at + rd * (1e-4 * max(1.0, hit.t));
210 continue;
211 }
212 if primary {
213 // Marked as the sky is, with no depth: the colour is
214 // the image's own and has no noise to take out, so the
215 // denoiser passes it through and mixes it into nothing
216 // — smoothed, the image's fine print is the first
217 // thing to go.
218 features[2u * idx] = vec4<f32>(0.0, 0.0, 0.0, 1e30);
219 features[2u * idx + 1u] = vec4<f32>(texel.rgb, 0.0);
220 }
221 radiance = radiance + throughput * texel.rgb;
222 break;
223 }
224 radiance = radiance + throughput * mat.emission.rgb;
225 var n = normalize(cross(tri.p1.xyz - tri.p0.xyz, tri.p2.xyz - tri.p0.xyz));
226 if dot(n, rd) > 0.0 {
227 n = -n;
228 }
229 if primary {
230 features[2u * idx] = vec4<f32>(n, length(at - eye));
231 features[2u * idx + 1u] = vec4<f32>(albedo, 0.0);
232 }
233 primary = false;
234 straight = false;
235 throughput = throughput * albedo;
236 ro = at + n * 1e-4;
237 rd = cosine_dir(n, rand(&rng), rand(&rng));
238 }
239 // Firefly clamp: rare sun-spike paths otherwise leave speckles the
240 // variance can't average out (and the denoiser's edge-stopping
241 // weights deliberately refuse to smear). Slight energy loss on
242 // extreme highlights, big variance win.
243 total = total + min(radiance, vec3<f32>(4.0));
244 }
245
246 var acc = accum[idx];
247 if params.sample_index == 0u {
248 acc = vec4<f32>(0.0);
249 }
250 acc = acc + vec4<f32>(total, f32(params.spp));
251 accum[idx] = acc;
252 let color = acc.rgb / max(acc.a, 1.0);
253 textureStore(
254 out_img,
255 vec2<i32>(i32(gid.x), i32(gid.y)),
256 vec4<f32>(clamp(color, vec3<f32>(0.0), vec3<f32>(1.0)), 1.0),
257 );
258 }