From df45ca5bd092911a2161f22d726fbb5587441173 Mon Sep 17 00:00:00 2001 From: yujiaxu-ops Date: Sat, 14 Nov 2026 03:58:30 +0800 Subject: [PATCH] The kernel's variants by WESL's conditional translation, and the repository's Rust pinned (1.100.0) The kernel was made for a device and a scene by text patches: the hardware traversal cut in at markers (kernel_base), a device's switches and a scene's source features found and replaced at anchors that were whole lines of code (kernel_switches, kernel_source), an edit of such a line breaking a patch at run time. Now the kernel is one WESL module (pathtrace.wgsl, trace_units.wgsl, trace_hw.wgsl, translucent.wgsl) with every variant in it under @if(flag) on a declaration or a statement, made by wesl-rs 0.6 (natively and in a browser): hwrt; the switches strict, cutout_lite, cutout_lite2, last_direct, light_pick; the source features vcolor, grid, tint, unlit, units, background (gpu.rs KERNEL_FILES, kernel_template, kernel_flags, kernel_wgsl). A flag the source names that is not given is an error. Still text, before the variant: a scene's functions in their slots, the merged layout, the uniform tables, the loop guard. rust-toolchain.toml pins 1.100.0 (wesl-rs 0.6 asks for 1.97.1; it was whatever stable a machine had), and every Rust job of CI runs for it. fetch_update is try_update. WGSL a browser refuses, found on the way: WESL parses by the specification's grammar, as Chrome's Tint does, and the scene library's surface-material hash (`a * b ^ c * d`: a bitwise operator's operands must be unary) did not compile through it; naga had taken it, so those materials and the particle passes (the same, three times) never compiled in a browser. Parenthesised in the order naga read them. A unit test holds the shader files outside the kernel to the grammar. Checks: for all 2,304 combinations of the flags the Metal source naga writes is the patched kernel's (two rewrites allowed for: the light pick's blocks take their condition from a local, and vertex and instance colours multiply in turn: the same products in the same order). On 1.100 the voxel and physics determinism checks give their recorded hashes, and main built with 1.96 and with 1.100 renders the same (render-identity, native and merged). Against main on 1.100, render-identity native, THREERT_MERGED=1, every switch on, the uniform tables: nothing changed. tools/web-hello.ts in Chrome: pass. cargo test (126), npm typecheck / test / npm-pack. Co-Authored-By: Claude Opus 5.5 --- .github/workflows/ci.yml | 11 +- docs/ROADMAP.md | 25 ++ docs/RUNTIME.md | 43 ++- packages/scene/src/materials.ts | 4 +- renderer/Cargo.lock | 379 +++++++++++++++++++++++- renderer/Cargo.toml | 2 + renderer/shaders/gpu_particles.wgsl | 4 +- renderer/shaders/particles.wgsl | 2 +- renderer/shaders/pathtrace.wgsl | 123 +++++++- renderer/shaders/trace_hw.wgsl | 32 +- renderer/shaders/trace_units.wgsl | 10 +- renderer/shaders/translucent.wgsl | 30 +- renderer/src/gpu.rs | 434 +++++++++++----------------- renderer/src/merged.rs | 30 +- renderer/src/scene.rs | 2 +- renderer/src/script.rs | 2 +- renderer/src/units.rs | 2 +- renderer/web/Cargo.lock | 375 ++++++++++++++++++++++++ rust-toolchain.toml | 6 + 19 files changed, 1166 insertions(+), 350 deletions(-) create mode 100644 rust-toolchain.toml diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index f7ebc3f..2ba6ed4 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -54,19 +54,20 @@ jobs: } # The renderer's release build: the crate, its path dependencies (voxel-core, physics-core, third_party) and # the presets it embeds (src/files.rs); not the browser module's crate (renderer/web), which it does not build. - job renderer '^(renderer/|presets/)' '^renderer/web/' + # Every Rust job runs for the toolchain (rust-toolchain.toml: the Rust they all build with). + job renderer '^(renderer/|presets/|rust-toolchain\.toml$)' '^renderer/web/' # voxel-core and its determinism check. - job voxel '^(renderer/voxel-core/|tools/voxel-determinism\.ts$)' + job voxel '^(renderer/voxel-core/|tools/voxel-determinism\.ts$|rust-toolchain\.toml$)' # physics-core, the patched crates it builds with, its determinism check, and what that check builds of the # renderer for wasm32 (examples/physics_bits.rs, in the renderer's crate graph: its manifest and lock file; # the voxel bodies' script it runs there: voxel-core, the renderer's voxel and physics adapters, and the api # and world they go through). - job physics '^(renderer/(physics-core/|voxel-core/|third_party/|Cargo\.toml$|Cargo\.lock$|examples/physics_bits\.rs$|src/(voxel|physics|api|world)\.rs$)|tools/physics-determinism\.ts$)' + job physics '^(renderer/(physics-core/|voxel-core/|third_party/|Cargo\.toml$|Cargo\.lock$|examples/physics_bits\.rs$|src/(voxel|physics|api|world)\.rs$)|tools/physics-determinism\.ts$|rust-toolchain\.toml$)' # The engine's library for wasm32: the same as the renderer's. - job wasm '^(renderer/|presets/)' '^renderer/web/' + job wasm '^(renderer/|presets/|rust-toolchain\.toml$)' '^renderer/web/' # The browser module (renderer/web, on the engine's library) and tools/web-engine.ts: the module in Node with # the physics hello world's bundle (games/, packages/, the bundler's and the build's tools, the npm manifests). - job web '^(renderer/|presets/|games/|packages/|tools/(web-engine\.ts|bundle-game\.ts|lib/)|package(-lock)?\.json$)' + job web '^(renderer/|presets/|games/|packages/|tools/(web-engine\.ts|bundle-game\.ts|lib/)|package(-lock)?\.json$|rust-toolchain\.toml$)' ts: name: "ts: typecheck, tests" diff --git a/docs/ROADMAP.md b/docs/ROADMAP.md index b7551e9..ea7f506 100644 --- a/docs/ROADMAP.md +++ b/docs/ROADMAP.md @@ -2235,6 +2235,31 @@ Checks: render-identity against the build before, native: 3 of 190 outputs diffe the passes after the trace 6.00-6.08 -> 6.19-6.24; `track` 0.27 less `compose`'s 0.08; the surfaces' copy in `refl` +0.02, a pass of its own read the samples again: 0.18). +**The kernel's variants by WESL, and the Rust pinned (2026-11-14, after it).** The kernel was made for a device and +a scene by text patches: the hardware traversal cut in at markers (`kernel_base`), a device's switches and a scene's +source features found and replaced at anchors that were whole lines of code (`kernel_switches`, `kernel_source`), an +edit of such a line breaking a patch at run time. Now the kernel is one WESL module (pathtrace.wgsl, then +trace_units.wgsl, trace_hw.wgsl, translucent.wgsl) and every variant is in it under WESL's conditional translation, +`@if(flag)` on a declaration or a statement (WESL: the community's WGSL extension, Bevy's and Vello's; wesl-rs 0.6, +natively and in a browser). The flags: `hwrt`, the switches (`strict`, `cutout_lite`, `cutout_lite2`, `last_direct`, +`light_pick`), the source features (`vcolor`, `grid`, `tint`, `unlit`, `units`, `background`); gpu.rs `kernel_flags`, +`kernel_wgsl`, docs/RUNTIME.md "The kernel's source". Checks: for all 2,304 combinations of the flags the Metal source +naga writes is the patched kernel's, two rewrites allowed for (the light pick's blocks take their condition from a +local; vertex and instance colours multiply in turn: the same products in the same order); render-identity against +main, native, merged, with every switch on and with the uniform tables (below). Left as text, before the variant: a +scene's functions in their slots, the merged layout, the uniform tables, the loop guard. + +The repository's Rust is pinned (`rust-toolchain.toml`, 1.100.0; CI's jobs run for it): wesl-rs 0.6 asks for 1.97.1, +and it was whatever stable a machine had (1.96 here). On 1.100 the determinism checks give their recorded hashes +(voxel, physics sequential and on 2 and 4 threads, the renderer's wasm32 graph), and main built with either compiler +renders the same (render-identity, native and merged: nothing changed). One deprecation: `fetch_update` is +`try_update`. + +On the way, WGSL a browser refuses: WESL parses by the specification's grammar, as Chrome's Tint does, and the +surface materials' hash (`a * b ^ c * d`: a bitwise operator's operands must be unary) did not compile through it. +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. + ## 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 e6f0499..5864991 100644 --- a/docs/RUNTIME.md +++ b/docs/RUNTIME.md @@ -908,11 +908,10 @@ procedural objects (`trace_procs`: the ray into the object's space, a tree of it waited at every unit for the slowest of its group (1.5 ms of the voxel level's 6.8). The top tree's root's box is 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. -- **`pathtrace.wgsl` is not edited.** The unit walk is `shaders/trace_units.wgsl`, put into the source by - `kernel_source` for scenes that have units under software traversal (`F_UNITS`, a source feature, as vertex - colours are): the file's text, a call at the end of the still walk, and the still walk's stack declaration taken - out ("The stack"). A scene without units is handed the source it was handed before, byte for byte; hardware - traversal's source never contains it. +- **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 + 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 of them than a top tree holds (2^18), and every instance of an offline render (`prepare`; a live session's scene goes through `prepare_live`), unless `THREERT_UNITS=1`. `THREERT_UNITS=0` flattens everything, as before this @@ -949,8 +948,7 @@ not yet used, and a tree given back, hold leaves that end (`units::FILL`), not z whose children are nodes 0 and 1). With no unit shown the count is 0 and the kernel does not enter the top range. **The stack.** One stack of 64 entries for an invocation of a kernel with units: a private variable of the module, -which the still walk uses too (`kernel_source` takes its own declaration out of the source for such a kernel; the -file is not edited). Why not a walk's own: an array a function declares is set to zero each time the function is +which the still walk uses too (its own declaration is under `@if(!units)`). Why not a walk's own: an array a function declares is set to zero each time the function is entered, whatever the ray then meets, and a ray's walk is entered many times a pixel. The unit walk's own 64 words cost hello world 0.3 ms a frame, Kabukichō 0.75 and the voxel level 0.55; the still walk's own 48 cost them 0.36, 1.1 and 1.5 more, **and cost every flattened scene the same** (a kernel without units is the file as it is, so it @@ -1171,6 +1169,35 @@ the CPU's update of the instances, not the trace. main; `cargo test --release` whole (79 pass); the atomic tests under hardware traversal (170 pass) and the 29 with units under software traversal; hello world natively and in Chrome (every check, the guard 0, no frame rebuilt). +## The kernel's source + +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`). +- 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 + emitters and environment take their conditions from a local (a statement is under `@if` whole, and those blocks are + long), and vertex and instance colours multiply the base colour in turn rather than in one expression (the same + products, in the same order). +- WESL's parser is the specification's grammar, as strict as a browser's (Chrome's Tint); naga, wgpu's front end, is + not. Through it the scene library's surface materials did not compile: their hash read `a * b ^ c * d`, which WGSL + forbids (a bitwise operator's operands are unary expressions) and which naga took. The same in the particle passes. + So they never compiled in a browser either; fixed 2026-11-14 with the parentheses in the order naga read them, and + the shader files outside the kernel are held to the grammar by a unit test (gpu.rs + `the_other_shaders_are_wgsl_as_the_specification_writes_it`). A scene's own functions go through WESL with the + kernel: WGSL a browser would refuse is refused natively too. +- The repository's Rust is pinned (`rust-toolchain.toml`, 1.100.0): wesl-rs 0.6 asks for 1.97.1 or newer, and the + determinism checks' recorded hashes are of one compiler. + ## Android (Adreno): what a phone's Vulkan driver needed The engine's library builds for `aarch64-linux-android` and runs under an Android client on Vulkan with ray queries. @@ -1214,7 +1241,7 @@ on the device", and differs where it says so. Two of these are what a device run should look at first: `prev_verts` is the one storage buffer left after the acceleration structure's binding (13) and feeds the motion vectors of moving objects and the vertex colours; `env_cdf` also holds the procedural geometry (voxel worlds). -- **The kernel's switches** (`gpu::KernelSwitches`; patches of the kernel's source, as a scene's features are): +- **The kernel's switches** (`gpu::KernelSwitches`; flags of the kernel's conditional translation, `@if` in its source, as a scene's features are: "The kernel's source" below): | switch | default | what it does | on the device | |---|---|---|---| diff --git a/packages/scene/src/materials.ts b/packages/scene/src/materials.ts index 2f1fe26..c32a05e 100644 --- a/packages/scene/src/materials.ts +++ b/packages/scene/src/materials.ts @@ -655,7 +655,7 @@ fn std_pcg(v: u32) -> u32 { return (w >> 22u) ^ w; } fn std_hashi(x: i32, y: i32, z: i32) -> f32 { - return f32(std_pcg(bitcast(x) * 73856093u ^ bitcast(y) * 19349663u ^ bitcast(z) * 83492791u) >> 8u) * (1.0 / 16777216.0); + return f32(std_pcg((bitcast(x) * 73856093u) ^ (bitcast(y) * 19349663u) ^ (bitcast(z) * 83492791u)) >> 8u) * (1.0 / 16777216.0); } fn std_vnoise(p: vec2) -> f32 { let i = floor(p); @@ -1013,7 +1013,7 @@ export class SurfaceMaterial extends LibraryMaterial { // text, and the Metal compiler failed on the translucent pass with SURFACE_HELPERS' window code in it.) const WET_HELPERS = ` fn std_wl_hash(x: i32, y: i32, z: i32) -> f32 { - let v = bitcast(x) * 73856093u ^ bitcast(y) * 19349663u ^ bitcast(z) * 83492791u; + let v = (bitcast(x) * 73856093u) ^ (bitcast(y) * 19349663u) ^ (bitcast(z) * 83492791u); let s = v * 747796405u + 2891336453u; let w = ((s >> ((s >> 28u) + 4u)) ^ s) * 277803737u; return f32(((w >> 22u) ^ w) >> 8u) * (1.0 / 16777216.0); diff --git a/renderer/Cargo.lock b/renderer/Cargo.lock index 24560ba..c53a925 100644 --- a/renderer/Cargo.lock +++ b/renderer/Cargo.lock @@ -37,6 +37,15 @@ dependencies = [ "zerocopy", ] +[[package]] +name = "aho-corasick" +version = "1.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c982642fa9e8606056828ee9a8505737230110bb1099153c79efe865c59d12ba" +dependencies = [ + "memchr", +] + [[package]] name = "allocator-api2" version = "0.2.21" @@ -77,6 +86,67 @@ dependencies = [ "libc", ] +[[package]] +name = "annotate-snippets" +version = "0.12.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0258346f92c78dbc5b336b9a5b27c64447aa10643bc64e651db2edb145361564" +dependencies = [ + "anstyle", + "memchr", + "unicode-width", +] + +[[package]] +name = "anstream" +version = "1.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "824a212faf96e9acacdbd09febd34438f8f711fb84e09a8916013cd7815ca28d" +dependencies = [ + "anstyle", + "anstyle-parse", + "anstyle-query", + "anstyle-wincon", + "colorchoice", + "is_terminal_polyfill", + "utf8parse", +] + +[[package]] +name = "anstyle" +version = "1.0.14" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "940b3a0ca603d1eade50a4846a2afffd5ef57a9feac2c0e2ec2e14f9ead76000" + +[[package]] +name = "anstyle-parse" +version = "1.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "52ce7f38b242319f7cabaa6813055467063ecdc9d355bbb4ce0c68908cd8130e" +dependencies = [ + "utf8parse", +] + +[[package]] +name = "anstyle-query" +version = "1.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "40c48f72fd53cd289104fc64099abca73db4166ad86ea0b4341abe65af83dadc" +dependencies = [ + "windows-sys 0.61.2", +] + +[[package]] +name = "anstyle-wincon" +version = "3.0.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "291e6a250ff86cd4a820112fb8898808a366d8f9f58ce16d1f538353ad55747d" +dependencies = [ + "anstyle", + "once_cell_polyfill", + "windows-sys 0.61.2", +] + [[package]] name = "anyhow" version = "1.0.104" @@ -328,6 +398,12 @@ dependencies = [ "unicode-width", ] +[[package]] +name = "colorchoice" +version = "1.0.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1d07550c9036bf2ae0c684c4297d503f838287c83c53686d05370d0e139ae570" + [[package]] name = "combine" version = "4.6.9" @@ -353,6 +429,15 @@ version = "0.10.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "be95d4a884358a21a199689323c29e9256b4a1b315e22f183d155422041d3a9a" +[[package]] +name = "convert_case" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "633458d4ef8c78b72454de2d54fd6ab2e60f9e02be22f3c6104cdc8a4e0fceb9" +dependencies = [ + "unicode-segmentation", +] + [[package]] name = "core-foundation" version = "0.9.4" @@ -482,6 +567,29 @@ version = "2.11.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "4583a4551df46e2792f82ceeac45e850d2e2d5debba0b91f102385cda5b11f06" +[[package]] +name = "derive_more" +version = "2.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d751e9e49156b02b44f9c1815bcb94b984cdcc4396ecc32521c739452808b134" +dependencies = [ + "derive_more-impl", +] + +[[package]] +name = "derive_more-impl" +version = "2.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "799a97264921d8623a957f6c3b9011f3b5492f557bbb7a5a19b7fa6d06ba8dcb" +dependencies = [ + "convert_case", + "proc-macro2", + "quote", + "rustc_version", + "syn 2.0.119", + "unicode-xid", +] + [[package]] name = "digest" version = "0.10.7" @@ -628,6 +736,12 @@ dependencies = [ "zlib-rs", ] +[[package]] +name = "fnv" +version = "1.0.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3f9eec918d3f24069decb9af1554cad7c880e2da24a9afd88aca000531ab82c1" + [[package]] name = "foldhash" version = "0.2.0" @@ -793,6 +907,12 @@ dependencies = [ "simba", ] +[[package]] +name = "glob" +version = "0.3.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e4eba85ea1d0a966a983acd07deee566e67395d2d96b6fb39e62b5a833f1eb0b" + [[package]] name = "glow" version = "0.17.0" @@ -986,6 +1106,21 @@ version = "1.1.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "a257582fdcde896fd96463bf2d40eefea0580021c0712a0e2b028b60b47a837a" +[[package]] +name = "is_terminal_polyfill" +version = "1.70.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a6cb138bb79a146c1bd460005623e142ef0181e3d0219cb493e02f7d08a35695" + +[[package]] +name = "itertools" +version = "0.15.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8b4baf93f58d4425749ca49a51c50ebab072c5df6994d08fed93541c331481dc" +dependencies = [ + "either", +] + [[package]] name = "itoa" version = "1.0.18" @@ -1088,6 +1223,15 @@ version = "3.1.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "e2db585e1d738fc771bf08a151420d3ed193d9d895a36df7f6f8a9456b911ddc" +[[package]] +name = "lalrpop-util" +version = "0.23.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "364b49d1c98c07972e16fe1346463becc08bcd663da04e6e37ab1bea068dce6a" +dependencies = [ + "regex-automata", +] + [[package]] name = "lazy_static" version = "1.5.1" @@ -1100,6 +1244,72 @@ version = "0.5.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "7a79a3332a6609480d7d0c9eab957bca6b455b91bb84e66d19f5ff66294b85b8" +[[package]] +name = "lexical" +version = "7.0.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1bc8a009b2ff1f419ccc62706f04fe0ca6e67b37460513964a3dfdb919bb37d6" +dependencies = [ + "lexical-core", +] + +[[package]] +name = "lexical-core" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7d8d125a277f807e55a77304455eb7b1cb52f2b18c143b60e766c120bd64a594" +dependencies = [ + "lexical-parse-float", + "lexical-parse-integer", + "lexical-util", + "lexical-write-float", + "lexical-write-integer", +] + +[[package]] +name = "lexical-parse-float" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "52a9f232fbd6f550bc0137dcb5f99ab674071ac2d690ac69704593cb4abbea56" +dependencies = [ + "lexical-parse-integer", + "lexical-util", +] + +[[package]] +name = "lexical-parse-integer" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9a7a039f8fb9c19c996cd7b2fcce303c1b2874fe1aca544edc85c4a5f8489b34" +dependencies = [ + "lexical-util", +] + +[[package]] +name = "lexical-util" +version = "1.0.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2604dd126bb14f13fb5d1bd6a66155079cb9fa655b37f875b3a742c705dbed17" + +[[package]] +name = "lexical-write-float" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "50c438c87c013188d415fbabbb1dceb44249ab81664efbd31b14ae55dabb6361" +dependencies = [ + "lexical-util", + "lexical-write-integer", +] + +[[package]] +name = "lexical-write-integer" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "409851a618475d2d5796377cad353802345cba92c867d9fbcde9cf4eac4e14df" +dependencies = [ + "lexical-util", +] + [[package]] name = "libc" version = "0.2.191" @@ -1167,6 +1377,38 @@ version = "0.4.35" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "5d15f86d12da44cd8d4ec3c73984537eb7f4cadcb9776a575a0a8a05ca11005c" +[[package]] +name = "logos" +version = "0.16.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "eb2c55a318a87600ea870ff8c2012148b44bf18b74fad48d0f835c38c7d07c5f" +dependencies = [ + "logos-derive", +] + +[[package]] +name = "logos-codegen" +version = "0.16.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "58b3ffaa284e1350d017a57d04ada118c4583cf260c8fb01e0fe28a2e9cf8970" +dependencies = [ + "fnv", + "proc-macro2", + "quote", + "regex-automata", + "regex-syntax", + "syn 2.0.119", +] + +[[package]] +name = "logos-derive" +version = "0.16.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "52d3a9855747c17eaf4383823f135220716ab49bea5fbea7dd42cc9a92f8aa31" +dependencies = [ + "logos-codegen", +] + [[package]] name = "matrixmultiply" version = "0.3.11" @@ -1704,6 +1946,12 @@ version = "1.21.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "9f7c3e4beb33f85d45ae3e3a1792185706c8e16d043238c593331cc7cd313b50" +[[package]] +name = "once_cell_polyfill" +version = "1.70.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "384b8ab6d37215f3c5301a95a4accb5d64aa607f1fcb26a11b5303878451b4fe" + [[package]] name = "orbclient" version = "0.3.55" @@ -2124,6 +2372,23 @@ dependencies = [ "bitflags 2.13.2", ] +[[package]] +name = "regex-automata" +version = "0.4.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ad8553b9b26413251cbf30e620595c7a41b3887f03da04579c0e6b0d6a06b4b2" +dependencies = [ + "aho-corasick", + "memchr", + "regex-syntax", +] + +[[package]] +name = "regex-syntax" +version = "0.8.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4" + [[package]] name = "relative-path" version = "2.0.3" @@ -2373,6 +2638,15 @@ dependencies = [ "zmij", ] +[[package]] +name = "serde_spanned" +version = "1.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9c187e0a369471c550af028cba56d4f4fa2a489ad28151ae95a4d340a2a8b782" +dependencies = [ + "serde_core", +] + [[package]] name = "sha1" version = "0.11.0" @@ -2632,6 +2906,7 @@ dependencies = [ "ureq", "voxel-core", "web-time", + "wesl", "wgpu", "winit", ] @@ -2670,11 +2945,26 @@ dependencies = [ "pin-project-lite", ] +[[package]] +name = "toml" +version = "1.2.1+spec-1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d5755c51863cdac711f1956baf8eb9588e060755f7d3c4c979e70ea1e1f32cd7" +dependencies = [ + "indexmap", + "serde_core", + "serde_spanned", + "toml_datetime", + "toml_parser", + "toml_writer", + "winnow", +] + [[package]] name = "toml_datetime" -version = "1.1.2+spec-1.1.0" +version = "1.1.3+spec-1.1.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "2b86d767906c6c42421dcba507eb9d203e779497710a47782a224bb871653053" +checksum = "173f25c883cd2176a2d8212155638c0a60d484a9d84546d270866f40f56b6e49" dependencies = [ "serde_core", ] @@ -2700,6 +2990,12 @@ dependencies = [ "winnow", ] +[[package]] +name = "toml_writer" +version = "1.2.0+spec-1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "009dbb7700f11c55ef83aebbe4c0e9cf1f1ac0844bd8d6f00dc38f302026e13a" + [[package]] name = "tracing" version = "0.1.44" @@ -2765,6 +3061,12 @@ version = "0.2.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "ca718471876992059f251e33aafc334346ada08042b7661dd1fe313a9bced557" +[[package]] +name = "unicode-xid" +version = "0.2.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ebc1c04c71510c7f702b52b7c350734c9ff1295c464a03335b00bb84fc54f853" + [[package]] name = "untrusted" version = "0.9.0" @@ -2812,6 +3114,12 @@ version = "0.8.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "b8c0a043c9540bae7c578c88f91dda8bd82e59ae27c21baca69c8b191aaf5a6e" +[[package]] +name = "utf8parse" +version = "0.2.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "06abde3611657adf66d383f00b093d7faecc7fa57071cce2578660c9f1010821" + [[package]] name = "version_check" version = "0.9.5" @@ -3049,6 +3357,46 @@ dependencies = [ "rustls-pki-types", ] +[[package]] +name = "wesl" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fdd8f9619baa4c7bde4df0899d6f626b70ee92d6e9c18c5f74473419b5bb6fa6" +dependencies = [ + "annotate-snippets", + "anstream", + "derive_more", + "glob", + "half", + "itertools", + "num-traits", + "serde", + "thiserror 2.0.21", + "toml", + "wesl-core", + "wesl-macros", + "wgsl-parse", + "wgsl-types", +] + +[[package]] +name = "wesl-core" +version = "0.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "84dd3e600497eda0b47afcd21bd84c8953a9dfce7e93ed27a38ba97ac0747e03" + +[[package]] +name = "wesl-macros" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ac8767ec26e75262a8e581d5f188b1f63a288baaf885bb33c9f4f73a693e06fc" +dependencies = [ + "itertools", + "proc-macro2", + "quote", + "syn 3.0.9", +] + [[package]] name = "wgpu" version = "30.0.1" @@ -3222,6 +3570,33 @@ dependencies = [ "web-sys", ] +[[package]] +name = "wgsl-parse" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "37cee524ba5cfabfd65d10a316b646f8a4c50f3663e7e787471f21b783261f47" +dependencies = [ + "annotate-snippets", + "derive_more", + "itertools", + "lalrpop-util", + "lexical", + "logos", + "thiserror 2.0.21", + "wgsl-types", +] + +[[package]] +name = "wgsl-types" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e119dd12e57d937397535201959ddddb5dbf26505402db5ef305ebb1ca9a7e5f" +dependencies = [ + "half", + "itertools", + "num-traits", +] + [[package]] name = "wide" version = "1.7.1" diff --git a/renderer/Cargo.toml b/renderer/Cargo.toml index 150b2bd..add0a69 100644 --- a/renderer/Cargo.toml +++ b/renderer/Cargo.toml @@ -9,6 +9,8 @@ license = "MIT" voxel-core = { path = "voxel-core" } physics-core = { path = "physics-core" } wgpu = "30" +# The kernel's variants (gpu.rs `kernel_wgsl`): WESL's conditional translation, by wesl-rs (WESL's compiler in Rust). +wesl = { version = "0.6", features = ["naga-ext"] } bytemuck = { version = "1", features = ["derive"] } serde = { version = "1", features = ["derive"] } serde_json = "1" diff --git a/renderer/shaders/gpu_particles.wgsl b/renderer/shaders/gpu_particles.wgsl index afe64ef..eb094f3 100644 --- a/renderer/shaders/gpu_particles.wgsl +++ b/renderer/shaders/gpu_particles.wgsl @@ -94,7 +94,7 @@ fn solid(x: vec3) -> bool { if (any(c < vec3(0.0)) || c.x >= f32(PPS.grid_n.x) || c.y >= f32(PPS.grid_n.y) || c.z >= f32(PPS.grid_n.z)) { return false; } let q = vec3(c); let row = (q.y * PPS.grid_n.z + q.z) * PPS.grid_n.w + q.x / 32u; - return (grid[row] >> (q.x % 32u) & 1u) != 0u; + return ((grid[row] >> (q.x % 32u)) & 1u) != 0u; } // One step of a particle that moves from `s.pos` by `s.vel` x dt, against the voxel grid and the plane (y = height): on @@ -191,7 +191,7 @@ fn gpu_particles(@builtin(global_invocation_id) gid: vec3) { var i: EIn; i.index = k; i.time = PPS.cam.w; - let h = pcg(sys.b.x * 2654435761u ^ pcg(k ^ pcg(sys.e.z))); + let h = pcg((sys.b.x * 2654435761u) ^ pcg(k ^ pcg(sys.e.z))); i.rand = vec4(unit(pcg(h)), unit(pcg(h + 1u)), unit(pcg(h + 2u)), unit(pcg(h + 3u))); i.pivot = sys.c.xyz; i.spawned = rel < sys.e.y; diff --git a/renderer/shaders/particles.wgsl b/renderer/shaders/particles.wgsl index 5855010..c2e431e 100644 --- a/renderer/shaders/particles.wgsl +++ b/renderer/shaders/particles.wgsl @@ -127,7 +127,7 @@ fn particles(@builtin(global_invocation_id) gid: vec3) { i.fwd = PPS.fwd.xyz; i.rain = PPS.env.x; i.wetness = PPS.env.y; - let h = pcg(s.b.x * 2654435761u ^ pcg(k)); + let h = pcg((s.b.x * 2654435761u) ^ pcg(k)); i.rand = vec4(unit(pcg(h)), unit(pcg(h + 1u)), unit(pcg(h + 2u)), unit(pcg(h + 3u))); i.pivot = s.c.xyz; i.custom = s.d; diff --git a/renderer/shaders/pathtrace.wgsl b/renderer/shaders/pathtrace.wgsl index c5e2144..e365585 100644 --- a/renderer/shaders/pathtrace.wgsl +++ b/renderer/shaders/pathtrace.wgsl @@ -4,6 +4,18 @@ // Material / atmosphere extensions (docs/PARITY.md): texture pages with ray-cone LOD, alpha cutout, world mapping, // wetness / puddles / ripples / grime, procedural facade windows, emissive animation, RIS many-light sampling, height // fog with volumetric single scattering. Each is compiled out when the scene does not use it (FEAT). +// +// The kernel is this file and trace_units.wgsl, trace_hw.wgsl and translucent.wgsl after it (gpu.rs `KERNEL_FILES`), +// compiled for a device and a scene by WESL's conditional translation (`@if`, gpu.rs `kernel_wgsl`): a declaration +// or a statement under `@if(x)` is there only where x holds. The flags: +// hwrt hardware traversal (ray queries) in the software BVH walk's place, and the translucency pass +// strict, cutout_lite, cutout_lite2, last_direct, light_pick +// a device's switches (gpu.rs `KernelSwitches`: what a GPU the kernel as written is wrong or +// too slow on needs) +// 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) +@if(hwrt) enable wgpu_ray_query; const PI: f32 = 3.14159265358979; const INF: f32 = 3.0e38; @@ -33,6 +45,7 @@ const EMIT_RIS_M: u32 = 6u; const VOL_RATE_MASK: u32 = 1u; const MAT_DOUBLE_SIDED: i32 = 1; const MAT_EMITTER_CARD: i32 = 2; +@if(unlit) const MAT_UNLIT: i32 = 4; // Live (real-time) mode: write this frame's 1-spp sample + first-hit distance instead of accumulating. override LIVE: bool = false; // Features the scene uses (gpu::F_*): the rest is compiled out, so plain scenes pay nothing for them. @@ -79,6 +92,7 @@ struct Params { ltree: vec4, // light tree (ltree.rs): first node in `emitters` (bits), its entry count (bits: the moving emitters' first), P(pick a moving one), candidates at camera hits (0 = off) procs: vec4, // procedural geometry (F_PROC): its data's first word in env_cdf, object count, first record in `instances`, - aerial: vec4, // the sky's aerial perspective (F_AERIAL): density multiple, Rayleigh scale, Mie scale, - + @if(background) background: vec4, // three.js scene.background (radiance), 1 }; struct TriAttr { n0: vec4, n1: vec4, n2: vec4, uv: vec4, t: vec4, x: vec4 }; @@ -127,6 +141,17 @@ struct Instance { m0: vec4, m1: vec4, m2: vec4, n0: vec4, n1 // 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; +// 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 +// vertices in `prev_verts` from vec4 P.procs.w. +@if(vcolor) fn vcolor_word(k: u32) -> u32 { return bitcast(prev_verts[P.procs.w + k / 4u][k % 4u]); } +@if(vcolor) fn vcolor(t: u32, u: f32, v: f32) -> vec3 { + let c0 = srgb_to_linear(unpack4x8unorm(vcolor_word(t * 3u)).rgb); + let c1 = srgb_to_linear(unpack4x8unorm(vcolor_word(t * 3u + 1u)).rgb); + let c2 = srgb_to_linear(unpack4x8unorm(vcolor_word(t * 3u + 2u)).rgb); + return c0 * (1.0 - u - v) + c1 * u + c2 * v; +} fn inst_vec(r: Instance, v: vec3) -> vec3 { return vec3(dot(r.m0.xyz, v), dot(r.m1.xyz, v), dot(r.m2.xyz, v)); } @@ -470,9 +495,9 @@ fn tri_hit(ti: u32, o: vec3, d: vec3) -> Hit { return Hit(dot(e2, q) * inv, ti, clamp(u, 0.0, 1.0), clamp(dot(d, q) * inv, 0.0, 1.0), NO_INST); } -// @trace-begin (software BVH; replaced by a ray-query implementation when hardware RT is available) +// The software BVH walk (without `hwrt`: trace_hw.wgsl's traversal is the hardware's). -fn ray_aabb(o: vec3, inv_d: vec3, bmin: vec3, bmax: vec3, tmax: f32) -> f32 { +@if(!hwrt) fn ray_aabb(o: vec3, inv_d: vec3, bmin: vec3, bmax: vec3, tmax: f32) -> f32 { let t0 = (bmin - o) * inv_d; let t1 = (bmax - o) * inv_d; let tmin_v = min(t0, t1); @@ -482,7 +507,7 @@ fn ray_aabb(o: vec3, inv_d: vec3, bmin: vec3, bmax: vec3, tm return select(INF, tn, tn <= tf); } -fn intersect_tri(o: vec3, d: vec3, i: u32, tmax: f32) -> vec3 { +@if(!hwrt) fn intersect_tri(o: vec3, d: vec3, i: u32, tmax: f32) -> vec3 { let v0 = tri_verts[i * 3u].xyz; let e1 = tri_verts[i * 3u + 1u].xyz; let e2 = tri_verts[i * 3u + 2u].xyz; @@ -502,11 +527,11 @@ fn intersect_tri(o: vec3, d: vec3, i: u32, tmax: f32) -> vec3 { } // mode: TRACE_CLOSEST, TRACE_SHADOW (any hit among shadow casters), TRACE_ANY (any hit on any surface). -fn trace(o: vec3, d: vec3, tmax_in: f32, mode: u32) -> Hit { +@if(!hwrt) fn trace(o: vec3, d: vec3, tmax_in: f32, mode: u32) -> Hit { var hit = Hit(tmax_in, 0xffffffffu, 0.0, 0.0, NO_INST); if (rays_off) { return hit; } let inv_d = 1.0 / select(d, vec3(1e-20), abs(d) < vec3(1e-20)); - var stack: array; + @if(!units) var stack: array; var sp = 0u; var node = 0u; loop { @@ -553,6 +578,7 @@ fn trace(o: vec3, d: vec3, tmax_in: f32, mode: u32) -> Hit { sp++; } } + @if(units) hit = trace_units(o, d, inv_d, hit, mode); if (has(F_PROC)) { return trace_procs(o, d, hit, mode); } return hit; } @@ -560,7 +586,7 @@ fn trace(o: vec3, d: vec3, tmax_in: f32, mode: u32) -> Hit { // Procedural geometry, software traversal (the reference the hardware one is checked against): each object's bounds, // then its boxes through the object's BVH (scene::ProcLayout with `sw`: nodes as bvh::GpuNode words, and the box of // each leaf slot, so the geometry function gets the box's own index). -fn trace_procs(o: vec3, d: vec3, hit_in: Hit, mode: u32) -> Hit { +@if(!hwrt) fn trace_procs(o: vec3, d: vec3, hit_in: Hit, mode: u32) -> Hit { var hit = hit_in; let first = P.procs.z; for (var k = 0u; k < P.procs.y; k++) { @@ -623,7 +649,6 @@ fn trace_procs(o: vec3, d: vec3, hit_in: Hit, mode: u32) -> Hit { return hit; } -// @trace-end // Shadow rays of the analytic lights: objects that cast no shadow (castShadow = false, glass) let them through. fn occluded(o: vec3, d: vec3, tmax: f32) -> bool { @@ -719,6 +744,8 @@ fn glass_at(h: Hit, d: vec3, t: f32) -> GlassAt { 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; } g.tint = m.base_color.rgb; + @if(tint) g.tint *= obj_tint(object_of(a)); + @if(vcolor) g.tint *= vcolor(h.tri, h.u, h.v); 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; @@ -969,7 +996,10 @@ fn emissive_material_function(f: i32, i: MatIn, m: Material, base: MatOut) -> Ma fn object_of(a: TriAttr) -> u32 { return bitcast(a.x.y); } // vec4s per object (scene.rs OBJECT_STRIDE): (pivot, power), (bounds size, time of the last power edit), per-instance // data (`MatIn.custom`). -const OBJECT_STRIDE: u32 = 3u; +@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; } // 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; } @@ -1061,6 +1091,8 @@ fn surface(o: vec3, d: vec3, h: Hit, width: f32) -> Surface { s.front = front; s.card = m.tex1.z == MAT_EMITTER_CARD; var base = m.base_color.rgb; + @if(tint) base *= obj_tint(object_of(a)); + @if(vcolor) base *= vcolor(i, h.u, h.v); var rough = m.params.y; var metal = m.params.x; var spec = m.params.z; @@ -1070,6 +1102,13 @@ fn surface(o: vec3, d: vec3, h: Hit, width: f32) -> Surface { if (m.tex0.w >= 0) { metal *= tex(m.tex0.w, uv, lod).b; } } s.emissive = emission(m, uv, lod) * obj_power(a); + // three.js MeshBasicMaterial: its colour to camera rays and mirror-like paths, lit by nothing, lighting nothing. + @if(unlit) if (m.tex1.z == MAT_UNLIT) { + s.emissive = select(vec3(0.0), base, cur_depth == 0u || mirror_path); + base = vec3(0.0); + spec = 0.0; + 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; // Detail that only matters where it is seen directly or in mirror-like reflections (puddles, glass, chrome). @@ -1141,6 +1180,7 @@ fn surface_proc(o: vec3, d: vec3, h: Hit, width: f32) -> Surface { s.front = front; s.card = false; var base = m.base_color.rgb; + @if(tint) base *= obj_tint(oi); var rough = m.params.y; var metal = m.params.x; var spec = m.params.z; @@ -1384,9 +1424,11 @@ fn env_pdf(d: vec3) -> f32 { let uv = env_uv(d); let x = min(u32(uv.x * f32(w)), w - 1u); let y = min(u32(uv.y * f32(h)), h - 1u); - let row_p = env_cdf[y] - select(0.0, env_cdf[y - 1u], y > 0u); + @if(!strict) let row_p = env_cdf[y] - select(0.0, env_cdf[y - 1u], y > 0u); + @if(strict) let row_p = env_cdf[y] - select(0.0, env_cdf[max(y, 1u) - 1u], y > 0u); let base = h + y * w; - let col_p = env_cdf[base + x] - select(0.0, env_cdf[base + x - 1u], x > 0u); + @if(!strict) let col_p = env_cdf[base + x] - select(0.0, env_cdf[base + x - 1u], x > 0u); + @if(strict) let col_p = env_cdf[base + x] - select(0.0, env_cdf[base + max(x, 1u) - 1u], x > 0u); let sin_t = sin((f32(y) + 0.5) / f32(h) * PI); return row_p * col_p * f32(w) * f32(h) / (2.0 * PI * PI * max(sin_t, 1e-4)); } @@ -1439,6 +1481,8 @@ fn offset_origin(p: vec3, n: vec3) -> vec3 { // What a BSDF-sampled ray (solid-angle pdf `pdf`; 0 = the camera ray, no MIS) that left the scene brings back: the // environment and the directional lights' disks, MIS-weighted against their sampling in `direct_light`. fn sky_mis(d: vec3, pdf: f32) -> vec3 { + // three.js scene.background: what camera rays (and paths bent straight through glass) see where nothing is. + @if(background) if (pdf <= 0.0) { return P.background.rgb; } var wgt = 1.0; if (pdf > 0.0 && P.counts2.y != 0u) { wgt = power_heuristic(pdf, env_pdf(d)); @@ -1846,11 +1890,35 @@ fn direct_light(s: Surface, lb: Lobes, v: vec3) -> vec3 { dl_det += sum; dl_det_spec += dl_spec; } - sum += local_lights(s, lb, v, o, s.p, false, vec3(0.0)); + // gpu.rs `KernelSwitches::light_pick`: past the camera hit, the directional lights and one of the other kinds a + // scene has (local lights, emitters, the sky), weighted by their number. + @if(light_pick) let has_local = P.counts.x > P.counts2.x; + @if(light_pick) let has_emit = P.counts.y > 0u; + @if(light_pick) let has_env = P.counts2.y != 0u && (P.env.x > 0.5 || luminance(P.sky_color.rgb) > 0.0); + @if(light_pick) var take_local = true; + @if(light_pick) var take_emit = true; + @if(light_pick) var take_env = true; + @if(light_pick) var pick_w = 1.0; + @if(light_pick) let kinds = u32(has_local) + u32(has_emit) + u32(has_env); + @if(light_pick) if (cur_depth > 0u && kinds > 1u) { + var k = min(u32(rnd() * f32(kinds)), kinds - 1u); + take_local = has_local && k == 0u; + if (has_local) { k -= 1u; } + take_emit = has_emit && k == 0u; + if (has_emit) { k -= 1u; } + take_env = has_env && k == 0u; + pick_w = f32(kinds); + } + @if(light_pick) let sum_sun = sum; + @if(light_pick) let spec_sun = dl_spec; + @if(!light_pick) sum += local_lights(s, lb, v, o, s.p, false, vec3(0.0)); + @if(light_pick) if (take_local) { sum += local_lights(s, lb, v, o, s.p, false, vec3(0.0)); } // Emissive triangles (incl. rect lights) mis_x = o; mis_n = s.n; - if (P.counts.y > 0u) { + var emitters = P.counts.y; + @if(light_pick) if (!take_emit) { emitters = 0u; } + if (emitters > 0u) { var es: ES; // With the light tree its candidates are already near the point: P.ltree.y of them at camera hits, one elsewhere. if (lt_on()) { es = sample_emitter_ris(o, s.n, select(1u, u32(P.ltree.w), cur_depth == 0u), s, lb, v); } else if (seen_directly()) { es = sample_emitter_ris(o, s.n, select(EMIT_RIS_M / 2u, EMIT_RIS_M, cur_depth == 0u), s, lb, v); } else if (P.light_lod.w > 0.0) { es = sample_emitter_ris(o, s.n, 2u, s, lb, v); } else { es = sample_emitter(o, s.n); } @@ -1865,7 +1933,9 @@ fn direct_light(s: Surface, lb: Lobes, v: vec3) -> vec3 { } } // Environment - if (P.counts2.y != 0u && (P.env.x > 0.5 || luminance(P.sky_color.rgb) > 0.0)) { + var env_on = P.counts2.y != 0u && (P.env.x > 0.5 || luminance(P.sky_color.rgb) > 0.0); + @if(light_pick) env_on = take_env && has_env; + if (env_on) { let l = sample_env(); let f = eval_bsdf(s, lb, v, l); if (any(f > vec3(0.0))) { @@ -1878,6 +1948,10 @@ fn direct_light(s: Surface, lb: Lobes, v: vec3) -> vec3 { } } } + @if(light_pick) if (pick_w > 1.0) { + sum = sum_sun + (sum - sum_sun) * pick_w; + dl_spec = spec_sun + (dl_spec - spec_sun) * pick_w; + } return sum; } @@ -1989,6 +2063,7 @@ struct VIn { lp: vec3, // position in the volume's own space: the box [-0.5, 0.5]^3 / the sphere of radius 0.5 time: f32, // scene time, s custom: vec4, // the node's customData + @if(grid) grid: vec2, // the node's density grid: its first layer in the atlas, its cells (10 bits an axis; 0: none) }; struct VOut { @@ -1997,6 +2072,15 @@ struct VOut { emission: vec3, // emitted radiance per metre (absorption x emitted radiance) }; +// The node's density grid at `lp` (the volume's own space, the box [-0.5, 0.5]^3): field A's four numbers, trilinear +// between cell centres, the border cells' beyond them; zeros for a node without a grid. +@if(grid) fn vgrid(i: VIn, lp: vec3) -> vec4 { + if (i.grid.y == 0u) { return vec4(0.0); } + let n = vec3(vec3(i.grid.y & 1023u, (i.grid.y >> 10u) & 1023u, (i.grid.y >> 20u) & 1023u)); + let c = clamp((lp + 0.5) * n, vec3(0.5), n - 0.5) + vec3(0.0, 0.0, f32(i.grid.x)); + return textureSampleLevel(grid_tex, samp, c / vec3(textureDimensions(grid_tex)), 0.0); +} + // @volume-functions-begin fn volume_function(f: u32, i: VIn, p: array, 8>) -> VOut { return VOut(0.0, vec3(1.0), vec3(0.0)); } // @volume-functions-end @@ -2034,6 +2118,7 @@ fn medium_eval(k: u32, x: vec3) -> VOut { i.lp = medium_local(m, x); i.time = P.time.x; i.custom = m.c; + @if(grid) i.grid = vec2(bitcast(m.b.z), bitcast(m.b.w)); var r = volume_function(m.a.x, i, m.p); r.extinction = clamp(r.extinction, 0.0, m.b.x); return r; @@ -3068,7 +3153,17 @@ fn main(@builtin(global_invocation_id) gid: vec3) { radiance += c; } if (s.card) { break; } - if (depth == P.counts.z) { break; } + @if(!last_direct) if (depth == P.counts.z) { break; } + @if(last_direct) if (depth == P.counts.z) { + // gpu.rs `KernelSwitches::last_direct`: the last vertex's direct light. + if (depth > 0u) { + var dz = thr * direct_light(s, lobes(s, max(dot(s.n, -d), 1e-4)), -d); + let lum = luminance(dz); + if (lum > clamp_v && !straight) { dz *= clamp_v / lum; } + radiance += dz; + } + break; + } // Live: the path ends in the radiance cache at its first bounce. Surfaces seen in mirror-like reflections // (specular lobe below roughness 0.4, where UE traces reflections) get exact direct light, as UE's hit // lighting does, and the cache's indirect part; others reflect the cache's total. Rotating tiles train the diff --git a/renderer/shaders/trace_hw.wgsl b/renderer/shaders/trace_hw.wgsl index 5f5c9c0..fd3ee89 100644 --- a/renderer/shaders/trace_hw.wgsl +++ b/renderer/shaders/trace_hw.wgsl @@ -2,9 +2,9 @@ // mask 2; opaque groups never report candidates, alpha-cutout groups are confirmed by `alpha_pass`. // instance_custom_data holds the instance's first global triangle index, or in scenes where a group starts beyond // 2^24 triangles (F_HUGE: custom data has 24 bits) its traversal group, whose first triangle is in P.group_first. -@group(0) @binding(13) var acc: acceleration_structure; +@if(hwrt) @group(0) @binding(13) var acc: acceleration_structure; -fn trace(o: vec3, d: vec3, tmax_in: f32, mode: u32) -> Hit { +@if(hwrt) fn trace(o: vec3, d: vec3, tmax_in: f32, mode: u32) -> Hit { // Ray tracing off (pathtrace.wgsl `rays_off`): past the camera ray, every ray misses. if (rays_off) { return Hit(tmax_in, 0xffffffffu, 0.0, 0.0, NO_INST); } let mask = select(0xffu, 1u, mode == TRACE_SHADOW); @@ -14,7 +14,8 @@ fn trace(o: vec3, d: vec3, tmax_in: f32, mode: u32) -> Hit { } // Live secondary rays, and live emitter shadow rays (TRACE_ANY_OPAQUE): cutouts opaque (pathtrace.wgsl // `cutouts_opaque`, `occluded_emitter`). - if (cutouts_opaque() || mode == TRACE_ANY_OPAQUE) { return trace_opaque(o, d, tmax_in, mask, mode != TRACE_CLOSEST); } + @if(!cutout_lite) if (cutouts_opaque() || mode == TRACE_ANY_OPAQUE) { return trace_opaque(o, d, tmax_in, mask, mode != TRACE_CLOSEST); } + @if(cutout_lite) if (cutouts_opaque() || mode != TRACE_CLOSEST) { return trace_opaque(o, d, tmax_in, mask, mode != TRACE_CLOSEST); } // Alpha-cutout triangles: every query forces opacity, and a cutout hit that fails its alpha test is stepped // through. The ray query's own candidate loop (non-opaque traversal) lost hits on Metal: ~0.1 % of rays near // cutout foliage missed even the opaque ground behind it or found it unlit (black pixels on the ground; the @@ -36,7 +37,8 @@ fn trace(o: vec3, d: vec3, tmax_in: f32, mode: u32) -> Hit { // In order, each query starting past the last cutout hit its alpha let through. Shadow rays take up to 64 of them // and count as blocked past that: with 16, as closest hits take, the rest counted as unblocked (light through dense // foliage: Bistro's offline render up to 2.3 % bright in a region against the software BVH). - let steps = select(16u, 64u, mode != TRACE_CLOSEST); + @if(!cutout_lite2) let steps = select(16u, 64u, mode != TRACE_CLOSEST); + @if(cutout_lite2) let steps = 4u; var t0 = 0.0; for (var k = 0u; k < steps; k++) { let h = trace_opaque(o + d * t0, d, tmax_in - t0, mask, false); @@ -51,7 +53,8 @@ fn trace(o: vec3, d: vec3, tmax_in: f32, mode: u32) -> Hit { // One query, not forced opaque; each cutout candidate that passes its alpha test is confirmed, which ends it (as does any // opaque triangle). Live only (see `trace`). -fn trace_any_cutout(o: vec3, d: vec3, tmax_in: f32, mask: u32) -> Hit { +@if(hwrt) fn trace_any_cutout(o: vec3, d: vec3, tmax_in: f32, mask: u32) -> Hit { + @if(strict) if (ray_bad(o, d, tmax_in)) { return Hit(tmax_in, 0xffffffffu, 0.0, 0.0, NO_INST); } var rq: ray_query; rayQueryInitialize(&rq, acc, RayDesc(RAY_FLAG_TERMINATE_ON_FIRST_HIT, mask, 0.0, tmax_in, o, d)); var pg = GOut(-1.0, vec3(0.0, 1.0, 0.0), 0u, vec2(0.0)); @@ -80,7 +83,8 @@ fn trace_any_cutout(o: vec3, d: vec3, tmax_in: f32, mask: u32) -> Hit } // The closest (or with `any`, the first found) hit, every triangle opaque. -fn trace_opaque(o: vec3, d: vec3, tmax_in: f32, mask: u32, any: bool) -> Hit { +@if(hwrt) fn trace_opaque(o: vec3, d: vec3, tmax_in: f32, mask: u32, any: bool) -> Hit { + @if(strict) if (ray_bad(o, d, tmax_in)) { return Hit(tmax_in, 0xffffffffu, 0.0, 0.0, NO_INST); } var rq: ray_query; let flags = select(RAY_FLAG_FORCE_OPAQUE, RAY_FLAG_FORCE_OPAQUE | RAY_FLAG_TERMINATE_ON_FIRST_HIT, any); rayQueryInitialize(&rq, acc, RayDesc(flags, mask, 0.0, tmax_in, o, d)); @@ -120,20 +124,20 @@ fn trace_opaque(o: vec3, d: vec3, tmax_in: f32, mask: u32, any: bool) // An alpha-cutout triangle: its material has an alpha test (as the build flags TRI_MASKED; model instances' pool // triangles carry no flags). -fn cutout(ti: u32) -> bool { +@if(hwrt) fn cutout(ti: u32) -> bool { return materials[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. -fn hit_inst(custom: u32) -> u32 { +@if(hwrt) fn hit_inst(custom: u32) -> u32 { return select(NO_INST, custom & (INSTANCE_BIT - 1u), has(F_INST) && (custom & INSTANCE_BIT) != 0u); } -const INSTANCE_BIT: u32 = 0x800000u; +@if(hwrt) const INSTANCE_BIT: u32 = 0x800000u; // Global triangle of a hit: a model instance's pool triangle (its cutout geometry 1 after the opaque one), or a // traversal group's. -fn hit_tri(custom: u32, geometry: u32, prim: u32) -> u32 { +@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; return r.x + select(0u, r.w, geometry == 1u) + prim; @@ -141,7 +145,13 @@ fn hit_tri(custom: u32, geometry: u32, prim: u32) -> u32 { return group_tri(custom, prim); } -fn group_tri(custom: u32, prim: u32) -> u32 { +@if(hwrt) fn group_tri(custom: u32, prim: u32) -> u32 { if (has(F_HUGE)) { return P.group_first[custom >> 2u][custom & 3u] + prim; } return custom + prim; } + +// gpu.rs `KernelSwitches::strict`: a ray a ray query must not be given (not finite, or a length that is negative or NaN). +@if(hwrt && strict) fn ray_bad(o: vec3, d: vec3, tmax: f32) -> bool { + let tb = bitcast(tmax); + return non_finite(o) || non_finite(d) || (tb & 0x80000000u) != 0u || (tb & 0x7fffffffu) > 0x7f800000u; +} diff --git a/renderer/shaders/trace_units.wgsl b/renderer/shaders/trace_units.wgsl index a82789b..eba43c2 100644 --- a/renderer/shaders/trace_units.wgsl +++ b/renderer/shaders/trace_units.wgsl @@ -1,6 +1,6 @@ // Units: software traversal's second level (gpu.rs F_UNITS, renderer/src/units.rs; docs/RUNTIME.md "Software traversal -// with units"). Spliced into pathtrace.wgsl's software traversal by `kernel_source` for scenes that have units; a scene -// without them, and hardware traversal, never see this file. +// with units"). Part of the kernel under `@if(units)` (pathtrace.wgsl's head): a scene with units under software +// traversal has it; a scene without them, and hardware traversal, do not. // // A unit is a model instance (`instances`, as hardware traversal's): its mesh's triangles in model space in // `tri_verts`, from the record's first triangle; its mesh's tree in `nodes`, from the root in the record's `n0.w` @@ -14,7 +14,7 @@ // `trace` walks the still scene first and calls this with its hit: the ray is already cut to it. // // The stack is the module's, one for an invocation of the kernel, and in a kernel with units the still walk uses it -// too (gpu.rs `kernel_source` takes its own away). An array a function declares is set to zero each time the +// too (`@if(!units)` takes its own away). An array a function declares is set to zero each time the // function is entered, whatever the ray then meets, and a ray's walk is entered many times a pixel: the unit walk's // own 64 words cost hello 0.3 ms a frame and the city 0.75, and the still walk's 48 cost every scene as much again, // flattened ones too. This one is set to zero once. A walk reads only what it has pushed: what an earlier ray left @@ -33,9 +33,9 @@ // // The root's box is tested before the walk: a ray that meets no unit reads one node, and a scene of one unit is a // root that is a leaf, whose record only the rays that meet its box read. -var stack: array; +@if(units) var stack: array; -fn trace_units(o: vec3, d: vec3, inv_d: vec3, hit_in: Hit, mode: u32) -> Hit { +@if(units) fn trace_units(o: vec3, d: vec3, inv_d: vec3, hit_in: Hit, mode: u32) -> Hit { var hit = hit_in; if (bitcast(P.motion.w) == 0u) { return hit; } let top = bitcast(P.motion.z); diff --git a/renderer/shaders/translucent.wgsl b/renderer/shaders/translucent.wgsl index 0bfe0fb..cd3e579 100644 --- a/renderer/shaders/translucent.wgsl +++ b/renderer/shaders/translucent.wgsl @@ -14,29 +14,29 @@ // Live, it runs on the composed HDR image after the temporal filter (no history: nothing to leave trails); offline it // accumulates (colour, transmittance) per sample, laid over the converged image. -struct TP { +@if(hwrt) struct TP { size: vec4, // output width, height, first translucent triangle, mode (0 offline: accumulate; 1 live: over the image) view: vec4, // live: traced width, height (the opaque first hits); training budget (points per frame); 1 = at the traced resolution misc: vec4, // live: share of hit pixels that train the volume voxels, 1 = train only (THREERT_TL_HIDE: what is behind), -, - }; -@group(2) @binding(0) var tacc: acceleration_structure; -@group(2) @binding(1) var TPS: TP; +@if(hwrt) @group(2) @binding(0) var tacc: acceleration_structure; +@if(hwrt) @group(2) @binding(1) var TPS: TP; // Live: the composed image (in place; w kept). Offline: per-pixel sums of (colour, transmittance). -@group(2) @binding(2) var timg: array>; +@if(hwrt) @group(2) @binding(2) var timg: array>; // Volume training points of this frame (live) and their count. -@group(2) @binding(3) var vpts: array>; -@group(2) @binding(4) var vcount: array>; +@if(hwrt) @group(2) @binding(3) var vpts: array>; +@if(hwrt) @group(2) @binding(4) var vcount: array>; // Live at a reduced render scale: (colour, transmittance) per traced pixel, laid over the image by translucent_up. -@group(2) @binding(5) var tlow: array>; +@if(hwrt) @group(2) @binding(5) var tlow: array>; -const TL_LAYERS: u32 = 4u; +@if(hwrt) const TL_LAYERS: u32 = 4u; // Volume voxels (cache_point, point_radiance): pathtrace.wgsl, shared with participating media. // Live: the nearest opaque first hit around output pixel (x, y): of the 2 x 2 traced pixels under it, the nearest // (sky = none), so nothing shows through at an edge. -fn opaque_depth(x: u32, y: u32) -> f32 { +@if(hwrt) fn opaque_depth(x: u32, y: u32) -> f32 { let iw = TPS.view.x; let ih = TPS.view.y; let fx = (f32(x) + 0.5) * f32(iw) / f32(TPS.size.x) - 0.5; @@ -56,7 +56,7 @@ fn opaque_depth(x: u32, y: u32) -> f32 { // Material of a translucent particle at triangle ti (barycentrics b): colour, opacity and emission after maps and its // material function. -fn particle_material(ti: u32, b: vec2, view: vec3, lod: f32) -> MatOut { +@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 uv = mesh_uv(a, m, b.x, b.y); @@ -74,7 +74,7 @@ fn particle_material(ti: u32, b: vec2, view: vec3, lod: f32) -> MatOut return r; } -@compute @workgroup_size(8, 8, 1) +@if(hwrt) @compute @workgroup_size(8, 8, 1) fn translucent(@builtin(global_invocation_id) gid: vec3) { let live = TPS.size.w == 1u; // Live below full resolution the particles are drawn per traced pixel (as UE draws translucency at the screen @@ -113,7 +113,9 @@ fn translucent(@builtin(global_invocation_id) gid: vec3) { var lb: array, TL_LAYERS>; var n = 0u; var rq: ray_query; - rayQueryInitialize(&rq, tacc, RayDesc(0u, 0xffu, 0.0, tmax, o, d)); + @if(!strict) rayQueryInitialize(&rq, tacc, RayDesc(0u, 0xffu, 0.0, tmax, o, d)); + @if(strict) let tbad = ray_bad(o, d, tmax); + @if(strict) rayQueryInitialize(&rq, tacc, RayDesc(0u, 0xffu, 0.0, select(tmax, 0.0, tbad), select(o, vec3(0.0), tbad), select(d, vec3(0.0, 0.0, 1.0), tbad))); while (rayQueryProceed(&rq)) { let c = rayQueryGetCandidateIntersection(&rq); if (n == TL_LAYERS && c.t >= lt[TL_LAYERS - 1u]) { continue; } @@ -174,7 +176,7 @@ fn translucent(@builtin(global_invocation_id) gid: vec3) { // Live, after `translucent` at the traced resolution: its (colour, transmittance) bilinearly up to each output pixel, // over the image. -@compute @workgroup_size(8, 8, 1) +@if(hwrt) @compute @workgroup_size(8, 8, 1) fn translucent_up(@builtin(global_invocation_id) gid: vec3) { let w = TPS.size.x; let h = TPS.size.y; @@ -197,7 +199,7 @@ fn translucent_up(@builtin(global_invocation_id) gid: vec3) { // Live, after `translucent`: one fresh estimate per training point into its volume voxels (both levels). The points // are capped per frame (TPS.view.z), so the cost is bounded whatever the particles cover. -@compute @workgroup_size(64, 1, 1) +@if(hwrt) @compute @workgroup_size(64, 1, 1) fn volume_train(@builtin(global_invocation_id) gid: vec3) { let k = gid.x; if (k >= min(atomicLoad(&vcount[0]), TPS.view.z)) { return; } diff --git a/renderer/src/gpu.rs b/renderer/src/gpu.rs index 214bd02..e482d25 100644 --- a/renderer/src/gpu.rs +++ b/renderer/src/gpu.rs @@ -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 source-patching features: `kernel_source`, the merged layout's lengths). + /// By (material-function code hash, the features that change the source: `kernel_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_source`): every other kernel's `Params` ends before it. + /// (`kernel_flags`): every other kernel's `Params` ends before it. pub background: [f32; 4], } @@ -698,81 +698,66 @@ pub fn instance_mask(cast_shadow: bool, hidden: bool) -> u8 { } } -/// The kernel source for features that change it rather than switch code in it (`F_SOURCE`). Kept out of -/// pathtrace.wgsl so that every scene without them runs the file as it is; each anchor must be there once. -/// - 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`). -/// - 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 (it is in no emitter list). -/// - 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. -/// - 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. -/// - F_VCOLOR, three.js `vertexColors` (`scene::tri_colors`): three sRGB 8-bit colours per triangle after the +/// The kernel's files, one module in this order (pathtrace.wgsl's head says what each holds and names the `@if` flags). +const KERNEL_FILES: [&str; 4] = [ + include_str!("../shaders/pathtrace.wgsl"), + include_str!("../shaders/trace_units.wgsl"), + include_str!("../shaders/trace_hw.wgsl"), + include_str!("../shaders/translucent.wgsl"), +]; + +/// The kernel's source: its files as one WESL module, every device's and scene's variant in it under `@if`. +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). -pub fn kernel_source(src: &str, features: u32) -> String { - let mut s = src.to_string(); - let mut patch = |anchor: &str, with: &str| { - assert_eq!(s.matches(anchor).count(), 1, "kernel_source: {anchor:?}"); - s = s.replace(anchor, with); - }; - if features & F_VCOLOR != 0 { - patch( - "@group(0) @binding(18) var media: array;\n", - "@group(0) @binding(18) var media: array;\n// Vertex colours (three.js `vertexColors`; gpu.rs F_VCOLOR): three sRGB 8-bit colours per triangle, after the previous\n// vertices in `prev_verts` from vec4 P.procs.w.\nfn vcolor_word(k: u32) -> u32 { return bitcast(prev_verts[P.procs.w + k / 4u][k % 4u]); }\nfn vcolor(t: u32, u: f32, v: f32) -> vec3 {\n let c0 = srgb_to_linear(unpack4x8unorm(vcolor_word(t * 3u)).rgb);\n let c1 = srgb_to_linear(unpack4x8unorm(vcolor_word(t * 3u + 1u)).rgb);\n let c2 = srgb_to_linear(unpack4x8unorm(vcolor_word(t * 3u + 2u)).rgb);\n return c0 * (1.0 - u - v) + c1 * u + c2 * v;\n}\n", - ); - patch(" s.card = m.tex1.z == MAT_EMITTER_CARD;\n var base = m.base_color.rgb;\n", " s.card = m.tex1.z == MAT_EMITTER_CARD;\n var base = m.base_color.rgb * vcolor(i, h.u, h.v);\n"); - patch(" g.tint = m.base_color.rgb;\n", " g.tint = m.base_color.rgb * vcolor(h.tri, h.u, h.v);\n"); - } - if features & F_GRID != 0 { - patch( - "@group(0) @binding(18) var media: array;\n", - "@group(0) @binding(18) var media: array;\n// Density grids (gpu.rs F_GRID, `GridsGpu`): field A of every grid, stacked along z.\n@group(0) @binding(19) var grid_tex: texture_3d;\n", - ); - patch( - " custom: vec4, // the node's customData\n};\n\nstruct VOut {\n", - " custom: vec4, // the node's customData\n grid: vec2, // the node's density grid: its first layer in the atlas, its cells (10 bits an axis; 0: none)\n};\n\nstruct VOut {\n", - ); - patch( - "// @volume-functions-begin\n", - "// The node's density grid at `lp` (the volume's own space, the box [-0.5, 0.5]^3): field A's four numbers, trilinear\n// between cell centres, the border cells' beyond them; zeros for a node without a grid.\nfn vgrid(i: VIn, lp: vec3) -> vec4 {\n if (i.grid.y == 0u) { return vec4(0.0); }\n let n = vec3(vec3(i.grid.y & 1023u, (i.grid.y >> 10u) & 1023u, (i.grid.y >> 20u) & 1023u));\n let c = clamp((lp + 0.5) * n, vec3(0.5), n - 0.5) + vec3(0.0, 0.0, f32(i.grid.x));\n return textureSampleLevel(grid_tex, samp, c / vec3(textureDimensions(grid_tex)), 0.0);\n}\n\n// @volume-functions-begin\n", - ); - patch(" i.custom = m.c;\n var r = volume_function(m.a.x, i, m.p);\n", " i.custom = m.c;\n i.grid = vec2(bitcast(m.b.z), bitcast(m.b.w));\n var r = volume_function(m.a.x, i, m.p);\n"); - } - if features & F_TINT != 0 { - patch( - "const OBJECT_STRIDE: u32 = 3u;\n", - "const OBJECT_STRIDE: u32 = 4u;\n// An instanced object's instance colour (three.js `instanceColor`), white for every other object.\nfn obj_tint(oi: u32) -> vec3 { return objects[oi * OBJECT_STRIDE + 3u].rgb; }\n", - ); - // (Anchored before the line's end: with F_VCOLOR the vertex colour follows.) - patch(" s.card = m.tex1.z == MAT_EMITTER_CARD;\n var base = m.base_color.rgb", " s.card = m.tex1.z == MAT_EMITTER_CARD;\n var base = m.base_color.rgb * obj_tint(object_of(a))"); - patch(" s.card = false;\n var base = m.base_color.rgb;\n", " s.card = false;\n var base = m.base_color.rgb * obj_tint(oi);\n"); - patch(" g.tint = m.base_color.rgb", " g.tint = m.base_color.rgb * obj_tint(object_of(a))"); - } - if features & F_UNLIT != 0 { - patch("const MAT_EMITTER_CARD: i32 = 2;\n", "const MAT_EMITTER_CARD: i32 = 2;\nconst MAT_UNLIT: i32 = 4;\n"); - patch( - " s.emissive = emission(m, uv, lod) * obj_power(a);\n", - " s.emissive = emission(m, uv, lod) * obj_power(a);\n // three.js MeshBasicMaterial: its colour to camera rays and mirror-like paths, lit by nothing, lighting nothing.\n if (m.tex1.z == MAT_UNLIT) {\n s.emissive = select(vec3(0.0), base, cur_depth == 0u || mirror_path);\n base = vec3(0.0);\n spec = 0.0;\n metal = 0.0;\n }\n", - ); - } - if features & F_UNITS != 0 { - // The still scene's walk ends by handing its hit to the units' (then to the procedural objects', as before). - // And it has no stack of its own here: the name is then the module's, which the units' file declares, set to - // zero once an invocation and not once a ray (shaders/trace_units.wgsl says what that is worth). - patch(STILL_STACK, STILL_STACK_GONE); - patch( - " if (has(F_PROC)) { return trace_procs(o, d, hit, mode); }\n return hit;\n}\n", - " hit = trace_units(o, d, inv_d, hit, mode);\n if (has(F_PROC)) { return trace_procs(o, d, hit, mode); }\n return hit;\n}\n", - ); - patch("\n// @trace-end\n", &format!("\n{}\n// @trace-end\n", include_str!("../shaders/trace_units.wgsl"))); - } - if features & F_BACKGROUND != 0 { - patch(" aerial: vec4, // the sky's aerial perspective (F_AERIAL): density multiple, Rayleigh scale, Mie scale, -\n};\n", " aerial: vec4, // the sky's aerial perspective (F_AERIAL): density multiple, Rayleigh scale, Mie scale, -\n background: vec4, // three.js scene.background (radiance), 1\n};\n"); - patch("fn sky_mis(d: vec3, pdf: f32) -> vec3 {\n", "fn sky_mis(d: vec3, pdf: f32) -> vec3 {\n // three.js scene.background: what camera rays (and paths bent straight through glass) see where nothing is.\n if (pdf <= 0.0) { return P.background.rgb; }\n"); - } - s +/// - 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)); + 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() } /// The loop guard's limit for every kernel, if one is set (merged.rs `guard`: every loop of a kernel bounded, so a @@ -782,11 +767,6 @@ pub fn loop_guard() -> Option { loop_guard_of(false) } -/// pathtrace.wgsl's `trace` declares its stack (software traversal); in a kernel with units it does not -/// (`kernel_source`). -const STILL_STACK: &str = "\n var stack: array;\n var sp = 0u;\n var node = 0u;\n"; -const STILL_STACK_GONE: &str = "\n var sp = 0u;\n var node = 0u;\n"; - /// The loop guard's limit for a kernel with units or without (`F_UNITS`). A kernel with units has the guard natively /// too, at the browser's limit, unless THREERT_LOOP_GUARD=0: its second level is the traversal that has run least, /// and what it walks is rewritten by edits while it runs. A kernel without units is as `loop_guard` says: natively, @@ -862,8 +842,9 @@ 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 patch of pathtrace.wgsl, trace_hw.wgsl or translucent.wgsl (`kernel_switches`, as `kernel_source` patches -/// for a scene's features); with none set the source is the files, which is every platform's default but Android's. +/// is a flag of the kernel's conditional translation (`kernel_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)] pub struct KernelSwitches { /// THREERT_KERNEL_STRICT=1|0 (Android: 1). The kernel gives the driver nothing undefined to act on: @@ -900,68 +881,6 @@ impl KernelSwitches { } } -/// The kernel's source (`kernel_base`) with a device's switches. Each anchor must be there once; hardware traversal's -/// are looked for only in its source. -fn kernel_switches(src: String, k: &KernelSwitches, hwrt: bool) -> String { - let mut s = src; - let mut patch = |anchor: &str, with: &str| { - assert_eq!(s.matches(anchor).count(), 1, "kernel_switches: {anchor:?}"); - s = s.replace(anchor, with); - }; - if k.strict { - patch(" let row_p = env_cdf[y] - select(0.0, env_cdf[y - 1u], y > 0u);\n", " let row_p = env_cdf[y] - select(0.0, env_cdf[max(y, 1u) - 1u], y > 0u);\n"); - patch(" let col_p = env_cdf[base + x] - select(0.0, env_cdf[base + x - 1u], x > 0u);\n", " let col_p = env_cdf[base + x] - select(0.0, env_cdf[base + max(x, 1u) - 1u], x > 0u);\n"); - } - if k.strict && hwrt { - let miss = " if (ray_bad(o, d, tmax_in)) { return Hit(tmax_in, 0xffffffffu, 0.0, 0.0, NO_INST); }\n"; - let any = " var rq: ray_query;\n rayQueryInitialize(&rq, acc, RayDesc(RAY_FLAG_TERMINATE_ON_FIRST_HIT, mask, 0.0, tmax_in, o, d));\n"; - patch(any, &format!("{miss}{any}")); - let opaque = " var rq: ray_query;\n let flags = select(RAY_FLAG_FORCE_OPAQUE, RAY_FLAG_FORCE_OPAQUE | RAY_FLAG_TERMINATE_ON_FIRST_HIT, any);\n"; - patch(opaque, &format!("{miss}{opaque}")); - // The translucency pass's query (its camera ray): a ray of no length in its place. - patch( - " rayQueryInitialize(&rq, tacc, RayDesc(0u, 0xffu, 0.0, tmax, o, d));\n", - " let tbad = ray_bad(o, d, tmax);\n rayQueryInitialize(&rq, tacc, RayDesc(0u, 0xffu, 0.0, select(tmax, 0.0, tbad), select(o, vec3(0.0), tbad), select(d, vec3(0.0, 0.0, 1.0), tbad)));\n", - ); - s += "\n// gpu.rs `KernelSwitches::strict`: a ray a ray query must not be given (not finite, or a length that is negative or NaN).\nfn ray_bad(o: vec3, d: vec3, tmax: f32) -> bool {\n let tb = bitcast(tmax);\n return non_finite(o) || non_finite(d) || (tb & 0x80000000u) != 0u || (tb & 0x7fffffffu) > 0x7f800000u;\n}\n"; - } - let mut patch = |anchor: &str, with: &str| { - assert_eq!(s.matches(anchor).count(), 1, "kernel_switches: {anchor:?}"); - s = s.replace(anchor, with); - }; - if k.cutout_lite >= 1 && hwrt { - patch( - " if (cutouts_opaque() || mode == TRACE_ANY_OPAQUE) { return trace_opaque(o, d, tmax_in, mask, mode != TRACE_CLOSEST); }\n", - " if (cutouts_opaque() || mode != TRACE_CLOSEST) { return trace_opaque(o, d, tmax_in, mask, mode != TRACE_CLOSEST); }\n", - ); - } - if k.cutout_lite >= 2 && hwrt { - patch(" let steps = select(16u, 64u, mode != TRACE_CLOSEST);\n", " let steps = 4u;\n"); - } - if k.last_direct { - patch( - " if (s.card) { break; }\n if (depth == P.counts.z) { break; }\n", - " if (s.card) { break; }\n if (depth == P.counts.z) {\n // gpu.rs `KernelSwitches::last_direct`: the last vertex's direct light.\n if (depth > 0u) {\n var dz = thr * direct_light(s, lobes(s, max(dot(s.n, -d), 1e-4)), -d);\n let lum = luminance(dz);\n if (lum > clamp_v && !straight) { dz *= clamp_v / lum; }\n radiance += dz;\n }\n break;\n }\n", - ); - } - if k.light_pick { - // One of the three kinds of light below the directional ones, drawn among those the scene has. - patch( - " sum += local_lights(s, lb, v, o, s.p, false, vec3(0.0));\n // Emissive triangles (incl. rect lights)\n mis_x = o;\n mis_n = s.n;\n if (P.counts.y > 0u) {\n", - " // gpu.rs `KernelSwitches::light_pick`.\n let has_local = P.counts.x > P.counts2.x;\n let has_emit = P.counts.y > 0u;\n let has_env = P.counts2.y != 0u && (P.env.x > 0.5 || luminance(P.sky_color.rgb) > 0.0);\n var take_local = true;\n var take_emit = true;\n var take_env = true;\n var pick_w = 1.0;\n let kinds = u32(has_local) + u32(has_emit) + u32(has_env);\n if (cur_depth > 0u && kinds > 1u) {\n var k = min(u32(rnd() * f32(kinds)), kinds - 1u);\n take_local = has_local && k == 0u;\n if (has_local) { k -= 1u; }\n take_emit = has_emit && k == 0u;\n if (has_emit) { k -= 1u; }\n take_env = has_env && k == 0u;\n pick_w = f32(kinds);\n }\n let sum_sun = sum;\n let spec_sun = dl_spec;\n if (take_local) { sum += local_lights(s, lb, v, o, s.p, false, vec3(0.0)); }\n // Emissive triangles (incl. rect lights)\n mis_x = o;\n mis_n = s.n;\n if (!take_emit) {\n } else if (P.counts.y > 0u) {\n", - ); - patch( - " // Environment\n if (P.counts2.y != 0u && (P.env.x > 0.5 || luminance(P.sky_color.rgb) > 0.0)) {\n let l = sample_env();\n let f = eval_bsdf(s, lb, v, l);\n", - " // Environment\n if (take_env && has_env) {\n let l = sample_env();\n let f = eval_bsdf(s, lb, v, l);\n", - ); - patch( - " dl_spec += env_radiance(l) * f * wgt / pdf_e * gtint(shadow_tint) * frac;\n }\n }\n }\n return sum;\n}\n", - " dl_spec += env_radiance(l) * f * wgt / pdf_e * gtint(shadow_tint) * frac;\n }\n }\n }\n if (pick_w > 1.0) {\n sum = sum_sun + (sum - sum_sun) * pick_w;\n dl_spec = spec_sun + (dl_spec - spec_sun) * pick_w;\n }\n return sum;\n}\n", - ); - } - s -} - /// The kernel's runtime checks: none (all indices come from our own BVH / scene builder). THREERT_CHECKED turns naga's /// on to diagnose a driver: 1 = every check, else a mask (2 bounds, 4 loop bounding, 8 ray-query initialization /// tracking). @@ -977,19 +896,6 @@ fn kernel_checks() -> wgpu::ShaderRuntimeChecks { c } -/// The kernel's source for a traversal: pathtrace.wgsl as it is (software), or with the hardware traversal in the -/// software one's place and the translucency pass after it. -pub fn kernel_base(hwrt: bool) -> String { - let base = include_str!("../shaders/pathtrace.wgsl"); - if hwrt { - let (head, rest) = base.split_once("// @trace-begin").unwrap(); - let (_, tail) = rest.split_once("// @trace-end").unwrap(); - format!("enable wgpu_ray_query;\n{head}{}{tail}{}", include_str!("../shaders/trace_hw.wgsl"), include_str!("../shaders/translucent.wgsl")) - } else { - base.to_string() - } -} - /// 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`). @@ -1269,14 +1175,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_source`; +/// Instance colours (`InstancedMesh.setColorAt`): object entries of 4 vec4s, the base colour times the 4th (`kernel_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_source`). +/// three.js MeshBasicMaterial (`scene::MAT_UNLIT`): seen by camera rays, lit by nothing, lighting nothing (`kernel_flags`). pub const F_UNLIT: u32 = 1 << 20; -/// three.js `scene.background`: camera rays that miss see `Params::background` (`kernel_source`). +/// three.js `scene.background`: camera rays that miss see `Params::background` (`kernel_flags`). pub const F_BACKGROUND: u32 = 1 << 21; -/// The features that patch the kernel's source (`kernel_source`) rather than switch code in it. +/// The features that change the kernel's declarations (`kernel_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 @@ -1474,14 +1380,12 @@ impl Gpu { let pipeline_cache = cache_feature.and_then(|_| crate::pipeline_cache::PipelineCache::open(&device, options.pipeline_cache.as_ref()?, &adapter_info)); let uniform_tables = options.uniform_tables_on(merged, cfg!(target_os = "android")); let switches = options.kernel; - let src = kernel_switches(kernel_base(hwrt), &switches, hwrt); + 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() }; - device.create_shader_module_trusted( - wgpu::ShaderModuleDescriptor { label: Some("pathtrace"), source: wgpu::ShaderSource::Wgsl(guarded(&src, false, false).into()) }, - kernel_checks(), - ) + let wgsl = kernel_wgsl(&guarded(&src, false, false), &kernel_flags(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 { binding, @@ -1672,19 +1576,15 @@ impl Gpu { let (_, tail) = rest.split_once(&format!("// @{marker}-end")).unwrap(); format!("{head}{code}{tail}") }; - let base = kernel_source(&self.source, patches); - let src = splice(&splice(&splice(&base, "material-functions", &k.functions), "volume-functions", &k.volume_functions), "geometry-functions", &k.geometry_functions); + 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 }; - unsafe { - self.device.create_shader_module_trusted( - wgpu::ShaderModuleDescriptor { label: Some("pathtrace+functions"), source: wgpu::ShaderSource::Wgsl(guarded(&src, self.merged, patches & F_UNITS != 0).into()) }, - kernel_checks(), - ) - } + // 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)); + unsafe { self.device.create_shader_module_trusted(wgpu::ShaderModuleDescriptor { label: Some("pathtrace+functions"), source: wgpu::ShaderSource::Wgsl(wgsl.into()) }, kernel_checks()) } }) .clone() } @@ -3283,36 +3183,44 @@ mod tests { } } - /// The source-patching features (`kernel_source`) find their anchors in pathtrace.wgsl, once each, alone and all - /// together; no feature leaves the source as it is. - /// The native default (no merged layout, no loop guard): what is handed to wgpu is the source as it was before - /// either existed. `guarded` gives a kernel back untouched; the live passes are their file. Prints the sources' - /// SHA-256 (`--nocapture`), which docs/WORKFLOW.md's gate compares with main's (PR #13). + fn naga_valid(src: &str) { + let m = naga::front::wgsl::parse_str(src).unwrap_or_else(|e| panic!("{}", e.emit_to_string(src))); + 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`). #[test] - fn the_default_sources_are_untouched() { + fn every_kernel_variant_is_valid() { use sha2::{Digest, Sha256}; if std::env::var("THREERT_LOOP_GUARD").is_ok() { return; } assert_eq!(loop_guard(), None); let sha = |s: &str| Sha256::digest(s.as_bytes()).iter().map(|b| format!("{b:02x}")).collect::(); - let live = include_str!("../shaders/live.wgsl"); - assert_eq!(crate::live::live_source(), live); - println!("source live {}", sha(&crate::live::live_source())); + 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(), + KernelSwitches { strict: true, ..Default::default() }, + KernelSwitches { cutout_lite: 1, ..Default::default() }, + KernelSwitches { cutout_lite: 2, ..Default::default() }, + KernelSwitches { last_direct: true, ..Default::default() }, + KernelSwitches { light_pick: true, ..Default::default() }, + all, + ]; for hwrt in [false, true] { - let base = kernel_base(hwrt); - assert_eq!(guarded(&base, false, false), base); - println!("source kernel hwrt={hwrt} {}", sha(&guarded(&base, false, false))); - // Every combination of the features that patch the source (`kernel_source`), each as compiled. - let bits = [F_TINT, F_UNLIT, F_BACKGROUND, F_VCOLOR, F_GRID]; - let mut all = Sha256::new(); - for m in 0..32u32 { - let f = bits.iter().enumerate().filter(|(k, _)| m >> k & 1 == 1).fold(0, |a, (_, b)| a | b); - let src = kernel_source(&base, f); - assert_eq!(guarded(&src, false, false), src); - all.update(guarded(&src, false, false).as_bytes()); + println!("source kernel hwrt={hwrt} {}", sha(&kernel_wgsl(&template, &kernel_flags(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(&crate::merged::uniform_tables(&template), &kernel_flags(hwrt, k, F_SOURCE & !F_UNITS))); } - println!("source kernel hwrt={hwrt} x32 features {}", all.finalize().iter().map(|b| format!("{b:02x}")).collect::()); } } @@ -3334,43 +3242,32 @@ 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_source(&kernel_base(false), F_UNITS); + 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"); - // (That it bounds every loop: merged.rs `the_guard_bounds_every_loop`. The unit walk's are among them.) - let stops = |s: &str| s.matches("if (loop_guard()) { break; }").count(); - let flat = crate::merged::guard(&format!("{}\n{}", kernel_base(false), crate::merged::guard_decl()), 65536, crate::merged::GUARD_TRIPS.name); - assert!(stops(&g) > stops(&flat) && stops(&flat) > 40); assert_eq!(guarded(&src, false, false), src, "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). - let m = naga::front::wgsl::parse_str(&g).unwrap_or_else(|e| panic!("{}", e.emit_to_string(&g))); - naga::valid::Validator::new(naga::valid::ValidationFlags::all(), naga::valid::Capabilities::all()).validate(&m).unwrap_or_else(|e| panic!("{}", e.emit_to_string(&g))); + // 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)); + 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()")); assert!(guard_room(false) && !guard_room(true)); } - /// Software traversal's units are in a kernel only when the scene has them (F_UNITS): the source without the - /// feature is the file; with it, the still scene's walk hands its hit to `trace_units`, whose text is there once, - /// inside the software traversal (the part hardware traversal's source replaces, which therefore cannot have it). + /// Software traversal's units are in a kernel only when the scene has them (F_UNITS, `units`): without, the still + /// walk keeps its stack and there is no unit walk; with, the still walk hands its hit to `trace_units` and its stack + /// is the module's (the procedural objects' walk keeps its own). Hardware traversal has neither software walk. #[test] - fn units_are_spliced_into_software_traversal_only() { - let base = kernel_base(false); - assert!(!base.contains("trace_units")); - let u = kernel_source(&base, F_UNITS); - assert_eq!(u.matches("fn trace_units(").count(), 1); - assert_eq!(u.matches("hit = trace_units(o, d, inv_d, hit, mode);").count(), 1); - let (begin, end, at) = (u.find("// @trace-begin").unwrap(), u.find("// @trace-end").unwrap(), u.find("fn trace_units(").unwrap()); - assert!(begin < at && at < end); - // One stack an invocation, the module's: the still walk declares none (the procedural objects' walk keeps - // its own), and the name it uses is the one the units' file declares. - assert_eq!((base.matches(STILL_STACK).count(), u.matches(STILL_STACK).count()), (1, 0)); - assert_eq!((u.matches("var stack: array;").count(), u.matches("var stack: array stack: array").count(), s.matches("var stack: array").count()); + assert_eq!((stacks(&sw), stacks(&u)), ((0, 2), (1, 1))); } /// With nothing set, every platform but Android opens the device as it was before the options existed: no @@ -3385,7 +3282,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_eq!(kernel_switches(kernel_base(hwrt), &desktop.kernel, hwrt), kernel_base(hwrt)); + assert!(kernel_flags(hwrt, &desktop.kernel, 0).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 }); @@ -3403,44 +3300,37 @@ mod tests { assert_eq!([crate::env_nee_of(None, false), crate::env_nee_of(None, true), crate::env_nee_of(Some("1"), true), crate::env_nee_of(Some("0"), false)], [true, false, true, false]); } - /// Each switch finds its anchors, alone and with the others, under both traversals, and what comes out is valid - /// WGSL (naga, wgpu's front end), with the uniform tables too. Strict guards every ray query the hardware kernel - /// makes; software traversal's source has none and takes only `env_pdf`'s index. + /// The other shader files are WGSL as the specification writes it: wesl's parser is as strict as a browser's + /// (Chrome's Tint) where naga is not. naga took `a * b ^ c`, which WGSL's grammar forbids (a bitwise operator's + /// operands are unary expressions): the scene library's surface materials and both particle passes would not + /// compile in a browser until 2026-11-14. (The kernel's files: `every_kernel_variant_is_valid`.) #[test] - fn the_kernel_switches_find_their_anchors() { - let valid = |src: &str| { - let m = naga::front::wgsl::parse_str(src).unwrap_or_else(|e| panic!("{}", e.emit_to_string(src))); - naga::valid::Validator::new(naga::valid::ValidationFlags::all(), naga::valid::Capabilities::all()).validate(&m).unwrap_or_else(|e| panic!("{}", e.emit_to_string(src))); - }; - let all = KernelSwitches { strict: true, cutout_lite: 2, last_direct: true, light_pick: true }; - let each = [ - KernelSwitches { strict: true, ..Default::default() }, - KernelSwitches { cutout_lite: 1, ..Default::default() }, - KernelSwitches { cutout_lite: 2, ..Default::default() }, - KernelSwitches { last_direct: true, ..Default::default() }, - KernelSwitches { light_pick: true, ..Default::default() }, - all, - ]; + 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, &[]); + } + } + + /// Each switch changes what it says, under both traversals (that every combination is valid WGSL: + /// `every_kernel_variant_is_valid`). Strict guards every ray query the hardware kernel makes; software traversal + /// has none and takes only `env_pdf`'s index; the cutout switches are hardware traversal's. + #[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)); for hwrt in [false, true] { - let base = kernel_base(hwrt); - for k in &each { - let src = kernel_switches(base.clone(), k, hwrt); - // (The cutout switches are hardware traversal's: software traversal's source is the file with them.) - assert_eq!(src == base, !hwrt && k.cutout_lite > 0 && *k == KernelSwitches { cutout_lite: k.cutout_lite, ..Default::default() }, "{k:?} hwrt={hwrt}"); - valid(&src); - // A scene's features patch the source after the switches, and the uniform tables after them. - valid(&crate::merged::uniform_tables(&kernel_source(&src, F_SOURCE & !F_UNITS))); - } - let strict = kernel_switches(base.clone(), &each[0], hwrt); + let plain = v(hwrt, KernelSwitches::default()); + let strict = v(hwrt, KernelSwitches { strict: true, ..Default::default() }); let queries = strict.matches("rayQueryInitialize(").count(); assert_eq!(queries, if hwrt { 3 } else { 0 }); assert_eq!(strict.matches("ray_bad(o, d, tmax").count(), queries, "every ray query is guarded"); assert_eq!(strict.matches("fn ray_bad(").count(), hwrt as usize); - assert!(!strict.contains("env_cdf[y - 1u]") && !strict.contains("env_cdf[base + x - 1u]") && base.contains("env_cdf[y - 1u]")); - let lit = kernel_switches(base.clone(), &all, hwrt); + assert!(!strict.contains("env_cdf[y - 1u]") && !strict.contains("env_cdf[base + x - 1u]") && plain.contains("env_cdf[y - 1u]")); + let lit = v(hwrt, KernelSwitches { strict: true, cutout_lite: 2, last_direct: true, light_pick: true }); assert_eq!(lit.matches("radiance += dz;").count(), 1); - assert_eq!(lit.matches("if (take_env && has_env) {").count(), 1); + assert_eq!(lit.matches("env_on = take_env && has_env;").count(), 1); assert_eq!(lit.contains("let steps = 4u;"), hwrt); + assert_eq!(v(hwrt, KernelSwitches { cutout_lite: 2, ..Default::default() }) == plain, !hwrt); } } @@ -3516,25 +3406,23 @@ 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 + /// is in a kernel without it. #[test] - fn kernel_source_finds_its_anchors() { - let base = include_str!("../shaders/pathtrace.wgsl"); - assert_eq!(kernel_source(base, 0), base); - let t = kernel_source(base, F_TINT); - assert!(t.contains("const OBJECT_STRIDE: u32 = 4u;") && !t.contains("const OBJECT_STRIDE: u32 = 3u;")); - assert_eq!(t.matches("* obj_tint(").count(), 3); - assert!(kernel_source(base, F_UNLIT).contains("if (m.tex1.z == MAT_UNLIT)")); - assert!(kernel_source(base, F_BACKGROUND).contains("background: vec4")); - let v = kernel_source(base, F_VCOLOR); - assert!(v.contains("fn vcolor(") && v.matches("* vcolor(").count() == 2 && !base.contains("vcolor")); - let gr = kernel_source(base, F_GRID); - assert!(gr.contains("@group(0) @binding(19) var grid_tex: texture_3d;") && gr.contains("fn vgrid(i: VIn, lp: vec3) -> vec4") && gr.contains(" grid: vec2,") && gr.contains("i.grid = vec2(") && !base.contains("grid_tex")); - assert_eq!(gr.matches("// @volume-functions-begin").count(), 1); - let all = kernel_source(base, F_SOURCE); - assert!(all.contains("obj_tint(object_of(a))") && all.contains("MAT_UNLIT") && all.contains("P.background.rgb")); - assert!(all.contains("var base = m.base_color.rgb * obj_tint(object_of(a)) * vcolor(i, h.u, h.v);") && all.contains("g.tint = m.base_color.rgb * obj_tint(object_of(a)) * vcolor(h.tri, h.u, h.v);")); - assert!(t.contains("var base = m.base_color.rgb * obj_tint(object_of(a));\n") && t.contains("g.tint = m.base_color.rgb * obj_tint(object_of(a));\n"), "F_TINT alone as before"); + 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 base = v(0); + let tint = v(F_TINT); + assert!(tint.contains("const OBJECT_STRIDE: u32 = 4u;") && !tint.contains("const OBJECT_STRIDE: u32 = 3u;")); + assert_eq!(tint.matches("*= obj_tint(").count(), 3); assert!(base.contains("const OBJECT_STRIDE: u32 = 3u;") && !base.contains("obj_tint")); + assert!(v(F_UNLIT).contains("MAT_UNLIT") && !base.contains("MAT_UNLIT")); + assert!(v(F_BACKGROUND).contains("P.background.rgb") && !base.contains("background")); + let vc = v(F_VCOLOR); + assert!(vc.contains("fn vcolor(") && vc.matches("*= vcolor(").count() == 2 && !base.contains("vcolor")); + let gr = v(F_GRID); + assert!(gr.contains("var grid_tex: texture_3d") && gr.contains("fn vgrid(") && gr.contains("i.grid = ") && !base.contains("grid_tex") && !base.contains("vgrid")); } /// Structural features other than F_INST (opaque particles, media) do not make a scene count as instanced: a group diff --git a/renderer/src/merged.rs b/renderer/src/merged.rs index d3a9d95..72dd485 100644 --- a/renderer/src/merged.rs +++ b/renderer/src/merged.rs @@ -3,7 +3,7 @@ //! live passes need none of this: each binds at most eight, live/layout.rs, and a page binds them as they are.) There //! is one source: pathtrace.wgsl, and the bindings gpu.rs makes for it, are what the native renderer runs, unchanged. //! This module writes the browser's from them, as text (as -//! `gpu::kernel_source` patches the kernel for a scene's features), and gives the Rust side the numbers that go with +//! `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"). //! @@ -301,7 +301,7 @@ fn merged_struct(name: &str, members: &[&Member], caps: &[u32]) -> String { format!("struct {name} {{\n{}\n}};\n", fields.join("\n")) } -/// The kernel (pathtrace.wgsl as a scene's variant: `gpu::kernel_source`, its functions spliced in) in the merged +/// 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 { @@ -461,6 +461,14 @@ mod tests { const KERNEL: &str = include_str!("../shaders/pathtrace.wgsl"); + /// 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. + 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)) + } + fn caps() -> Caps { // A cache of one entry, a view of 64 x 36: lengths that are not round. let n = 64 * 36; @@ -494,13 +502,13 @@ mod tests { /// structs' arrays where `offsets` puts them. The source as written has 19, so the count is the generator's. #[test] fn the_merged_kernel_fits_eight_storage_buffers() { - let (m, info) = parse(KERNEL); + let (m, info) = parse(&variant(0, None)); let (used, declared) = storage_buffers(&m, &info); assert_eq!(declared, 19, "the kernel as written"); assert_eq!(used.iter().map(|u| u.1).max(), Some(19), "{used:?}"); 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 = kernel(&crate::gpu::kernel_source(KERNEL, features), &caps); + let src = variant(features, Some(&caps)); let (m, info) = parse(&src); let (used, declared) = storage_buffers(&m, &info); assert!(declared as u32 <= STORAGE_BUFFERS, "features {features:#x}: {declared} declared"); @@ -539,10 +547,10 @@ mod tests { out.sort_by_key(|g| g.1); out }; - let (written, _) = parse(KERNEL); + 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 = crate::gpu::kernel_source(KERNEL, features); - let src = uniform_tables(&base); + 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 (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(); @@ -598,12 +606,14 @@ mod tests { words("loop {") + words("for (") + words("while (") }; // The kernel as written (its count a buffer of its own) and merged (its count in `MergedLive`). - let merged = kernel(&crate::gpu::kernel_source(KERNEL, crate::gpu::F_SOURCE), &caps()); - for (src, trips) in [(format!("{KERNEL}{}", guard_decl()), "loop_guard_trips"), (merged, "ML.loop_guard_trips")] { + 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)); - let (m, info) = parse(&g); + // (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 (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:?}"); diff --git a/renderer/src/scene.rs b/renderer/src/scene.rs index 47a7f7d..4a5b2c1 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, gpu.rs `kernel_source`). +/// three.js MeshBasicMaterial (`desc::Material::unlit`; the kernel's F_UNLIT, `@if(unlit)`: gpu.rs `kernel_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/script.rs b/renderer/src/script.rs index 8463e5b..1905dcd 100644 --- a/renderer/src/script.rs +++ b/renderer/src/script.rs @@ -140,7 +140,7 @@ impl Drop for Lent<'_> { let took = self.1.elapsed().as_micros() as u64; let give = took.min(LEND_US.saturating_sub(self.0.lent.load(Ordering::Relaxed))); // (0: no deadline is set, and none is made.) - if give > 0 && self.0.deadline.fetch_update(Ordering::Relaxed, Ordering::Relaxed, |d| (d != 0).then_some(d + give)).is_ok() { + if give > 0 && self.0.deadline.try_update(Ordering::Relaxed, Ordering::Relaxed, |d| (d != 0).then_some(d + give)).is_ok() { self.0.lent.fetch_add(give, Ordering::Relaxed); } } diff --git a/renderer/src/units.rs b/renderer/src/units.rs index cca2af8..3f46fed 100644 --- a/renderer/src/units.rs +++ b/renderer/src/units.rs @@ -1025,7 +1025,7 @@ pub(crate) mod cpu { } /// pathtrace.wgsl `trace`, the still scene's tree (no cutouts: the scenes here have none), then the units as - /// `kernel_source` splices them in. + /// the kernel has them under `@if(units)`. pub(crate) fn trace(&self, o: Vec3A, d: Vec3A, tmax_in: f32, mode: u32) -> Walk { let mut hit = Hit { t: tmax_in, tri: 0xffff_ffff, inst: NO_INST }; let inv_d = Self::inv(d); diff --git a/renderer/web/Cargo.lock b/renderer/web/Cargo.lock index 814be5f..ded1b0b 100644 --- a/renderer/web/Cargo.lock +++ b/renderer/web/Cargo.lock @@ -37,6 +37,15 @@ dependencies = [ "zerocopy", ] +[[package]] +name = "aho-corasick" +version = "1.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c982642fa9e8606056828ee9a8505737230110bb1099153c79efe865c59d12ba" +dependencies = [ + "memchr", +] + [[package]] name = "allocator-api2" version = "0.2.21" @@ -77,6 +86,67 @@ dependencies = [ "libc", ] +[[package]] +name = "annotate-snippets" +version = "0.12.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0258346f92c78dbc5b336b9a5b27c64447aa10643bc64e651db2edb145361564" +dependencies = [ + "anstyle", + "memchr", + "unicode-width", +] + +[[package]] +name = "anstream" +version = "1.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "824a212faf96e9acacdbd09febd34438f8f711fb84e09a8916013cd7815ca28d" +dependencies = [ + "anstyle", + "anstyle-parse", + "anstyle-query", + "anstyle-wincon", + "colorchoice", + "is_terminal_polyfill", + "utf8parse", +] + +[[package]] +name = "anstyle" +version = "1.0.14" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "940b3a0ca603d1eade50a4846a2afffd5ef57a9feac2c0e2ec2e14f9ead76000" + +[[package]] +name = "anstyle-parse" +version = "1.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "52ce7f38b242319f7cabaa6813055467063ecdc9d355bbb4ce0c68908cd8130e" +dependencies = [ + "utf8parse", +] + +[[package]] +name = "anstyle-query" +version = "1.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "40c48f72fd53cd289104fc64099abca73db4166ad86ea0b4341abe65af83dadc" +dependencies = [ + "windows-sys 0.61.2", +] + +[[package]] +name = "anstyle-wincon" +version = "3.0.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "291e6a250ff86cd4a820112fb8898808a366d8f9f58ce16d1f538353ad55747d" +dependencies = [ + "anstyle", + "once_cell_polyfill", + "windows-sys 0.61.2", +] + [[package]] name = "anyhow" version = "1.0.104" @@ -302,6 +372,12 @@ dependencies = [ "unicode-width", ] +[[package]] +name = "colorchoice" +version = "1.0.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1d07550c9036bf2ae0c684c4297d503f838287c83c53686d05370d0e139ae570" + [[package]] name = "combine" version = "4.6.9" @@ -321,6 +397,15 @@ dependencies = [ "crossbeam-utils", ] +[[package]] +name = "convert_case" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "633458d4ef8c78b72454de2d54fd6ab2e60f9e02be22f3c6104cdc8a4e0fceb9" +dependencies = [ + "unicode-segmentation", +] + [[package]] name = "core-foundation" version = "0.9.4" @@ -426,6 +511,29 @@ version = "1.2.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "f27ae1dd37df86211c42e150270f82743308803d90a6f6e6651cd730d5e1732f" +[[package]] +name = "derive_more" +version = "2.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d751e9e49156b02b44f9c1815bcb94b984cdcc4396ecc32521c739452808b134" +dependencies = [ + "derive_more-impl", +] + +[[package]] +name = "derive_more-impl" +version = "2.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "799a97264921d8623a957f6c3b9011f3b5492f557bbb7a5a19b7fa6d06ba8dcb" +dependencies = [ + "convert_case", + "proc-macro2", + "quote", + "rustc_version", + "syn 2.0.119", + "unicode-xid", +] + [[package]] name = "digest" version = "0.10.7" @@ -560,6 +668,12 @@ dependencies = [ "zlib-rs", ] +[[package]] +name = "fnv" +version = "1.0.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3f9eec918d3f24069decb9af1554cad7c880e2da24a9afd88aca000531ab82c1" + [[package]] name = "foldhash" version = "0.2.0" @@ -713,6 +827,12 @@ dependencies = [ "simba", ] +[[package]] +name = "glob" +version = "0.3.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e4eba85ea1d0a966a983acd07deee566e67395d2d96b6fb39e62b5a833f1eb0b" + [[package]] name = "glow" version = "0.17.0" @@ -881,6 +1001,21 @@ version = "1.1.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "a257582fdcde896fd96463bf2d40eefea0580021c0712a0e2b028b60b47a837a" +[[package]] +name = "is_terminal_polyfill" +version = "1.70.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a6cb138bb79a146c1bd460005623e142ef0181e3d0219cb493e02f7d08a35695" + +[[package]] +name = "itertools" +version = "0.15.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8b4baf93f58d4425749ca49a51c50ebab072c5df6994d08fed93541c331481dc" +dependencies = [ + "either", +] + [[package]] name = "itoa" version = "1.0.19" @@ -983,6 +1118,15 @@ version = "3.1.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "e2db585e1d738fc771bf08a151420d3ed193d9d895a36df7f6f8a9456b911ddc" +[[package]] +name = "lalrpop-util" +version = "0.23.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "364b49d1c98c07972e16fe1346463becc08bcd663da04e6e37ab1bea068dce6a" +dependencies = [ + "regex-automata", +] + [[package]] name = "lazy_static" version = "1.5.1" @@ -995,6 +1139,72 @@ version = "0.5.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "7a79a3332a6609480d7d0c9eab957bca6b455b91bb84e66d19f5ff66294b85b8" +[[package]] +name = "lexical" +version = "7.0.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1bc8a009b2ff1f419ccc62706f04fe0ca6e67b37460513964a3dfdb919bb37d6" +dependencies = [ + "lexical-core", +] + +[[package]] +name = "lexical-core" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7d8d125a277f807e55a77304455eb7b1cb52f2b18c143b60e766c120bd64a594" +dependencies = [ + "lexical-parse-float", + "lexical-parse-integer", + "lexical-util", + "lexical-write-float", + "lexical-write-integer", +] + +[[package]] +name = "lexical-parse-float" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "52a9f232fbd6f550bc0137dcb5f99ab674071ac2d690ac69704593cb4abbea56" +dependencies = [ + "lexical-parse-integer", + "lexical-util", +] + +[[package]] +name = "lexical-parse-integer" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9a7a039f8fb9c19c996cd7b2fcce303c1b2874fe1aca544edc85c4a5f8489b34" +dependencies = [ + "lexical-util", +] + +[[package]] +name = "lexical-util" +version = "1.0.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2604dd126bb14f13fb5d1bd6a66155079cb9fa655b37f875b3a742c705dbed17" + +[[package]] +name = "lexical-write-float" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "50c438c87c013188d415fbabbb1dceb44249ab81664efbd31b14ae55dabb6361" +dependencies = [ + "lexical-util", + "lexical-write-integer", +] + +[[package]] +name = "lexical-write-integer" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "409851a618475d2d5796377cad353802345cba92c867d9fbcde9cf4eac4e14df" +dependencies = [ + "lexical-util", +] + [[package]] name = "libc" version = "0.2.191" @@ -1062,6 +1272,38 @@ version = "0.4.35" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "5d15f86d12da44cd8d4ec3c73984537eb7f4cadcb9776a575a0a8a05ca11005c" +[[package]] +name = "logos" +version = "0.16.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "eb2c55a318a87600ea870ff8c2012148b44bf18b74fad48d0f835c38c7d07c5f" +dependencies = [ + "logos-derive", +] + +[[package]] +name = "logos-codegen" +version = "0.16.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "58b3ffaa284e1350d017a57d04ada118c4583cf260c8fb01e0fe28a2e9cf8970" +dependencies = [ + "fnv", + "proc-macro2", + "quote", + "regex-automata", + "regex-syntax", + "syn 2.0.119", +] + +[[package]] +name = "logos-derive" +version = "0.16.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "52d3a9855747c17eaf4383823f135220716ab49bea5fbea7dd42cc9a92f8aa31" +dependencies = [ + "logos-codegen", +] + [[package]] name = "matrixmultiply" version = "0.3.11" @@ -1609,6 +1851,12 @@ version = "1.21.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "9f7c3e4beb33f85d45ae3e3a1792185706c8e16d043238c593331cc7cd313b50" +[[package]] +name = "once_cell_polyfill" +version = "1.70.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "384b8ab6d37215f3c5301a95a4accb5d64aa607f1fcb26a11b5303878451b4fe" + [[package]] name = "orbclient" version = "0.3.55" @@ -1968,6 +2216,23 @@ dependencies = [ "bitflags 2.13.2", ] +[[package]] +name = "regex-automata" +version = "0.4.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ad8553b9b26413251cbf30e620595c7a41b3887f03da04579c0e6b0d6a06b4b2" +dependencies = [ + "aho-corasick", + "memchr", + "regex-syntax", +] + +[[package]] +name = "regex-syntax" +version = "0.8.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4" + [[package]] name = "relative-path" version = "2.0.3" @@ -2168,6 +2433,15 @@ dependencies = [ "zmij", ] +[[package]] +name = "serde_spanned" +version = "1.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9c187e0a369471c550af028cba56d4f4fa2a489ad28151ae95a4d340a2a8b782" +dependencies = [ + "serde_core", +] + [[package]] name = "sha2" version = "0.10.9" @@ -2412,6 +2686,7 @@ dependencies = [ "sha2", "voxel-core", "web-time", + "wesl", "wgpu", "winit", ] @@ -2465,6 +2740,21 @@ dependencies = [ "pin-project-lite", ] +[[package]] +name = "toml" +version = "1.2.1+spec-1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d5755c51863cdac711f1956baf8eb9588e060755f7d3c4c979e70ea1e1f32cd7" +dependencies = [ + "indexmap", + "serde_core", + "serde_spanned", + "toml_datetime", + "toml_parser", + "toml_writer", + "winnow", +] + [[package]] name = "toml_datetime" version = "1.1.3+spec-1.1.0" @@ -2495,6 +2785,12 @@ dependencies = [ "winnow", ] +[[package]] +name = "toml_writer" +version = "1.2.0+spec-1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "009dbb7700f11c55ef83aebbe4c0e9cf1f1ac0844bd8d6f00dc38f302026e13a" + [[package]] name = "tracing" version = "0.1.44" @@ -2541,12 +2837,24 @@ version = "0.2.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "ca718471876992059f251e33aafc334346ada08042b7661dd1fe313a9bced557" +[[package]] +name = "unicode-xid" +version = "0.2.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ebc1c04c71510c7f702b52b7c350734c9ff1295c464a03335b00bb84fc54f853" + [[package]] name = "urlencoding" version = "2.1.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "daf8dba3b7eb870caf1ddeed7bc9d2a049f3cfdfae7cb521b087cc33ae4c49da" +[[package]] +name = "utf8parse" +version = "0.2.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "06abde3611657adf66d383f00b093d7faecc7fa57071cce2578660c9f1010821" + [[package]] name = "version_check" version = "0.9.5" @@ -2760,6 +3068,46 @@ dependencies = [ "wasm-bindgen", ] +[[package]] +name = "wesl" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fdd8f9619baa4c7bde4df0899d6f626b70ee92d6e9c18c5f74473419b5bb6fa6" +dependencies = [ + "annotate-snippets", + "anstream", + "derive_more", + "glob", + "half", + "itertools", + "num-traits", + "serde", + "thiserror 2.0.21", + "toml", + "wesl-core", + "wesl-macros", + "wgsl-parse", + "wgsl-types", +] + +[[package]] +name = "wesl-core" +version = "0.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "84dd3e600497eda0b47afcd21bd84c8953a9dfce7e93ed27a38ba97ac0747e03" + +[[package]] +name = "wesl-macros" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ac8767ec26e75262a8e581d5f188b1f63a288baaf885bb33c9f4f73a693e06fc" +dependencies = [ + "itertools", + "proc-macro2", + "quote", + "syn 3.0.10", +] + [[package]] name = "wgpu" version = "30.0.1" @@ -2933,6 +3281,33 @@ dependencies = [ "web-sys", ] +[[package]] +name = "wgsl-parse" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "37cee524ba5cfabfd65d10a316b646f8a4c50f3663e7e787471f21b783261f47" +dependencies = [ + "annotate-snippets", + "derive_more", + "itertools", + "lalrpop-util", + "lexical", + "logos", + "thiserror 2.0.21", + "wgsl-types", +] + +[[package]] +name = "wgsl-types" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e119dd12e57d937397535201959ddddb5dbf26505402db5ef305ebb1ca9a7e5f" +dependencies = [ + "half", + "itertools", + "num-traits", +] + [[package]] name = "wide" version = "1.7.1" diff --git a/rust-toolchain.toml b/rust-toolchain.toml new file mode 100644 index 0000000..a39e087 --- /dev/null +++ b/rust-toolchain.toml @@ -0,0 +1,6 @@ +# The Rust every build of this repository uses: CI's and every machine's the same, so a build and the determinism +# checks' recorded hashes (tools/physics-determinism.ts, tools/voxel-determinism.ts) do not change with whatever stable +# a machine happens to have. A newer one is a PR of its own (or of the change that needs it), with those checks run. +[toolchain] +channel = "1.100.0" +targets = ["wasm32-unknown-unknown"]