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"]