From 14b5f3995349a5b5136ac4fba2a306cc414bd12e Mon Sep 17 00:00:00 2001 From: yujiaxu-ops Date: Sat, 14 Nov 2026 02:02:48 +0800 Subject: [PATCH] The live passes in a browser as natively: every pass at most eight storage buffers, merged.rs's live half gone The three passes over a browser's eight, rearranged; the browser's layout is now the native one, a layout a pass. - post (13): the view modes are passes of their own before it, not run in the lit picture (view_raw, view_hdr, view_sample); an HDR snapshot in a view mode is now that signal. - compose (11): the upscaler's state of each sample (Flicker, Lock) is a pass of its own after it, track. - dtemporal (10): the surfaces' compact copy (sb) is written by refl, which reads each sample whole anyway. - The sample's parts are read apart (sample_diffuse, _specular, _emission). - merged.rs: the slabs, the offsets' uniform and the live source's rewrite are gone; the kernel's outputs the passes read start at multiples of 256 bytes in its merged buffer (@align(256); the struct's size rounded to it), so a page binds those ranges. A unit test holds every pass to eight, and a view in the merged layout checks it as it starts. Checks: cargo test 125 passed. render-identity against the build before, native: 3 of 190 outputs differ by rounding (particles-motion-repose 277 pixels, at most ~3e-6 nits on dark ones; wind-pivot-sway-live, voxel-cut-body-push ~3e-5 relative); THREERT_MERGED=1: 3 (at most 7.5e-5 nits on values of 1 to 4). View modes 1-8: every PNG identical (the material tests). tools/web-hello.ts in Chrome: pass. Cost +0.18 ms (Bistro, 1440p at 0.375, alternated: track 0.27 less compose's 0.08; the surfaces' copy in refl +0.02). Co-Authored-By: Claude Opus 5.5 --- docs/ROADMAP.md | 22 +++++ docs/RUNTIME.md | 20 ++--- renderer/shaders/live.wgsl | 169 +++++++++++++++++++++++------------ renderer/src/gpu.rs | 6 +- renderer/src/live/frame.rs | 13 ++- renderer/src/live/layout.rs | 26 +++++- renderer/src/live/targets.rs | 133 +++++++-------------------- renderer/src/merged.rs | 140 ++++------------------------- 8 files changed, 229 insertions(+), 300 deletions(-) diff --git a/docs/ROADMAP.md b/docs/ROADMAP.md index a714dd8..b7551e9 100644 --- a/docs/ROADMAP.md +++ b/docs/ROADMAP.md @@ -2213,6 +2213,28 @@ is still the merged one, for all passes, until those three are rearranged. Check build before, native and merged, nothing changed; the passes after the trace 6.03-6.09 ms on both (Bistro, 1440p at 0.375, alternated). +**The live passes in a browser as natively (2026-11-14, after it).** The three passes over a browser's eight storage +buffers, rearranged so that each binds at most eight, and the browser's layout is now the native one, a layout a +pass: +- `post` (13): the view modes are passes of their own before it, which a frame in the lit picture does not run + (`view_raw`: base colour, normal, history length into the output as they are; `view_hdr` and `view_sample`: the + signals in place of the image `post` exposes). An HDR snapshot in a view mode is now that signal (it was the + lit image). +- `compose` (11): the upscaler's state of each sample (`Flicker`, `Lock`) is a pass of its own after it, `track`. +- `dtemporal` (10): the surfaces' compact copy (`sb`) is written by `refl`, which reads each sample whole anyway. +- The sample's parts are read apart (`sample_diffuse`, `_specular`, `_emission`): a pass binds the resolved reflection + only if it takes the specular part. +- merged.rs's live half is gone (the slabs, the offsets' uniform, the live source's rewrite); the kernel's outputs the + passes read start at multiples of 256 bytes in its merged buffer (`@align(256)`), so a page binds those ranges. + `every_pass_fits_a_browser` holds every pass to eight. + +Checks: render-identity against the build before, native: 3 of 190 outputs differ by rounding +(`particles-motion-repose` 277 pixels, at most ~3e-6 nits on dark ones; `wind-pivot-sway-live` and +`voxel-cut-body-push` ~3e-5 relative); `THREERT_MERGED=1`: 3 (at most 7.5e-5 nits on values of 1 to 4). The view modes +1-8: every PNG identical. `tools/web-hello.ts` in Chrome: pass. Cost +0.18 ms (Bistro, 1440p at 0.375, alternated: +the passes after the trace 6.00-6.08 -> 6.19-6.24; `track` 0.27 less `compose`'s 0.08; the surfaces' copy in `refl` ++0.02, a pass of its own read the samples again: 0.18). + ## 2. Temporal fixes Targets (2026-11-08): bench-temporal, bench-ghosting and their UE MRQ references move to Bistro night (§ 0 step 7); the diff --git a/docs/RUNTIME.md b/docs/RUNTIME.md index 32f8ea0..e6f0499 100644 --- a/docs/RUNTIME.md +++ b/docs/RUNTIME.md @@ -688,22 +688,20 @@ tree (ROADMAP § 6). the node host as `GameSession` does, and sends the canvas' clicks down the same path. `node tools/build-web.ts [game.ts] --serve` builds and serves it (needs the wasm-bindgen CLI of the crate's version). - **The merged binding layout** (`renderer/src/merged.rs`). WebGPU gives a stage 8 storage buffers by default; the - kernel binds 19 and the live passes 25. (Natively each live pass has a layout of its own, of the bindings its - entry point uses: `renderer/src/live/layout.rs`, from live.wgsl as build.rs reads it; three passes use more than - eight, so the browser's layout is still one for all, merged.) There is one source: the browser's shaders are written from - `pathtrace.wgsl` and `live.wgsl` as text when a module is compiled, and the native ones are untouched. + kernel binds 19. (The live passes need none of it: each has a layout of its own, of the bindings its entry point + uses, `renderer/src/live/layout.rs` from live.wgsl as build.rs reads it, and none binds more than eight, a unit + test's check; a page binds them as natively.) There is one source: the browser's kernel is written from + `pathtrace.wgsl` as text when a module is compiled, and the native one is untouched. - The kernel's seven tables become the arrays of one struct in one buffer; their lengths are in the source, each a power of two above a floor, so a kernel is compiled again only when a table doubles. The renderer keeps writing its own buffers and a frame copies them into the struct's. - The kernel's six group-1 buffers (the radiance cache, what the live passes read of a trace) are ranges of one - buffer, the arrays of a second struct. - - The live passes' buffers of one element type are ranges of a slab each (the two atomic ones, the exposure - histogram and the carried thin-feature locks, share one), indexed from offsets in a uniform. A - lasting buffer's two halves (this frame's and the last, `LiveParams.state`) lie in its range, as natively. + buffer, the arrays of a second struct; what the live passes read of it starts at multiples of 256 bytes, so they + bind those ranges as buffers of their own. - A use of a merged name the generator does not know, a declaration not found exactly once, a name it would add that is taken: each stops the renderer, with the line. Locals that shadow a name are followed through their - scopes. A test reads the result with naga: valid WGSL, 8 storage buffers for the kernel and for the live passes - (19 and 25 as written), the structs' arrays where the Rust side puts them. + scopes. A test reads the result with naga: valid WGSL, 8 storage buffers for the kernel (19 as written), the + structs' arrays where the Rust side puts them. - `THREERT_MERGED=1` runs it natively (software traversal, the device held to 8). - **The loop guard** (`merged::guard`). The kernel's traversals follow indices they read from buffers and end only on sound data. On 2026-11-11 two dispatches that did not return took this machine's display down (WindowServer's @@ -1211,7 +1209,7 @@ on the device", and differs where it says so. | `accum` (11), group 1 (radiance cache, the trace's outputs, the loop guard's count) | per pixel, per cache entry | written by the kernel: cannot be uniform | | group 2, the translucency pass (`timg`, `vpts`, `vcount`, `tlow`) | per pixel | written by the pass | | the particle passes' `systems`, `spawns`, `grid`; the grid programs' `srcs` | small, read-only | storage: other shaders, of eleven bindings at most and no acceleration structure; not known to misread (the fix on the device left them alone) | - | the live passes (`live.wgsl`: 25 storage buffers, `ae_bins` and `ae_state` among them) | per pixel; two tiny | storage: written by the passes; the device showed their output right | + | the live passes (`live.wgsl`: 25 storage buffers, `ae_bins` and `ae_state` among them; at most eight a pass) | per pixel; two tiny | storage: written by the passes; the device showed their output right | Two of these are what a device run should look at first: `prev_verts` is the one storage buffer left after the acceleration structure's binding (13) and feeds the motion vectors of moving objects and the vertex colours; diff --git a/renderer/shaders/live.wgsl b/renderer/shaders/live.wgsl index c097f0c..e999c25 100644 --- a/renderer/shaders/live.wgsl +++ b/renderer/shaders/live.wgsl @@ -43,7 +43,8 @@ struct Pixel { color: vec4, albedo: vec4, normal: vec4, emit: vec // ---- The live passes' buffers, by what they hold and how long. // // The declarations below are the passes' layout: build.rs reads them (with naga) when the renderer is built, and each -// pass binds what its entry point uses; live/layout.rs says what each binding is bound to, how large a buffer of the +// pass binds what its entry point uses, natively and in a browser alike (at most eight storage buffers a pass, a +// browser's: live/layout.rs's test); live/layout.rs says what each binding is bound to, how large a buffer of the // passes' own is, and the passes' order and grids. A buffer more: its declaration here, a line there. // // What lasts from one frame to the next has a buffer of its own, in two halves: this frame's, which the passes write @@ -56,23 +57,23 @@ struct Pixel { color: vec4, albedo: vec4, normal: vec4, emit: vec // gbuffer octahedral normal, view depth (0 = sky), roughness (`stemporal`; `GB`); once the temporal passes // have read the previous half, `lock_carry` puts its locks' depths into it (`judge` reads it) // color_hist the output grid's colour history: rgb, samples (`upscale`) -// up_state the upscaler's state per traced sample (`compose`): x, y `Flicker`; z, w `Lock` (its life and the +// up_state the upscaler's state per traced sample (`track`): x, y `Flicker`; z, w `Lock` (its life and the // view depth of what it is for; where in the traced pixel that is, f16 each) // What lives within a frame shares buffers whose uses do not overlap (the passes run in live/layout.rs `PASSES` order): // sample_vel per traced sample: screen velocity (output pixels this frame), view depth, animated shading // (`compose`; `judge`, `upscale`, depth of field and motion blur read it) -// judged `judge`'s findings per traced sample (`Judged`), in up_state's previous half: `compose` has read it +// judged `judge`'s findings per traced sample (`Judged`), in up_state's previous half: `track` has read it // by then, and the next frame writes the other (`judged_slot`) // fa, fb the irradiance filter's ping-pong on the traced grid; then fb the composed colour of the traced grid // (`compose`; w: first-hit distance) and fa the image on the output grid (what the post chain reads); // fb depth of field's output when motion blur follows -// sb the traced surfaces for the diffuse filter (`dtemporal` writes, the à-trous passes read), then the +// sb the traced surfaces for the diffuse filter (`refl` writes, the à-trous passes read), then the // upscaled image when depth of field / motion blur follow (w: the circle of confusion) // spec_pp the specular filter's ping-pong (`stemporal` -> `sfilt0..2`: two f16 slots, the result in slot 1), // then in slot 0 `reject`'s verdict // lumq depth of field's tiles (`dof_tiles`, `dof`), then the quarter-resolution log2 luminance // (`lum_down`, `lum_blur`) -// lock_fwd per traced sample: the previous frame's lock that lands on it (`lock_carry` writes, `compose` reads +// lock_fwd per traced sample: the previous frame's lock that lands on it (`lock_carry` writes, `track` reads // and clears it for the next frame) @group(0) @binding(2) var diffuse_hist: array>; @group(0) @binding(3) var color_hist: array>; @@ -155,21 +156,30 @@ fn lum(c: vec3) -> f32 { // ------------------------------------------------------------------------------------------------ temporal -// One traced sample's parts (averaged over this frame's samples): diffuse lighting, specular lighting, and what the -// camera ray saw directly (emission, sky). Lighting is clamped at `sample_clamp()` (an exposure-relative ceiling: -// fireflies far above white). Returns (diffuse, specular, emission). -fn sample_parts(q: u32) -> array, 3> { - let n = max(L.misc2.w, 1.0); - var d = sample_px[q].color.rgb / n; - var s = sample_px[q].spec.rgb / n + refl_res[q].rgb; - var e = sample_px[q].emit.rgb / n; +// One traced sample's parts (averaged over this frame's samples): diffuse lighting, specular lighting (with the +// resolved reflection, `refl`), and what the camera ray saw directly (emission, sky). Lighting is clamped at +// `sample_clamp()` (an exposure-relative ceiling: fireflies far above white). Each part on its own: a pass binds only +// what the part it takes reads. +fn sample_diffuse(q: u32) -> vec3 { + var d = sample_px[q].color.rgb / max(L.misc2.w, 1.0); if (non_finite(d)) { d = vec3(0.0); } - if (non_finite(s)) { s = vec3(0.0); } - if (non_finite(e)) { e = vec3(0.0); } let cl = sample_clamp(); if (lum(d) > cl) { d *= cl / lum(d); } + return d; +} + +fn sample_specular(q: u32) -> vec3 { + var s = sample_px[q].spec.rgb / max(L.misc2.w, 1.0) + refl_res[q].rgb; + if (non_finite(s)) { s = vec3(0.0); } + let cl = sample_clamp(); if (lum(s) > cl) { s *= cl / lum(s); } - return array, 3>(d, s, e); + return s; +} + +fn sample_emission(q: u32) -> vec3 { + var e = sample_px[q].emit.rgb / max(L.misc2.w, 1.0); + if (non_finite(e)) { e = vec3(0.0); } + return e; } // Roughness below which a surface counts as a near mirror: its reflection hit distance is the resolved one (pass `refl`). @@ -240,6 +250,10 @@ fn refl(@builtin(global_invocation_id) gid: vec3) { var rt = 0.0; let px = sample_px[i]; let t = px.color.w; + // (And a compact copy of the sample's surface for the diffuse filter's edge stops, `sb`, free until the upscaler: + // normal, first-hit distance, negative on animated emitters. Here because this pass reads the sample whole: the + // 80-byte samples cost the filter twice the bandwidth, and a pass of its own read them again, 0.18 ms.) + sb[i] = vec4(normalize(px.normal.xyz + vec3(1e-6)), select(t, -t, px.albedo.w > 1.5 && t > 0.0)); if (t > 0.0 && arrayLength(&refl_in) >= 2u * iw * ih) { let e = refl_in[i * 2u + 1u].rgb; if (any(e > vec3(0.0))) { @@ -325,8 +339,7 @@ fn fast_load(j: u32) -> vec3 { } fn sample_color(q: u32) -> vec3 { - let s = sample_parts(q); - return (s[0] + s[1]) * medium_tr(q) + s[2]; + return (sample_diffuse(q) + sample_specular(q)) * medium_tr(q) + sample_emission(q); } // The traced sample under an output pixel (nearest by the pixels' centres). @@ -415,13 +428,10 @@ fn dtemporal(@builtin(global_invocation_id) gid: vec3) { let t = px.color.w; let n = normalize(px.normal.xyz + vec3(1e-6)); let anim = px.albedo.w > 1.5; - // A compact copy of the sample's surface for the denoiser's and the upsample's edge stops (`sb` is free until - // `spec0`; the 80-byte samples cost twice the bandwidth): normal, first-hit distance, negative on animated emitters. - sb[q] = vec4(n, select(t, -t, anim && t > 0.0)); let sp = sample_pos(i32(gid.x), i32(gid.y)); let dir = ray_dir(L.cur, sp.x, sp.y); let p = select(L.cur.pos.xyz + dir * 1.0e6, L.cur.pos.xyz + dir * t, t > 0.0); - var irr = sample_parts(q)[0] / max(px.albedo.rgb, vec3(0.02)); + var irr = sample_diffuse(q) / max(px.albedo.rgb, vec3(0.02)); var l = lum(irr); var hd = vec4(0.0); var hm = vec4(0.0); @@ -762,7 +772,7 @@ fn stemporal(@builtin(global_invocation_id) gid: vec3) { p = L.cur.pos.xyz + dir * t; depth = dot(p - L.cur.pos.xyz, L.cur.fwd.xyz); } - let cur = sample_parts(q)[1]; + let cur = sample_specular(q); // This frame's neighbourhood mean (same surface), in YCoCg. var m1 = vec3(0.0); var mn = 0.0; @@ -779,7 +789,7 @@ fn stemporal(@builtin(global_invocation_id) gid: vec3) { ok = s.color.w == 0.0; } if (!ok) { continue; } - m1 += ycocg(sample_parts(j)[1]); + m1 += ycocg(sample_specular(j)); mn += 1.0; } let smean = m1 / max(mn, 1.0); @@ -911,7 +921,7 @@ fn compose(@builtin(global_invocation_id) gid: vec3) { let q = gid.y * iw + gid.x; let px = sample_px[q]; let t = px.color.w; - let c = (fa[q].rgb * max(px.albedo.rgb, vec3(0.02)) + spec_load(1u, q).rgb) * medium_tr(q) + sample_parts(q)[2] + direct_in[q].rgb; + let c = (fa[q].rgb * max(px.albedo.rgb, vec3(0.02)) + spec_load(1u, q).rgb) * medium_tr(q) + sample_emission(q) + direct_in[q].rgb; let sp = sample_pos(i32(gid.x), i32(gid.y)); let dir = ray_dir(L.cur, sp.x, sp.y); var p = L.cur.pos.xyz + dir * 1.0e6; @@ -931,8 +941,26 @@ fn compose(@builtin(global_invocation_id) gid: vec3) { } let shown = finite_or_zero(max(c, vec3(0.0))); fb[q] = vec4(shown, t); + // (None on animated shading: it shows this frame's picture, and its history is not kept for it to stand by.) + let reactive = max(select(0.0, 1.0, px.albedo.w > 1.5), clamp(medium_share(q) / UP_MEDIUM, 0.0, 1.0)); + sample_vel[q] = vec4(vel, depth, reactive); +} + +// The upscaler's state of each traced sample, after `compose`: how much it flickers (`Flicker`) and the thin feature's +// lock (`Lock`), into this frame's half of `up_state`. +@compute @workgroup_size(8, 8, 1) +fn track(@builtin(global_invocation_id) gid: vec3) { + let iw = L.isz.x; + let ih = L.isz.y; + if (gid.x >= iw || gid.y >= ih) { return; } + let q = gid.y * iw + gid.x; + let sp = sample_pos(i32(gid.x), i32(gid.y)); + let sv = sample_vel[q]; + let vel = sv.xy; + let depth = sv.z; + let reactive = sv.w; // How much this sample flickers (`Flicker`), from where its surface was in the previous frame's traced grid. - let now = min(lum(shown), 6.0e4); + let now = min(lum(fb[q].rgb), 6.0e4); var f = Flicker(vec2(0.0), 0.0, 0.0); let ps = vec2(L.isz.zw); if (L.size.w == 0u && ps.x > 0) { @@ -966,8 +994,6 @@ fn compose(@builtin(global_invocation_id) gid: vec3) { across = across || (e9[4] > THIN_LIKE * max(a, b)) || (THIN_LIKE * e9[4] < min(a, b)); } let thin = across && (like & 27u) != 27u && (like & 54u) != 54u && (like & 216u) != 216u && (like & 432u) != 432u; - // (None on animated shading: it shows this frame's picture, and its history is not kept for it to stand by.) - let reactive = max(select(0.0, 1.0, px.albedo.w > 1.5), clamp(medium_share(q) / UP_MEDIUM, 0.0, 1.0)); let fw = atomicLoad(&lock_fwd[q]); atomicStore(&lock_fwd[q], 0u); var was = Carried(0.0, 0.0, 0.0, vec2(0.0), false); @@ -982,7 +1008,6 @@ fn compose(@builtin(global_invocation_id) gid: vec3) { } life *= 1.0 - reactive; up_state[this_half(q)] = vec4(pack2x16float(vec2(now, select(now, f.last.x, f.n >= 1.0))), pack2x16float(vec2(sd * still, min(max(f.n + still, 1.0), 9.0))), pack2x16float(vec2(life, min(lz, 6.0e4))), pack2x16float(at)); - sample_vel[q] = vec4(vel, depth, reactive); } // The upscaler works on linear radiance and its means are plain ones. (Weighted by 1 / (1 + luminance x exposure), so @@ -1043,7 +1068,7 @@ const UP_STALE: f32 = 0.05; // traced sample at this place changes from one frame to the next under the jitter, which is what a temporal // upscaler's flicker history is for (UE's TSR keeps the input's luminance over time for the same decision). -// `Flicker`, per traced sample, in `up_state`'s x and y (f16; `compose` keeps it): the composed sample's luminance of +// `Flicker`, per traced sample, in `up_state`'s x and y (f16; `track` keeps it): the composed sample's luminance of // the last frame and of the one before, the deviation of its frame-to-frame change as a share of its level, and of how // many frames that is (0: none; the deviation is then 1: unknown, as much as the level), all read where the sample's // surface was a frame ago (in motion the place on screen would read the motion as flicker). @@ -1117,7 +1142,7 @@ fn carry(p: u32) -> Carried { // Each lock of the previous frame onto the traced sample its point falls in now (`Lock`), the longest-lived where two // fall in one (the later sample where they tie): an atomic max of (life in 7 bits, the sample + 1 in 25: an 8K grid), -// which `compose` reads. And the previous frame's g-buffer learns what was locked there: the nearest thing in that +// which `track` reads. And the previous frame's g-buffer learns what was locked there: the nearest thing in that // traced pixel was the lamp, which its sample most frames missed (`judge`'s test for something gone from in front // reads it: the sky a lamp has passed over keeps the lamp in its history otherwise, which no sample there contradicts, // since the history's pixels around a sparse lamp are dark too). Over the previous frame's traced grid (live/frame.rs). @@ -1857,6 +1882,60 @@ fn mblur(@builtin(global_invocation_id) gid: vec3) { fa[i] = vec4(outv, 0.0); } +// ------------------------------------------------------------------------------------------------ view modes + +// Unreal-style view modes (`L.mode.x`), passes of their own before `post` that run only in them (live/layout.rs +// `PASSES`), so that `post` binds only what the picture needs. `view_raw`: base colour (1), world normal (2) and the +// diffuse history's length (5) of the traced sample under each output pixel, as they are, into the output (`post` +// leaves it). +@compute @workgroup_size(8, 8, 1) +fn view_raw(@builtin(global_invocation_id) gid: vec3) { + let mode = L.mode.x; + if (gid.x >= L.size.x || gid.y >= L.size.y || (mode != 1u && mode != 2u && mode != 5u)) { return; } + let tq = sample_under(gid.x, gid.y); + var d = vec3(0.0); + if (mode == 1u) { + let a = clamp(sample_px[tq].albedo.rgb, vec3(0.0), vec3(1.0)); + d = select(pow((a + 0.055) / 1.055, vec3(1.0 / 2.4)) * 1.055 - 0.055, a * 12.92, a <= vec3(0.0031308)); + } else if (mode == 2u) { + let g = gb_unpack(gbuffer[this_half(tq)]); + d = select(g.n * 0.5 + 0.5, vec3(0.0), g.depth == 0.0); + } else if (mode == 5u) { + let t = clamp(log2(max(diffuse_hist[this_half(tq)].w, 1.0)) / 12.0, 0.0, 1.0); // 1 .. 4096 frames (diffuse) + d = clamp(vec3(1.5 - abs(4.0 * t - 3.0), 1.5 - abs(4.0 * t - 2.0), 1.5 - abs(4.0 * t - 1.0)), vec3(0.0), vec3(1.0)); + } else { + return; + } + let q = vec3(clamp(d, vec3(0.0), vec3(1.0)) * 255.0 + 0.5); + out_px[gid.y * L.size.x + gid.x] = q.z | (q.y << 8u) | (q.x << 16u) | (255u << 24u); +} + +// `view_hdr`: in place of the image `post` exposes and tonemaps (`fa`), the traced sample under the pixel: its +// accumulations without the spatial filters (3), and the denoised signals apart: diffuse (6: the composed colour less +// the rest), specular (7), deterministic direct light (8). (6 does not hold with depth of field and motion blur both +// on: the composed colour's buffer is theirs then.) +@compute @workgroup_size(8, 8, 1) +fn view_hdr(@builtin(global_invocation_id) gid: vec3) { + let mode = L.mode.x; + if (gid.x >= L.size.x || gid.y >= L.size.y || (mode != 3u && mode != 6u && mode != 7u && mode != 8u)) { return; } + let i = gid.y * L.size.x + gid.x; + let tq = sample_under(gid.x, gid.y); + var src = fa[i].rgb; + if (mode == 3u) { src = (diffuse_hist[this_half(tq)].rgb * max(sample_px[tq].albedo.rgb, vec3(0.02)) + spec_hist[this_half(tq)].rgb) * medium_tr(tq) + sample_emission(tq) + direct_in[tq].rgb; } + if (mode == 6u) { src = max(fb[tq].rgb - spec_load(1u, tq).rgb * medium_tr(tq) - sample_emission(tq) - direct_in[tq].rgb, vec3(0.0)); } + if (mode == 8u) { src = direct_in[tq].rgb; } + if (mode == 7u) { src = spec_load(1u, tq).rgb * medium_tr(tq); } + fa[i] = vec4(src, fa[i].w); +} + +// `view_sample`: the same with the raw sample under the pixel (4). +@compute @workgroup_size(8, 8, 1) +fn view_sample(@builtin(global_invocation_id) gid: vec3) { + if (gid.x >= L.size.x || gid.y >= L.size.y || L.mode.x != 4u) { return; } + let i = gid.y * L.size.x + gid.x; + fa[i] = vec4(sample_color(sample_under(gid.x, gid.y)), fa[i].w); +} + // ------------------------------------------------------------------------------------------------ post // UE LinToLog (PostProcessCombineLUTs): LUT input encoding for the 32^3 tonemap LUT. @@ -1873,33 +1952,9 @@ fn post(@builtin(global_invocation_id) gid: vec3) { let x = f32(gid.x) + 0.5; let y = f32(gid.y) + 0.5; let mode = L.mode.x; - // Unreal-style view modes. Base color / normal / history bypass exposure + tonemapping. - if (mode == 1u || mode == 2u || mode == 5u) { - var d = vec3(0.0); - if (mode == 1u) { - let a = clamp(sample_px[sample_under(gid.x, gid.y)].albedo.rgb, vec3(0.0), vec3(1.0)); - d = select(pow((a + 0.055) / 1.055, vec3(1.0 / 2.4)) * 1.055 - 0.055, a * 12.92, a <= vec3(0.0031308)); - } else if (mode == 2u) { - let g = gb_unpack(gbuffer[this_half(sample_under(gid.x, gid.y))]); - d = select(g.n * 0.5 + 0.5, vec3(0.0), g.depth == 0.0); - } else { - let t = clamp(log2(max(diffuse_hist[this_half(sample_under(gid.x, gid.y))].w, 1.0)) / 12.0, 0.0, 1.0); // 1 .. 4096 frames (diffuse) - d = clamp(vec3(1.5 - abs(4.0 * t - 3.0), 1.5 - abs(4.0 * t - 2.0), 1.5 - abs(4.0 * t - 1.0)), vec3(0.0), vec3(1.0)); - } - let q = vec3(clamp(d, vec3(0.0), vec3(1.0)) * 255.0 + 0.5); - out_px[i] = q.z | (q.y << 8u) | (q.x << 16u) | (255u << 24u); - return; - } + // (The view modes that show a buffer as it is are `view_raw`'s picture.) + if (mode == 1u || mode == 2u || mode == 5u) { return; } var src = fa[i].rgb; - // The traced sample under the pixel: its accumulations without the spatial filters (3), the raw sample (4), and the - // denoised signals apart: diffuse (6: the composed colour less the rest), specular (7), deterministic direct light (8). - // (6 does not hold with depth of field and motion blur both on: the composed colour's buffer is theirs then.) - let tq = sample_under(gid.x, gid.y); - if (mode == 3u) { src = (diffuse_hist[this_half(tq)].rgb * max(sample_px[tq].albedo.rgb, vec3(0.02)) + spec_hist[this_half(tq)].rgb) * medium_tr(tq) + sample_parts(tq)[2] + direct_in[tq].rgb; } - if (mode == 4u) { src = sample_color(tq); } - if (mode == 6u) { src = max(fb[tq].rgb - spec_load(1u, tq).rgb * medium_tr(tq) - sample_parts(tq)[2] - direct_in[tq].rgb, vec3(0.0)); } - if (mode == 8u) { src = direct_in[tq].rgb; } - if (mode == 7u) { src = spec_load(1u, tq).rgb * medium_tr(tq); } if (L.post.x > 0.0 && mode == 0u) { // Scene fringe: red sampled radially outward, blue inward (nearest texel). let k = L.post.x * 0.0025; diff --git a/renderer/src/gpu.rs b/renderer/src/gpu.rs index 5901716..214bd02 100644 --- a/renderer/src/gpu.rs +++ b/renderer/src/gpu.rs @@ -1447,7 +1447,7 @@ impl Gpu { // layout with more does not validate). let merged = cfg!(target_arch = "wasm32") || std::env::var("THREERT_MERGED").is_ok_and(|v| v == "1"); limits.max_storage_buffers_per_shader_stage = if merged { crate::merged::STORAGE_BUFFERS } else { al.max_storage_buffers_per_shader_stage.min(31) }; - // Storage + uniform buffers together (the kernel binds 19 storage buffers, the live passes 25, each beside its + // Storage + uniform buffers together (the kernel binds 19 storage buffers, a live pass at most 8, each beside its // uniforms; Metal allows 31). limits.max_buffers_and_acceleration_structures_per_shader_stage = al.max_buffers_and_acceleration_structures_per_shader_stage.min(31); limits.max_acceleration_structures_per_shader_stage = al.max_acceleration_structures_per_shader_stage; @@ -3297,8 +3297,8 @@ mod tests { assert_eq!(loop_guard(), None); let sha = |s: &str| Sha256::digest(s.as_bytes()).iter().map(|b| format!("{b:02x}")).collect::(); let live = include_str!("../shaders/live.wgsl"); - assert_eq!(crate::live::live_source(false), live); - println!("source live {}", sha(&crate::live::live_source(false))); + assert_eq!(crate::live::live_source(), live); + println!("source live {}", sha(&crate::live::live_source())); for hwrt in [false, true] { let base = kernel_base(hwrt); assert_eq!(guarded(&base, false, false), base); diff --git a/renderer/src/live/frame.rs b/renderer/src/live/frame.rs index 299c62a..b927b38 100644 --- a/renderer/src/live/frame.rs +++ b/renderer/src/live/frame.rs @@ -556,10 +556,12 @@ impl LiveCore { for (k, (name, grid, pipe)) in self.passes.pipes.iter().enumerate() { let q = (k as u32 + 1) * 2; let mut pass = e.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some(name), timestamp_writes: Some(wgpu::ComputePassTimestampWrites { query_set: &p.queries, beginning_of_pass_write_index: Some(q), end_of_pass_write_index: Some(q + 1) }) }); - pass.set_pipeline(pipe); - pass.set_bind_group(0, &self.sc.bgs[k], &[]); - let (gw, gh) = grid.size((w, h), (iw, ih), self.prev_isz); - pass.dispatch_workgroups(gw.div_ceil(8), gh.div_ceil(8), 1); + if self.view_mode != 0 || !layout::VIEW_MODE_PASSES.contains(name) { + pass.set_pipeline(pipe); + pass.set_bind_group(0, &self.sc.bgs[k], &[]); + let (gw, gh) = grid.size((w, h), (iw, ih), self.prev_isz); + pass.dispatch_workgroups(gw.div_ceil(8), gh.div_ceil(8), 1); + } drop(pass); if *name == "upscale" && tl.is_some() { for (j, train) in [(0u32, false), (1, true)] { @@ -576,6 +578,9 @@ impl LiveCore { let tw = self.timer.as_ref().zip(timed).map(|(t, slot)| wgpu::ComputePassTimestampWrites { query_set: &t.queries, beginning_of_pass_write_index: Some(slot as u32 * 8 + 2), end_of_pass_write_index: Some(slot as u32 * 8 + 3) }); let mut pass = e.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("live post"), timestamp_writes: tw }); for (k, (name, grid, pipe)) in self.passes.pipes.iter().enumerate() { + if self.view_mode == 0 && layout::VIEW_MODE_PASSES.contains(name) { + continue; + } pass.set_pipeline(pipe); pass.set_bind_group(0, &self.sc.bgs[k], &[]); let (gw, gh) = grid.size((w, h), (iw, ih), self.prev_isz); diff --git a/renderer/src/live/layout.rs b/renderer/src/live/layout.rs index d0979fa..18e2f26 100644 --- a/renderer/src/live/layout.rs +++ b/renderer/src/live/layout.rs @@ -2,8 +2,9 @@ //! `BINDINGS`, each binding's name, number, kind and element stride; `ENTRY_POINTS`, the bindings each entry point //! uses); here is only what the shader cannot say: what each binding is bound to and how many elements a buffer of //! the passes' own holds (`RESOURCES`), and the passes' order and the grid each runs over (`PASSES`). `check` holds -//! the two to each other (a unit test, and every view as it starts). A buffer or a pass more: its declaration or -//! entry point in live.wgsl, a line here. +//! the two to each other (a unit test, and every view as it starts). Each pass binds what its entry point uses, and +//! none more than a browser's eight storage buffers (the test below), so a page binds them as they are natively. A buffer +//! or a pass more: its declaration or entry point in live.wgsl, a line here. /// A binding live.wgsl declares. pub(super) struct Binding { @@ -117,6 +118,7 @@ pub(super) const PASSES: &[(&str, Grid)] = &[ ("sfilt2", Grid::Traced), ("lock_carry", Grid::PrevTraced), ("compose", Grid::Traced), + ("track", Grid::Traced), ("judge", Grid::Traced), ("reject", Grid::Traced), ("upscale", Grid::Output), @@ -143,9 +145,16 @@ pub(super) const PASSES: &[(&str, Grid)] = &[ ("bloom_v1", Grid::Quarter), ("bloom_h0", Grid::Quarter), ("bloom_v0", Grid::Quarter), + ("view_raw", Grid::Output), + ("view_hdr", Grid::Output), + ("view_sample", Grid::Output), ("post", Grid::Output), ]; +/// The passes of the view modes (`LiveParams::mode`.x above 0): a frame in the lit picture does not run them (the +/// per-pass profile times an empty pass for each, to keep its timestamps in `PASSES`' order). +pub(super) const VIEW_MODE_PASSES: &[&str] = &["view_raw", "view_hdr", "view_sample"]; + impl Grid { /// Invocations, for an output of `out`, a traced grid of `traced` and the previous frame's `prev`. pub(super) fn size(self, out: (u32, u32), traced: (u32, u32), prev: (u32, u32)) -> (u32, u32) { @@ -217,10 +226,23 @@ pub(super) fn check() { assert!(problems.is_empty(), "the live passes' layout (live/layout.rs) and live.wgsl disagree:\n {}", problems.join("\n ")); } +/// The storage buffers pass `name` binds. +pub(super) fn storage_count(name: &str) -> usize { + uses(name).iter().filter(|&&b| BINDINGS.iter().any(|x| x.binding == b && matches!(x.kind, Kind::Storage { .. }))).count() +} + #[cfg(test)] mod tests { #[test] fn the_table_and_the_shader_agree() { super::check(); } + + /// Every pass binds at most the storage buffers a browser gives a stage (WebGPU's default: merged.rs + /// `STORAGE_BUFFERS`), so that a page binds them as they are. + #[test] + fn every_pass_fits_a_browser() { + let over: Vec = super::PASSES.iter().map(|p| (p.0, super::storage_count(p.0))).filter(|p| p.1 as u32 > crate::merged::STORAGE_BUFFERS).map(|p| format!("{} {}", p.0, p.1)).collect(); + assert!(over.is_empty(), "passes over a browser's storage buffers: {over:?}"); + } } diff --git a/renderer/src/live/targets.rs b/renderer/src/live/targets.rs index 16e660a..ea3535e 100644 --- a/renderer/src/live/targets.rs +++ b/renderer/src/live/targets.rs @@ -47,18 +47,14 @@ pub(super) fn tonemap_lut(p: &crate::desc::PostPreset) -> Vec { out } -/// The live passes' source: live.wgsl as it is, or in the merged layout its buffers merged into slabs (merged.rs). -pub(crate) fn live_source(merged: bool) -> std::borrow::Cow<'static, str> { - let src = include_str!("../../shaders/live.wgsl"); - if merged { - crate::merged::live(src).into() - } else { - src.into() - } +/// The live passes' source: live.wgsl, natively and in a browser alike (each pass binds at most a browser's storage +/// buffers: `layout`). +pub(crate) fn live_source() -> &'static str { + include_str!("../../shaders/live.wgsl") } pub(super) struct Passes { - /// Each pass's bind group layout: the bindings it uses. In the merged layout (merged.rs) one for all. + /// Each pass's bind group layout: the bindings it uses. pub(super) layouts: Vec, /// The passes in `layout::PASSES`' order, with the grid each runs over. pub(super) pipes: Vec<(&'static str, layout::Grid, wgpu::ComputePipeline)>, @@ -78,42 +74,33 @@ fn layout_entry(b: &layout::Binding) -> wgpu::BindGroupLayoutEntry { impl Passes { pub(super) fn new(g: &gpu::Gpu) -> Passes { layout::check(); + if g.merged { + // (A browser's stage binds at most eight storage buffers; the device here is held to them.) + let over: Vec = layout::PASSES.iter().map(|p| (p.0, layout::storage_count(p.0))).filter(|p| p.1 as u32 > crate::merged::STORAGE_BUFFERS).map(|p| format!("{} binds {}", p.0, p.1)).collect(); + assert!(over.is_empty(), "live passes over a browser's storage buffers: {}", over.join(", ")); + } let d = &g.device; - let module = d.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("live"), source: wgpu::ShaderSource::Wgsl(live_source(g.merged).into()) }); + let module = d.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("live"), source: wgpu::ShaderSource::Wgsl(live_source().into()) }); let pipe = |name: &str, layout: &wgpu::BindGroupLayout| { let pl = d.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: None, bind_group_layouts: &[Some(layout)], immediate_size: 0 }); d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some(name), layout: Some(&pl), module: &module, entry_point: Some(name), compilation_options: Default::default(), cache: g.cache() }) }; - let (layouts, pipes) = if g.merged { - // The merged layout: one for all passes, the buffers that stay, the slabs, and the offsets' uniform. - use crate::merged::{LIVE_KEPT, OFFSETS_BINDING, SLABS}; - let mut entries: Vec = layout::BINDINGS.iter().map(layout_entry).collect(); - let storage = |e: &wgpu::BindGroupLayoutEntry| matches!(e.ty, wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { .. }, .. }); - entries.retain(|e| !storage(e) || LIVE_KEPT.iter().any(|k| k.0 == e.binding)); - entries.extend(SLABS.iter().map(|s| layout_entry(&layout::Binding { name: s.name, binding: s.binding, kind: layout::Kind::Storage { read_only: s.read_only, stride: s.stride, len: 0 } }))); - entries.push(wgpu::BindGroupLayoutEntry { binding: OFFSETS_BINDING, ..layout_entry(layout::binding("L")) }); - let all = d.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: Some("live"), entries: &entries }); - let pipes = layout::PASSES.iter().map(|&(name, grid)| (name, grid, pipe(name, &all))).collect(); - (vec![all], pipes) - } else { - // A layout a pass, of the bindings it uses. - let mut layouts = vec![]; - let mut pipes = vec![]; - for &(name, grid) in layout::PASSES { - let entries: Vec = layout::uses(name).iter().map(|&b| layout_entry(layout::BINDINGS.iter().find(|x| x.binding == b).expect("a binding of live.wgsl"))).collect(); - let l = d.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: Some(name), entries: &entries }); - pipes.push((name, grid, pipe(name, &l))); - layouts.push(l); - } - (layouts, pipes) - }; + // A layout a pass, of the bindings it uses. + let mut layouts = vec![]; + let mut pipes = vec![]; + for &(name, grid) in layout::PASSES { + let entries: Vec = layout::uses(name).iter().map(|&b| layout_entry(layout::BINDINGS.iter().find(|x| x.binding == b).expect("a binding of live.wgsl"))).collect(); + let l = d.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: Some(name), entries: &entries }); + pipes.push((name, grid, pipe(name, &l))); + layouts.push(l); + } g.save_pipeline_cache(); Passes { layouts, pipes } } /// The layout pass `k` binds. pub(super) fn layout(&self, k: usize) -> &wgpu::BindGroupLayout { - &self.layouts[if self.layouts.len() == 1 { 0 } else { k }] + &self.layouts[k] } } @@ -129,8 +116,6 @@ pub struct Targets { /// a trace: the specular-lobe rays (`refl`), motion vectors per traced pixel (pathtrace.wgsl `motion_out`) and the /// deterministic direct light per traced pixel (`direct_out`). pub kernel: gpu::KernelBufs, - /// The merged layout (merged.rs): the slabs the buffers above are ranges of (`SLABS` but the kernel's). - pub(super) slabs: Option<[wgpu::Buffer; 5]>, pub(super) lut_tex: wgpu::Texture, pub(super) lut_view: wgpu::TextureView, pub(super) sampler: wgpu::Sampler, @@ -139,7 +124,6 @@ pub struct Targets { impl Targets { /// For a view of `w` x `h` and a radiance cache of `cache` entries. pub(super) fn new(g: &gpu::Gpu, w: u32, h: u32, lut: &[u16], cache: u64) -> Targets { - use crate::merged::{LIVE_BUFS, SLABS, SLAB_K}; let d = &g.device; let n = (w * h) as u64; let sb = |label: &str, bytes: u64| d.create_buffer(&wgpu::BufferDescriptor { label: Some(label), size: bytes, usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, mapped_at_creation: false }); @@ -162,10 +146,8 @@ impl Targets { wgpu::TexelCopyBufferLayout { offset: 0, bytes_per_row: Some(LUT * 8), rows_per_image: Some(LUT) }, wgpu::Extent3d { width: LUT, height: LUT, depth_or_array_layers: LUT }, ); - // The passes' own buffers, in binding order: their bytes, and in the merged layout the slab of their element - // type (merged.rs `LIVE_BUFS`; none: a buffer that stays, `out_px`). Each range in a slab starts at a multiple - // of 16 bytes: of every slab's element. - let own: Vec<(&layout::Binding, u64, Option)> = layout::BINDINGS + // The passes' own buffers, in binding order. + let bufs: Vec<(u32, Buf)> = layout::BINDINGS .iter() .filter_map(|b| { let layout::Res::Own(count) = layout::resource(b) else { return None }; @@ -176,26 +158,7 @@ impl Targets { layout::Count::Bloom => bloom_texels, layout::Count::Declared => len, }; - Some((b, elements * stride, LIVE_BUFS.iter().find(|m| m.name == b.name).map(|m| m.slab))) - }) - .collect(); - let slabs: Option<[wgpu::Buffer; 5]> = g.merged.then(|| { - debug_assert_eq!(SLAB_K, 5, "the owned slabs are the first five"); - std::array::from_fn(|k| sb(SLABS[k].name, own.iter().filter(|o| o.2 == Some(k)).map(|o| o.1.next_multiple_of(16)).sum())) - }); - let mut at = [0u64; 5]; - let bufs: Vec<(u32, Buf)> = own - .iter() - .map(|&(b, bytes, slab)| { - let buf = match (&slabs, slab) { - (Some(s), Some(k)) => { - let r = Buf { buffer: s[k].clone(), offset: at[k], size: bytes }; - at[k] += bytes.next_multiple_of(16); - r - } - _ => Buf::whole(sb(b.name, bytes)), - }; - (b.binding, buf) + Some((b.binding, Buf::whole(sb(b.name, elements * stride)))) }) .collect(); let out = bufs.iter().find(|b| b.0 == layout::binding("out_px").binding).expect("out_px").1.buffer.clone(); @@ -204,7 +167,6 @@ impl Targets { bufs, out, kernel: gpu::KernelBufs::new(d, g.merged, g.guard_room(), cache, n * 32, n * 16, n * 16), - slabs, lut_view: lut_tex.create_view(&Default::default()), lut_tex, sampler: d.create_sampler(&wgpu::SamplerDescriptor { mag_filter: wgpu::FilterMode::Linear, min_filter: wgpu::FilterMode::Linear, ..Default::default() }), @@ -226,8 +188,7 @@ impl Targets { } } - /// The bind group of each pass, in `Passes::pipes`' order (`sample`: the trace's samples). In the merged layout - /// (merged.rs) one for all, the slabs and the offsets the passes read (`LO`). + /// The bind group of each pass, in `Passes::pipes`' order (`sample`: the trace's samples). pub(super) fn bind_groups(&self, d: &wgpu::Device, passes: &Passes, sample: &wgpu::Buffer) -> Vec { let entry = |b: &'_ layout::Binding| -> wgpu::BindGroupEntry<'_> { let resource = match layout::resource(b) { @@ -239,40 +200,14 @@ impl Targets { }; wgpu::BindGroupEntry { binding: b.binding, resource } }; - match (&self.slabs, &self.kernel.merged) { - (Some(slabs), Some((live, _))) => { - use crate::merged::{LIVE_BUFS, LIVE_KEPT, OFFSETS_BINDING, SLABS, SLAB_K}; - // Each merged buffer's (first element, length) in its slab, in `MergedOffsets`' order, and the - // struct's unused last pair (`merged::OFFSETS_BYTES`). - let mut lo: Vec<[u32; 2]> = LIVE_BUFS - .iter() - .map(|m| { - let buf = self.buf(m.name); - let stride = SLABS[m.slab].stride; - debug_assert!(buf.offset % stride == 0 && buf.size % stride == 0); - [(buf.offset / stride) as u32, (buf.size / stride) as u32] - }) - .collect(); - lo.push([0, 0]); - debug_assert_eq!(lo.len() * 8, crate::merged::OFFSETS_BYTES); - let offsets = wgpu::util::DeviceExt::create_buffer_init(d, &wgpu::util::BufferInitDescriptor { label: Some("live offsets"), contents: bytemuck::cast_slice(&lo), usage: wgpu::BufferUsages::UNIFORM }); - let storage = |b: &layout::Binding| matches!(b.kind, layout::Kind::Storage { .. }); - let mut entries: Vec = layout::BINDINGS.iter().filter(|b| !storage(b) || LIVE_KEPT.iter().any(|k| k.0 == b.binding)).map(entry).collect(); - entries.extend(slabs.iter().zip(&SLABS).map(|(s, m)| wgpu::BindGroupEntry { binding: m.binding, resource: s.as_entire_binding() })); - entries.push(wgpu::BindGroupEntry { binding: SLABS[SLAB_K].binding, resource: live.as_entire_binding() }); - entries.push(wgpu::BindGroupEntry { binding: OFFSETS_BINDING, resource: offsets.as_entire_binding() }); - let all = d.create_bind_group(&wgpu::BindGroupDescriptor { label: Some("live"), layout: passes.layout(0), entries: &entries }); - vec![all; passes.pipes.len()] - } - _ => passes - .pipes - .iter() - .enumerate() - .map(|(k, (name, ..))| { - let entries: Vec = layout::uses(name).iter().map(|&b| entry(layout::BINDINGS.iter().find(|x| x.binding == b).expect("a binding of live.wgsl"))).collect(); - d.create_bind_group(&wgpu::BindGroupDescriptor { label: Some(name), layout: passes.layout(k), entries: &entries }) - }) - .collect(), - } + passes + .pipes + .iter() + .enumerate() + .map(|(k, (name, ..))| { + let entries: Vec = layout::uses(name).iter().map(|&b| entry(layout::BINDINGS.iter().find(|x| x.binding == b).expect("a binding of live.wgsl"))).collect(); + d.create_bind_group(&wgpu::BindGroupDescriptor { label: Some(name), layout: passes.layout(k), entries: &entries }) + }) + .collect() } } diff --git a/renderer/src/merged.rs b/renderer/src/merged.rs index 9ae0c3f..d3a9d95 100644 --- a/renderer/src/merged.rs +++ b/renderer/src/merged.rs @@ -1,7 +1,8 @@ -//! The merged binding layout: the engine's shaders with at most eight storage buffers a stage, WebGPU's default limit -//! (`maxStorageBuffersPerShaderStage`). The kernel binds 19 as it is written and the live passes 25: more than a -//! browser gives a page. There is one source: pathtrace.wgsl and live.wgsl, and the bindings gpu.rs and live/targets.rs make for -//! them, are what the native renderer runs, unchanged. This module writes the browser's from them, as text (as +//! The merged binding layout: the kernel with at most eight storage buffers a stage, WebGPU's default limit +//! (`maxStorageBuffersPerShaderStage`). The kernel binds 19 as it is written: more than a browser gives a page. (The +//! live passes need none of this: each binds at most eight, live/layout.rs, and a page binds them as they are.) There +//! is one source: pathtrace.wgsl, and the bindings gpu.rs makes for it, are what the native renderer runs, unchanged. +//! This module writes the browser's from them, as text (as //! `gpu::kernel_source` patches the kernel for a scene's features), and gives the Rust side the numbers that go with //! it. `THREERT_MERGED=1` runs the same layout natively (software traversal), so the tests that check the renderer //! check it too (docs/RUNTIME.md "The engine in a browser"). @@ -14,10 +15,9 @@ //! the struct's (`Gpu::sync_tables`; they are small). //! - The kernel's six group-1 buffers (`KERNEL_LIVE`: the radiance cache, what the live passes read //! of a trace) become the arrays of `ML: MergedLive`, one buffer that they are ranges of (`gpu::KernelBufs`). -//! - The live passes' buffers of one element type are ranges of one slab each (`SLABS`), indexed from an offset -//! the pass reads from a uniform (`LO`): the source does not change with the size of the view. (The lasting -//! buffers' two halves alternate by a uniform, live.wgsl `this_half`.) `MK` is the kernel's `ML` buffer read as -//! vec4s. +//! What the live passes read of it starts at a multiple of 256 bytes (`@align(256)`): the passes bind those +//! ranges as buffers of their own, and a binding's offset is a multiple of that (WebGPU's +//! `minStorageBufferOffsetAlignment`). //! //! A use of a merged name the generator does not know (anything but `name[...]`, `&name[...]` and //! `arrayLength(&name)`), a declaration that is not found exactly once, a name the generator adds that the source @@ -31,7 +31,8 @@ pub const STORAGE_BUFFERS: u32 = 8; /// One of the kernel's buffers that becomes an array of a merged struct. pub struct Member { pub name: &'static str, - /// Its element type, the element's size and alignment in a storage buffer. + /// Its element type and the element's size in a storage buffer; where its array may start in the struct: its + /// element's alignment, or more (written `@align`). pub ty: &'static str, pub stride: u64, pub align: u64, @@ -58,9 +59,9 @@ pub const KERNEL_LIVE: [Member; 6] = [ Member { name: "cache_keys", ty: "atomic", stride: 4, align: 4, decl: "@group(1) @binding(1) var cache_keys: array>;" }, Member { name: "cache_acc", ty: "atomic", stride: 4, align: 4, decl: "@group(1) @binding(2) var cache_acc: array>;" }, Member { name: "cache_val", ty: "vec4", stride: 16, align: 16, decl: "@group(1) @binding(3) var cache_val: array>;" }, - Member { name: "refl", ty: "vec4", stride: 16, align: 16, decl: "@group(1) @binding(5) var refl: array>;" }, - Member { name: "motion_out", ty: "vec4", stride: 16, align: 16, decl: "@group(1) @binding(6) var motion_out: array>;" }, - Member { name: "direct_out", ty: "vec4", stride: 16, align: 16, decl: "@group(1) @binding(7) var direct_out: array>;" }, + Member { name: "refl", ty: "vec4", stride: 16, align: 256, decl: "@group(1) @binding(5) var refl: array>;" }, + Member { name: "motion_out", ty: "vec4", stride: 16, align: 256, decl: "@group(1) @binding(6) var motion_out: array>;" }, + Member { name: "direct_out", ty: "vec4", stride: 16, align: 256, decl: "@group(1) @binding(7) var direct_out: array>;" }, ]; pub const KERNEL_LIVE_BINDING: u32 = 1; /// `MergedLive`'s last array, after `KERNEL_LIVE`'s: the loop guard's count (`guard`; its first word: the invocations @@ -81,7 +82,8 @@ pub fn live_layout(caps: &[u32; 6]) -> (Vec, u64) { out.push(at); at += m.stride * c as u64; } - (out, at.next_multiple_of(16)) + // (A struct's size is a multiple of its largest member's alignment: 256 where the live passes' ranges are.) + (out, at.next_multiple_of(members.iter().map(|m| m.align).max().unwrap_or(16))) } /// The lengths of the merged structs' arrays: what a kernel's source is written for (`gpu::Kernel::caps`). @@ -113,74 +115,6 @@ pub fn offsets(members: &[Member], caps: &[u32]) -> (Vec, u64) { (out, at.next_multiple_of(members.iter().map(|m| m.align).max().unwrap_or(1))) } -/// A slab of the live passes: buffers of one element type as ranges of one buffer. -pub struct Slab { - pub name: &'static str, - pub ty: &'static str, - pub stride: u64, - pub binding: u32, - pub read_only: bool, -} - -pub const SLABS: [Slab; 6] = [ - Slab { name: "MV4", ty: "vec4", stride: 16, binding: 2, read_only: false }, - Slab { name: "MU2", ty: "vec2", stride: 8, binding: 4, read_only: false }, - Slab { name: "MF1", ty: "f32", stride: 4, binding: 15, read_only: false }, - Slab { name: "MU4", ty: "vec4", stride: 16, binding: 17, read_only: false }, - Slab { name: "MA", ty: "atomic", stride: 4, binding: 13, read_only: false }, - // The kernel's `ML` buffer (what the live passes read of a trace), as vec4s: every array of it starts at a - // multiple of 16 bytes. - Slab { name: "MK", ty: "vec4", stride: 16, binding: 25, read_only: true }, -]; -pub const SLAB_V4: usize = 0; -pub const SLAB_U2: usize = 1; -pub const SLAB_F1: usize = 2; -pub const SLAB_U4: usize = 3; -pub const SLAB_A: usize = 4; -pub const SLAB_K: usize = 5; - -/// One of the live passes' buffers that becomes a range of a slab: its binding in live.wgsl as written (what -/// `live::Targets` binds there), its slab, its declaration up to the `;`. -pub struct LiveBuf { - pub name: &'static str, - pub binding: u32, - pub slab: usize, - pub decl: &'static str, -} - -/// In `MergedOffsets`' order (group 0, binding `OFFSETS_BINDING`: an (offset, length) pair of u32 each, in elements). -pub const LIVE_BUFS: [LiveBuf; 23] = [ - LiveBuf { name: "diffuse_hist", binding: 2, slab: SLAB_V4, decl: "@group(0) @binding(2) var diffuse_hist: array>;" }, - LiveBuf { name: "color_hist", binding: 3, slab: SLAB_V4, decl: "@group(0) @binding(3) var color_hist: array>;" }, - LiveBuf { name: "gbuffer", binding: 4, slab: SLAB_U2, decl: "@group(0) @binding(4) var gbuffer: array>;" }, - LiveBuf { name: "spec_hist", binding: 5, slab: SLAB_V4, decl: "@group(0) @binding(5) var spec_hist: array>;" }, - LiveBuf { name: "diffuse_mom", binding: 6, slab: SLAB_V4, decl: "@group(0) @binding(6) var diffuse_mom: array>;" }, - LiveBuf { name: "diffuse_fast", binding: 7, slab: SLAB_U2, decl: "@group(0) @binding(7) var diffuse_fast: array>;" }, - LiveBuf { name: "up_state", binding: 8, slab: SLAB_U4, decl: "@group(0) @binding(8) var up_state: array>;" }, - LiveBuf { name: "sample_vel", binding: 9, slab: SLAB_V4, decl: "@group(0) @binding(9) var sample_vel: array>;" }, - LiveBuf { name: "ae_bins", binding: 13, slab: SLAB_A, decl: "@group(0) @binding(13) var ae_bins: array, 64>;" }, - LiveBuf { name: "ae_state", binding: 14, slab: SLAB_V4, decl: "@group(0) @binding(14) var ae_state: array, 2>;" }, - LiveBuf { name: "lumq", binding: 15, slab: SLAB_F1, decl: "@group(0) @binding(15) var lumq: array;" }, - LiveBuf { name: "lumb", binding: 16, slab: SLAB_F1, decl: "@group(0) @binding(16) var lumb: array;" }, - LiveBuf { name: "fa", binding: 17, slab: SLAB_V4, decl: "@group(0) @binding(17) var fa: array>;" }, - LiveBuf { name: "fb", binding: 18, slab: SLAB_V4, decl: "@group(0) @binding(18) var fb: array>;" }, - LiveBuf { name: "bloom_a", binding: 19, slab: SLAB_V4, decl: "@group(0) @binding(19) var bloom_a: array>;" }, - LiveBuf { name: "bloom_b", binding: 20, slab: SLAB_V4, decl: "@group(0) @binding(20) var bloom_b: array>;" }, - LiveBuf { name: "spec_pp", binding: 21, slab: SLAB_U2, decl: "@group(0) @binding(21) var spec_pp: array>;" }, - LiveBuf { name: "sb", binding: 22, slab: SLAB_V4, decl: "@group(0) @binding(22) var sb: array>;" }, - LiveBuf { name: "refl_in", binding: 23, slab: SLAB_K, decl: "@group(0) @binding(23) var refl_in: array>;" }, - LiveBuf { name: "refl_res", binding: 24, slab: SLAB_V4, decl: "@group(0) @binding(24) var refl_res: array>;" }, - LiveBuf { name: "motion_in", binding: 25, slab: SLAB_K, decl: "@group(0) @binding(25) var motion_in: array>;" }, - LiveBuf { name: "direct_in", binding: 26, slab: SLAB_K, decl: "@group(0) @binding(26) var direct_in: array>;" }, - LiveBuf { name: "lock_fwd", binding: 27, slab: SLAB_A, decl: "@group(0) @binding(27) var lock_fwd: array>;" }, -]; -pub const OFFSETS_BINDING: u32 = 31; -/// `MergedOffsets`' size: a pair for each of `LIVE_BUFS`, and one unused that makes it a multiple of 16 bytes. -pub const OFFSETS_BYTES: usize = (LIVE_BUFS.len() + 1) * 8; - -/// The live passes' storage bindings that stay as they are: (binding, read only). -pub const LIVE_KEPT: [(u32, bool); 2] = [(1, true), (10, false)]; - /// How a merged name's uses are rewritten. struct Rule { name: &'static str, @@ -363,7 +297,7 @@ fn fresh(src: &str, names: &[&str]) { } fn merged_struct(name: &str, members: &[&Member], caps: &[u32]) -> String { - let fields: Vec = members.iter().zip(caps).map(|(m, c)| format!(" {}: array<{}, {c}>,", m.name, m.ty)).collect(); + let fields: Vec = members.iter().zip(caps).map(|(m, c)| format!(" {}{}: array<{}, {c}>,", if m.align > 16 { format!("@align({}) ", m.align) } else { String::new() }, m.name, m.ty)).collect(); format!("struct {name} {{\n{}\n}};\n", fields.join("\n")) } @@ -433,30 +367,6 @@ pub fn uniform_tables(src: &str) -> String { rewrite(&s, &rules) } -/// The live passes (live.wgsl) in the merged layout: sample_px, out_px and ae_bins as they are, the slabs, and the -/// offsets `LO`. -pub fn live(src: &str) -> String { - let mut names: Vec<&str> = SLABS.iter().map(|s| s.name).collect(); - names.extend(["LO", "MergedOffsets", "merged_pad"]); - fresh(src, &names); - let mut s = src.to_string(); - for b in LIVE_BUFS.iter() { - take(&mut s, b.decl, ""); - } - // (A last pair unused: the struct is 208 bytes, a multiple of 16. WebGPU does not ask for that, but wgpu's - // downlevel targets do of a uniform binding's size, and it costs nothing; `OFFSETS_BYTES`.) - let fields: Vec = LIVE_BUFS.iter().map(|b| format!(" {}: vec2,", b.name)).chain([" merged_pad: vec2,".to_string()]).collect(); - let mut head = format!("// The merged layout (renderer/src/merged.rs): each buffer's (first element, length) in its slab.\nstruct MergedOffsets {{\n{}\n}};\n@group(0) @binding({OFFSETS_BINDING}) var LO: MergedOffsets;\n", fields.join("\n")); - for sl in &SLABS { - head += &format!("@group(0) @binding({}) var {}: array<{}>;\n", sl.binding, if sl.read_only { "read" } else { "read_write" }, sl.name, sl.ty); - } - // After the uniforms' declaration (the file's first binding). - let anchor = "@group(0) @binding(0) var L: LiveParams;\n"; - take(&mut s, anchor, &format!("{anchor}{head}")); - let rules: Vec = LIVE_BUFS.iter().map(|b| Rule { name: b.name, open: format!("{}[LO.{}.x + (", SLABS[b.slab].name, b.name), close: ")".into(), len: format!("LO.{}.y", b.name) }).collect(); - rewrite(&s, &rules) -} - /// A bound on every loop of a shader: each `loop`, `for` and `while` body begins with a count of the invocation's /// loop iterations so far, and leaves the loop past `limit`. The kernel's traversals (the BVH, the light tree, the /// procedural boxes) follow indices they read from buffers and end only if the data is sound: on data that is not, a @@ -550,7 +460,6 @@ mod tests { use super::*; const KERNEL: &str = include_str!("../shaders/pathtrace.wgsl"); - const LIVE: &str = include_str!("../shaders/live.wgsl"); fn caps() -> Caps { // A cache of one entry, a view of 64 x 36: lengths that are not round. @@ -658,23 +567,6 @@ mod tests { assert!(out.ends_with("{ return objects[min(u32(a[1] * 3u + 1u), 4095u)].w + tex_info[min(u32(u32(objects[min(u32(k), 4095u)].x)), 4095u)].y + media[min(u32(0), 31u)].b.x; }\n"), "{out}"); } - /// The same for the live passes (25 as written). - #[test] - fn the_merged_live_passes_fit_eight_storage_buffers() { - let (m, info) = parse(LIVE); - assert_eq!(storage_buffers(&m, &info).1, 25, "the live passes as written"); - let passes = m.entry_points.len(); - let (m, info) = parse(&live(LIVE)); - let (used, declared) = storage_buffers(&m, &info); - assert_eq!(declared as u32, STORAGE_BUFFERS); - assert!(used.iter().all(|u| u.1 as u32 <= STORAGE_BUFFERS), "{used:?}"); - assert_eq!(used.len(), passes, "every pass is still there"); - // The offsets' struct: a pair of u32 each, in order (what `live::Targets` writes). - assert_eq!(layout(&m, "MergedOffsets"), ((0..24).map(|k| k * 8).collect(), OFFSETS_BYTES as u64)); - assert_eq!(OFFSETS_BYTES % 16, 0); - assert_eq!(LIVE_BUFS.len(), 23); - } - /// The Rust side's records are the size the merged structs give their elements. #[test] fn the_tables_strides_are_the_records() {