diff --git a/docs/ROADMAP.md b/docs/ROADMAP.md index ea7f506..c5f454e 100644 --- a/docs/ROADMAP.md +++ b/docs/ROADMAP.md @@ -2260,6 +2260,34 @@ surface materials' hash (`a * b ^ c * d`: a bitwise operator's operands must be naga had taken it, so those materials (and the particle passes, the same three times) never compiled in a browser. Fixed with parentheses in the order naga read them; a unit test holds the other shader files to the grammar. +**The kernel's layouts in its source, and the loop guard on its syntax tree (2026-11-14, after it).** The browser's +merged layout, Android's uniform tables and the loop guard were still made by rewriting the kernel's text: merged.rs +scanned the WGSL for every use of thirteen buffer names, followed the locals that shadowed them through their scopes, +rewrote each use (`MT.materials[`, `min(u32(i), last)`), and put a counter after the `{` of every loop it found by its +keyword. Now: +- The kernel reads its seven tables and uses its six group-1 buffers only through functions (`material_at`, + `light_at`, ..., `cache_key_claim`, `cache_acc_add`, `refl_store`, `direct_len`, ...: 87 reads of the tables, 40 uses + of group 1), and each layout is a block of declarations in pathtrace.wgsl that declares the buffers and those + functions: `@if(merged) { .. } @else { @if(uniform_tables) { .. } @else { .. } }` (WESL's blocks; scene data + behind an accessor module, as engines keep it). +- The numbers a variant needs are WESL constants the source imports (`import constants::..`): the merged arrays' + lengths, the uniform arrays', the loop guard's limit. gpu.rs `KernelVariant` gives a variant's flags and constants. +- The loop guard's count is put into every loop by a pass over the variant's syntax tree (gpu.rs `guard_loops`), the + scene's functions' loops too; its declarations are in the source under `@if(guard)`. +- merged.rs's scanner, its rewriter and their tests are gone (637 lines to 302); the numbers the Rust side needs for + its buffers stay. What is left as text: a scene's functions in their slots. + +A read or write that goes past the functions does not compile in the merged variant, whose block declares none of +the buffers by name, and the tests read the merged variants with naga. Checks: render-identity against main, all +with the loop guard first (natively, merged, with the uniform tables, with every switch on), then native without it: +0 of 190 outputs changed in any (172-175 byte-identical, the rest within their run-to-run spread), so the tree pass +and the text scan guard the same loops; `tools/web-hello.ts` in Chrome: pass. Cost (1440p at 0.375, alternated, both +orders): native hardware traversal equal (Bistro 8.28-8.97 ms against 8.30-8.96, Kabukichō 15.46-16.17 against +15.53-16.25); software traversal -0.2 ms in the layout as written (46.8 against 47.0) and +0.2 ms merged (47.3 against +47.1). The Metal source naga writes is the same in meaning both ways (compared function by function); neither +one-field accessors nor main's parameter order moved the merged +0.2, which comes and goes with the layout: the +compiler's arrangement around the accessors, +-0.5 %. + ## 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 5864991..e116868 100644 --- a/docs/RUNTIME.md +++ b/docs/RUNTIME.md @@ -690,25 +690,32 @@ tree (ROADMAP § 6). - **The merged binding layout** (`renderer/src/merged.rs`). WebGPU gives a stage 8 storage buffers by default; the 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. + test's check; a page binds them as natively.) There is one source: pathtrace.wgsl declares the kernel's tables and + its group-1 buffers in an `@if(merged) { .. } @else { .. }` block each (WESL's conditional translation, "The + kernel's source"), and the kernel reads and writes them only through functions each block declares + (`material_at`, `cache_acc_add`, `refl_store`, ...). + - The kernel's seven tables become the arrays of one struct in one buffer; their lengths are WESL constants the + source imports (gpu.rs `KernelVariant::constants`), 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; 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 (19 as written), the - structs' arrays where the Rust side puts them. + - The merged block declares none of the buffers by its own name, so a read or write that goes past the functions + does not compile in it (naga names the line). A test reads the merged variants with naga: valid WGSL, 8 storage + buffers for the kernel (19 as written), the structs' arrays where the Rust side puts them. + - Until 2026-11-14 the browser's kernel was written from the native one as text: a scanner of WGSL in merged.rs + found every use of the thirteen names, followed the locals that shadowed them through their scopes, and rewrote + each use (the uniform tables and the loop guard were made the same way). - `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 +- **The loop guard** (`@if(guard)` in pathtrace.wgsl, gpu.rs `guard_loops`). 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 watchdog; software traversal under the atomic tests, at two scenes without a triangle; the cause below is likely, not confirmed). A page's shader runs on the player's GPU beside everything else, so **a browser's kernel always has the guard**: every loop body counts the - invocation's iterations and leaves past a limit (65,536; a page can set it). No formula changes, but the image is + invocation's iterations and leaves past a limit (65,536; a page can set it). The count is put at the start of + every loop body, the scene's functions' too, by a pass over the variant's syntax tree (WESL's, as a compiler + instruments code); its declarations and the limit (a WESL constant) are in the source. No formula changes, but the image is not bit-identical to the kernel's without the guard: the compiler arranges the arithmetic of the functions it is in differently (hello world, one sample: 354 of 43,200 numbers differ, by at most 2 × 10⁻⁷ of their value). What it stops it counts: the count is read back without waiting (`LiveCore::guard_trips`, the page's `guardTrips()`, the frame log), an offline @@ -909,7 +916,7 @@ procedural objects (`trace_procs`: the ray into the object's space, a tree of it tested first: a ray that meets no unit reads one node, and a scene of one unit has no tree to walk. The hit it returns is the hit hardware traversal returns (the triangle, the record), so nothing after the traversal differs. - **The unit walk is `shaders/trace_units.wgsl`**, in the kernel under `@if(units)` (`F_UNITS`, a source feature, - as vertex colours are; gpu.rs `kernel_flags`): the file's walk, a call at the end of the still walk, and the still + as vertex colours are; gpu.rs `KernelVariant::flags`): the file's walk, a call at the end of the still walk, and the still walk's stack declaration taken out ("The stack"). A scene without units has none of it; hardware traversal's kernel never does. - Still flattened, as before: meshes that deform (a pose shared by a crowd), every instance of a scene with more @@ -989,7 +996,7 @@ to it. kernel without units is the source it was (`a_kernel_with_units_has_the_loop_guard`; naga validates the guarded kernel with units as written). -**The loop guard.** A browser's kernel always has it (`merged::guard`). **Natively a kernel with units has it with +**The loop guard.** A browser's kernel always has it (gpu.rs `guard_loops`). **Natively a kernel with units has it with nothing set** (65,536 iterations an invocation; `THREERT_LOOP_GUARD=0` takes it out, `=` sets another): what hung on 2026-11-11 is not confirmed, the unit walk is new, and no kernel with units has run without it. A native kernel without units has none unless `THREERT_LOOP_GUARD` says so, as before (its source and every existing byte @@ -1173,15 +1180,22 @@ units under software traversal; hello world natively and in Chrome (every check, The path-tracing kernel is one WESL module: `pathtrace.wgsl`, then `trace_units.wgsl`, `trace_hw.wgsl` and `translucent.wgsl` (gpu.rs `KERNEL_FILES`), every device's and scene's variant in it under WESL's conditional -translation (`@if(flag)` on a declaration or a statement). wesl-rs (WESL's compiler in Rust, natively and in a -browser alike) makes a kernel's WGSL from the flags (gpu.rs `kernel_flags`, `kernel_wgsl`): the traversal -(`hwrt`), a device's switches (`strict`, `cutout_lite`, `cutout_lite2`, `last_direct`, `light_pick`) and a scene's -features that change declarations (`vcolor`, `grid`, `tint`, `unlit`, `units`, `background`; the other features switch -code under the pipeline-overridable constant `FEAT`). pathtrace.wgsl's head names them; each is said where it changes -something, and a flag the source names that a kernel is not given is an error. -- Before that, still as text: a scene's material, volume and geometry functions between their markers in - pathtrace.wgsl, the merged layout (merged.rs `kernel`, the browser), the uniform tables (merged.rs `uniform_tables`, - Android) and the loop guard (merged.rs `guard`). +translation (`@if(flag)` on a declaration or a statement, `@if(flag) { .. } @else { .. }` on a block of +declarations). wesl-rs (WESL's compiler in Rust, natively and in a browser alike) makes a kernel's WGSL from a +variant (gpu.rs `KernelVariant`, `kernel_wgsl`): the traversal (`hwrt`), a device's switches (`strict`, +`cutout_lite`, `cutout_lite2`, `last_direct`, `light_pick`), a scene's features that change declarations (`vcolor`, +`grid`, `tint`, `unlit`, `units`, `background`; the other features switch code under the pipeline-overridable constant +`FEAT`), the binding layout (`merged`, the browser's; `uniform_tables`, Android's) and the loop guard (`guard`). +pathtrace.wgsl's head names them; each is said where it changes something, and a flag the source names that a +kernel is not given is an error. +- The numbers a variant is compiled for are WESL constants, which the source imports (`import constants::..`; gpu.rs + `KernelVariant::constants`): the merged layout's array lengths, the uniform tables', the loop guard's limit. One + the source imports and is not given is an error. +- The kernel reads its tables and its group-1 buffers only through functions (`material_at`, `instance_at`, + `cache_acc_add`, `refl_store`, ...), declared once in each layout's block, so a layout is its declarations and + nothing else of the kernel knows which it is in. +- The loop guard's count goes into every loop by a pass over the variant's syntax tree (gpu.rs `guard_loops`). +- Still as text: a scene's material, volume and geometry functions, between their markers in pathtrace.wgsl. - Until 2026-11-14 the variants were patches of the source at anchors, whole lines of code found and replaced (gpu.rs `kernel_base`, `kernel_switches`, `kernel_source`). For every combination of the flags, 2,304, the Metal source naga wrote from the patched kernel and from the one with `@if` was the same, two rewrites allowed for: the light pick's @@ -1215,9 +1229,9 @@ on the device", and differs where it says so. - **Uniform tables** (`GpuOptions::uniform_tables`, `THREERT_UNIFORM_TABLES=1|0`; default on Android; never with the merged layout). The driver gave the kernel other buffers' contents for four of its read-only storage buffers: `tex_info` (binding 14), `objects` (15), `instances` (17), `media` (18). As uniform buffers they read right. - - The kernel's source is rewritten when a module is compiled, as the merged layout's is (`merged::uniform_tables`): - the four declarations become sized uniform arrays at the same bindings, and every index into them is held to the - array (`min(u32(i), last)`), so nothing is read outside a uniform buffer. + - The kernel's source has them under `@if(uniform_tables)` (pathtrace.wgsl): the four declarations as sized + uniform arrays at the same bindings (their lengths WESL constants, merged.rs `UNIFORM_LEN`), and their accessors + hold every index to the array (`min(i, last)`), so nothing is read outside a uniform buffer. - The renderer keeps writing its storage buffers. `Gpu::sync_tables` copies the four into their uniform buffers in a frame's encoder, after the particle passes (which write their objects' entries on the GPU) and before the trace. - A uniform binding holds 64 KiB, so the arrays are bounded: 2,048 textures, 4,096 object entries (1,365 objects; diff --git a/renderer/shaders/pathtrace.wgsl b/renderer/shaders/pathtrace.wgsl index e365585..481b225 100644 --- a/renderer/shaders/pathtrace.wgsl +++ b/renderer/shaders/pathtrace.wgsl @@ -15,6 +15,16 @@ // vcolor, grid, tint, unlit, units, background // a scene's features that change the kernel's declarations rather than switch code in it // (gpu.rs `F_SOURCE`; the others are `FEAT`, a pipeline-overridable constant) +// merged the merged binding layout (merged.rs: at most eight storage buffers a stage, a browser's) +// uniform_tables four tables in uniform buffers (merged.rs `UNIFORM_TABLES`: Android's drivers) +// guard the loop guard (gpu.rs `loop_guard_of`, `guard_loops`) +// The numbers a variant is compiled for are WESL constants (`import constants::`, gpu.rs `KernelVariant::constants`): the +// merged layout's array lengths, the uniform tables', the loop guard's limit. +@if(merged) import constants::{MT_MATERIALS, MT_LIGHTS, MT_EMITTERS, MT_TEX_INFO, MT_OBJECTS, MT_INSTANCES, MT_MEDIA, + ML_CACHE_KEYS, ML_CACHE_ACC, ML_CACHE_VAL, ML_REFL, ML_MOTION_OUT, ML_DIRECT_OUT}; +@if(uniform_tables) import constants::{UT_TEX_INFO, UT_OBJECTS, UT_INSTANCES, UT_MEDIA}; +@if(merged || guard) import constants::GUARD_WORDS; +@if(guard) import constants::LOOP_GUARD_LIMIT; @if(hwrt) enable wgpu_ray_query; const PI: f32 = 3.14159265358979; @@ -114,11 +124,6 @@ struct Pixel { color: vec4, albedo: vec4, normal: vec4, emit: vec @group(0) @binding(1) var tri_verts: array>; @group(0) @binding(2) var attrs: array; @group(0) @binding(3) var nodes: array; -@group(0) @binding(4) var materials: array; -@group(0) @binding(5) var lights: array; -// emitters[i] = (triangle index bits, cumulative weight (normalized), weight, area) for i < counts.y, then the guide -// table: EMITTER_GUIDE + 1 entry indices (bits), four per vec4 (lib.rs emitter_cdf). -@group(0) @binding(6) var emitters: array>; @group(0) @binding(7) var textures: texture_2d_array; @group(0) @binding(8) var samp: sampler; @group(0) @binding(9) var env_tex: texture_2d; @@ -127,20 +132,73 @@ struct Pixel { color: vec4, albedo: vec4, normal: vec4, emit: vec @group(0) @binding(11) var accum: array; // GGX energy LUT (E.x, E.y) for 32 roughness x 32 NoV, packed 2 per vec4 @group(0) @binding(12) var energy_lut: array, 512>; -// Per texture: (u0, v0, du, dv) of its rect in the page, (layer, wrap, width, height). See texpack.rs. -@group(0) @binding(14) var tex_info: array>; -// Per object (scene node): pivot (world position), world bounds size (UE ObjectPositionWS / ObjectBounds). -@group(0) @binding(15) var objects: array>; // The dynamic triangles' vertices (kernel layout, from triangle P.group_first[1].x on) as they were the previous frame. @group(0) @binding(16) var prev_verts: array>; // Model instances (scene::GpuInstance): world matrix rows, normal matrix rows, (first triangle, object, LOD offset bits, -). struct Instance { m0: vec4, m1: vec4, m2: vec4, n0: vec4, n1: vec4, n2: vec4, info: vec4, p0: vec4, p1: vec4, p2: vec4 }; -@group(0) @binding(17) var instances: array; // Bounded participating media (`VolumeMesh`, gpu.rs MediumGpu; F_MEDIA), P.time.w of them: world -> local affine rows, // (volume function, shape: 0 box / 1 sphere / 2 none while its node is hidden, homogeneous, 1 = it changes with time: `VolumeFunction` `animated`, a // simulated grid), (majorant 1/m, HG anisotropy, -, -), customData, parameters. struct Medium { w0: vec4, w1: vec4, w2: vec4, a: vec4, b: vec4, c: vec4, p: array, 8> }; -@group(0) @binding(18) var media: array; + +// The scene's tables, read through `material_at`, `light_at`, `emitter_at`, `tex_info_at`, `object_at`, `instance_at` +// and `medium_at` in every layout: +// - materials (scene::GpuMaterial), lights (scene::GpuLight); +// - emitters: (triangle index bits, cumulative weight (normalized), weight, area) for i < counts.y, then the guide +// table: EMITTER_GUIDE + 1 entry indices (bits), four per vec4 (lib.rs emitter_cdf), then the light tree; +// - tex_info, per texture: (u0, v0, du, dv) of its rect in the page, (layer, wrap, width, height). See texpack.rs; +// - objects, per object (scene node): pivot (world position), world bounds size (UE ObjectPositionWS / +// ObjectBounds), then what `OBJECT_STRIDE` says; +// - instances (Instance), media (Medium). +@if(merged) { + // The merged layout (merged.rs): the seven as the arrays of one struct in one buffer, their lengths the renderer's + // (`merged::Caps`: each a power of two, so a scene that grows compiles a kernel again only when a table doubles). + struct MergedTables { + materials: array, + lights: array, + emitters: array, MT_EMITTERS>, + tex_info: array, MT_TEX_INFO>, + objects: array, MT_OBJECTS>, + instances: array, + media: array, + }; + @group(0) @binding(4) var MT: MergedTables; + fn material_at(i: u32) -> Material { return MT.materials[i]; } + fn light_at(i: u32) -> Light { return MT.lights[i]; } + fn emitter_at(i: u32) -> vec4 { return MT.emitters[i]; } + fn tex_info_at(i: u32) -> vec4 { return MT.tex_info[i]; } + fn object_at(i: u32) -> vec4 { return MT.objects[i]; } + fn instance_at(i: u32) -> Instance { return MT.instances[i]; } + fn medium_at(i: u32) -> Medium { return MT.media[i]; } +} @else { + @group(0) @binding(4) var materials: array; + @group(0) @binding(5) var lights: array; + @group(0) @binding(6) var emitters: array>; + fn material_at(i: u32) -> Material { return materials[i]; } + fn light_at(i: u32) -> Light { return lights[i]; } + fn emitter_at(i: u32) -> vec4 { return emitters[i]; } + @if(uniform_tables) { + // The uniform tables (merged.rs `UNIFORM_TABLES`): sized uniform arrays, every index held to the array (a + // scene with more entries than one holds reads the last for the rest; nothing outside the buffer is read). + @group(0) @binding(14) var tex_info: array, UT_TEX_INFO>; + @group(0) @binding(15) var objects: array, UT_OBJECTS>; + @group(0) @binding(17) var instances: array; + @group(0) @binding(18) var media: array; + fn tex_info_at(i: u32) -> vec4 { return tex_info[min(i, UT_TEX_INFO - 1u)]; } + fn object_at(i: u32) -> vec4 { return objects[min(i, UT_OBJECTS - 1u)]; } + fn instance_at(i: u32) -> Instance { return instances[min(i, UT_INSTANCES - 1u)]; } + fn medium_at(i: u32) -> Medium { return media[min(i, UT_MEDIA - 1u)]; } + } @else { + @group(0) @binding(14) var tex_info: array>; + @group(0) @binding(15) var objects: array>; + @group(0) @binding(17) var instances: array; + @group(0) @binding(18) var media: array; + fn tex_info_at(i: u32) -> vec4 { return tex_info[i]; } + fn object_at(i: u32) -> vec4 { return objects[i]; } + fn instance_at(i: u32) -> Instance { return instances[i]; } + fn medium_at(i: u32) -> Medium { return media[i]; } + } +} // Density grids (gpu.rs F_GRID, `GridsGpu`): field A of every grid, stacked along z. @if(grid) @group(0) @binding(19) var grid_tex: texture_3d; // Vertex colours (three.js `vertexColors`; gpu.rs F_VCOLOR): three sRGB 8-bit colours per triangle, after the previous @@ -221,7 +279,7 @@ fn rnd() -> f32 { const EMITTER_GUIDE: u32 = 65536u; fn emitter_guide(k: u32) -> u32 { - return bitcast(emitters[P.counts.y + k / 4u][k % 4u]); + return bitcast(emitter_at(P.counts.y + k / 4u)[k % 4u]); } // One emitter CDF entry: the first whose cumulative weight is not below u. The guide table bounds the binary search to @@ -237,9 +295,9 @@ fn emitter_pick_u(u: f32) -> vec4 { var hi = emitter_guide(k + 1u); while (lo < hi) { let mid = (lo + hi) / 2u; - if (emitters[mid].y < u) { lo = mid + 1u; } else { hi = mid; } + if (emitter_at(mid).y < u) { lo = mid + 1u; } else { hi = mid; } } - return emitters[min(lo, P.counts.y - 1u)]; + return emitter_at(min(lo, P.counts.y - 1u)); } // Light tree (ltree.rs): nodes after the CDF in `emitters`, one vec4 each: pack2x16float(centre x, y), @@ -254,7 +312,7 @@ const LT_FAR: f32 = 120.0; fn lt_on() -> bool { return P.ltree.w > 0.5; } -fn lt_node(i: u32) -> vec4 { return emitters[bitcast(P.ltree.x) + i]; } +fn lt_node(i: u32) -> vec4 { return emitter_at(bitcast(P.ltree.x) + i); } // Importance of a tree node seen from x on a surface with normal n (n = 0: no hemisphere, a fog point): power x the // best cosine any point of its sphere can have at x / max(distance², radius²). @@ -282,17 +340,17 @@ fn lt_imp(nd: vec4, x: vec3, n: vec3) -> f32 { // A CDF entry by power within entries [lo, hi) (cumulative p0 .. p1): one guide-table lookup. fn emitter_pick_range(lo: u32, hi: u32) -> vec4 { - let p0 = select(0.0, emitters[max(lo, 1u) - 1u].y, lo > 0u); - let p1 = emitters[hi - 1u].y; + let p0 = select(0.0, emitter_at(max(lo, 1u) - 1u).y, lo > 0u); + let p1 = emitter_at(hi - 1u).y; let u = min(p0 + rnd() * (p1 - p0), p1); let k = min(u32(u * f32(EMITTER_GUIDE)), EMITTER_GUIDE - 1u); var l = emitter_guide(k); var h = emitter_guide(k + 1u); while (l < h) { let mid = (l + h) / 2u; - if (emitters[mid].y < u) { l = mid + 1u; } else { h = mid; } + if (emitter_at(mid).y < u) { l = mid + 1u; } else { h = mid; } } - return emitters[clamp(l, lo, hi - 1u)]; + return emitter_at(clamp(l, lo, hi - 1u)); } // An emitter entry for a point x (normal n), its selection probability `p` and its weight `w` (area x luminance x @@ -325,11 +383,11 @@ fn emitter_draw(x: vec3, n: vec3) -> EPick { let lo = w & 0x1fffffffu; let hi = lo + ((w >> 29u) & 3u) + 1u; var wsum = 0.0; - for (var j = lo; j < hi; j++) { wsum += emitters[j].z; } + for (var j = lo; j < hi; j++) { wsum += emitter_at(j).z; } var t = rnd() * wsum; - var e = emitters[hi - 1u]; + var e = emitter_at(hi - 1u); for (var j = lo; j < hi; j++) { - let c = emitters[j]; + let c = emitter_at(j); if (t < c.z) { e = c; break; } t -= c.z; } @@ -345,7 +403,7 @@ fn emitter_sel(ti: u32, x: vec3, n: vec3) -> f32 { if (P.misc.y <= 0.0 || (has(F_PEMIT) && a.x.z < 0.0)) { return 0.0; } var area = 0.5 * length(cross(tri_verts[ti * 3u + 1u].xyz, tri_verts[ti * 3u + 2u].xyz)); if (has(F_PEMIT) && a.x.z > 0.0) { area = a.x.z; } - return luminance(materials[bitcast(a.uv.w)].emissive.rgb) * obj_power(a) * area / P.misc.y; + return luminance(material_at(bitcast(a.uv.w)).emissive.rgb) * obj_power(a) * area / P.misc.y; } let depth = (code >> LT_PATH_BITS) - 1u; var p = 1.0 - P.ltree.z; @@ -365,7 +423,7 @@ fn emitter_sel(ti: u32, x: vec3, n: vec3) -> f32 { var wsum = 0.0; var we = 0.0; for (var j = lo; j < hi; j++) { - let c = emitters[j]; + let c = emitter_at(j); wsum += c.z; if (bitcast(c.x) == ti) { we = c.z; } } @@ -440,11 +498,11 @@ fn proc_point(r: Instance, p: vec3) -> vec3 { return proc_dir(r, p - v // Box `prim` of record `rec` against the object-space ray: the geometry function's hit in (tmin, tmax), else t = -1. fn proc_box(rec: u32, prim: u32, o: vec3, d: vec3, tmin: f32, tmax: f32) -> GOut { - let info = instances[rec].info; + let info = instance_at(rec).info; let b = pword(info.x + 1u) + prim * 6u; gdata_first = pword(info.x + 2u); gdata_count = pword(info.x + 11u); - let g = geometry_function(info.z, GIn(o, d, tmin, tmax, prim, proc_vec(b), proc_vec(b + 3u), gdata_first, gdata_count, objects[info.y * OBJECT_STRIDE + 2u], P.time.x)); + let g = geometry_function(info.z, GIn(o, d, tmin, tmax, prim, proc_vec(b), proc_vec(b + 3u), gdata_first, gdata_count, object_at(info.y * OBJECT_STRIDE + 2u), P.time.x)); if (!(g.t > tmin && g.t < tmax)) { return GOut(-1.0, g.normal, 0u, g.attr); } return g; } @@ -472,7 +530,7 @@ fn proc_attr(h: Hit) -> vec2 { // A procedural hit's world normal (geometric = shading) through its record's normal rows. fn proc_normal(h: Hit) -> vec3 { - let r = instances[h.inst]; + let r = instance_at(h.inst); let p = proc_bits(h.u); let n = proc_unoct(vec2(f32(p & 0x7fffu), f32(p >> 15u)) / 32767.0 * 2.0 - 1.0); return normalize(vec3(dot(r.n0.xyz, n), dot(r.n1.xyz, n), dot(r.n2.xyz, n))); @@ -591,7 +649,7 @@ fn tri_hit(ti: u32, o: vec3, d: vec3) -> Hit { let first = P.procs.z; for (var k = 0u; k < P.procs.y; k++) { let rec = first + k; - let r = instances[rec]; + let r = instance_at(rec); if (mode == TRACE_SHADOW && r.info.w == 0u) { continue; } let lo = proc_point(r, o); let ld = proc_dir(r, d); @@ -734,12 +792,12 @@ fn glass_at(h: Hit, d: vec3, t: f32) -> GlassAt { var e1 = tri_verts[h.tri * 3u + 1u].xyz; var e2 = tri_verts[h.tri * 3u + 2u].xyz; if (has(F_INST) && h.inst != NO_INST) { - let r = instances[h.inst]; + let r = instance_at(h.inst); e1 = inst_vec(r, e1); e2 = inst_vec(r, e2); a = inst_attr(r, a); } - let m = materials[bitcast(a.uv.w)]; + let m = material_at(bitcast(a.uv.w)); var g = GlassAt(normalize(cross(e1, e2)), 0.0, false, vec3(1.0)); var tr = clamp(m.layer.x, 0.0, 1.0) * (1.0 - clamp(m.params.x, 0.0, 1.0)); if (tr <= 0.0) { return g; } @@ -749,7 +807,7 @@ fn glass_at(h: Hit, d: vec3, t: f32) -> GlassAt { if (has(F_TEX) && (m.tex0.x >= 0 || m.uv.z > 0.5)) { let uv = mesh_uv(a, m, h.u, h.v); var lod = mesh_lod(a, m) + log2(max(t * 2.0 * P.cam_pos.w / lod_height(), 1e-7)) - LOD_BIAS; - if (has(F_INST) && h.inst != NO_INST) { lod += bitcast(instances[h.inst].info.z); } + if (has(F_INST) && h.inst != NO_INST) { lod += bitcast(instance_at(h.inst).info.z); } if (m.tex0.x >= 0) { g.tint *= srgb_to_linear(tex(m.tex0.x, uv, lod).rgb); } if (m.uv.z > 0.5) { tr *= tex(i32(m.uv.z) - 1, uv, lod).r; } } @@ -787,7 +845,7 @@ fn media_homogeneous(o: vec3, d: vec3, tmax: f32, light: bool) -> Marc if (!light) { // Transmittance alone: the media's product, in any order. for (var k = 0u; k < n; k++) { - let m = media[k]; + let m = medium_at(k); if (m.a.z == 0u) { continue; } let sp = medium_span(m, o, d); let a = max(sp.x, 0.0); @@ -806,7 +864,7 @@ fn media_homogeneous(o: vec3, d: vec3, tmax: f32, light: bool) -> Marc var end = tmax; var inside = false; for (var k = 0u; k < n; k++) { - let m = media[k]; + let m = medium_at(k); if (m.a.z == 0u) { continue; } let sp = medium_span(m, o, d); let a = max(sp.x, 0.0); @@ -903,14 +961,14 @@ fn srgb_to_linear(c: vec3) -> vec3 { // Texture pages: `lod` is log2 of the ray footprint in texture-UV units; the mip level adds the texture's resolution. fn tex_level(id: i32, uv: vec2, level: f32) -> vec4 { - let a = tex_info[u32(id) * 2u]; - let b = tex_info[u32(id) * 2u + 1u]; + let a = tex_info_at(u32(id) * 2u); + let b = tex_info_at(u32(id) * 2u + 1u); let f = select(clamp(uv, vec2(0.0), vec2(1.0)), fract(uv), b.y > 0.5); return textureSampleLevel(textures, samp, a.xy + f * a.zw, i32(b.x), clamp(level, 0.0, TEX_MAX_LEVEL)); } fn tex(id: i32, uv: vec2, lod: f32) -> vec4 { - let b = tex_info[u32(id) * 2u + 1u]; + let b = tex_info_at(u32(id) * 2u + 1u); return tex_level(id, uv, lod + 0.5 * log2(b.z * b.w)); } @@ -999,23 +1057,23 @@ fn object_of(a: TriAttr) -> u32 { return bitcast(a.x.y); } @if(!tint) const OBJECT_STRIDE: u32 = 3u; @if(tint) const OBJECT_STRIDE: u32 = 4u; // An instanced object's instance colour (three.js `instanceColor`), white for every other object. -@if(tint) fn obj_tint(oi: u32) -> vec3 { return objects[oi * OBJECT_STRIDE + 3u].rgb; } +@if(tint) fn obj_tint(oi: u32) -> vec3 { return object_at(oi * OBJECT_STRIDE + 3u).rgb; } // Instance power (UE material "Power", scripts animate it): scales the emission of everything the node draws. -fn obj_power(a: TriAttr) -> f32 { return objects[object_of(a) * OBJECT_STRIDE].w; } +fn obj_power(a: TriAttr) -> f32 { return object_at(object_of(a) * OBJECT_STRIDE).w; } // Inputs at a point of triangle ti (barycentrics u, v) for material functions outside full shading (alpha, NEE). fn mat_in_at(ti: u32, u: f32, v: f32, view: vec3, lod: f32) -> MatIn { let a = attrs[ti]; - let m = materials[bitcast(a.uv.w)]; + let m = material_at(bitcast(a.uv.w)); let w = 1.0 - u - v; let p = tri_verts[ti * 3u].xyz + tri_verts[ti * 3u + 1u].xyz * u + tri_verts[ti * 3u + 2u].xyz * v; let ng = normalize(cross(tri_verts[ti * 3u + 1u].xyz, tri_verts[ti * 3u + 2u].xyz)); var n = normalize(a.n0.xyz * w + a.n1.xyz * u + a.n2.xyz * v); let t = normalize(a.t.xyz - n * dot(n, a.t.xyz) + vec3(1e-9)); let oi = object_of(a); - return MatIn(p, n, ng, t, cross(n, t) * a.t.w, mesh_uv(a, m, u, v), mesh_lod(a, m), view, objects[oi * OBJECT_STRIDE].xyz, objects[oi * OBJECT_STRIDE + 1u].xyz, - P.time.x, P.time.y, P.time.z, lod, cur_depth, objects[oi * OBJECT_STRIDE].w, cur_depth == 0u || mirror_path, - objects[oi * OBJECT_STRIDE + 2u]); + return MatIn(p, n, ng, t, cross(n, t) * a.t.w, mesh_uv(a, m, u, v), mesh_lod(a, m), view, object_at(oi * OBJECT_STRIDE).xyz, object_at(oi * OBJECT_STRIDE + 1u).xyz, + P.time.x, P.time.y, P.time.z, lod, cur_depth, object_at(oi * OBJECT_STRIDE).w, cur_depth == 0u || mirror_path, + object_at(oi * OBJECT_STRIDE + 2u)); } fn mat_base(m: Material) -> MatOut { @@ -1036,10 +1094,10 @@ fn cutouts_opaque() -> bool { // the model instance a pool triangle belongs to (its scale shifts the texel density, as in `inst_attr`), or NO_INST. fn alpha_pass(ti: u32, inst: u32, u: f32, v: f32, t: f32) -> bool { let a = attrs[ti]; - let m = materials[bitcast(a.uv.w)]; + let m = material_at(bitcast(a.uv.w)); var alpha = m.base_color.a; var lod = mesh_lod(a, m) + log2(max(t * 2.0 * P.cam_pos.w / lod_height(), 1e-7)) - LOD_BIAS; - if (has(F_INST) && inst != NO_INST) { lod += bitcast(instances[inst].info.z); } + if (has(F_INST) && inst != NO_INST) { lod += bitcast(instance_at(inst).info.z); } if (has(F_TEX) && m.tex0.x >= 0) { alpha *= tex(m.tex0.x, mesh_uv(a, m, u, v), lod).a; } @@ -1068,7 +1126,7 @@ fn surface(o: vec3, d: vec3, h: Hit, width: f32) -> Surface { var e2 = tri_verts[i * 3u + 2u].xyz; var a = attrs[i]; if (has(F_INST) && h.inst != NO_INST) { - let r = instances[h.inst]; + let r = instance_at(h.inst); e1 = inst_vec(r, e1); e2 = inst_vec(r, e2); a = inst_attr(r, a); @@ -1077,7 +1135,7 @@ fn surface(o: vec3, d: vec3, h: Hit, width: f32) -> Surface { var ng = normalize(cross(e1, e2)); var n = normalize(a.n0.xyz * w + a.n1.xyz * h.u + a.n2.xyz * h.v); let mi = bitcast(a.uv.w); - let m = materials[mi]; + let m = material_at(mi); let front = dot(ng, d) < 0.0; if (!front) { ng = -ng; } if (dot(n, ng) < 0.0) { n = -n; } @@ -1110,7 +1168,7 @@ fn surface(o: vec3, d: vec3, h: Hit, width: f32) -> Surface { metal = 0.0; } // Power just changed by a script / gameplay (objects[1].w = scene time of the edit): short history. - s.anim = P.time.x - objects[object_of(a) * OBJECT_STRIDE + 1u].w < 0.5; + s.anim = P.time.x - object_at(object_of(a) * OBJECT_STRIDE + 1u).w < 0.5; // Detail that only matters where it is seen directly or in mirror-like reflections (puddles, glass, chrome). let near = cur_depth == 0u || mirror_path; if (has(F_TEX) && m.tex0.y >= 0 && near) { @@ -1128,9 +1186,9 @@ fn surface(o: vec3, d: vec3, h: Hit, width: f32) -> Surface { var t = normalize(a.t.xyz - n * dot(n, a.t.xyz) + vec3(1e-9)); let nv = normalize(a.n0.xyz * w + a.n1.xyz * h.u + a.n2.xyz * h.v) * select(-1.0, 1.0, front || dot(a.n0.xyz, ng) > 0.0); let i = MatIn(p, select(-nv, nv, dot(nv, ng) >= 0.0), ng, t, cross(n, t) * a.t.w, mesh_uv(a, m, h.u, h.v), mesh_lod(a, m), -d, - objects[oi * OBJECT_STRIDE].xyz, objects[oi * OBJECT_STRIDE + 1u].xyz, P.time.x, P.time.y, P.time.z, - log2(max(width, 1e-7)) - 0.5 * log2(max(abs(dot(ng, d)), 1e-3)), cur_depth, objects[oi * OBJECT_STRIDE].w, near, - objects[oi * OBJECT_STRIDE + 2u]); + object_at(oi * OBJECT_STRIDE).xyz, object_at(oi * OBJECT_STRIDE + 1u).xyz, P.time.x, P.time.y, P.time.z, + log2(max(width, 1e-7)) - 0.5 * log2(max(abs(dot(ng, d)), 1e-3)), cur_depth, object_at(oi * OBJECT_STRIDE).w, near, + object_at(oi * OBJECT_STRIDE + 2u)); let r = material_function(m.tex1.w - 1, i, m, MatOut(base, rough, metal, spec, s.emissive, n, m.base_color.a, false)); base = r.base; rough = r.roughness; @@ -1165,10 +1223,10 @@ fn surface(o: vec3, d: vec3, h: Hit, width: f32) -> Surface { // two-sided), the object's power and customData. No textures (no uv: material functions get uv 0 and can work from // `i.wp`), no transmission. fn surface_proc(o: vec3, d: vec3, h: Hit, width: f32) -> Surface { - let r = instances[h.inst]; + let r = instance_at(h.inst); let hd = r.info.x; let mi = pword(pword(hd + 7u) + min((h.tri >> 23u) & 0xffu, pword(hd + 3u) - 1u)); - let m = materials[mi]; + let m = material_at(mi); let oi = r.info.y; var ng = proc_normal(h); let front = dot(ng, d) < 0.0; @@ -1184,15 +1242,15 @@ fn surface_proc(o: vec3, d: vec3, h: Hit, width: f32) -> Surface { var rough = m.params.y; var metal = m.params.x; var spec = m.params.z; - s.emissive = m.emissive.rgb * objects[oi * OBJECT_STRIDE].w; - s.anim = P.time.x - objects[oi * OBJECT_STRIDE + 1u].w < 0.5; + s.emissive = m.emissive.rgb * object_at(oi * OBJECT_STRIDE).w; + s.anim = P.time.x - object_at(oi * OBJECT_STRIDE + 1u).w < 0.5; if (has(F_FUNC) && m.tex1.w > 0) { let near = cur_depth == 0u || mirror_path; let a = select(vec3(1.0, 0.0, 0.0), vec3(0.0, 1.0, 0.0), abs(n.x) > 0.9); let t = normalize(cross(a, n)); - let i = MatIn(p, n, ng, t, cross(n, t), proc_attr(h), 0.0, -d, objects[oi * OBJECT_STRIDE].xyz, objects[oi * OBJECT_STRIDE + 1u].xyz, + let i = MatIn(p, n, ng, t, cross(n, t), proc_attr(h), 0.0, -d, object_at(oi * OBJECT_STRIDE).xyz, object_at(oi * OBJECT_STRIDE + 1u).xyz, P.time.x, P.time.y, P.time.z, log2(max(width, 1e-7)) - 0.5 * log2(max(abs(dot(ng, d)), 1e-3)), cur_depth, - objects[oi * OBJECT_STRIDE].w, near, objects[oi * OBJECT_STRIDE + 2u]); + object_at(oi * OBJECT_STRIDE).w, near, object_at(oi * OBJECT_STRIDE + 2u)); let f = material_function(m.tex1.w - 1, i, m, MatOut(base, rough, metal, spec, s.emissive, n, m.base_color.a, false)); base = f.base; rough = f.roughness; @@ -1492,7 +1550,7 @@ fn sky_mis(d: vec3, pdf: f32) -> vec3 { // BSDF rays that fall inside a directional light's cone see its disk (MIS vs cone sampling). if (pdf > 0.0) { for (var i = 0u; i < P.counts2.x; i++) { - let li = lights[i]; + let li = light_at(i); if (dot(d, -li.dir.xyz) >= li.dir.w) { let omega = 2.0 * PI * (1.0 - li.dir.w); c += (li.color.rgb / omega) * power_heuristic(pdf, 1.0 / omega); @@ -1522,7 +1580,7 @@ fn emit_mis(hit: Hit, d: vec3, pdf: f32) -> f32 { let area = 0.5 * length(cr); let cos_l = abs(dot(normalize(cr), d)); let ha = attrs[hit.tri]; - let mat = materials[bitcast(ha.uv.w)]; + let mat = material_at(bitcast(ha.uv.w)); let pdf_a = luminance(mat.emissive.rgb) * obj_power(ha) / P.emit_power.x; let pdf_w = pdf_a * hit.t * hit.t / max(cos_l, 1e-6); if (!has(F_PEMIT) || ha.x.z == 0.0) { return power_heuristic(pdf, pdf_w); } @@ -1642,14 +1700,14 @@ fn local_lights(s: Surface, lb: Lobes, v: vec3, o: vec3, x: vec3, var i = nd; loop { if (i >= end) { break; } - let gs = lights[P.counts.x + (i - nd) / LIGHT_GROUP].pos; + let gs = light_at(P.counts.x + (i - nd) / LIGHT_GROUP).pos; let gend = min(i + LIGHT_GROUP, end); if (distance(x, gs.xyz) > gs.w) { i = gend; continue; } for (; i < gend; i++) { - let li = lights[i]; + let li = light_at(i); // Cheap rejects before sampling the light: black (off, or hidden), beyond its attenuation range (the // window is zero there), or entirely behind the surface. if (all(li.color.rgb <= vec3(0.0))) { continue; } @@ -1681,9 +1739,9 @@ fn local_lights(s: Surface, lb: Lobes, v: vec3, o: vec3, x: vec3, var pick_ls: LS; for (var k = 0u; k < m_cand; k++) { let idx = first + min(u32(rnd() * f32(count)), count - 1u); - let ls = local_light(lights[idx], x); + let ls = local_light(light_at(idx), x); var ph = luminance(ls.c); - if (vol) { ph *= lights[idx].x.x; } else { ph *= max(dot(s.n, ls.l), 0.0) + 0.05 * max(dot(s.ng, ls.l), 0.0); } + if (vol) { ph *= light_at(idx).x.x; } else { ph *= max(dot(s.n, ls.l), 0.0) + 0.05 * max(dot(s.ng, ls.l), 0.0); } if (ph <= 0.0) { continue; } wsum += ph; if (rnd() * wsum < ph) { @@ -1694,10 +1752,10 @@ fn local_lights(s: Surface, lb: Lobes, v: vec3, o: vec3, x: vec3, } if (pick_ph > 0.0) { var c: vec3; - let lk = lights[pick].x; + let lk = light_at(pick).x; if (vol) { c = pick_ls.c * hg_phase(P.vol.y, dot(pick_ls.l, d)) * lk.x; } else { c = pick_ls.c * eval_bsdf_k(s, lb, v, pick_ls.l, lk.y); } if (any(c > vec3(0.0))) { - let vis = light_vis(lights[pick].ext.w >= 0.5, o, pick_ls.sl, pick_ls.dist * 0.999); + let vis = light_vis(light_at(pick).ext.w >= 0.5, o, pick_ls.sl, pick_ls.dist * 0.999); if (any(vis > vec3(0.0))) { let r = c * gtint(vis) * (wsum * f32(count) / (f32(m_cand) * pick_ph)); sum += r; @@ -1716,7 +1774,7 @@ struct ES { le: vec3, l: vec3, dist: f32, pdf_w: f32, w: f32 }; fn emitter_radiance(ti: u32, b1: f32, b2: f32, l: vec3) -> vec3 { let a = attrs[ti]; let mi = bitcast(a.uv.w); - let m = materials[mi]; + let m = material_at(mi); // 5 cm footprint: a blurred texel is plenty for lighting and keeps NEE cheap. let pw = obj_power(a); var le = emission(m, mesh_uv(a, m, b1, b2), mesh_lod(a, m) + log2(0.05)) * pw; @@ -1797,7 +1855,7 @@ fn sample_emitter_ris(x: vec3, n: vec3, m: u32, sf: Surface, lb: Lobes let dist = sqrt(d2); let l = dl / dist; var cos_l = dot(ln, -l); - let m_card = materials[bitcast(attrs[ti].uv.w)].tex1.z == MAT_EMITTER_CARD; + let m_card = material_at(bitcast(attrs[ti].uv.w)).tex1.z == MAT_EMITTER_CARD; if (!m_card) { cos_l = abs(cos_l); } if (cos_l <= 0.0) { return r; } r.le = emitter_radiance(ti, pick_b.x, pick_b.y, l); @@ -1833,7 +1891,7 @@ fn sample_emitter(x: vec3, n: vec3) -> ES { let l = dl / dist; let a = attrs[ti]; let mi = bitcast(a.uv.w); - let m = materials[mi]; + let m = material_at(mi); var cos_l = dot(ln, -l); if (m.tex1.z != MAT_EMITTER_CARD) { cos_l = abs(cos_l); } if (cos_l <= 0.0 || (has(F_PEMIT) && !(cos_l > 0.0))) { return r; } // also a particle not shown (no area: NaN) @@ -1863,7 +1921,7 @@ fn direct_light(s: Surface, lb: Lobes, v: vec3) -> vec3 { let o = offset_origin(s.p, s.ng); // Directional lights: uniformly sampled cone (soft shadows), MIS against BSDF sampling. for (var i = 0u; i < P.counts2.x; i++) { - let li = lights[i]; + let li = light_at(i); let cos_max = li.dir.w; let to_light = -li.dir.xyz; let cos_t = 1.0 - rnd() * (1.0 - cos_max); @@ -2011,7 +2069,7 @@ fn fog_scatter(o: vec3, d: vec3, t0: f32, t1: f32, tr: f32) -> vec3(0.0); for (var i = 0u; i < P.counts2.x; i++) { - let li = lights[i]; + let li = light_at(i); let l = -li.dir.xyz; if (all(li.color.rgb <= vec3(0.0))) { continue; } // black: no shadow ray (direct_light) if (li.ext.w < 0.5 || !occluded(x, l, INF)) { ls += li.color.rgb * hg_phase(P.vol.y, dot(l, d)) * li.x.x; } @@ -2112,7 +2170,7 @@ fn medium_span(m: Medium, o: vec3, d: vec3) -> vec2 { } fn medium_eval(k: u32, x: vec3) -> VOut { - let m = media[k]; + let m = medium_at(k); var i: VIn; i.p = x; i.lp = medium_local(m, x); @@ -2141,7 +2199,7 @@ fn media_track(o: vec3, d: vec3, tmax: f32) -> Track { var mj = 0.0; var end = tmax; for (var k = 0u; k < n; k++) { - let sp = medium_span(media[k], o, d); + let sp = medium_span(medium_at(k), o, d); let a = max(sp.x, 0.0); let b = min(sp.y, tmax); if (a >= b) { continue; } @@ -2149,7 +2207,7 @@ fn media_track(o: vec3, d: vec3, tmax: f32) -> Track { end = min(end, a); } else if (b > t) { mask |= 1u << k; - mj += media[k].b.x; + mj += medium_at(k).b.x; end = min(end, b); } } @@ -2167,7 +2225,7 @@ fn media_track(o: vec3, d: vec3, tmax: f32) -> Track { // The medium a collision here is in: each by its extinction (the first without a draw). if (v.extinction > 0.0 && (st <= v.extinction || rnd() * st < v.extinction)) { r.albedo = v.albedo; - r.g = media[k].b.y; + r.g = medium_at(k).b.y; } } r.emission += em / mj; @@ -2190,9 +2248,9 @@ fn media_tr(o: vec3, d: vec3, tmax: f32) -> f32 { var tr = 1.0; let n = min(u32(P.time.w), 32u); for (var k = 0u; k < n; k++) { - let sp = medium_span(media[k], o, d); + let sp = medium_span(medium_at(k), o, d); let b = min(sp.y, tmax); - let mj = media[k].b.x; + let mj = medium_at(k).b.x; if (max(sp.x, 0.0) >= b || mj <= 0.0) { continue; } var t = max(sp.x, 0.0); loop { @@ -2233,7 +2291,7 @@ fn media_march(o: vec3, d: vec3, tmax: f32, camera: bool) -> March { var dt = 1.0e30; var gridded = false; for (var k = 0u; k < n; k++) { - let sp = medium_span(media[k], o, d); + let sp = medium_span(medium_at(k), o, d); let a = max(sp.x, 0.0); let b = min(sp.y, tmax); if (a >= b) { continue; } @@ -2242,7 +2300,7 @@ fn media_march(o: vec3, d: vec3, tmax: f32, camera: bool) -> March { } else if (b > t) { mask |= 1u << k; end = min(end, b); - let grid = bitcast(media[k].b.w) != 0u; + let grid = bitcast(medium_at(k).b.w) != 0u; gridded = gridded || grid; dt = min(dt, (b - a) / (per * select(1.0, 8.0, grid))); } @@ -2265,7 +2323,7 @@ fn media_march(o: vec3, d: vec3, tmax: f32, camera: bool) -> March { st += v.extinction; sc += v.albedo * v.extinction; em += v.emission; - if (media[k].a.w != 0u) { + if (medium_at(k).a.w != 0u) { st_m += v.extinction; sc_m += v.albedo * v.extinction; em_m += v.emission; @@ -2315,7 +2373,7 @@ fn sample_hg(d: vec3, g: f32) -> vec3 { fn phase_direct(x: vec3, d: vec3, g: f32) -> vec3 { var sum = vec3(0.0); for (var i = 0u; i < P.counts2.x; i++) { - let li = lights[i]; + let li = light_at(i); let cos_max = li.dir.w; let cos_t = 1.0 - rnd() * (1.0 - cos_max); let sin_t = sqrt(max(0.0, 1.0 - cos_t * cos_t)); @@ -2337,7 +2395,7 @@ fn phase_direct(x: vec3, d: vec3, g: f32) -> vec3 { var pick_ls: LS; for (var k = 0u; k < 3u; k++) { let idx = nd + min(u32(rnd() * f32(nl)), nl - 1u); - let ls = local_light(lights[idx], x); + let ls = local_light(light_at(idx), x); let wk = luminance(ls.c) * hg_phase(g, dot(ls.l, d)); if (wk <= 0.0) { continue; } wsum += wk; @@ -2348,7 +2406,7 @@ fn phase_direct(x: vec3, d: vec3, g: f32) -> vec3 { } } if (pick_w > 0.0) { - let vis = light_vis(lights[pick].ext.w >= 0.5, x, pick_ls.sl, pick_ls.dist * 0.999); + let vis = light_vis(light_at(pick).ext.w >= 0.5, x, pick_ls.sl, pick_ls.dist * 0.999); if (any(vis > vec3(0.0))) { sum += pick_ls.c * gtint(vis) * hg_phase(g, dot(pick_ls.l, d)) * (wsum * f32(nl) / (3.0 * pick_w)); } } } @@ -2395,11 +2453,11 @@ fn cache_key_point(p: vec3, coarser: f32) -> vec2 { fn cache_point(p: vec3) -> vec4 { let f = cache_find(cache_key_point(p, VOLUME_FINE), false); var v = vec4(0.0); - if (f != 0xffffffffu) { v = cache_val[f * 2u]; } + if (f != 0xffffffffu) { v = cache_val_load(f * 2u); } if (v.w >= CACHE_WARM) { return v; } let c = cache_find(cache_key_point(p, VOLUME_COARSE), false); if (c != 0xffffffffu) { - let cv = cache_val[c * 2u]; + let cv = cache_val_load(c * 2u); if (cv.w > v.w) { return cv; } } return v; @@ -2412,7 +2470,7 @@ fn point_direct(x: vec3) -> vec3 { let ph = 1.0 / (4.0 * PI); var sum = vec3(0.0); for (var i = 0u; i < P.counts2.x; i++) { - let li = lights[i]; + let li = light_at(i); let cos_max = li.dir.w; let cos_t = 1.0 - rnd() * (1.0 - cos_max); let sin_t = sqrt(max(0.0, 1.0 - cos_t * cos_t)); @@ -2430,7 +2488,7 @@ fn point_direct(x: vec3) -> vec3 { var pick_ls: LS; for (var k = 0u; k < 3u; k++) { let idx = nd + min(u32(rnd() * f32(nl)), nl - 1u); - let ls = local_light(lights[idx], x); + let ls = local_light(light_at(idx), x); let wk = luminance(ls.c); if (wk <= 0.0) { continue; } wsum += wk; @@ -2440,7 +2498,7 @@ fn point_direct(x: vec3) -> vec3 { pick_ls = ls; } } - if (pick_w > 0.0 && (lights[pick].ext.w < 0.5 || !occluded(x, pick_ls.sl, pick_ls.dist * 0.999))) { + if (pick_w > 0.0 && (light_at(pick).ext.w < 0.5 || !occluded(x, pick_ls.sl, pick_ls.dist * 0.999))) { sum += pick_ls.c * ph * (wsum * f32(nl) / (3.0 * pick_w)); } } @@ -2551,17 +2609,87 @@ fn sphere_hit(x: vec3, l: vec3, hit: Hit) -> vec3 { // ------------------------------------------------------------------------------------------------ // What a live trace leaves the live passes (live.wgsl). -// Reflection resolve inputs (live; live.wgsl `refl`), 2 per traced pixel: the camera hit's specular-lobe ray -// (radiance it brought back, its lobe pdf; pdf 0 = the pixel sampled another lobe), (specular albedo E at the hit, -// octahedral ray direction). Neighbours reuse each other's rays weighted by their own BRDF (Lumen / SSSR). -@group(1) @binding(5) var refl: array>; -// Live motion vectors per traced pixel: where the camera hit was the previous frame, as a world-space offset (w = 1: -// on a moving triangle). -@group(1) @binding(6) var motion_out: array>; -// Live deterministic direct light per traced pixel (`dl_det` at the camera hit). -@group(1) @binding(7) var direct_out: array>; - -fn resolve_bound() -> bool { return LIVE && arrayLength(&refl) >= 2u * P.size.x * P.size.y; } +// Group 1, the buffers the kernel writes, used through the functions below in every layout: +// - refl: reflection resolve inputs (live; live.wgsl `refl`), 2 per traced pixel: the camera hit's specular-lobe ray +// (radiance it brought back, its lobe pdf; pdf 0 = the pixel sampled another lobe), (specular albedo E at the hit, +// octahedral ray direction). Neighbours reuse each other's rays weighted by their own BRDF (Lumen / SSSR); +// - motion_out: live motion vectors per traced pixel: where the camera hit was the previous frame, as a world-space +// offset (w = 1: on a moving triangle); +// - direct_out: live deterministic direct light per traced pixel (`dl_det` at the camera hit); +// - cache_keys, cache_acc, cache_val: the radiance cache (below, at `cache_on`); +// - loop_guard_trips: the loop guard's count (first word: the invocations a loop of which was stopped). +@if(merged) { + // The merged layout (merged.rs `KERNEL_LIVE`): the arrays of one struct, one buffer they are ranges of + // (gpu.rs `KernelBufs`). What the live passes read of it starts at a multiple of 256 bytes: they bind those + // ranges as buffers of their own (WebGPU's `minStorageBufferOffsetAlignment`). The loop guard's count is always + // there, written only by a kernel with the guard. + struct MergedLive { + cache_keys: array, ML_CACHE_KEYS>, + cache_acc: array, ML_CACHE_ACC>, + cache_val: array, ML_CACHE_VAL>, + @align(256) refl: array, ML_REFL>, + @align(256) motion_out: array, ML_MOTION_OUT>, + @align(256) direct_out: array, ML_DIRECT_OUT>, + loop_guard_trips: array, GUARD_WORDS>, + }; + @group(1) @binding(1) var ML: MergedLive; + fn refl_len() -> u32 { return ML_REFL; } + fn refl_store(i: u32, v: vec4) { ML.refl[i] = v; } + fn motion_len() -> u32 { return ML_MOTION_OUT; } + fn motion_store(i: u32, v: vec4) { ML.motion_out[i] = v; } + fn direct_len() -> u32 { return ML_DIRECT_OUT; } + fn direct_load(i: u32) -> vec4 { return ML.direct_out[i]; } + fn direct_store(i: u32, v: vec4) { ML.direct_out[i] = v; } + fn cache_key_load(i: u32) -> u32 { return atomicLoad(&ML.cache_keys[i]); } + fn cache_key_claim(i: u32, k: u32) -> bool { let r = atomicCompareExchangeWeak(&ML.cache_keys[i], 0u, k); return r.exchanged || r.old_value == k; } + fn cache_key_clear(i: u32) { atomicStore(&ML.cache_keys[i], 0u); } + fn cache_acc_add(i: u32, v: u32) -> u32 { return atomicAdd(&ML.cache_acc[i], v); } + fn cache_acc_load(i: u32) -> u32 { return atomicLoad(&ML.cache_acc[i]); } + fn cache_acc_clear(i: u32) { atomicStore(&ML.cache_acc[i], 0u); } + fn cache_val_load(i: u32) -> vec4 { return ML.cache_val[i]; } + fn cache_val_store(i: u32, v: vec4) { ML.cache_val[i] = v; } + @if(guard) fn loop_guard_stopped() { atomicAdd(&ML.loop_guard_trips[0], 1u); } +} @else { + @group(1) @binding(1) var cache_keys: array>; + @group(1) @binding(2) var cache_acc: array>; + @group(1) @binding(3) var cache_val: array>; + @group(1) @binding(5) var refl: array>; + @group(1) @binding(6) var motion_out: array>; + @group(1) @binding(7) var direct_out: array>; + fn refl_len() -> u32 { return arrayLength(&refl); } + fn refl_store(i: u32, v: vec4) { refl[i] = v; } + fn motion_len() -> u32 { return arrayLength(&motion_out); } + fn motion_store(i: u32, v: vec4) { motion_out[i] = v; } + fn direct_len() -> u32 { return arrayLength(&direct_out); } + fn direct_load(i: u32) -> vec4 { return direct_out[i]; } + fn direct_store(i: u32, v: vec4) { direct_out[i] = v; } + fn cache_key_load(i: u32) -> u32 { return atomicLoad(&cache_keys[i]); } + fn cache_key_claim(i: u32, k: u32) -> bool { let r = atomicCompareExchangeWeak(&cache_keys[i], 0u, k); return r.exchanged || r.old_value == k; } + fn cache_key_clear(i: u32) { atomicStore(&cache_keys[i], 0u); } + fn cache_acc_add(i: u32, v: u32) -> u32 { return atomicAdd(&cache_acc[i], v); } + fn cache_acc_load(i: u32) -> u32 { return atomicLoad(&cache_acc[i]); } + fn cache_acc_clear(i: u32) { atomicStore(&cache_acc[i], 0u); } + fn cache_val_load(i: u32) -> vec4 { return cache_val[i]; } + fn cache_val_store(i: u32, v: vec4) { cache_val[i] = v; } + @if(guard) { + @group(1) @binding(8) var loop_guard_trips: array, GUARD_WORDS>; + fn loop_guard_stopped() { atomicAdd(&loop_guard_trips[0], 1u); } + } +} + +// The loop guard (gpu.rs `loop_guard_of`; docs/RUNTIME.md "The engine in a browser"): every loop of the kernel, the +// scene's functions' too, begins with `if (loop_guard()) { break; }` (gpu.rs `guard_loops`, a pass over the variant's +// syntax tree). The kernel's traversals follow indices they read from buffers and end only on sound data; with the +// guard a dispatch always ends: an invocation's iterations are counted over all its loops together, past +// LOOP_GUARD_LIMIT it leaves them, and the invocations stopped are counted for the renderer to read back. +@if(guard) var LOOP_GUARD_N: u32 = 0u; +@if(guard) fn loop_guard() -> bool { + LOOP_GUARD_N += 1u; + if (LOOP_GUARD_N == LOOP_GUARD_LIMIT + 1u) { loop_guard_stopped(); } + return LOOP_GUARD_N > LOOP_GUARD_LIMIT; +} + +fn resolve_bound() -> bool { return LIVE && refl_len() >= 2u * P.size.x * P.size.y; } // THREERT_CACHE_FLAGS bit 16 turns the resolve off (rays then stay in their pixel's specular, the resolve reads zeros). fn resolve_on() -> bool { return resolve_bound() && (u32(P.cache.w) & 16u) == 0u; } @@ -2602,9 +2730,7 @@ const CACHE_EVICT: f32 = 120.0; // frames without samples before an entry i // cache_keys[i]: checksum of the voxel key (0 = free). cache_acc[8i ..]: this frame's sums (fixed point) of the total // r, g, b, of its indirect part r, g, b, and the sample count. cache_val[2i] = (irradiance / π, history samples), // cache_val[2i + 1] = (its indirect part: all but the voxel's own direct light, frames since the last sample). -@group(1) @binding(1) var cache_keys: array>; -@group(1) @binding(2) var cache_acc: array>; -@group(1) @binding(3) var cache_val: array>; +// (Declared with the trace's outputs, group 1, above `resolve_bound`.) fn cache_on() -> bool { return LIVE && P.cache.x > 0.5; } @@ -2632,13 +2758,12 @@ fn cache_find(k: vec2, insert: bool) -> u32 { var free = 0xffffffffu; for (var i = 0u; i < CACHE_PROBE; i++) { let s = (k.x + i) & (CACHE_SIZE - 1u); - let c = atomicLoad(&cache_keys[s]); + let c = cache_key_load(s); if (c == k.y) { return s; } if (c == 0u && free == 0xffffffffu) { free = s; } } if (!insert || free == 0xffffffffu) { return 0xffffffffu; } - let r = atomicCompareExchangeWeak(&cache_keys[free], 0u, k.y); - if (r.exchanged || r.old_value == k.y) { return free; } + if (cache_key_claim(free, k.y)) { return free; } return 0xffffffffu; } @@ -2660,8 +2785,8 @@ fn cache_get(s: Surface, indirect: bool) -> vec4 { fn cache_entry(k: vec2, indirect: bool) -> vec4 { let i = cache_find(k, false); if (i == 0xffffffffu) { return vec4(0.0); } - let v = cache_val[i * 2u]; - if (indirect) { return vec4(cache_val[i * 2u + 1u].rgb, v.w); } + let v = cache_val_load(i * 2u); + if (indirect) { return vec4(cache_val_load(i * 2u + 1u).rgb, v.w); } return v; } @@ -2672,15 +2797,15 @@ fn cache_fixed(e: vec3) -> vec3 { fn cache_add(i: u32, total: vec3, indirect: vec3) { if (i == 0xffffffffu || non_finite(total) || non_finite(indirect)) { return; } - if (atomicAdd(&cache_acc[i * 8u + 6u], 1u) >= CACHE_FRAME_CAP) { return; } + if (cache_acc_add(i * 8u + 6u, 1u) >= CACHE_FRAME_CAP) { return; } let a = cache_fixed(total); let b = cache_fixed(indirect); - atomicAdd(&cache_acc[i * 8u], a.x); - atomicAdd(&cache_acc[i * 8u + 1u], a.y); - atomicAdd(&cache_acc[i * 8u + 2u], a.z); - atomicAdd(&cache_acc[i * 8u + 3u], b.x); - atomicAdd(&cache_acc[i * 8u + 4u], b.y); - atomicAdd(&cache_acc[i * 8u + 5u], b.z); + cache_acc_add(i * 8u, a.x); + cache_acc_add(i * 8u + 1u, a.y); + cache_acc_add(i * 8u + 2u, a.z); + cache_acc_add(i * 8u + 3u, b.x); + cache_acc_add(i * 8u + 4u, b.y); + cache_acc_add(i * 8u + 5u, b.z); } // White Lambertian lobes: lighting gathered with them is irradiance / π, and a cosine-sampled ray has weight 1. @@ -2824,16 +2949,16 @@ fn media_train(@builtin(global_invocation_id) gid: vec3) { let n = min(u32(P.time.w), 32u); var total = 0.0; for (var k = 0u; k < n; k++) { - let sp = medium_span(media[k], o, d); - total += max(min(sp.y, select(tmax, tmax_h, media[k].a.z != 0u)) - max(sp.x, 0.0), 0.0); + let sp = medium_span(medium_at(k), o, d); + total += max(min(sp.y, select(tmax, tmax_h, medium_at(k).a.z != 0u)) - max(sp.x, 0.0), 0.0); } if (total <= 0.0) { return; } var u = rnd() * total; var x = o; for (var k = 0u; k < n; k++) { - let sp = medium_span(media[k], o, d); + let sp = medium_span(medium_at(k), o, d); let a = max(sp.x, 0.0); - let len = max(min(sp.y, select(tmax, tmax_h, media[k].a.z != 0u)) - a, 0.0); + let len = max(min(sp.y, select(tmax, tmax_h, medium_at(k).a.z != 0u)) - a, 0.0); if (u < len) { x = o + d * (a + u); break; @@ -2876,7 +3001,7 @@ fn media_composite(@builtin(global_invocation_id) gid: vec3) { px.normal = vec4(px.normal.xyz, floor(clamp(share, 0.0, 1.0) * 255.0) + 0.5 * clamp(mm.tr, 0.0, 1.0)); px.emit = vec4(px.emit.rgb * mm.tr + mm.light * samples, px.emit.w); accum[pix] = px; - if (arrayLength(&direct_out) >= w * h) { direct_out[pix] = vec4(direct_out[pix].rgb * mm.tr, direct_out[pix].w); } + if (direct_len() >= w * h) { direct_store(pix, vec4(direct_load(pix).rgb * mm.tr, direct_load(pix).w)); } } // Once per frame after the trace: folds the frame's training sums into each entry (history capped at P.cache.y @@ -2884,29 +3009,29 @@ fn media_composite(@builtin(global_invocation_id) gid: vec3) { @compute @workgroup_size(256, 1, 1) fn cache_resolve(@builtin(global_invocation_id) gid: vec3) { let i = gid.x; - if (i >= CACHE_SIZE || atomicLoad(&cache_keys[i]) == 0u) { return; } - let n = min(atomicLoad(&cache_acc[i * 8u + 6u]), CACHE_FRAME_CAP); - var v = cache_val[i * 2u]; - var ind = cache_val[i * 2u + 1u]; + if (i >= CACHE_SIZE || cache_key_load(i) == 0u) { return; } + let n = min(cache_acc_load(i * 8u + 6u), CACHE_FRAME_CAP); + var v = cache_val_load(i * 2u); + var ind = cache_val_load(i * 2u + 1u); if (n > 0u) { var sum: array; - for (var c = 0u; c < 6u; c++) { sum[c] = f32(atomicLoad(&cache_acc[i * 8u + c])) * (P.misc.x / 2097152.0) / f32(n); } + for (var c = 0u; c < 6u; c++) { sum[c] = f32(cache_acc_load(i * 8u + c)) * (P.misc.x / 2097152.0) / f32(n); } let hist = min(v.w + f32(n), P.cache.y); // The frame cap counts frames, also those without samples since the last update (sparse voxels). let a = max(f32(n) / hist, min((ind.w + 1.0) / P.cache.z, 1.0)); v = vec4(mix(v.rgb, vec3(sum[0], sum[1], sum[2]), a), hist); ind = vec4(mix(ind.rgb, vec3(sum[3], sum[4], sum[5]), a), 0.0); - for (var c = 0u; c < 8u; c++) { atomicStore(&cache_acc[i * 8u + c], 0u); } + for (var c = 0u; c < 8u; c++) { cache_acc_clear(i * 8u + c); } } else { ind.w += 1.0; if (ind.w > CACHE_EVICT) { - atomicStore(&cache_keys[i], 0u); + cache_key_clear(i); v = vec4(0.0); ind = vec4(0.0); } } - cache_val[i * 2u] = v; - cache_val[i * 2u + 1u] = ind; + cache_val_store(i * 2u, v); + cache_val_store(i * 2u + 1u, ind); } // ------------------------------------------------------------------------------------------------ @@ -2919,7 +3044,7 @@ fn main(@builtin(global_invocation_id) gid: vec3) { let pix = gid.y * w + gid.x; // Every pixel's deterministic direct light is written each frame (0 unless the camera hits a surface): a pixel that // turned to sky must not keep what a surface left there. - if (LIVE && P.misc.w < 0.5 && arrayLength(&direct_out) >= w * h) { direct_out[pix] = vec4(0.0); } + if (LIVE && P.misc.w < 0.5 && direct_len() >= w * h) { direct_store(pix, vec4(0.0)); } rng_state = pcg(pix ^ pcg(P.size.z * 9781u + P.size.w * 6271u)); // Raster primary: one visibility sample per pixel at the frame's jitter (lens sampling would need its own rays). @@ -3113,7 +3238,7 @@ fn main(@builtin(global_invocation_id) gid: vec3) { let kp = (ti - P.group_first[1].x) * 3u; if (LIVE && P.motion.x > 0.5 && is_proc(hit)) { // Procedural geometry: its object-space hit point through the previous frame's rows. - let r = instances[hit.inst]; + let r = instance_at(hit.inst); if (any(r.p0 != r.m0) || any(r.p1 != r.m1) || any(r.p2 != r.m2)) { let lp = vec4(proc_point(r, s.p), 1.0); mv = vec4(vec3(dot(r.p0, lp), dot(r.p1, lp), dot(r.p2, lp)) - s.p, 1.0); @@ -3121,7 +3246,7 @@ fn main(@builtin(global_invocation_id) gid: vec3) { } else if (LIVE && P.motion.x > 0.5 && has(F_INST) && hit.inst != NO_INST) { // A movable instance: its model-space hit point (the previous pose's for a deforming mesh) through the // previous frame's transform (p rows; the same as m when it did not move). - let r = instances[hit.inst]; + let r = instance_at(hit.inst); if (deforms || any(r.p0 != r.m0) || any(r.p1 != r.m1) || any(r.p2 != r.m2)) { var lp = vec4(tri_verts[ti * 3u].xyz + tri_verts[ti * 3u + 1u].xyz * hit.u + tri_verts[ti * 3u + 2u].xyz * hit.v, 1.0); if (deforms) { lp = vec4(prev_verts[kp].xyz + prev_verts[kp + 1u].xyz * hit.u + prev_verts[kp + 2u].xyz * hit.v, 1.0); } @@ -3206,10 +3331,10 @@ fn main(@builtin(global_invocation_id) gid: vec3) { // Live: the deterministic direct light leaves the radiance (and its specular part the specular) for a signal of // its own, written here so nothing more stays live through the path. It carries the transmittance to the camera // (fog, media) as the radiance did: left out, the diffuse signal's clamp at 0 dropped the attenuation. - if (LIVE && depth == 0u && arrayLength(&direct_out) >= P.size.x * P.size.y && !non_finite(dl_det) && !non_finite(dl_det_spec)) { + if (LIVE && depth == 0u && direct_len() >= P.size.x * P.size.y && !non_finite(dl_det) && !non_finite(dl_det_spec)) { radiance -= thr * dl_det; dl_spec = max(dl_spec - dl_det_spec, vec3(0.0)); - if (P.misc.w < 0.5) { direct_out[gid.y * P.size.x + gid.x] = vec4(thr * dl_det, 0.0); } + if (P.misc.w < 0.5) { direct_store(gid.y * P.size.x + gid.x, vec4(thr * dl_det, 0.0)); } } let l = sample_bsdf(s, lb, v); @@ -3288,8 +3413,8 @@ fn main(@builtin(global_invocation_id) gid: vec3) { if (non_finite(dropped)) { dropped = vec3(0.0); } spec = clamp(spec, vec3(0.0), max(radiance - e0, vec3(0.0))); let diff = max(radiance - e0 - spec - dropped, vec3(0.0)); - if (P.misc.w < 0.5 && arrayLength(&motion_out) >= w * h) { - motion_out[pix] = select(vec4(0.0), mv, !non_finite(mv.xyz)); + if (P.misc.w < 0.5 && motion_len() >= w * h) { + motion_store(pix, select(vec4(0.0), mv, !non_finite(mv.xyz))); } if (resolve_bound() && P.misc.w < 0.5) { // Radiance along the specular ray: what the path gathered after it, over the ray's weight. @@ -3301,8 +3426,8 @@ fn main(@builtin(global_invocation_id) gid: vec3) { // A sample that took the diffuse lobe leaves its specular share (dropped) for the resolve to return when no // specular-lobe ray around it can stand in (w = -1 marks it). if (resolve && !(refl_pdf > 0.0)) { li = dropped; } - refl[pix * 2u] = vec4(li, select(select(0.0, -1.0, resolve), min(refl_pdf, 1e6), resolve && refl_pdf > 0.0)); - refl[pix * 2u + 1u] = vec4(select(vec3(0.0), refl_e, resolve), bitcast(oct_enc(select(vec3(0.0, 0.0, 1.0), refl_l, refl_pdf > 0.0)))); + refl_store(pix * 2u, vec4(li, select(select(0.0, -1.0, resolve), min(refl_pdf, 1e6), resolve && refl_pdf > 0.0))); + refl_store(pix * 2u + 1u, vec4(select(vec3(0.0), refl_e, resolve), bitcast(oct_enc(select(vec3(0.0, 0.0, 1.0), refl_l, refl_pdf > 0.0))))); } // misc.w = index of this sample within the frame: the first one also writes the g-buffer. if (P.misc.w < 0.5) { diff --git a/renderer/shaders/trace_hw.wgsl b/renderer/shaders/trace_hw.wgsl index fd3ee89..1569452 100644 --- a/renderer/shaders/trace_hw.wgsl +++ b/renderer/shaders/trace_hw.wgsl @@ -125,7 +125,7 @@ // An alpha-cutout triangle: its material has an alpha test (as the build flags TRI_MASKED; model instances' pool // triangles carry no flags). @if(hwrt) fn cutout(ti: u32) -> bool { - return materials[bitcast(attrs[ti].uv.w)].uv2.w > 0.0; + return material_at(bitcast(attrs[ti].uv.w)).uv2.w > 0.0; } // The model-instance record a hit's custom data names (F_INST), or NO_INST for a traversal group. @@ -139,7 +139,7 @@ // traversal group's. @if(hwrt) fn hit_tri(custom: u32, geometry: u32, prim: u32) -> u32 { if (has(F_INST) && (custom & INSTANCE_BIT) != 0u) { - let r = instances[custom & (INSTANCE_BIT - 1u)].info; + let r = instance_at(custom & (INSTANCE_BIT - 1u)).info; return r.x + select(0u, r.w, geometry == 1u) + prim; } return group_tri(custom, prim); diff --git a/renderer/shaders/trace_units.wgsl b/renderer/shaders/trace_units.wgsl index eba43c2..be6af43 100644 --- a/renderer/shaders/trace_units.wgsl +++ b/renderer/shaders/trace_units.wgsl @@ -77,7 +77,7 @@ } else if (rec == NO_INST) { // A top leaf: one unit. The ray in the unit's space (its direction not renormalised: t is the world's): // `proc_point`, `proc_dir`, or the origin moved alone. - let u = instances[n.left_or_first]; + let u = instance_at(n.left_or_first); if (mode != TRACE_SHADOW || bitcast(u.n1.w) != 0u) { rec = n.left_or_first; ro = o - vec3(u.m0.w, u.m1.w, u.m2.w); diff --git a/renderer/shaders/translucent.wgsl b/renderer/shaders/translucent.wgsl index cd3e579..fc9e522 100644 --- a/renderer/shaders/translucent.wgsl +++ b/renderer/shaders/translucent.wgsl @@ -58,7 +58,7 @@ // material function. @if(hwrt) fn particle_material(ti: u32, b: vec2, view: vec3, lod: f32) -> MatOut { let a = attrs[ti]; - let m = materials[bitcast(a.uv.w)]; + let m = material_at(bitcast(a.uv.w)); let uv = mesh_uv(a, m, b.x, b.y); let tl = mesh_lod(a, m) + lod - LOD_BIAS; var r = mat_base(m); diff --git a/renderer/src/gpu.rs b/renderer/src/gpu.rs index e482d25..fc999e3 100644 --- a/renderer/src/gpu.rs +++ b/renderer/src/gpu.rs @@ -14,7 +14,7 @@ pub struct Gpu { pub device: wgpu::Device, pub queue: wgpu::Queue, pub adapter_info: wgpu::AdapterInfo, - /// The kernel's bounded tables are uniform buffers (`GpuOptions::uniform_tables`; merged.rs `uniform_tables`). + /// The kernel's bounded tables are uniform buffers (`GpuOptions::uniform_tables`; merged.rs `UNIFORM_TABLES`). pub uniform_tables: bool, /// The kernel's switches this device was opened with (`KernelSwitches`): they are in `source`. pub switches: KernelSwitches, @@ -27,7 +27,7 @@ pub struct Gpu { /// The merged binding layout (merged.rs): a browser's, where a stage has eight storage buffers. Natively with /// THREERT_MERGED=1 (software traversal, the device held to eight too). pub merged: bool, - /// By (material-function code hash, the features that change the source: `kernel_flags`, the merged layout's lengths). + /// By (material-function code hash, the features that change the source: `KernelVariant::flags`, the merged layout's lengths). modules: std::sync::Mutex), wgpu::ShaderModule>>, pipeline_layout: wgpu::PipelineLayout, /// Kernel variants by (feature mask, material-function code hash, LIVE): `FEAT` compiles out the features a scene @@ -114,7 +114,7 @@ pub struct Params { /// The sky's aerial perspective (F_AERIAL; `scene::aerial_params`): density multiple, Rayleigh scale, Mie scale, -. pub aerial: [f32; 4], /// three.js `scene.background` (F_BACKGROUND): its radiance, 1. Last, and read only by the kernel patched for it - /// (`kernel_flags`): every other kernel's `Params` ends before it. + /// (`KernelVariant::flags`): every other kernel's `Params` ends before it. pub background: [f32; 4], } @@ -711,56 +711,182 @@ pub fn kernel_template() -> String { KERNEL_FILES.join("\n") } -/// The flags of the kernel's conditional translation (pathtrace.wgsl's head): a device's traversal and switches, and -/// a scene's features that change the kernel's declarations (`F_SOURCE`; the others switch code under `override FEAT`). -/// - vcolor (F_VCOLOR), three.js `vertexColors` (`scene::tri_colors`): three sRGB 8-bit colours per triangle after the -/// previous vertices in `prev_verts` (from vec4 `Params::procs.w`; the kernel has no binding to spare), interpolated -/// in linear at the hit, multiply the base colour of surfaces and glass (white for triangles without them). -/// - grid (F_GRID), density grids (`scene::GridBuild`): the grid atlas at binding 19 (a texture: the buffers are at -/// Metal's limit), `VIn::grid` and `vgrid` for volume functions. -/// - tint (F_TINT), instance colours (`scene::object_stride`): object entries of 4 vec4s, the 4th the instance colour, -/// which multiplies the base colour of surfaces, procedural surfaces and glass (three.js `instanceColor`). -/// - unlit (F_UNLIT), three.js MeshBasicMaterial (`scene::MAT_UNLIT`): its colour (x map, x tint) is emission seen by -/// camera rays and mirror-like paths only, its albedo 0: lit by nothing, lighting nothing (in no emitter list). -/// - units (F_UNITS), software traversal's second level (trace_units.wgsl): the still scene's walk hands its hit to -/// the units' and has no stack of its own (the module's, set to zero once an invocation and not once a ray). -/// - background (F_BACKGROUND), three.js `scene.background`: camera rays that miss, and paths bent straight on through -/// glass (the rays `sky_mis` gets with pdf 0), see `Params::background`, which `Params` ends with in this kernel only. -pub fn kernel_flags(hwrt: bool, k: &KernelSwitches, features: u32) -> [(&'static str, bool); 12] { - let f = |bit: u32| features & bit != 0; - [ - ("hwrt", hwrt), - ("strict", k.strict), - ("cutout_lite", k.cutout_lite >= 1), - ("cutout_lite2", k.cutout_lite >= 2), - ("last_direct", k.last_direct), - ("light_pick", k.light_pick), - ("vcolor", f(F_VCOLOR)), - ("grid", f(F_GRID)), - ("tint", f(F_TINT)), - ("unlit", f(F_UNLIT)), - ("units", f(F_UNITS)), - ("background", f(F_BACKGROUND)), - ] -} - -/// A kernel's WGSL: `src` (the template, with what is put into its text first: the scene's functions, the merged -/// layout, the uniform tables, the loop guard) through WESL's conditional translation for `flags` (wesl-rs, WESL's -/// reference compiler in Rust). Every flag the source names must be given: one that is not is an error, not off. -pub fn kernel_wgsl(src: &str, flags: &[(&str, bool)]) -> String { - use wesl::{resolver::VirtualResolver, syntax::ModulePath, CompileOptions, Feature, Features}; - let path: ModulePath = "package::kernel".parse().expect("a module path"); - let mut resolver = VirtualResolver::new(); - resolver.add_module(path.clone(), std::borrow::Cow::Borrowed(src)); +/// What a kernel is compiled for (pathtrace.wgsl's head names the flags): the device's traversal, switches, binding +/// layout and loop guard, and the scene's features that change the kernel's declarations (`F_SOURCE`; the others +/// switch code under `override FEAT`). +#[derive(Clone, Copy, Debug, Default)] +pub struct KernelVariant { + pub hwrt: bool, + pub switches: KernelSwitches, + /// The scene's features: `F_SOURCE`'s are flags. + pub features: u32, + /// The merged layout (merged.rs), with its arrays' lengths. + pub merged: Option, + /// Four tables in uniform buffers (merged.rs `UNIFORM_TABLES`; `GpuOptions::uniform_tables`). Never with `merged`. + pub uniform_tables: bool, + /// The loop guard's limit, in a kernel with the guard (`loop_guard_of`, `guard_loops`). + pub guard: Option, +} + +impl KernelVariant { + /// A device's traversal and switches, a scene's features; the layout as written, no guard. + pub fn new(hwrt: bool, switches: KernelSwitches, features: u32) -> KernelVariant { + KernelVariant { hwrt, switches, features, ..Default::default() } + } + + /// The flags of the kernel's conditional translation. A scene's: + /// - vcolor (F_VCOLOR), three.js `vertexColors` (`scene::tri_colors`): three sRGB 8-bit colours per triangle after + /// the previous vertices in `prev_verts` (from vec4 `Params::procs.w`; the kernel has no binding to spare), + /// interpolated in linear at the hit, multiply the base colour of surfaces and glass (white for triangles + /// without them). + /// - grid (F_GRID), density grids (`scene::GridBuild`): the grid atlas at binding 19 (a texture: the buffers are + /// at Metal's limit), `VIn::grid` and `vgrid` for volume functions. + /// - tint (F_TINT), instance colours (`scene::object_stride`): object entries of 4 vec4s, the 4th the instance + /// colour, which multiplies the base colour of surfaces, procedural surfaces and glass (three.js `instanceColor`). + /// - unlit (F_UNLIT), three.js MeshBasicMaterial (`scene::MAT_UNLIT`): its colour (x map, x tint) is emission seen + /// by camera rays and mirror-like paths only, its albedo 0: lit by nothing, lighting nothing (in no emitter list). + /// - units (F_UNITS), software traversal's second level (trace_units.wgsl): the still scene's walk hands its hit + /// to the units' and has no stack of its own (the module's, set to zero once an invocation and not once a ray). + /// - background (F_BACKGROUND), three.js `scene.background`: camera rays that miss, and paths bent straight on + /// through glass (the rays `sky_mis` gets with pdf 0), see `Params::background`, which `Params` ends with in + /// this kernel only. + pub fn flags(&self) -> [(&'static str, bool); 15] { + let (k, f) = (&self.switches, |bit: u32| self.features & bit != 0); + [ + ("hwrt", self.hwrt), + ("strict", k.strict), + ("cutout_lite", k.cutout_lite >= 1), + ("cutout_lite2", k.cutout_lite >= 2), + ("last_direct", k.last_direct), + ("light_pick", k.light_pick), + ("vcolor", f(F_VCOLOR)), + ("grid", f(F_GRID)), + ("tint", f(F_TINT)), + ("unlit", f(F_UNLIT)), + ("units", f(F_UNITS)), + ("background", f(F_BACKGROUND)), + ("merged", self.merged.is_some()), + ("uniform_tables", self.uniform_tables), + ("guard", self.guard.is_some()), + ] + } + + /// The numbers the variant is compiled for, WESL constants (`import constants::`): the merged layout's array + /// lengths (`MT_`, `ML_`), the uniform tables' (`UT_
`), the loop guard's count's words and + /// its limit. + pub fn constants(&self) -> Vec<(String, u32)> { + use crate::merged::{GUARD_WORDS, KERNEL_LIVE, TABLES, UNIFORM_LEN, UNIFORM_TABLES}; + let mut out = vec![]; + if let Some(c) = &self.merged { + out.extend(TABLES.iter().zip(c.tables).map(|(m, n)| (format!("MT_{}", m.name.to_uppercase()), n))); + out.extend(KERNEL_LIVE.iter().zip(c.live).map(|(m, n)| (format!("ML_{}", m.name.to_uppercase()), n))); + } + if self.uniform_tables { + out.extend(UNIFORM_TABLES.iter().zip(UNIFORM_LEN).map(|(&k, n)| (format!("UT_{}", TABLES[k].name.to_uppercase()), n))); + } + if self.merged.is_some() || self.guard.is_some() { + out.push(("GUARD_WORDS".into(), GUARD_WORDS)); + } + if let Some(n) = self.guard { + out.push(("LOOP_GUARD_LIMIT".into(), n)); + } + out + } +} + +/// A kernel's WGSL: `src` (the template, the scene's functions spliced in) for a variant: WESL's conditional +/// translation for its flags, with its constants, then the loop guard if it has one (`guard_loops`). +pub fn kernel_wgsl(src: &str, v: &KernelVariant) -> String { + assert!(!(v.merged.is_some() && v.uniform_tables), "the uniform tables are never in the merged layout"); + let mut out = wesl_wgsl(src, &v.flags(), &v.constants()); + if v.guard.is_some() { + guard_loops(&mut out); + } + out.to_string() +} + +/// `src` through wesl-rs (WESL's reference compiler in Rust): conditional translation for `flags`, imports of the +/// `constants` package from `constants`. Every flag the source names must be given (one that is not is an error, not +/// off), and every constant it imports. +pub fn wesl_wgsl(src: &str, flags: &[(&str, bool)], constants: &[(String, u32)]) -> wesl::syntax::TranslationUnit { + use wesl::{syntax::ModulePath, CompileOptions, Feature, Features}; let mut features = Features { default: Feature::Error, ..Default::default() }; for &(name, on) in flags { features.set(name, on); } - let options = CompileOptions { imports: false, condcomp: true, strip: false, validate: false, sourcemap: false, mangler: wesl::ManglerKind::None, features, ..Default::default() }; - wesl::compile(&path, &options, &resolver).unwrap_or_else(|e| panic!("the kernel's WESL: {e}")).to_string() + let constants = constants.iter().map(|(name, n)| format!("public const {name} = {n}u;\n")).collect(); + let resolver = KernelResolver { src, constants }; + let options = CompileOptions { imports: true, condcomp: true, strip: false, validate: false, sourcemap: false, mangler: wesl::ManglerKind::None, features, ..Default::default() }; + let path: ModulePath = "package::kernel".parse().expect("a module path"); + wesl::compile(&path, &options, &resolver).unwrap_or_else(|e| panic!("the kernel's WESL: {e}")).syntax +} + +/// The modules a kernel's WESL sees: itself, and WESL's `constants` package (`wesl_wgsl`). +struct KernelResolver<'a> { + src: &'a str, + constants: String, +} + +impl wesl::resolver::Resolver for KernelResolver<'_> { + fn resolve_source<'a>(&'a self, path: &wesl::syntax::ModulePath) -> Result, wesl::error::ResolveError> { + use wesl::syntax::PathOrigin; + match &path.origin { + PathOrigin::Package(p) if p == "constants" => Ok(std::borrow::Cow::Borrowed(&self.constants)), + PathOrigin::Absolute if path.components == ["kernel"] => Ok(std::borrow::Cow::Borrowed(self.src)), + _ => Err(wesl::error::ResolveError::ModuleNotFound(path.clone(), "the kernel imports nothing but `constants`".into())), + } + } +} + +/// The loop guard (docs/RUNTIME.md "The engine in a browser"): every `loop`, `for` and `while` of a variant, the +/// scene's functions' too, begins with `if (loop_guard()) { break; }` (pathtrace.wgsl declares `loop_guard`, under +/// `@if(guard)`). A pass over the variant's syntax tree, as a compiler instruments code: the tree knows what a loop +/// and its body are, so nothing the text says (comments, strings, a loop without braces) can mislead it. +pub fn guard_loops(tu: &mut wesl::syntax::TranslationUnit) { + use wesl::syntax::{CompoundStatement, GlobalDeclaration, Statement, StatementNode}; + fn body(b: &mut CompoundStatement, guard: &StatementNode) { + for s in b.statements.iter_mut() { + statement(s, guard); + } + } + fn opened(b: &mut CompoundStatement, guard: &StatementNode) { + body(b, guard); + b.statements.insert(0, guard.clone()); + } + fn statement(s: &mut StatementNode, guard: &StatementNode) { + match s.node_mut() { + Statement::Compound(c) => body(c, guard), + Statement::If(i) => { + body(&mut i.if_clause.body, guard); + for e in i.else_if_clauses.iter_mut() { + body(&mut e.body, guard); + } + if let Some(e) = i.else_clause.as_mut() { + body(&mut e.body, guard); + } + } + Statement::Switch(sw) => sw.clauses.iter_mut().for_each(|c| body(&mut c.body, guard)), + Statement::Loop(l) => { + opened(&mut l.body, guard); + if let Some(c) = l.continuing.as_mut() { + body(&mut c.body, guard); + } + } + Statement::For(f) => opened(&mut f.body, guard), + Statement::While(w) => opened(&mut w.body, guard), + _ => {} + } + } + let guard: Statement = "if (loop_guard()) { break; }".parse().expect("the guard's statement"); + let guard = StatementNode::from(guard); + for d in tu.global_declarations.iter_mut() { + if let GlobalDeclaration::Function(f) = d.node_mut() { + body(&mut f.body, &guard); + } + } } -/// The loop guard's limit for every kernel, if one is set (merged.rs `guard`: every loop of a kernel bounded, so a +/// The loop guard's limit for every kernel, if one is set (`guard_loops`: every loop of a kernel bounded, so a /// dispatch ends whatever the data): THREERT_LOOP_GUARD= natively, where a kernel is /// otherwise as written; always in a browser, where a dispatch that does not end is the player's whole GPU. pub fn loop_guard() -> Option { @@ -811,7 +937,7 @@ fn guard_room(hwrt: bool) -> bool { #[derive(Clone, Debug, Default, PartialEq)] pub struct GpuOptions { /// The kernel's bounded tables (texture pages, objects, instances, media) as uniform buffers instead of storage - /// buffers (merged.rs `uniform_tables`). None: where a driver needs it, which is Android. Never with the merged + /// buffers (merged.rs `UNIFORM_TABLES`). None: where a driver needs it, which is Android. Never with the merged /// layout, whose tables are one buffer of its own. pub uniform_tables: Option, /// A directory to keep compiled pipelines in between runs (pipeline_cache.rs). None: nothing is kept. @@ -842,7 +968,7 @@ impl GpuOptions { } /// The kernel's switches: changes to its source for a GPU that the source as written is wrong or too slow on. Each -/// is a flag of the kernel's conditional translation (`kernel_flags`; pathtrace.wgsl, trace_hw.wgsl and +/// is a flag of the kernel's conditional translation (`KernelVariant::flags`; pathtrace.wgsl, trace_hw.wgsl and /// translucent.wgsl say where, under `@if`); with none set the kernel is as written, every platform's default but /// Android's. #[derive(Clone, Copy, Debug, Default, PartialEq, Eq)] @@ -896,17 +1022,6 @@ fn kernel_checks() -> wgpu::ShaderRuntimeChecks { c } -/// A kernel source as it is compiled: with the loop guard if it has one (`loop_guard_of`: `units`, whether the source -/// has the unit walk), its count in `MergedLive` (the merged layout) or in a buffer of its own (group 1, -/// `merged::GUARD_BINDING`). -fn guarded(src: &str, merged: bool, units: bool) -> String { - match loop_guard_of(units) { - None => src.to_string(), - Some(n) if merged => crate::merged::guard(src, n, "ML.loop_guard_trips"), - Some(n) => crate::merged::guard(&format!("{src}\n{}", crate::merged::guard_decl()), n, crate::merged::GUARD_TRIPS.name), - } -} - pub fn tlas_rows(w: glam::Mat4) -> [f32; 12] { let row = |r: usize| [w.x_axis[r], w.y_axis[r], w.z_axis[r], w.w_axis[r]]; let [r0, r1, r2] = [row(0), row(1), row(2)]; @@ -975,7 +1090,7 @@ pub struct KernelBufs { pub refl: Buf, pub motion: Buf, pub direct: Buf, - /// The loop guard's count (merged.rs `guard`: its first word, the invocations a loop of which was stopped): with + /// The loop guard's count (`guard_loops`: its first word, the invocations a loop of which was stopped): with /// the guard, and always in the merged layout (the last array of `MergedLive`). pub guard: Option, /// The merged layout: the one buffer, and its arrays' lengths. @@ -1175,14 +1290,14 @@ pub const F_HMEDIA: u32 = 131072; /// Procedural geometry (`ProceduralMesh`; pathtrace.wgsl `proc_*`): AABB candidates in every ray query, its hits' /// own surface path. Compiled out where no scene has any (`THREERT_FEATURES_ON` forces it in, to measure that cost). pub const F_PROC: u32 = 262144; -/// Instance colours (`InstancedMesh.setColorAt`): object entries of 4 vec4s, the base colour times the 4th (`kernel_flags`; +/// Instance colours (`InstancedMesh.setColorAt`): object entries of 4 vec4s, the base colour times the 4th (`KernelVariant::flags`; /// `scene::object_stride`). pub const F_TINT: u32 = 524288; -/// three.js MeshBasicMaterial (`scene::MAT_UNLIT`): seen by camera rays, lit by nothing, lighting nothing (`kernel_flags`). +/// three.js MeshBasicMaterial (`scene::MAT_UNLIT`): seen by camera rays, lit by nothing, lighting nothing (`KernelVariant::flags`). pub const F_UNLIT: u32 = 1 << 20; -/// three.js `scene.background`: camera rays that miss see `Params::background` (`kernel_flags`). +/// three.js `scene.background`: camera rays that miss see `Params::background` (`KernelVariant::flags`). pub const F_BACKGROUND: u32 = 1 << 21; -/// The features that change the kernel's declarations (`kernel_flags`, `@if` in its source) rather than switch code in it. +/// The features that change the kernel's declarations (`KernelVariant::flags`, `@if` in its source) rather than switch code in it. /// three.js `vertexColors` (`scene::tri_colors`): a colour per triangle vertex multiplies the base colour (F_SOURCE). pub const F_VCOLOR: u32 = 1 << 22; /// Density grids (`scene::GridBuild`): the grid atlas (binding 19) and `vgrid` for volume functions (F_SOURCE). (Bit @@ -1383,8 +1498,7 @@ impl Gpu { let src = kernel_template(); // Bounds checks are not needed: all indices come from our own BVH / scene builder. let module = (!merged).then(|| unsafe { - let src = if uniform_tables { crate::merged::uniform_tables(&src) } else { src.clone() }; - let wgsl = kernel_wgsl(&guarded(&src, false, false), &kernel_flags(hwrt, &switches, 0)); + let wgsl = kernel_wgsl(&src, &KernelVariant { uniform_tables, guard: loop_guard_of(false), ..KernelVariant::new(hwrt, switches, 0) }); device.create_shader_module_trusted(wgpu::ShaderModuleDescriptor { label: Some("pathtrace"), source: wgpu::ShaderSource::Wgsl(wgsl.into()) }, kernel_checks()) }); let storage = |binding: u32, read_only: bool| wgpu::BindGroupLayoutEntry { @@ -1577,13 +1691,10 @@ impl Gpu { format!("{head}{code}{tail}") }; let src = splice(&splice(&splice(&self.source, "material-functions", &k.functions), "volume-functions", &k.volume_functions), "geometry-functions", &k.geometry_functions); - // The merged layout: the same source, its buffers merged for this kernel's arrays (merged.rs). - let src = k.caps.as_ref().map_or(src.clone(), |c| crate::merged::kernel(&src, c)); - // The uniform tables (never with the merged layout): the same source, four tables' declarations and - // indices rewritten. - let src = if self.uniform_tables { crate::merged::uniform_tables(&src) } else { src }; - // The loop guard, then the variant: this device's traversal and switches, this scene's features. - let wgsl = kernel_wgsl(&guarded(&src, self.merged, patches & F_UNITS != 0), &kernel_flags(self.hwrt, &self.switches, patches)); + // The variant: this device's traversal, switches, layout (the merged one with this kernel's arrays' + // lengths) and loop guard, this scene's features. + let v = KernelVariant { merged: k.caps, uniform_tables: self.uniform_tables, guard: loop_guard_of(patches & F_UNITS != 0), ..KernelVariant::new(self.hwrt, self.switches, patches) }; + let wgsl = kernel_wgsl(&src, &v); unsafe { self.device.create_shader_module_trusted(wgpu::ShaderModuleDescriptor { label: Some("pathtrace+functions"), source: wgpu::ShaderSource::Wgsl(wgsl.into()) }, kernel_checks()) } }) .clone() @@ -2082,7 +2193,7 @@ impl Gpu { bg_entries.retain(|e| !MERGED_TABLE_BINDINGS.contains(&e.binding)); bg_entries.push(e(crate::merged::TABLES_BINDING, t)); } - // The uniform tables: four tables' uniform copies in their place (merged.rs `uniform_tables`). + // The uniform tables: four tables' uniform copies in their place (merged.rs `UNIFORM_TABLES`). let uniform = self.uniform_tables.then(|| { // What the scene has of each, against what its array holds (the buffers have the edits' room after it). let used = [b_texinfo.size() / 16, scene.objects.len() as u64, records.len() as u64, scene.media.len() as u64]; @@ -3188,10 +3299,10 @@ mod tests { naga::valid::Validator::new(naga::valid::ValidationFlags::all(), naga::valid::Capabilities::all()).validate(&m).unwrap_or_else(|e| panic!("{}", e.emit_to_string(src))); } - /// The native default (no merged layout, no loop guard): `guarded` gives the kernel back untouched and the live - /// passes are their file. Every variant the kernel's flags make is valid WGSL (naga, wgpu's front end): each - /// traversal with no switch, each alone and all of them, by no scene feature, each alone and all of them (units are - /// software traversal's), and with the uniform tables. Prints the default variants' SHA-256 (`--nocapture`). + /// The native default has no loop guard, and the live passes are their file. Every variant the kernel's flags + /// make is valid WGSL (naga, wgpu's front end): each traversal with no switch, each alone and all of them, by no + /// scene feature, each alone and all of them (units are software traversal's), with the uniform tables, and with + /// the loop guard. (The merged layout's: merged.rs.) Prints the default variants' SHA-256 (`--nocapture`). #[test] fn every_kernel_variant_is_valid() { use sha2::{Digest, Sha256}; @@ -3202,7 +3313,6 @@ mod tests { let sha = |s: &str| Sha256::digest(s.as_bytes()).iter().map(|b| format!("{b:02x}")).collect::(); assert_eq!(crate::live::live_source(), include_str!("../shaders/live.wgsl")); let template = kernel_template(); - assert_eq!(guarded(&template, false, false), template); let all = KernelSwitches { strict: true, cutout_lite: 2, last_direct: true, light_pick: true }; let switches = [ KernelSwitches::default(), @@ -3214,12 +3324,14 @@ mod tests { all, ]; for hwrt in [false, true] { - println!("source kernel hwrt={hwrt} {}", sha(&kernel_wgsl(&template, &kernel_flags(hwrt, &KernelSwitches::default(), 0)))); - for k in &switches { + println!("source kernel hwrt={hwrt} {}", sha(&kernel_wgsl(&template, &KernelVariant::new(hwrt, KernelSwitches::default(), 0)))); + for &k in &switches { for f in [0, F_VCOLOR, F_GRID, F_TINT, F_UNLIT, F_UNITS, F_BACKGROUND, F_SOURCE] { - naga_valid(&kernel_wgsl(&template, &kernel_flags(hwrt, k, if hwrt { f & !F_UNITS } else { f }))); + naga_valid(&kernel_wgsl(&template, &KernelVariant::new(hwrt, k, if hwrt { f & !F_UNITS } else { f }))); } - naga_valid(&kernel_wgsl(&crate::merged::uniform_tables(&template), &kernel_flags(hwrt, k, F_SOURCE & !F_UNITS))); + let all_features = KernelVariant::new(hwrt, k, if hwrt { F_SOURCE & !F_UNITS } else { F_SOURCE }); + naga_valid(&kernel_wgsl(&template, &KernelVariant { uniform_tables: true, ..all_features })); + naga_valid(&kernel_wgsl(&template, &KernelVariant { guard: Some(1 << 16), ..all_features })); } } } @@ -3242,13 +3354,12 @@ mod tests { } // As compiled: the source with units has every loop bounded and the count's buffer; hardware traversal, which // never has units, has room for the count only when a limit is set. - let src = kernel_template(); - let g = guarded(&src, false, true); - assert!(g.contains("return LOOP_GUARD_N > 65536u;") && g.contains(&crate::merged::guard_decl()), "the guard and its count"); - assert_eq!(guarded(&src, false, false), src, "units told to have no guard: the source as written"); + let v = KernelVariant { guard: loop_guard_of(true), ..KernelVariant::new(false, KernelSwitches::default(), F_UNITS) }; + assert_eq!((v.guard, loop_guard_of(false)), (Some(65536), None), "units told to have no guard: the source as written"); // What a native session with units compiles with nothing set, to naga (wgpu's front end); the unit walk's loops // are bounded too (that every loop is: merged.rs `the_guard_bounds_every_loop`). - let wgsl = kernel_wgsl(&g, &kernel_flags(false, &KernelSwitches::default(), F_UNITS)); + let wgsl = kernel_wgsl(&kernel_template(), &v); + assert!(wgsl.contains("const LOOP_GUARD_LIMIT = 65536u;") && wgsl.contains("var loop_guard_trips: array, GUARD_WORDS>;"), "the guard and its count"); naga_valid(&wgsl); let walk = &wgsl[wgsl.find("fn trace_units(").expect("the unit walk")..]; assert!(walk[..walk[1..].find("\nfn ").unwrap_or(walk.len() - 1)].contains("loop_guard()")); @@ -3261,9 +3372,9 @@ mod tests { #[test] fn units_are_software_traversals_only() { let t = kernel_template(); - let sw = kernel_wgsl(&t, &kernel_flags(false, &KernelSwitches::default(), 0)); - let u = kernel_wgsl(&t, &kernel_flags(false, &KernelSwitches::default(), F_UNITS)); - let hw = kernel_wgsl(&t, &kernel_flags(true, &KernelSwitches::default(), 0)); + let sw = kernel_wgsl(&t, &KernelVariant::new(false, KernelSwitches::default(), 0)); + let u = kernel_wgsl(&t, &KernelVariant::new(false, KernelSwitches::default(), F_UNITS)); + let hw = kernel_wgsl(&t, &KernelVariant::new(true, KernelSwitches::default(), 0)); assert!(!sw.contains("trace_units") && !hw.contains("trace_units") && !hw.contains("fn trace_procs(")); assert_eq!((u.matches("fn trace_units(").count(), u.matches("trace_units(o, d, inv_d, hit, mode)").count()), (1, 1)); let stacks = |s: &str| (s.matches("var stack: array").count(), s.matches("var stack: array").count()); @@ -3282,7 +3393,7 @@ mod tests { assert_eq!(desktop.kernel, KernelSwitches { strict: false, cutout_lite: 0, last_direct: false, light_pick: false }); assert!(!desktop.uniform_tables_on(false, false)); for hwrt in [false, true] { - assert!(kernel_flags(hwrt, &desktop.kernel, 0).iter().all(|&(name, on)| on == (name == "hwrt" && hwrt))); + assert!(KernelVariant::new(hwrt, desktop.kernel, 0).flags().iter().all(|&(name, on)| on == (name == "hwrt" && hwrt))); } let android = GpuOptions::from_vars(&vars(&[]), true); assert_eq!(android.kernel, KernelSwitches { strict: true, cutout_lite: 1, last_direct: true, light_pick: false }); @@ -3307,7 +3418,7 @@ mod tests { #[test] fn the_other_shaders_are_wgsl_as_the_specification_writes_it() { for src in [include_str!("../shaders/live.wgsl"), include_str!("../shaders/particles.wgsl"), include_str!("../shaders/gpu_particles.wgsl"), include_str!("../shaders/grid_sim.wgsl")] { - kernel_wgsl(src, &[]); + wesl_wgsl(src, &[], &[]); } } @@ -3317,7 +3428,7 @@ mod tests { #[test] fn the_kernel_switches_change_what_they_say() { let t = kernel_template(); - let v = |hwrt: bool, k: KernelSwitches| kernel_wgsl(&t, &kernel_flags(hwrt, &k, 0)); + let v = |hwrt: bool, k: KernelSwitches| kernel_wgsl(&t, &KernelVariant::new(hwrt, k, 0)); for hwrt in [false, true] { let plain = v(hwrt, KernelSwitches::default()); let strict = v(hwrt, KernelSwitches { strict: true, ..Default::default() }); @@ -3344,7 +3455,7 @@ mod tests { assert_eq!(bindings, [14, 15, 17, 18]); for (k, &t) in UNIFORM_TABLES.iter().enumerate() { let m = &TABLES[t]; - assert!(m.decl.starts_with(&format!("@group(0) @binding({}) var {}:", bindings[k], m.name)), "{}", m.decl); + assert!(kernel_template().contains(&format!("@group(0) @binding({}) var {}: array<", bindings[k], m.name)), "{}", m.name); let bytes = m.stride * UNIFORM_LEN[k] as u64; assert!(bytes <= 65536 && bytes % 16 == 0, "{}: {bytes} bytes", m.name); } @@ -3406,12 +3517,12 @@ mod tests { assert_eq!(words[..2], [11, 22]); } - /// Each feature that changes the kernel's declarations (`kernel_flags`) changes what it says, and nothing of it + /// Each feature that changes the kernel's declarations (`KernelVariant::flags`) changes what it says, and nothing of it /// is in a kernel without it. #[test] fn the_source_features_change_what_they_say() { let t = kernel_template(); - let v = |f: u32| kernel_wgsl(&t, &kernel_flags(false, &KernelSwitches::default(), f)); + let v = |f: u32| kernel_wgsl(&t, &KernelVariant::new(false, KernelSwitches::default(), f)); let base = v(0); let tint = v(F_TINT); assert!(tint.contains("const OBJECT_STRIDE: u32 = 4u;") && !tint.contains("const OBJECT_STRIDE: u32 = 3u;")); diff --git a/renderer/src/live/mod.rs b/renderer/src/live/mod.rs index 2eeda4f..1b1366f 100644 --- a/renderer/src/live/mod.rs +++ b/renderer/src/live/mod.rs @@ -363,7 +363,7 @@ pub struct LiveCore { pub cpu_wait_ms: f32, /// The cloud cover the sky and the sun light were made with (`move_sun`). sky_cover: f32, - /// The loop guard (merged.rs `guard`), if a kernel of this session may have it: the invocations it has stopped at + /// The loop guard (gpu.rs `guard_loops`), if a kernel of this session may have it: the invocations it has stopped at /// its limit since the view was made, as last read back (a frame or two behind, without waiting). Above zero: /// frames with wrong pixels, a traversal that did not end on the scene's data. None without the guard. pub guard_trips: Option, diff --git a/renderer/src/merged.rs b/renderer/src/merged.rs index 72dd485..e3efbb8 100644 --- a/renderer/src/merged.rs +++ b/renderer/src/merged.rs @@ -1,29 +1,28 @@ //! 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_wgsl` makes the kernel's variants), 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"). +//! live passes need none of this: each binds at most eight, live/layout.rs, and a page binds them as they are.) //! -//! What is merged, and how: +//! The kernel's source has both layouts: pathtrace.wgsl declares its tables and its group-1 buffers in an +//! `@if(merged) { .. } @else { .. }` block each (WESL's conditional translation, `gpu::KernelVariant`), and reads and +//! writes them only through functions declared in each (`material_at`, `cache_acc_add`, `refl_store`, ...). This +//! module is the Rust side's half: which buffers are merged, the lengths of the merged structs' arrays (WESL constants +//! the source imports, `gpu::KernelVariant::constants`) and where the arrays are, for the buffers the renderer makes +//! and copies. `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"). +//! +//! What is merged: //! - The kernel's seven tables (`TABLES`: materials, lights, emitters, texture pages, objects, instances, media) -//! become the arrays of one struct, `MT: MergedTables`, in one buffer. Their lengths are written into the source -//! (a struct's arrays are sized), each a power of two no less than `TABLE_MIN`, so a scene that grows compiles a -//! kernel again only when a table doubles. The renderer keeps writing its own buffers; each frame copies them into -//! 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`). -//! 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 +//! become the arrays of one struct, `MT: MergedTables`, in one buffer. Their lengths are each a power of two no +//! less than `TABLE_MIN`, so a scene that grows compiles a kernel again only when a table doubles. The renderer +//! keeps writing its own buffers; each frame copies them into 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) and +//! the loop guard's count become the arrays of `ML: MergedLive`, one buffer that they are ranges of +//! (`gpu::KernelBufs`). 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 -//! already has: each stops the renderer with the name and the line. Locals and parameters that shadow a merged name -//! are followed through their scopes and left alone. The test below reads what comes out with naga: it is valid WGSL, -//! no stage has more than `STORAGE_BUFFERS`, and the structs' arrays are where the Rust side puts them. +//! The tests below read every variant with naga: valid WGSL, no stage with more than `STORAGE_BUFFERS`, and the +//! structs' arrays where `offsets` and `live_layout` put them. /// WebGPU's default `maxStorageBuffersPerShaderStage`: what the merged layout fits. pub const STORAGE_BUFFERS: u32 = 8; @@ -31,24 +30,21 @@ 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 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, + /// Its element's size in a storage buffer; where its array may start in the struct: its element's alignment, or + /// more (`@align` in pathtrace.wgsl's struct). pub stride: u64, pub align: u64, - /// Its declaration in pathtrace.wgsl, as written. - pub decl: &'static str, } /// The kernel's tables, in `MergedTables`' order (group 0, binding `TABLES_BINDING`). pub const TABLES: [Member; 7] = [ - Member { name: "materials", ty: "Material", stride: 288, align: 16, decl: "@group(0) @binding(4) var materials: array;" }, - Member { name: "lights", ty: "Light", stride: 80, align: 16, decl: "@group(0) @binding(5) var lights: array;" }, - Member { name: "emitters", ty: "vec4", stride: 16, align: 16, decl: "@group(0) @binding(6) var emitters: array>;" }, - Member { name: "tex_info", ty: "vec4", stride: 16, align: 16, decl: "@group(0) @binding(14) var tex_info: array>;" }, - Member { name: "objects", ty: "vec4", stride: 16, align: 16, decl: "@group(0) @binding(15) var objects: array>;" }, - Member { name: "instances", ty: "Instance", stride: 160, align: 16, decl: "@group(0) @binding(17) var instances: array;" }, - Member { name: "media", ty: "Medium", stride: 224, align: 16, decl: "@group(0) @binding(18) var media: array;" }, + Member { name: "materials", stride: 288, align: 16 }, + Member { name: "lights", stride: 80, align: 16 }, + Member { name: "emitters", stride: 16, align: 16 }, + Member { name: "tex_info", stride: 16, align: 16 }, + Member { name: "objects", stride: 16, align: 16 }, + Member { name: "instances", stride: 160, align: 16 }, + Member { name: "media", stride: 224, align: 16 }, ]; pub const TABLES_BINDING: u32 = 4; /// The least length of each table's array: scenes below it share one kernel source. @@ -56,17 +52,17 @@ pub const TABLE_MIN: [u32; 7] = [64, 64, 1024, 256, 1024, 64, 16]; /// The kernel's group-1 buffers, in `MergedLive`'s order (group 1, binding `KERNEL_LIVE_BINDING`). 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: 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>;" }, + Member { name: "cache_keys", stride: 4, align: 4 }, + Member { name: "cache_acc", stride: 4, align: 4 }, + Member { name: "cache_val", stride: 16, align: 16 }, + Member { name: "refl", stride: 16, align: 256 }, + Member { name: "motion_out", stride: 16, align: 256 }, + Member { name: "direct_out", stride: 16, align: 256 }, ]; 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 +/// `MergedLive`'s last array, after `KERNEL_LIVE`'s: the loop guard's count (`gpu::guard_loops`; its first word: the invocations /// a loop of which was stopped). Always there in the merged layout, written only by a kernel with the guard. -pub const GUARD_TRIPS: Member = Member { name: "loop_guard_trips", ty: "atomic", stride: 4, align: 4, decl: "" }; +pub const GUARD_TRIPS: Member = Member { name: "loop_guard_trips", stride: 4, align: 4 }; /// The same count in the layout as written, when the kernel has the guard (group 1, binding `GUARD_BINDING`). pub const GUARD_BINDING: u32 = 8; pub const GUARD_WORDS: u32 = 4; @@ -115,237 +111,17 @@ pub fn offsets(members: &[Member], caps: &[u32]) -> (Vec, u64) { (out, at.next_multiple_of(members.iter().map(|m| m.align).max().unwrap_or(1))) } -/// How a merged name's uses are rewritten. -struct Rule { - name: &'static str, - /// What `name[` becomes, and what goes before the `]` that closes it. - open: String, - close: String, - /// What `arrayLength(&name)` becomes. - len: String, -} - -fn fail(what: &str, src: &str, at: usize) -> ! { - panic!("merged layout: {what} (line {})", src[..at].matches('\n').count() + 1) -} - -/// The identifier at `i` (bytes), if one starts there. -fn ident(b: &[u8], i: usize) -> Option { - let start = |c: u8| c.is_ascii_alphabetic() || c == b'_'; - if i >= b.len() || !start(b[i]) { - return None; - } - let mut j = i + 1; - while j < b.len() && (start(b[j]) || b[j].is_ascii_digit()) { - j += 1; - } - Some(j) -} - -fn skip_space(b: &[u8], mut i: usize) -> usize { - while i < b.len() && b[i].is_ascii_whitespace() { - i += 1; - } - i -} - -/// Rewrites every use of the rules' names in `src` (WGSL): a scan that knows comments, blocks and the locals and -/// parameters that shadow a name (from the end of their declaration to the end of their block). -fn rewrite(src: &str, rules: &[Rule]) -> String { - let b = src.as_bytes(); - let rule = |name: &[u8]| rules.iter().position(|r| r.name.as_bytes() == name); - let mut out: Vec = Vec::with_capacity(b.len() + b.len() / 16); - // Blocks and brackets open; the brackets a rewritten index closes with more than `]`. - let (mut depth, mut brackets, mut parens) = (0usize, 0usize, 0usize); - let mut closes: Vec<(usize, &str)> = vec![]; - // Names shadowed, each with the block depth it lives in; those declared in the statement being read. - let mut shadows: Vec<(usize, usize)> = vec![]; - let mut pending: Vec<(usize, usize)> = vec![]; - // The last byte of code before here (not space, not comment); after `var` / `let` / `const` in a function; in a - // function's header (its parameters). - let (mut last, mut declares, mut header) = (0u8, false, false); - let mut i = 0; - while i < b.len() { - if b[i..].starts_with(b"//") { - let j = b[i..].iter().position(|&c| c == b'\n').map_or(b.len(), |k| i + k); - out.extend_from_slice(&b[i..j]); - i = j; - continue; - } - if b[i..].starts_with(b"/*") { - let j = src[i + 2..].find("*/").map_or(b.len(), |k| i + 2 + k + 2); - out.extend_from_slice(&b[i..j]); - i = j; - continue; - } - let Some(j) = ident(b, i) else { - let c = b[i]; - match c { - b'{' => { - depth += 1; - header = false; - } - b'}' => { - depth = depth.checked_sub(1).unwrap_or_else(|| fail("an unmatched }", src, i)); - shadows.retain(|s| s.1 <= depth); - } - b';' => shadows.append(&mut pending), - b'(' => parens += 1, - b')' => parens = parens.saturating_sub(1), - b'[' => brackets += 1, - b']' => { - if closes.last().is_some_and(|c| c.0 == brackets) { - out.extend_from_slice(closes.pop().unwrap().1.as_bytes()); - } - brackets = brackets.checked_sub(1).unwrap_or_else(|| fail("an unmatched ]", src, i)); - } - _ => {} - } - if !c.is_ascii_whitespace() { - last = c; - } - out.push(c); - i += 1; - continue; - }; - let word = &b[i..j]; - let after = skip_space(b, j); - let next = b.get(after).copied().unwrap_or(0); - let (was, declared) = (last, std::mem::take(&mut declares)); - last = b'a'; - if depth > 0 && matches!(word, b"var" | b"let" | b"const") { - declares = true; - } else if depth == 0 && word == b"fn" { - header = true; - } else if word == b"arrayLength" && next == b'(' { - // `arrayLength(&name)` of a merged name that nothing shadows: its length. - let amp = skip_space(b, after + 1); - if b.get(amp) == Some(&b'&') { - let n0 = skip_space(b, amp + 1); - if let Some(n1) = ident(b, n0) { - let end = skip_space(b, n1); - if let Some(r) = rule(&b[n0..n1]).filter(|r| !shadows.iter().any(|s| s.0 == *r)) { - if b.get(end) != Some(&b')') { - fail(&format!("arrayLength of something inside `{}`", rules[r].name), src, i); - } - out.extend_from_slice(rules[r].len.as_bytes()); - last = b')'; - i = end + 1; - continue; - } - } - } - } else if let Some(r) = rule(word) { - let name = rules[r].name; - if declared { - // A local of that name: it shadows the buffer from the end of this statement to the end of its block - // (a `for`'s own variable: the loop's block). - pending.push((r, depth + (parens > 0) as usize)); - } else if was == b'.' { - // A member of something else. - } else if next == b':' && b.get(after + 1) != Some(&b':') { - // A parameter of that name shadows the buffer in the function's body; a struct's member does nothing. - if header { - shadows.push((r, 1)); - } - } else if shadows.iter().any(|s| s.0 == r) { - // The local. - } else if next == b'[' { - out.extend_from_slice(rules[r].open.as_bytes()); - brackets += 1; - if !rules[r].close.is_empty() { - closes.push((brackets, &rules[r].close)); - } - last = b'['; - i = after + 1; - continue; - } else { - fail(&format!("`{name}` is used in a way the generator does not know"), src, i); - } - } - out.extend_from_slice(word); - i = j; - } - if depth != 0 || brackets != 0 || !closes.is_empty() { - fail("unbalanced braces or brackets", src, b.len()); - } - String::from_utf8(out).expect("the source's own bytes") -} - -/// Takes the declaration `decl` out of `s` (it is there exactly once), leaving `with` in its place. -fn take(s: &mut String, decl: &str, with: &str) { - assert_eq!(s.matches(decl).count(), 1, "merged layout: the declaration {decl:?}"); - *s = s.replacen(decl, with, 1); -} - -/// Stops if the source already has a name the generator adds. -fn fresh(src: &str, names: &[&str]) { - let b = src.as_bytes(); - let mut i = 0; - while i < b.len() { - if b[i..].starts_with(b"//") { - i = b[i..].iter().position(|&c| c == b'\n').map_or(b.len(), |k| i + k); - } else if let Some(j) = ident(b, i) { - if names.iter().any(|n| n.as_bytes() == &b[i..j]) { - fail(&format!("the source already has a `{}`", &src[i..j]), src, i); - } - i = j; - } else { - i += 1; - } - } -} - -fn merged_struct(name: &str, members: &[&Member], caps: &[u32]) -> String { - 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")) -} - -/// The kernel (its source, `gpu::kernel_template`, its functions spliced in; before its variant is made) in the merged -/// layout for arrays of `caps`' lengths: group 0 binds tri_verts, attrs, nodes, `MT`, env_cdf, accum and prev_verts, -/// group 1 `ML`. -pub fn kernel(src: &str, caps: &Caps) -> String { - fresh(src, &["MT", "ML", "MergedTables", "MergedLive", GUARD_TRIPS.name]); - let mut s = src.to_string(); - for (k, m) in TABLES.iter().enumerate() { - let with = if k == 0 { - format!("{}@group(0) @binding({TABLES_BINDING}) var MT: MergedTables;", merged_struct("MergedTables", &TABLES.iter().collect::>(), &caps.tables)) - } else { - String::new() - }; - take(&mut s, m.decl, &with); - } - for (k, m) in KERNEL_LIVE.iter().enumerate() { - let with = if k == 0 { - format!("{}@group(1) @binding({KERNEL_LIVE_BINDING}) var ML: MergedLive;", merged_struct("MergedLive", &KERNEL_LIVE.iter().chain([&GUARD_TRIPS]).collect::>(), &caps.live.iter().copied().chain([GUARD_WORDS]).collect::>())) - } else { - String::new() - }; - take(&mut s, m.decl, &with); - } - let rule = |var: &str, m: &Member, cap: u32| Rule { name: m.name, open: format!("{var}.{}[", m.name), close: String::new(), len: format!("{cap}u") }; - let rules: Vec = TABLES.iter().zip(caps.tables).map(|(m, c)| rule("MT", m, c)).chain(KERNEL_LIVE.iter().zip(caps.live).map(|(m, c)| rule("ML", m, c))).collect(); - rewrite(&s, &rules) -} - /// The uniform tables (`gpu::GpuOptions::uniform_tables`; docs/RUNTIME.md "Android"): which of `TABLES` the kernel /// reads from uniform buffers instead of storage buffers, in the layout as written. Qualcomm's Adreno 740 driver /// (Vulkan, Android 13) gave the kernel other buffers' contents for these four: a model's textures came out as /// another table's numbers, instance records sent hits to the wrong triangles, an emissive object's power read as /// something else's and lit nothing. As uniform buffers they read right. They are the four of the kernel's read-only -/// tables that are bounded (the others: `uniform_tables`). -pub const UNIFORM_TABLES: [usize; 4] = [3, 4, 5, 6]; -/// Their arrays' lengths (a uniform buffer's array is sized): what 64 KiB hold, the size every device gives a uniform -/// binding (WebGPU's `maxUniformBufferBindingSize`, Vulkan's least `maxUniformBufferRange` on the phones this is -/// for), and for the media what a scene may have (`scene::MAX_MEDIA`). In a scene's terms: 2,048 textures, 1,365 -/// objects (1,024 with instance colours), 409 instance records, every medium. -pub const UNIFORM_LEN: [u32; 4] = [4096, 4096, 409, 32]; - -/// The kernel (as `kernel` takes it: a scene's variant, its functions spliced in) with the uniform tables: the four -/// declarations as sized uniform arrays at their own bindings, and every index into them held to the array -/// (`min(u32(i), last)`). A scene with more entries than an array holds reads the last one for the rest (the renderer -/// says so when it uploads one, `Gpu::upload`); nothing is read outside the buffer, which the kernel's unchecked -/// indexing would otherwise do. Everything else is the source as written. +/// tables that are bounded (the others: below). +/// +/// In the kernel (pathtrace.wgsl, `@if(uniform_tables)`) the four are sized uniform arrays at their own bindings, and +/// every index into them is held to the array (`tex_info_at` and the others: `min(i, last)`). A scene with more +/// entries than an array holds reads the last one for the rest (the renderer says so when it uploads one, +/// `Gpu::upload`); nothing is read outside the buffer, which the kernel's unchecked indexing would otherwise do. /// /// The kernel's other storage buffers stay storage buffers, each for its reason: /// - `materials` (288 B each), `lights` (80 B): the device drew right with them as they are, and 64 KiB would hold @@ -354,119 +130,22 @@ pub const UNIFORM_LEN: [u32; 4] = [4096, 4096, 409, 32]; /// geometry's data), `tri_verts`, `attrs`, `nodes`, `prev_verts` (and the vertex colours after them): megabytes, /// which no uniform binding holds; /// - `accum` and group 1 (the radiance cache, the trace's outputs): written by the kernel. -pub fn uniform_tables(src: &str) -> String { - let mut s = src.to_string(); - let mut rules = vec![]; - for (&k, &len) in UNIFORM_TABLES.iter().zip(&UNIFORM_LEN) { - let m = &TABLES[k]; - let sized = m.decl.replace("var", "var").replace(&format!("array<{}>", m.ty), &format!("array<{}, {len}>", m.ty)); - assert!(sized.contains("var") && sized.ends_with(&format!(", {len}>;")), "uniform tables: the declaration {:?}", m.decl); - take(&mut s, m.decl, &sized); - rules.push(Rule { name: m.name, open: format!("{}[min(u32(", m.name), close: format!("), {}u)", len - 1), len: format!("{len}u") }); - } - 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 -/// dispatch never ends, and a GPU that never gets a dispatch back takes the whole machine's display with it (it did: -/// docs/RUNTIME.md "The engine in a browser"). With the bound a dispatch always ends: wrong pixels instead of no -/// frame. No formula changes, but the image is not bit-identical to the kernel's without the guard: the compiler -/// arranges the arithmetic of the functions it is in differently (hello world, one sample a pixel: 354 of 43,200 -/// numbers differ, by at most 2 x 10^-7 of their value). Between the two layouts, guard or not, the bytes are equal. -/// `THREERT_LOOP_GUARD=` natively (`gpu::loop_guard_of`: a kernel with units has it with nothing set); a -/// browser's kernel always has it. -/// -/// A stopped loop is counted, not hidden: the invocation that reaches the limit adds one to the first word of `trips` -/// (an `array, 4>` the source can name: `ML.loop_guard_trips` in the merged layout, the buffer `guard_decl` -/// declares in the layout as written), which the renderer reads back (`live::LiveCore::guard_trips`, an offline -/// render's error). A frame with a count above zero is a frame with wrong pixels, and says so. -pub fn guard(src: &str, limit: u32, trips: &str) -> String { - fresh(src, &["loop_guard", "LOOP_GUARD_N"]); - let b = src.as_bytes(); - let mut out: Vec = Vec::with_capacity(b.len() + 4096); - let mut i = 0; - while i < b.len() { - if b[i..].starts_with(b"//") { - let j = b[i..].iter().position(|&c| c == b'\n').map_or(b.len(), |k| i + k); - out.extend_from_slice(&b[i..j]); - i = j; - continue; - } - if b[i..].starts_with(b"/*") { - let j = src[i + 2..].find("*/").map_or(b.len(), |k| i + 2 + k + 2); - out.extend_from_slice(&b[i..j]); - i = j; - continue; - } - let Some(j) = ident(b, i) else { - out.push(b[i]); - i += 1; - continue; - }; - let word = &b[i..j]; - let mut body = None; - let after = skip_space(b, j); - if word == b"loop" && b.get(after) == Some(&b'{') { - body = Some(after); - } else if matches!(word, b"for" | b"while") && b.get(after) == Some(&b'(') { - // The header's closing parenthesis, then the body. - let (mut depth, mut k) = (0usize, after); - while k < b.len() { - match b[k] { - b'(' => depth += 1, - b')' => { - depth -= 1; - if depth == 0 { - break; - } - } - _ => {} - } - k += 1; - } - let open = skip_space(b, k + 1); - if b.get(open) != Some(&b'{') { - fail("a loop without a body the guard can open", src, i); - } - body = Some(open); - } - match body { - Some(open) => { - out.extend_from_slice(&b[i..=open]); - out.extend_from_slice(b" if (loop_guard()) { break; }"); - i = open + 1; - } - None => { - out.extend_from_slice(word); - i = j; - } - } - } - let mut s = String::from_utf8(out).expect("the source's own bytes"); - s += &format!("\n// The loop guard (renderer/src/merged.rs `guard`): an invocation's loop iterations, all its loops together; the\n// invocations stopped at the limit are counted.\nvar LOOP_GUARD_N: u32 = 0u;\nfn loop_guard() -> bool {{\n LOOP_GUARD_N += 1u;\n if (LOOP_GUARD_N == {limit}u + 1u) {{ atomicAdd(&{trips}[0], 1u); }}\n return LOOP_GUARD_N > {limit}u;\n}}\n"); - s -} - -/// The loop guard's count as a buffer of its own, for a kernel in the layout as written (the merged layout has it in -/// `MergedLive`): the declaration `guard`'s `trips` names. -pub fn guard_decl() -> String { - format!("@group(1) @binding({GUARD_BINDING}) var {}: array, {GUARD_WORDS}>;\n", GUARD_TRIPS.name) -} +pub const UNIFORM_TABLES: [usize; 4] = [3, 4, 5, 6]; +/// Their arrays' lengths (a uniform buffer's array is sized): what 64 KiB hold, the size every device gives a uniform +/// binding (WebGPU's `maxUniformBufferBindingSize`, Vulkan's least `maxUniformBufferRange` on the phones this is +/// for), and for the media what a scene may have (`scene::MAX_MEDIA`). In a scene's terms: 2,048 textures, 1,365 +/// objects (1,024 with instance colours), 409 instance records, every medium. +pub const UNIFORM_LEN: [u32; 4] = [4096, 4096, 409, 32]; #[cfg(test)] mod tests { use super::*; - const KERNEL: &str = include_str!("../shaders/pathtrace.wgsl"); + use crate::gpu::{kernel_template, kernel_wgsl, KernelVariant}; - /// A scene's kernel (software traversal, no switch) as compiled: the merged layout put into its text first when - /// given, as `gpu::Gpu::kernel_module` does, then its variant. + /// A scene's kernel (software traversal, no switch) as compiled, in the merged layout when given its lengths. fn variant(features: u32, merged: Option<&Caps>) -> String { - let t = crate::gpu::kernel_template(); - let t = merged.map_or(t.clone(), |c| kernel(&t, c)); - crate::gpu::kernel_wgsl(&t, &crate::gpu::kernel_flags(false, &Default::default(), features)) + kernel_wgsl(&kernel_template(), &KernelVariant { merged: merged.copied(), ..KernelVariant::new(false, Default::default(), features) }) } fn caps() -> Caps { @@ -497,9 +176,11 @@ mod tests { } } - /// Every kernel variant the source patches make, in the merged layout: valid WGSL, no entry point with more than - /// eight storage buffers, none declared beyond them (a pipeline layout counts what is declared), and the merged - /// structs' arrays where `offsets` puts them. The source as written has 19, so the count is the generator's. + /// Kernel variants in the merged layout: valid WGSL, no entry point with more than eight storage buffers, none + /// declared beyond them (a pipeline layout counts what is declared), and the merged structs' arrays where `offsets` + /// and `live_layout` put them. The layout as written has 19. Because the merged block declares none of the + /// buffers by their own names, a read or write of one past its accessor does not compile here (and naga says + /// where), whatever the other layouts take. #[test] fn the_merged_kernel_fits_eight_storage_buffers() { let (m, info) = parse(&variant(0, None)); @@ -509,6 +190,7 @@ mod tests { let caps = caps(); for features in [0, crate::gpu::F_VCOLOR, crate::gpu::F_GRID, crate::gpu::F_TINT | crate::gpu::F_UNLIT | crate::gpu::F_BACKGROUND, crate::gpu::F_SOURCE] { let src = variant(features, Some(&caps)); + assert!(src.contains("const MT_MEDIA = 16u;") && src.contains(&format!("const ML_REFL = {}u;", caps.live[3])), "the lengths, as constants"); let (m, info) = parse(&src); let (used, declared) = storage_buffers(&m, &info); assert!(declared as u32 <= STORAGE_BUFFERS, "features {features:#x}: {declared} declared"); @@ -521,15 +203,10 @@ mod tests { /// Every kernel variant with the uniform tables: valid WGSL; the four tables are uniform buffers at their own /// bindings, sized as the Rust side makes their buffers; the other fifteen storage buffers are as written; and - /// every index into the four is held to its array, so nothing is read outside a uniform buffer. + /// each of the four is indexed once, in its accessor, held to its array, so nothing is read outside a uniform + /// buffer. #[test] fn the_uniform_tables_are_four_sized_uniform_arrays() { - // A name's index sites in code (not comments): `name[` that is not part of a longer name or a member. - let sites = |src: &str, name: &str| { - let code: String = src.lines().map(|l| l.split("//").next().unwrap_or("")).collect::>().join("\n"); - let open = format!("{name}["); - code.match_indices(&open).filter(|(k, _)| !code[..*k].ends_with(|c: char| c.is_alphanumeric() || c == '_' || c == '.')).count() - }; let globals = |m: &naga::Module, uniform: bool| -> Vec<(String, u32, u32)> { let mut out: Vec<(String, u32, u32)> = m .global_variables @@ -549,8 +226,7 @@ mod tests { }; let (written, _) = parse(&variant(0, None)); for features in [0, crate::gpu::F_VCOLOR, crate::gpu::F_GRID, crate::gpu::F_TINT | crate::gpu::F_UNLIT | crate::gpu::F_BACKGROUND, crate::gpu::F_SOURCE] { - let base = variant(features, None); - let src = crate::gpu::kernel_wgsl(&uniform_tables(&crate::gpu::kernel_template()), &crate::gpu::kernel_flags(false, &Default::default(), features)); + let src = kernel_wgsl(&kernel_template(), &KernelVariant { uniform_tables: true, ..KernelVariant::new(false, Default::default(), features) }); let (m, info) = parse(&src); assert_eq!(storage_buffers(&m, &info).1, 15, "features {features:#x}"); let tables: Vec<(String, u32, u32)> = globals(&m, true).into_iter().filter(|g| g.0 != "P" && g.0 != "energy_lut").collect(); @@ -559,20 +235,13 @@ mod tests { // The storage buffers that stay: the same names at the same bindings as the kernel as written. let kept = |m: &naga::Module| globals(m, false).into_iter().map(|g| (g.0, g.1)).filter(|g| !want.iter().any(|w| w.0 == g.0)).collect::>(); assert_eq!(kept(&m), kept(&written), "features {features:#x}"); - for (&k, len) in UNIFORM_TABLES.iter().zip(UNIFORM_LEN) { + for &k in &UNIFORM_TABLES { let name = TABLES[k].name; - let held = format!("{name}[min(u32("); - assert!(sites(&base, name) >= 3, "{name}: the kernel indexes it"); - assert_eq!(src.matches(&held).count(), sites(&base, name), "{name}: every index is held (features {features:#x})"); - assert_eq!(sites(&src, name), src.matches(&held).count(), "{name}: no index is left as written"); - assert!(src.contains(&format!("), {}u)]", len - 1))); + let upper = name.to_uppercase(); + assert_eq!(src.matches(&format!("{name}[")).count(), 1, "{name}: indexed in its accessor only (features {features:#x})"); + assert!(src.contains(&format!("{name}[min(i, UT_{upper} - 1u)]")), "{name}: held to its array"); } } - // One index, as it is rewritten: an expression with brackets of its own, and a table inside another's index. - let src = "@group(0) @binding(14) var tex_info: array>;\n@group(0) @binding(15) var objects: array>;\n@group(0) @binding(17) var instances: array;\n@group(0) @binding(18) var media: array;\nfn f(a: array, k: i32) -> f32 { return objects[a[1] * 3u + 1u].w + tex_info[u32(objects[k].x)].y + media[0].b.x; }\n"; - let out = uniform_tables(src); - assert!(out.starts_with("@group(0) @binding(14) var tex_info: array, 4096>;\n@group(0) @binding(15) var objects: array, 4096>;\n@group(0) @binding(17) var instances: array;\n@group(0) @binding(18) var media: array;\n"), "{out}"); - 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 Rust side's records are the size the merged structs give their elements. @@ -586,52 +255,48 @@ mod tests { assert_eq!(table_cap(2, 512), 1024); } - /// The scan follows scopes: a local or a parameter of a merged name is left alone, from the end of its - /// declaration to the end of its block; comments are not read. - #[test] - fn locals_that_shadow_a_buffer_are_left_alone() { - let rules = [Rule { name: "refl", open: "ML.refl[".into(), close: String::new(), len: "7u".into() }, Rule { name: "fa", open: "MV4[LO.fa.x + (".into(), close: ")".into(), len: "LO.fa.y".into() }]; - let src = "struct S { refl: f32 };\nfn f(i: u32) -> f32 {\n let a = refl[i].x; // refl[0] }\n {\n let refl = refl[i + 1u].y;\n if (refl > 0.0) { return refl; }\n }\n var s: S;\n return s.refl + refl[arrayLength(&refl) - 1u].x + fa[fa[i].x].y;\n}\nfn g(refl: array) -> f32 { return refl[1]; }\nfn h() -> f32 { return refl[2].x; }\n"; - let want = "struct S { refl: f32 };\nfn f(i: u32) -> f32 {\n let a = ML.refl[i].x; // refl[0] }\n {\n let refl = ML.refl[i + 1u].y;\n if (refl > 0.0) { return refl; }\n }\n var s: S;\n return s.refl + ML.refl[7u - 1u].x + MV4[LO.fa.x + (MV4[LO.fa.x + (i)].x)].y;\n}\nfn g(refl: array) -> f32 { return refl[1]; }\nfn h() -> f32 { return ML.refl[2].x; }\n"; - assert_eq!(rewrite(src, &rules), want); - } - - /// The loop guard opens every loop of the kernel and of the live passes, in the layout as written and in the - /// merged one, and what comes out is valid WGSL. + /// The loop guard opens every loop of the kernel, the scene's functions' too, in the layout as written and in the + /// merged one (its count a buffer of its own, or in `MergedLive`), and what comes out is valid WGSL. #[test] fn the_guard_bounds_every_loop() { - let loops = |s: &str| { - let code: String = s.lines().map(|l| l.split("//").next().unwrap_or("")).collect::>().join("\n"); - let words = |w: &str| code.match_indices(w).filter(|(k, _)| !code[..*k].ends_with(|c: char| c.is_alphanumeric() || c == '_')).count(); - words("loop {") + words("for (") + words("while (") - }; - // The kernel as written (its count a buffer of its own) and merged (its count in `MergedLive`). - let template = crate::gpu::kernel_template(); - let merged = kernel(&template, &caps()); - for (src, trips) in [(format!("{template}{}", guard_decl()), "loop_guard_trips"), (merged, "ML.loop_guard_trips")] { - let g = guard(&src, 1 << 16, trips); - assert!(loops(&src) > 40, "the kernel's loops"); - assert_eq!(g.matches("if (loop_guard()) { break; }").count(), loops(&src)); - // (The variant every loop of which is there: software traversal with the scene's features.) - let (m, info) = parse(&crate::gpu::kernel_wgsl(&g, &crate::gpu::kernel_flags(false, &Default::default(), crate::gpu::F_SOURCE))); + let loops = |s: &str| s.matches("loop {").count() + s.matches("for (").count() + s.matches("while (").count(); + for merged in [None, Some(caps())] { + let v = KernelVariant { merged, ..KernelVariant::new(false, Default::default(), crate::gpu::F_SOURCE) }; + let plain = kernel_wgsl(&kernel_template(), &v); + let g = kernel_wgsl(&kernel_template(), &KernelVariant { guard: Some(1 << 16), ..v }); + assert!(loops(&plain) > 40, "the kernel's loops"); + assert_eq!(g.matches("if (loop_guard()) {").count(), loops(&plain)); + let (m, info) = parse(&g); let (used, declared) = storage_buffers(&m, &info); // (As written: one buffer more, where the device has room. Merged: still eight.) - assert!(declared == 20 || declared as u32 == STORAGE_BUFFERS, "{declared}: {used:?}"); + assert_eq!(declared as u32, if merged.is_some() { STORAGE_BUFFERS } else { 20 }, "{used:?}"); } - let g = guard("fn f() {\n for (var i = 0u; i < n(2u); i++) { x(); }\n loop { // {\n break;\n }\n while (a) { }\n}\n", 9, "T"); - assert!(g.starts_with("fn f() {\n for (var i = 0u; i < n(2u); i++) { if (loop_guard()) { break; } x(); }\n loop { if (loop_guard()) { break; } // {\n break;\n }\n while (a) { if (loop_guard()) { break; } }\n}\n"), "{g}"); - assert!(g.ends_with("fn loop_guard() -> bool {\n LOOP_GUARD_N += 1u;\n if (LOOP_GUARD_N == 9u + 1u) { atomicAdd(&T[0], 1u); }\n return LOOP_GUARD_N > 9u;\n}\n"), "{g}"); + // The pass on its own: nested loops, a `continuing` block, comments and braces in them, a loop in an `if`. + let mut tu = crate::gpu::wesl_wgsl("fn f() {\n for (var i = 0u; i < 2u; i++) { // {\n loop { if (a) { break; } continuing { while (b) { } } }\n }\n if (c) { while (d) { } } else { x(); }\n}\n", &[], &[]); + crate::gpu::guard_loops(&mut tu); + let g = tu.to_string(); + assert_eq!(g.matches("if (loop_guard()) {").count(), 4, "{g}"); + let first = |w: &str| g[g.find(w).unwrap()..].split('\n').nth(1).unwrap().trim().to_string(); + assert_eq!([first("for ("), first("loop {"), first("while (b)"), first("while (d)")], ["if (loop_guard()) {"; 4], "{g}"); } + /// The constants a variant imports are the ones it is given, by name: a merged kernel's lengths, a guard's limit. #[test] - #[should_panic(expected = "`refl` is used in a way the generator does not know (line 2)")] - fn a_use_the_generator_does_not_know_stops_it() { - rewrite("fn f() {\n let p = &refl;\n}\n", &[Rule { name: "refl", open: "ML.refl[".into(), close: String::new(), len: "7u".into() }]); + fn the_variant_gives_what_it_imports() { + let v = KernelVariant { merged: Some(caps()), guard: Some(9), ..KernelVariant::new(false, Default::default(), 0) }; + let c = v.constants(); + assert_eq!(c.len(), 7 + 6 + 2); + assert!(c.contains(&("MT_TEX_INFO".into(), TABLE_MIN[3])) && c.contains(&("ML_DIRECT_OUT".into(), 64 * 36)) && c.contains(&("LOOP_GUARD_LIMIT".into(), 9)) && c.contains(&("GUARD_WORDS".into(), GUARD_WORDS))); + let u = KernelVariant { uniform_tables: true, ..KernelVariant::new(false, Default::default(), 0) }.constants(); + assert_eq!(u, vec![("UT_TEX_INFO".into(), 4096), ("UT_OBJECTS".into(), 4096), ("UT_INSTANCES".into(), 409), ("UT_MEDIA".into(), 32)]); + assert!(KernelVariant::default().constants().is_empty()); } #[test] - #[should_panic(expected = "the declaration")] - fn a_declaration_that_moved_stops_it() { - kernel(&KERNEL.replace("var lights: array;", "var lights: array;"), &caps()); + #[should_panic(expected = "MT_MEDIA")] + fn a_constant_the_source_imports_must_be_given() { + let mut c = KernelVariant { merged: Some(caps()), ..KernelVariant::new(false, Default::default(), 0) }.constants(); + c.retain(|(n, _)| n != "MT_MEDIA"); + crate::gpu::wesl_wgsl(&kernel_template(), &KernelVariant { merged: Some(caps()), ..Default::default() }.flags(), &c); } } diff --git a/renderer/src/scene.rs b/renderer/src/scene.rs index 4a5b2c1..461360d 100644 --- a/renderer/src/scene.rs +++ b/renderer/src/scene.rs @@ -13,7 +13,7 @@ use std::sync::Arc; pub const MAT_DOUBLE_SIDED: i32 = 1; pub const MAT_EMITTER_CARD: i32 = 2; -/// three.js MeshBasicMaterial (`desc::Material::unlit`; the kernel's F_UNLIT, `@if(unlit)`: gpu.rs `kernel_flags`). +/// three.js MeshBasicMaterial (`desc::Material::unlit`; the kernel's F_UNLIT, `@if(unlit)`: gpu.rs `KernelVariant::flags`). pub const MAT_UNLIT: i32 = 4; pub const TRI_NO_SHADOW: u32 = 1; /// Alpha-tested triangle (cutout material): traversal runs an any-hit test. diff --git a/renderer/src/units.rs b/renderer/src/units.rs index 3f46fed..c47475b 100644 --- a/renderer/src/units.rs +++ b/renderer/src/units.rs @@ -1220,7 +1220,7 @@ mod tests { "stack[sp] = far;", "continue;", "} else if (rec == NO_INST) {", - "let u = instances[n.left_or_first];", + "let u = instance_at(n.left_or_first);", "if (mode != TRACE_SHADOW || bitcast(u.n1.w) != 0u) {", "rec = n.left_or_first;", "ro = o - vec3(u.m0.w, u.m1.w, u.m2.w);", diff --git a/tools/web-hello.ts b/tools/web-hello.ts index 577f136..d848610 100644 --- a/tools/web-hello.ts +++ b/tools/web-hello.ts @@ -16,7 +16,7 @@ // --cold a fresh browser profile: the shaders compiled again (the first visit's times) // // A GPU job that does not end takes this machine's display with it (2026-11-11), and a page's shader runs on the same -// GPU as everything else. So: the kernel in the page has the loop guard (renderer/src/merged.rs `guard`), and its +// GPU as everything else. So: the kernel in the page has the loop guard (renderer/src/gpu.rs `guard_loops`), and its // count is a check here (0, or the frames are wrong); the first step is small; every wait on the page has a timeout; // and a timeout stops the run where it is: nothing is retried, no second page is opened (exit 3; see the message). // The GPU lock is held.