git.lucas.co / cce-ui
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 }