diff --git a/docs/RUNTIME.md b/docs/RUNTIME.md index 1515376..3190106 100644 --- a/docs/RUNTIME.md +++ b/docs/RUNTIME.md @@ -1170,6 +1170,158 @@ 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). +## 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. +On Qualcomm's Adreno 740 (Snapdragon 8 Gen 2, Android 13; an AYN Thor handheld) the renderer as written did not +draw correctly: a model's textures came out as another table's numbers, emissive objects lit nothing, Light Room was +black. What follows is what fixed it there, on this engine's current source, and it has run on the device ("Checked +on the device" below). All of it is off on every other platform unless asked for: with nothing set the device, its +layout and its kernels' source are what they were (`gpu.rs` `the_options_default_to_the_device_as_written`, +`the_default_sources_are_untouched`). + +**The figures in the two tables are the original author's, measured on that handheld with the engine at `b646a20`,** +on scenes that are not all at hand. What was measured again on this source, on the same handheld, is under "Checked +on the device", and differs where it says so. + +- **Uniform tables** (`GpuOptions::uniform_tables`, `THREERT_UNIFORM_TABLES=1|0`; default on Android; never with the + merged layout). The driver gave the kernel other buffers' contents for four of its read-only storage buffers: + `tex_info` (binding 14), `objects` (15), `instances` (17), `media` (18). As uniform buffers they read right. + - The kernel's source is rewritten when a module is compiled, as the merged layout's is (`merged::uniform_tables`): + the four declarations become sized uniform arrays at the same bindings, and every index into them is held to the + array (`min(u32(i), last)`), so nothing is read outside a uniform buffer. + - The renderer keeps writing its storage buffers. `Gpu::sync_tables` copies the four into their uniform buffers in + a frame's encoder, after the particle passes (which write their objects' entries on the GPU) and before the trace. + - A uniform binding holds 64 KiB, so the arrays are bounded: 2,048 textures, 4,096 object entries (1,365 objects; + 1,024 with instance colours), 409 instance records (model instances, procedural objects, `dynamicGeometry` + meshes), all 32 media. A scene with more reads the last entry for the rest, and the renderer says so on stderr + when it uploads the scene and once when in-place edits pass a bound. `THREERT_UNIFORM_TABLES=0` binds storage + buffers again. +- **The kernel's other storage buffers stay as they are**, each for its reason: + + | buffer | size | decision | + |---|---|---| + | `materials` (4), `lights` (5) | 288 B and 80 B an entry | storage: the device drew right with them as they are; 64 KiB would hold 227 materials | + | `emitters` (6) | the emitter CDF and the light tree | storage: megabytes in a lit scene | + | `env_cdf` (10) | the environment's CDF, then the procedural geometry's data | storage (megabytes); its light sampling is off on Android, below | + | `tri_verts` (1), `attrs` (2), `nodes` (3), `prev_verts` (16) | per triangle | storage (megabytes) | + | `accum` (11), group 1 (radiance cache, reservoirs, the trace's outputs, the loop guard's count) | per pixel, per cache entry | written by the kernel: cannot be uniform | + | group 2, the translucency pass (`timg`, `vpts`, `vcount`, `tlow`) | per pixel | written by the pass | + | the particle passes' `systems`, `spawns`, `grid`; the grid programs' `srcs` | small, read-only | storage: other shaders, of eleven bindings at most and no acceleration structure; not known to misread (the fix on the device left them alone) | + | the live passes (`live.wgsl`: 28 storage buffers, `ae_bins` and `ae_state` among them) | per pixel; two tiny | storage: written by the passes; the device showed their output right | + + Two of these are what a device run should look at first: `prev_verts` is the one storage buffer left after the + acceleration structure's binding (13) and feeds the motion vectors of moving objects and the vertex colours; + `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): + + | switch | default | what it does | on the device | + |---|---|---|---| + | `THREERT_KERNEL_STRICT` | Android: 1 | a ray query with a ray that is not finite, or a length that is negative or NaN, is not made (the ray misses): the kernel's two and the translucency pass's; `env_pdf` does not form the index before the first entry | the odd such ray faulted the GPU's MMU thousands of times a second and doubled the trace | + | `THREERT_ENV_NEE` | Android: 0 | 0 = a baked environment map (an HDRI, a `SkyAtmosphere`) is not light-sampled, at load and when a dynamic sky is baked again; the sky lights through the paths' own rays, the sun is a directional light as ever | its reads of `env_cdf` went to the wrong memory: tens of thousands of GPU page faults a second, a trace up to twice as slow | + | `THREERT_CUTOUT_LITE` | Android: 1 | hardware traversal: 1 = shadow and any-hit rays take alpha cutouts as opaque; 2 = camera rays also step through at most 4 cutouts | a map with cutout foliage: the trace 235 ms, 41 ms without cutouts | + | `THREERT_LAST_DIRECT` | Android: 1 | the path's last vertex (at the bounce limit, past the camera hit) takes its direct light | with one bounce and no radiance cache (which cost 30-60 % more trace) sunlit ground lit nothing and interiors were black | + | `THREERT_LIGHT_PICK` | 0 | past the camera hit, the directional lights and one of the other kinds (local lights, emitters, the sky), weighted by their number: two shadow rays a vertex instead of four; unbiased, noisier | a handheld's host sets it with `THREERT_LAST_DIRECT` | + + `THREERT_CHECKED=1` (or a mask: 2 bounds, 4 loop bounding, 8 ray-query initialization) compiles the kernel with + naga's runtime checks, to tell a driver's fault from the kernel's. +- **The pipeline cache** (`GpuOptions::pipeline_cache`, `THREERT_PIPELINE_CACHE=`; `pipeline_cache.rs`). + The driver compiles a kernel variant for about ten seconds. With a directory given, the device has a wgpu + `PipelineCache` (Vulkan only), every compute pipeline of the engine and the raster pass's are made with it, and it + is saved after a kernel variant, a scene's passes and the live passes are compiled. Opening a game a second time + went from 10.6 s to 0.4 s (the author's; on this source a scene's pipelines compile for 16 to 31 s the first time + and are there in under a second from the cache, below). + - One file per adapter, named by wgpu's key (backend, vendor, device). wgpu's header in the data holds the driver's + cache identity: after a driver update the data is not used and the next save replaces the file. + - The data is sealed (a magic, its length, its SHA-256): a file cut short or damaged is not handed to the driver. + - Two processes may share the directory: a save writes a file of its own name and renames it over the old one. The + last save wins; what the other compiled meanwhile is compiled again by whoever misses it. + - The directory may be deleted at any time. A file past 256 MiB is not read. +- **What a client sets.** `Gpu::new()` reads the environment, so an Android client that sets + `THREERT_PIPELINE_CACHE` to a directory of its own before opening the device has everything above. In code: + `Gpu::new_with(GpuOptions { pipeline_cache: Some(dir), ..GpuOptions::from_env() })` (`from_env` is the platform's + defaults under the environment; `GpuOptions::default()` has no switch on). +- **What the device affords.** The author's figure for `b646a20` was 640 × 360 at 30 fps with one or two bounces + (a tenth of the rays an iPhone 15 Pro traces). This source does not reach it: the trace costs 0.42 to 0.65 µs a + traced pixel at one bounce whatever the scene, which is 249 × 140 to 333 × 187 traced at 30 fps ("Checked on the + device"). +- **Dynamic resolution by the clock** (`THREERT_GOVERN_CLOCK=1|0`, default on Android; `live.rs` + `govern_resolution`). The render scale was governed in frames (a check every 30th, the urgent one after the first + 120), which is twice a second and 2 s only at 60 frames a second. A scene opens at its full scale, on this GPU at + about 6 frames a second, so it was a slide show for 25 s. With the clock it is at its frame rate 3 s after its + first frame. +- **Not carried over from the original change:** it also sorted bind group layout entries and bind group entries by + binding, for a driver thought to mind their order. wgpu puts both in binding order itself before its backend sees + them (wgpu-core 30 `bgl::EntryMap::from_entries`, `Device::create_bind_group`), so the sort changed nothing; a test + holds wgpu to it (`wgpu_takes_bind_group_entries_in_any_order`). + +**Checked on the device** (AYN Thor: Snapdragon 8 Gen 2, Adreno 740, Android 13, driver 0676.53; 2026-11-13, this +source under ThreeRT for Android 0.4.0, which sets `THREERT_PIPELINE_CACHE` and nothing else): + +- **It draws right.** The live view's frame (the client's `POST /screenshot`, 1280 × 720 or 960 × 540, 32 to 88 + frames of history) against the same scene's picture from a Mac (the published cover, same camera): mean absolute + difference per channel, of 255: Material balls 0.8, Light Room 1.3 (it was black), Damaged Helmet 1.6 (its + textures were another table's numbers), Boutique Showroom 3.0 (its shelf strips and spots light the room). + OpenStrikeRT's hall from the place its page's picture was taken: the same walls, pillars, crates, signs, tube + lights, the sun through the skylight on the floor, the furnace's glow; soldiers and their weapons (instanced, + posed) are whole and in their teams' colours. No GPU fault in the kernel log (`dmesg`, kgsl) in any run. +- **The bounds.** OpenStrikeRT (104,898 triangles, 271 objects) with nine bots stays inside the uniform tables: + the renderer's warning does not appear. +- **The pipeline cache works.** A first open, after the package is on the device: OpenStrikeRT 31 s, Boutique + Showroom 18 s, Material balls 16 s (kernel variants and the live passes compiling). The same after a restart of + the app: 0.9 s, 1.7 s (1.1 s of it the scene's setup), 0.5 s. +- **What a frame costs.** GPU time at 640 × 360 traced (scale 1), one bounce, Android's defaults: Hello (31,362 + triangles) 101 ms, Material balls (217,730) 96 ms, Sunlit room (6,400) 149 ms, Boutique Showroom (131,200) + 139 ms, OpenStrikeRT 146 ms; with ray tracing off (the camera ray alone) OpenStrikeRT still takes 90 ms. Of + Material balls' frame at 1280 × 720 out, 416 × 234 traced: the trace 40.7 ms, every other pass together 9 ms + (upscale 1.7 ms, the denoisers 4 ms); at 960 × 540 out the others take about 5 ms. So the trace is the frame, + it goes with the pixels traced and not with the scene, and 30 fps is 249 × 140 (OpenStrikeRT, the showroom) to + 333 × 187 (Material balls, two bounces) under a 960 × 540 view; 40 fps is 249 × 140 to 291 × 164 under + 853 × 480 (36 to 37 fps reached); 60 fps is 192 × 108 to 224 × 126 under 640 × 360 (54 to 58 reached). +- **The switches, one at a time** (OpenStrikeRT, 640 × 360 traced, one bounce; 146 ms as shipped; one run each): + `THREERT_LIGHT_PICK=1` 137 ms; `THREERT_LAST_DIRECT=0` 116 ms; `THREERT_UNIFORM_TABLES=0` 135 ms (so the uniform + tables and their index clamps cost about 8 %); `THREERT_KERNEL_STRICT=0` 151 ms (no slower with the guard, and no + fault without it in these scenes); `THREERT_CACHE=0` with the light pick 135 ms (Material balls: 96.9 against + 96.5 ms: the radiance cache costs nothing measurable here). None of them is the factor of four to the author's + figure. +- **Hardware against software traversal** (`THREERT_HWRT=0`, same size and bounce): Sunlit room 78 ms against + 149 ms and Material balls 74 against 96 are faster without ray queries; OpenStrikeRT is 336 against 146. A ray + query has a high fixed cost on this driver; ray queries stay the default. `THREERT_RASTER=1` (raster primary + visibility): Sunlit room 136 against 149 ms. +- **Not looked at on the device:** voxel worlds and procedural geometry (`env_cdf`'s second half), particles, media, + a map with cutout foliage (`THREERT_CUTOUT_LITE`'s figures are the author's), `THREERT_ENV_NEE=1`, and the + motion vectors of moving objects (`prev_verts`): soldiers were looked at standing, not frame by frame in motion. + +**Checked on a Mac** (Apple M3 Max, Metal, 2026-11-13; main `f69b1bc` against this change): +- The atomic tests (`npm run test:render`, 187), run in full on main's build, on this build, and on this build + with `THREERT_UNIFORM_TABLES=1` (the last two twice): every test passes in every run but `geometry/grow-live` + (below). Their HDR outputs, byte for byte: all 126 offline renders are identical in all five runs. Of the 60 other + live tests, 41 gave the same bytes in every run, with the uniform tables and without. The other 19 differ between + two runs without the uniform tables, and by as much as with them (18 of them run three or four more times both + ways: about as many differing numbers and as large a relative difference, 10⁻⁶ to 0.6 by test). +- Under software traversal (`THREERT_HWRT=0`) the same 24 tests fail with and without the uniform tables (particles, + translucency: hardware traversal's), and 173 of 187 outputs are identical. The other 14 are live tests: 13 of them + differ from run to run without the uniform tables too, and `animation/shared-pose-live` was identical both ways + when run again. +- With the uniform tables and `THREERT_KERNEL_STRICT=1`: the 126 offline renders are still main's bytes. With all of + Android's defaults: 187 pass, and 117 of the 126 are main's bytes (nine differ as the last vertex's direct light + and the unsampled sky intend). +- `geometry/grow-live` adds 1,100 `dynamicGeometry` objects in a frame and often stops at the script's deadline + ("interrupted") on this machine while other jobs run: on main's build in 4 of the 7 times it was run, on this + build more often than not, whatever is set. Where it ran, it passed under hardware traversal in every configuration. Its 1,102 instance + records are past the uniform tables' 409, which the renderer says; under software traversal with the uniform + tables its centre reads 1 (the sky) instead of 0.5, both times it was run. That is the bound above, at work: a + scene past it is not drawn right with the uniform tables. +- `cargo test --release` passes as written, with `THREERT_UNIFORM_TABLES=1`, and with all of Android's defaults. + `live::the_uniform_tables_render_what_the_tables_do` renders a scene offline on a device with and without the + uniform tables to the same bytes and reads the uniform buffers back against their tables, before and after an + edit in place. +- The library checks for `aarch64-linux-android` (NDK 27.1, `cargo ndk -t arm64-v8a --platform 29 check --release + --lib --features "winit/android-native-activity rquickjs/bindgen"`, with `LIBCLANG_PATH` and the NDK sysroot for + bindgen) and for wasm32. +- Not checked here: the pipeline cache against a driver (Metal has none: its file handling is tested, its wgpu + calls only compile), and anything on an Adreno. + ## Correctness * `threert-render edit-check` builds the GPU tables both ways for a list of events: edited in place, and built fresh diff --git a/renderer/src/gpu.rs b/renderer/src/gpu.rs index f4dcd7d..7e53aba 100644 --- a/renderer/src/gpu.rs +++ b/renderer/src/gpu.rs @@ -14,6 +14,12 @@ pub struct Gpu { pub device: wgpu::Device, pub queue: wgpu::Queue, pub adapter_info: wgpu::AdapterInfo, + /// The kernel's bounded tables are uniform buffers (`GpuOptions::uniform_tables`; merged.rs `uniform_tables`). + pub uniform_tables: bool, + /// The kernel's switches this device was opened with (`KernelSwitches`): they are in `source`. + pub switches: KernelSwitches, + /// Compiled pipelines kept between runs (`GpuOptions::pipeline_cache`), where the backend has such a cache. + pipeline_cache: Option, /// Kernel source (hardware or software traversal spliced in); scenes with material functions get their own module. source: String, /// The kernel as written (none in the merged layout: its sources are written for their arrays' lengths). @@ -181,6 +187,33 @@ impl Uploaded { } } +/// The uniform tables' buffers among the seven (`table_buffers`), in merged.rs `UNIFORM_TABLES`' order. +fn uniform_sources(tables: [&wgpu::Buffer; 7]) -> [&wgpu::Buffer; 4] { + crate::merged::UNIFORM_TABLES.map(|k| tables[k]) +} + +/// A note for each uniform table that a scene has more entries of (`used`, merged.rs `UNIFORM_TABLES`' order) than +/// its array holds: the kernel reads the last entry for the rest. +pub(crate) fn uniform_overflow(used: [u64; 4]) -> Vec { + use crate::merged::{TABLES, UNIFORM_LEN, UNIFORM_TABLES}; + (0..4) + .filter(|&k| used[k] > UNIFORM_LEN[k] as u64) + .map(|k| format!("uniform tables: the scene has {} entries of `{}`, a uniform buffer holds {} (the rest read as the last one; THREERT_UNIFORM_TABLES=0 binds storage buffers)", used[k], TABLES[UNIFORM_TABLES[k]].name, UNIFORM_LEN[k])) + .collect() +} + +/// The bytes of a table that its uniform copy takes: what both hold, in whole words (a copy's unit). +fn uniform_copy_bytes(table: u64, copy: u64) -> u64 { + table.min(copy) & !3 +} + +/// A scene bind group's entries with the uniform tables' copies in their tables' places. +fn bind_uniform_tables<'a>(entries: &mut Vec>, copies: &'a [wgpu::Buffer; 4]) { + let bindings = crate::merged::UNIFORM_TABLES.map(|k| MERGED_TABLE_BINDINGS[k]); + entries.retain(|e| !bindings.contains(&e.binding)); + entries.extend(copies.iter().zip(bindings).map(|(b, binding)| wgpu::BindGroupEntry { binding, resource: b.as_entire_binding() })); +} + /// What the scene bind group holds besides `Uploaded`'s buffers (`Gpu::rebind_scene` makes it again). pub struct BindParts { pub nodes: wgpu::Buffer, @@ -199,6 +232,9 @@ pub struct Uploaded { /// The merged layout (merged.rs): the buffer of the kernel's `MergedTables`, the seven tables below as its arrays /// (`Gpu::sync_tables` copies them there each frame; their lengths: `kernel.caps`). pub tables: Option, + /// The uniform tables (`Gpu::uniform_tables`): the uniform buffers the kernel reads `bg_parts.texinfo`, `objects`, + /// `instances` and `media` from (merged.rs `UNIFORM_TABLES`' order; `Gpu::sync_tables` copies them there each frame). + pub uniform: Option<[wgpu::Buffer; 4]>, pub accum: wgpu::Buffer, pub params: wgpu::Buffer, pub width: u32, @@ -794,6 +830,158 @@ fn guard_room(hwrt: bool) -> bool { loop_guard_of(!hwrt).is_some() } +/// What a host decides about a device before it is opened (`Gpu::new_with`). `Gpu::new` takes it from the environment +/// (`from_env`), where every platform but Android gets `default()`: the device, its layout and its kernels as they +/// are without any of this. What a phone's GPU needed and why: docs/RUNTIME.md "Android (Adreno)". +#[derive(Clone, Debug, Default, PartialEq)] +pub struct GpuOptions { + /// The kernel's bounded tables (texture pages, objects, instances, media) as uniform buffers instead of storage + /// buffers (merged.rs `uniform_tables`). None: where a driver needs it, which is Android. Never with the merged + /// layout, whose tables are one buffer of its own. + pub uniform_tables: Option, + /// A directory to keep compiled pipelines in between runs (pipeline_cache.rs). None: nothing is kept. + pub pipeline_cache: Option, + pub kernel: KernelSwitches, +} + +impl GpuOptions { + /// THREERT_UNIFORM_TABLES=1|0, THREERT_PIPELINE_CACHE=, and the kernel's switches + /// (`KernelSwitches::from_env`). + pub fn from_env() -> GpuOptions { + Self::from_vars(&|name| std::env::var(name).ok(), cfg!(target_os = "android")) + } + + /// `from_env`, of these variables on this platform. + fn from_vars(var: &dyn Fn(&str) -> Option, android: bool) -> GpuOptions { + GpuOptions { + uniform_tables: var("THREERT_UNIFORM_TABLES").map(|v| v == "1"), + pipeline_cache: var("THREERT_PIPELINE_CACHE").filter(|d| !d.is_empty()).map(std::path::PathBuf::from), + kernel: KernelSwitches::from_vars(var, android), + } + } + + /// Whether the device has the uniform tables: as asked, else on Android; not in the merged layout. + fn uniform_tables_on(&self, merged: bool, android: bool) -> bool { + !merged && self.uniform_tables.unwrap_or(android) + } +} + +/// 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. +#[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: + /// - a ray query whose ray is not finite (origin, direction) or whose length is negative or NaN is not made: the + /// ray misses. Such a description is undefined behaviour; on Adreno 740 the odd one faulted the GPU's MMU + /// thousands of times a second and doubled the trace. (naga's `ray_query_initialization_tracking` would do + /// it with bookkeeping around every ray-query call.) + /// - `env_pdf` no longer forms the index before the first entry (`select` evaluates both sides: `y - 1` at + /// y = 0 is 0xffffffff, read and then not used). + pub strict: bool, + /// THREERT_CUTOUT_LITE=0|1|2 (Android: 1). Hardware traversal: 1 = shadow rays and the other any-hit rays take + /// alpha cutouts as opaque; 2 = camera rays also step through at most 4 cutouts (16 as written). Alpha-tested + /// traversal is what costs most on Adreno 740 (a map with cutout foliage: the trace 235 ms, 41 ms without). + pub cutout_lite: u32, + /// THREERT_LAST_DIRECT=1|0 (Android: 1). The path's last vertex (past the camera hit, at the bounce limit) still + /// takes its direct light. A device that affords one bounce and no radiance cache otherwise lights that bounce + /// by sky and emitters alone: sunlit ground lit nothing, interiors were black. + pub last_direct: bool, + /// THREERT_LIGHT_PICK=1|0 (off everywhere unless set). Past the camera hit, `direct_light` takes the directional + /// lights and one of the other kinds a scene has (local lights, emitters, the sky), weighted by their number, + /// instead of all of them: two shadow rays a vertex instead of four. Unbiased; the indirect light is noisier. + pub light_pick: bool, +} + +impl KernelSwitches { + pub fn from_env() -> KernelSwitches { + Self::from_vars(&|name| std::env::var(name).ok(), cfg!(target_os = "android")) + } + + fn from_vars(var: &dyn Fn(&str) -> Option, android: bool) -> KernelSwitches { + let num = |name: &str, default: u32| var(name).and_then(|v| v.parse::().ok()).unwrap_or(default); + let on = android as u32; + KernelSwitches { strict: num("THREERT_KERNEL_STRICT", on) != 0, cutout_lite: num("THREERT_CUTOUT_LITE", on), last_direct: num("THREERT_LAST_DIRECT", on) != 0, light_pick: num("THREERT_LIGHT_PICK", 0) != 0 } + } +} + +/// 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 && restir_here()) {\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 && restir_here()) {\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). +fn kernel_checks() -> wgpu::ShaderRuntimeChecks { + let v: u32 = std::env::var("THREERT_CHECKED").ok().and_then(|v| v.parse().ok()).unwrap_or(0); + if v == 1 { + return wgpu::ShaderRuntimeChecks::checked(); + } + let mut c = wgpu::ShaderRuntimeChecks::unchecked(); + c.bounds_checks = v & 2 != 0; + c.force_loop_bounding = v & 4 != 0; + c.ray_query_initialization_tracking = v & 8 != 0; + 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 { @@ -1239,8 +1427,19 @@ impl Gpu { pollster::block_on(Self::new_async()) } + /// `new`, with a host's own options instead of the environment's. + #[cfg(not(target_arch = "wasm32"))] + pub fn new_with(options: GpuOptions) -> Result { + pollster::block_on(Self::new_async_with(options)) + } + /// The adapter and the device come from futures: a browser resolves them from its event loop. pub async fn new_async() -> Result { + Self::new_async_with(GpuOptions::from_env()).await + } + + /// `new_async`, with a host's own options instead of the environment's. + pub async fn new_async_with(options: GpuOptions) -> Result { let backends = wgpu::Backends::from_env().unwrap_or(wgpu::Backends::PRIMARY); let mut desc = wgpu::InstanceDescriptor::new_without_display_handle(); desc.backends = backends; @@ -1275,6 +1474,9 @@ impl Gpu { if timestamps { features |= wgpu::Features::TIMESTAMP_QUERY; } + // Compiled pipelines kept between runs, where a host gave a directory and the backend has a cache. + let cache_feature = options.pipeline_cache.as_ref().and_then(|_| crate::pipeline_cache::PipelineCache::feature(&adapter, &adapter_info)); + features |= cache_feature.unwrap_or(wgpu::Features::empty()); let (device, queue) = adapter.request_device(&wgpu::DeviceDescriptor { label: Some("threert"), required_features: features, @@ -1283,12 +1485,16 @@ impl Gpu { experimental_features: if hwrt { unsafe { wgpu::ExperimentalFeatures::enabled() } } else { wgpu::ExperimentalFeatures::disabled() }, ..Default::default() }).await?; - let src = kernel_base(hwrt); + 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); // 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()) }, - wgpu::ShaderRuntimeChecks::unchecked(), + kernel_checks(), ) }); let storage = |binding: u32, read_only: bool| wgpu::BindGroupLayoutEntry { @@ -1303,6 +1509,8 @@ impl Gpu { ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None }, count: None, }; + // (The uniform tables: these four bindings hold uniform buffers; merged.rs `UNIFORM_TABLES`.) + let table = |binding: u32| if uniform_tables { uniform(binding) } else { storage(binding, true) }; let mut entries = vec![ uniform(0), storage(1, true), @@ -1327,11 +1535,11 @@ impl Gpu { storage(10, true), storage(11, false), uniform(12), - storage(14, true), - storage(15, true), + table(14), + table(15), storage(16, true), - storage(17, true), - storage(18, true), + table(17), + table(18), // The density grids' atlas (F_GRID; a 1-cell texture in scenes without grids, whose kernel does not read it). wgpu::BindGroupLayoutEntry { binding: 19, @@ -1417,7 +1625,7 @@ impl Gpu { } else { (None, None) }; - Ok(Gpu { instance, adapter, timestamps, device, queue, adapter_info, source: src, module, merged, modules: Default::default(), pipeline_layout: pl, pipelines: Default::default(), layout, vis_layout, no_vis, no_vis_bufs, no_vis_view, no_grid_view, grid_layout, hwrt, pages: Default::default(), tl_layout, tl_pipeline_layout, tl_pipelines: Default::default(), media_passes: Default::default(), proc_cache: Default::default(), proc_blases: Default::default() }) + Ok(Gpu { instance, adapter, timestamps, device, queue, adapter_info, uniform_tables, switches, pipeline_cache, source: src, module, merged, modules: Default::default(), pipeline_layout: pl, pipelines: Default::default(), layout, vis_layout, no_vis, no_vis_bufs, no_vis_view, no_grid_view, grid_layout, hwrt, pages: Default::default(), tl_layout, tl_pipeline_layout, tl_pipelines: Default::default(), media_passes: Default::default(), proc_cache: Default::default(), proc_blases: Default::default() }) } /// Whether this device's kernel buffers have the loop guard's count (`guard_room`; always in the merged layout). @@ -1425,6 +1633,20 @@ impl Gpu { guard_room(self.hwrt) } + /// The pipeline cache for the descriptor of a pipeline about to be made (None: nothing is kept between runs). + pub fn cache(&self) -> Option<&wgpu::PipelineCache> { + self.pipeline_cache.as_ref().map(|c| c.for_pipeline()) + } + + /// The pipelines made since the last call, to the cache's file (pipeline_cache.rs: nothing when none was made, or + /// the cache already had them). The renderer calls it after it compiles a kernel variant, a scene's passes and + /// the live passes; a host that makes pipelines of its own with `cache` calls it after them. + pub fn save_pipeline_cache(&self) { + if let Some(c) = &self.pipeline_cache { + c.save(); + } + } + /// The path-tracing kernel specialised for a scene (feature mask + material functions), compiled on first use. pub fn pipeline(&self, k: &Kernel, live: bool) -> wgpu::ComputePipeline { let mut cache = self.pipelines.lock().unwrap(); @@ -1433,14 +1655,16 @@ impl Gpu { .or_insert_with(|| { let module = self.kernel_module(k); let constants = [("LIVE", if live { 1.0 } else { 0.0 }), ("FEAT", k.features as f64)]; - self.device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + let p = self.device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some(if live { "pathtrace-live" } else { "pathtrace" }), layout: Some(&self.pipeline_layout), module: &module, entry_point: Some("main"), compilation_options: wgpu::PipelineCompilationOptions { constants: &constants, ..Default::default() }, - cache: None, - }) + cache: self.cache(), + }); + self.save_pipeline_cache(); + p }) .clone() } @@ -1467,10 +1691,13 @@ impl Gpu { let src = splice(&splice(&splice(&base, "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()) }, - wgpu::ShaderRuntimeChecks::unchecked(), + kernel_checks(), ) } }) @@ -1486,14 +1713,16 @@ impl Gpu { .or_insert_with(|| { let module = self.kernel_module(k); let constants = [("LIVE", 1.0), ("FEAT", k.features as f64)]; - self.device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + let p = self.device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some(entry), layout: Some(&self.pipeline_layout), module: &module, entry_point: Some(entry), compilation_options: wgpu::PipelineCompilationOptions { constants: &constants, ..Default::default() }, - cache: None, - }) + cache: self.cache(), + }); + self.save_pipeline_cache(); + p }) .clone() } @@ -1513,10 +1742,12 @@ impl Gpu { module: &module, entry_point: Some(entry), compilation_options: wgpu::PipelineCompilationOptions { constants: &constants, ..Default::default() }, - cache: None, + cache: self.cache(), }) }; - TranslucentPipelines { draw: make("translucent"), train: make("volume_train"), up: make("translucent_up") } + let made = TranslucentPipelines { draw: make("translucent"), train: make("volume_train"), up: make("translucent_up") }; + self.save_pipeline_cache(); + made }) .clone() } @@ -1541,14 +1772,16 @@ impl Gpu { (None, Some((_, live))) => self.kernel_module(&Kernel::plain(crate::merged::Caps { tables: crate::merged::TABLE_MIN, live: *live })), (None, None) => unreachable!("the merged layout's group 1 is one buffer"), }; - self.device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + let p = self.device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some("cache resolve"), layout: Some(&self.pipeline_layout), module: &module, entry_point: Some("cache_resolve"), compilation_options: Default::default(), - cache: None, - }) + cache: self.cache(), + }); + self.save_pipeline_cache(); + p } /// Texture pages for the scene's texture set (see `texpack`), cached across uploads of the same set. @@ -1585,8 +1818,8 @@ impl Gpu { } } let view = tex.create_view(&wgpu::TextureViewDescriptor { dimension: Some(wgpu::TextureViewDimension::D2Array), ..Default::default() }); - // (The merged layout copies it into the tables' buffer, `sync_tables`.) - let info = d.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("tex_info"), contents: bytemuck::cast_slice(&packed.info), usage: if self.merged { wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC } else { wgpu::BufferUsages::STORAGE } }); + // (The merged layout copies it into the tables' buffer, the uniform tables into its uniform copy: `sync_tables`.) + let info = d.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("tex_info"), contents: bytemuck::cast_slice(&packed.info), usage: if self.merged || self.uniform_tables { wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC } else { wgpu::BufferUsages::STORAGE } }); timings.push(("texture_pages_ms".into(), t.elapsed().as_secs_f64() * 1000.0)); *cache = Some((key, view.clone(), info.clone())); (view, info) @@ -1707,7 +1940,7 @@ impl Gpu { let pipeline = grids.iter().any(|g| g.sim.is_some()).then(|| { let module = d.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("grid sim"), source: wgpu::ShaderSource::Wgsl(grid_sim_source(scene).into()) }); let layout = d.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: Some("grid sim"), bind_group_layouts: &[Some(&self.grid_layout)], immediate_size: 0 }); - d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some("grid sim"), layout: Some(&layout), module: &module, entry_point: Some("grid_sim"), compilation_options: Default::default(), cache: None }) + d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some("grid sim"), layout: Some(&layout), module: &module, entry_point: Some("grid_sim"), compilation_options: Default::default(), cache: self.cache() }) }); GridsGpu { atlas, pipeline, grids, step: std::cell::Cell::new(0) } }); @@ -1958,11 +2191,24 @@ impl Gpu { bg_entries.push(wgpu::BindGroupEntry { binding: 13, resource: a.tlas.as_binding() }); } // The merged layout: the tables' one buffer in their place (merged.rs). - let tables = self.merged.then(|| self.tables_buffer([&b_mats, &b_lights, &b_emit, &b_texinfo, &b_objects, &b_instances, &b_media])); + let all_tables = [&b_mats, &b_lights, &b_emit, &b_texinfo, &b_objects, &b_instances, &b_media]; + let tables = self.merged.then(|| self.tables_buffer(all_tables)); if let Some((t, _)) = &tables { bg_entries.retain(|e| !MERGED_TABLE_BINDINGS.contains(&e.binding)); bg_entries.push(e(crate::merged::TABLES_BINDING, t)); } + // The uniform tables: four tables' uniform copies in their place (merged.rs `uniform_tables`). + let uniform = self.uniform_tables.then(|| { + // What the scene has of each, against what its array holds (the buffers have the edits' room after it). + let used = [b_texinfo.size() / 16, scene.objects.len() as u64, records.len() as u64, scene.media.len() as u64]; + for note in uniform_overflow(used) { + eprintln!("threert: {note}"); + } + self.uniform_copies(uniform_sources(all_tables)) + }); + if let Some(u) = &uniform { + bind_uniform_tables(&mut bg_entries, u); + } let bg = d.create_bind_group(&wgpu::BindGroupDescriptor { label: None, layout: &self.layout, entries: &bg_entries }); let mut kernel = Kernel::new(scene, groups); kernel.caps = tables.as_ref().map(|t| crate::merged::Caps { tables: t.1, live: self.no_vis_bufs.merged.as_ref().expect("merged").1 }); @@ -1996,9 +2242,9 @@ impl Gpu { let particles = match &accel { Some(a) if !scene.particles.is_empty() => { let module = d.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("particles"), source: wgpu::ShaderSource::Wgsl(particle_source(scene).into()) }); - let pipeline = d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some("particles"), layout: None, module: &module, entry_point: Some("particles"), compilation_options: Default::default(), cache: None }); + let pipeline = d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some("particles"), layout: None, module: &module, entry_point: Some("particles"), compilation_options: Default::default(), cache: self.cache() }); let sm = d.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("stamp"), source: wgpu::ShaderSource::Wgsl(STAMP_WGSL.into()) }); - let sp = d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some("stamp"), layout: None, module: &sm, entry_point: Some("stamp"), compilation_options: Default::default(), cache: None }); + let sp = d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some("stamp"), layout: None, module: &sm, entry_point: Some("stamp"), compilation_options: Default::default(), cache: self.cache() }); // The stamps depend on what is rebuilt last: the translucent TLAS when there is one. let (sv, stl) = translucent.as_ref().map_or((&a.verts, &a.tlas), |t| (&t.verts, &t.tlas)); let sbg = d.create_bind_group(&wgpu::BindGroupDescriptor { label: Some("stamp"), layout: &sp.get_bind_group_layout(0), entries: &[e(0, sv), wgpu::BindGroupEntry { binding: 1, resource: stl.as_binding() }] }); @@ -2038,7 +2284,7 @@ impl Gpu { let gpu_particles = match &accel { Some(a) if !scene.gpu_particles.is_empty() => { let module = d.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("gpu particles"), source: wgpu::ShaderSource::Wgsl(gpu_particle_source(scene).into()) }); - let pipeline = d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some("gpu particles"), layout: None, module: &module, entry_point: Some("gpu_particles"), compilation_options: Default::default(), cache: None }); + let pipeline = d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some("gpu particles"), layout: None, module: &module, entry_point: Some("gpu_particles"), compilation_options: Default::default(), cache: self.cache() }); let systems: Vec = scene .gpu_particles .iter() @@ -2069,7 +2315,7 @@ impl Gpu { }); let host = scene.gpu_particles.iter().map(|ps| (0u32, 0.0f32, ps.rate)).collect(); let sm = d.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("stamp"), source: wgpu::ShaderSource::Wgsl(STAMP_WGSL.into()) }); - let sp = d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some("stamp"), layout: None, module: &sm, entry_point: Some("stamp"), compilation_options: Default::default(), cache: None }); + let sp = d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some("stamp"), layout: None, module: &sm, entry_point: Some("stamp"), compilation_options: Default::default(), cache: self.cache() }); let sbg = d.create_bind_group(&wgpu::BindGroupDescriptor { label: Some("stamp"), layout: &sp.get_bind_group_layout(0), entries: &[e(0, &a.verts), wgpu::BindGroupEntry { binding: 1, resource: a.tlas.as_binding() }] }); Some(GpuParticlesGpu { pipeline, stamp: (sp, sbg), bg, params, table, spawns, grid, systems: std::cell::RefCell::new(systems), host: std::cell::RefCell::new(host), nodes: scene.gpu_particles.iter().map(|p| p.node).collect(), max_cap: scene.gpu_particles.iter().map(|p| p.capacity).max().unwrap_or(0), frame: std::cell::Cell::new(0), fresh: std::cell::Cell::new(true), grid_o, grid_n }) } @@ -2079,7 +2325,9 @@ impl Gpu { self.device.poll(wgpu::PollType::wait_indefinitely())?; timings.push(("upload_ms".into(), t0.elapsed().as_secs_f64() * 1000.0)); let bg_parts = BindParts { nodes: b_nodes, energy: b_energy, texinfo: b_texinfo, tex_view, env_view, sampler, grid_view }; - Ok(Uploaded { bg, bg_parts, tables, accum: b_accum, params: b_params, width: w, height: h, kernel, materials: b_mats, lights: b_lights, objects: b_objects, emitters: b_emit, verts: b_verts, attrs: b_attrs, prev_verts: b_prev, dynamic_first, accel, units: scene.units.clone(), particles, gpu_particles, grids, translucent, media: b_media, vcolor_base: (!scene.gpu_vcolors.is_empty()).then_some(scene.vcolor_base as u64 * 16), media_nodes: scene.media_nodes.clone(), media_hidden: std::cell::RefCell::new(scene.media.iter().map(|m| m.a[1] == crate::scene::MEDIUM_HIDDEN).collect()), instances: b_instances, room, cdf: b_cdf, cdf_words: env_cdf.len(), region_words, env_tex }) + // (The scene's own pipelines, made above: particles, grid programs.) + self.save_pipeline_cache(); + Ok(Uploaded { bg, bg_parts, tables, uniform, accum: b_accum, params: b_params, width: w, height: h, kernel, materials: b_mats, lights: b_lights, objects: b_objects, emitters: b_emit, verts: b_verts, attrs: b_attrs, prev_verts: b_prev, dynamic_first, accel, units: scene.units.clone(), particles, gpu_particles, grids, translucent, media: b_media, vcolor_base: (!scene.gpu_vcolors.is_empty()).then_some(scene.vcolor_base as u64 * 16), media_nodes: scene.media_nodes.clone(), media_hidden: std::cell::RefCell::new(scene.media.iter().map(|m| m.a[1] == crate::scene::MEDIUM_HIDDEN).collect()), instances: b_instances, room, cdf: b_cdf, cdf_words: env_cdf.len(), region_words, env_tex }) } /// Bytes Metal allocates for an AABB BLAS over `count` boxes in `boxes` (its own size query through wgpu-hal; None @@ -2675,6 +2923,10 @@ impl Gpu { entries.retain(|e| !MERGED_TABLE_BINDINGS.contains(&e.binding)); entries.push(e(crate::merged::TABLES_BINDING, t)); } + // The uniform tables: the same copies (their size does not follow a table's: merged.rs `UNIFORM_LEN`). + if let Some(u) = &up.uniform { + bind_uniform_tables(&mut entries, u); + } let bg = self.device.create_bind_group(&wgpu::BindGroupDescriptor { label: None, layout: &self.layout, entries: &entries }); drop(entries); up.bg = bg; @@ -2707,10 +2959,39 @@ impl Gpu { (self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("merged tables"), size, usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false }), caps) } + /// The uniform tables' buffers for these four tables (merged.rs `UNIFORM_TABLES`' order), filled from them as they + /// are now. + fn uniform_copies(&self, tables: [&wgpu::Buffer; 4]) -> [wgpu::Buffer; 4] { + use crate::merged::{TABLES, UNIFORM_LEN, UNIFORM_TABLES}; + let copies: [wgpu::Buffer; 4] = std::array::from_fn(|k| { + let m = &TABLES[UNIFORM_TABLES[k]]; + // (COPY_SRC: read back by the test that holds them to the tables, and by whoever diagnoses a driver.) + self.device.create_buffer(&wgpu::BufferDescriptor { label: Some(m.name), size: m.stride * UNIFORM_LEN[k] as u64, usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::COPY_SRC, mapped_at_creation: false }) + }); + let mut e = self.device.create_command_encoder(&Default::default()); + Self::copy_uniform(&mut e, tables, &copies); + self.queue.submit([e.finish()]); + copies + } + + fn copy_uniform(e: &mut wgpu::CommandEncoder, tables: [&wgpu::Buffer; 4], copies: &[wgpu::Buffer; 4]) { + for (src, dst) in tables.into_iter().zip(copies) { + let bytes = uniform_copy_bytes(src.size(), dst.size()); + if bytes > 0 { + e.copy_buffer_to_buffer(src, 0, dst, 0, bytes); + } + } + } + /// The merged layout: the tables as they are now, copied into the struct's buffer. The renderer writes the tables /// where it always did; the passes that read the scene read this copy, so it is made in their encoder, before - /// them, each frame (the tables are small: a few copies of kilobytes). + /// them, each frame (the tables are small: a few copies of kilobytes). The uniform tables: the same, into the + /// four uniform buffers. After whatever writes a table on the GPU in that encoder (the particle passes write + /// their objects' entries). pub fn sync_tables(&self, up: &Uploaded, e: &mut wgpu::CommandEncoder) { + if let Some(u) = &up.uniform { + Self::copy_uniform(e, uniform_sources(up.table_buffers()), u); + } let (Some(t), Some(caps)) = (&up.tables, &up.kernel.caps) else { return }; let (at, _) = crate::merged::offsets(&crate::merged::TABLES, &caps.tables); for (k, src) in up.table_buffers().into_iter().enumerate() { @@ -3107,6 +3388,149 @@ mod tests { assert!(!hw.contains("// @trace-end") && !hw.contains("return trace_procs(o, d, hit, mode);")); } + /// With nothing set, every platform but Android opens the device as it was before the options existed: no + /// uniform tables, no pipeline cache, the kernel's source the files (`the_default_sources_are_untouched` has its + /// hashes). Android's defaults are the uniform tables and the switches its GPUs needed; the environment overrides + /// either way, and the merged layout never has the uniform tables. + #[test] + fn the_options_default_to_the_device_as_written() { + let vars = |set: &'static [(&'static str, &'static str)]| move |name: &str| set.iter().find(|(n, _)| *n == name).map(|(_, v)| v.to_string()); + let desktop = GpuOptions::from_vars(&vars(&[]), false); + assert_eq!(desktop, GpuOptions::default()); + 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)); + } + let android = GpuOptions::from_vars(&vars(&[]), true); + assert_eq!(android.kernel, KernelSwitches { strict: true, cutout_lite: 1, last_direct: true, light_pick: false }); + assert_eq!((android.uniform_tables, android.pipeline_cache.clone()), (None, None)); + assert!(android.uniform_tables_on(false, true) && !android.uniform_tables_on(true, true)); + // Set: what is asked, on either platform. + let off = GpuOptions::from_vars(&vars(&[("THREERT_UNIFORM_TABLES", "0"), ("THREERT_KERNEL_STRICT", "0"), ("THREERT_CUTOUT_LITE", "0"), ("THREERT_LAST_DIRECT", "0")]), true); + assert_eq!((off.kernel, off.uniform_tables_on(false, true)), (KernelSwitches::default(), false)); + let on = GpuOptions::from_vars(&vars(&[("THREERT_UNIFORM_TABLES", "1"), ("THREERT_CUTOUT_LITE", "2"), ("THREERT_LIGHT_PICK", "1"), ("THREERT_PIPELINE_CACHE", "/data/cache/pipelines")]), false); + assert!(on.uniform_tables_on(false, false) && !on.uniform_tables_on(true, false)); + assert_eq!(on.kernel, KernelSwitches { strict: false, cutout_lite: 2, last_direct: false, light_pick: true }); + assert_eq!(on.pipeline_cache.as_deref(), Some(std::path::Path::new("/data/cache/pipelines"))); + assert_eq!(GpuOptions::from_vars(&vars(&[("THREERT_PIPELINE_CACHE", "")]), true).pipeline_cache, None); + // (And the environment map's light sampling, which is the scene's preparation's: lib.rs `env_nee_on`.) + 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. + #[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, + ]; + 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 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_eq!(lit.matches("radiance += dz;").count(), 1); + assert_eq!(lit.matches("if (take_env && has_env) {").count(), 1); + assert_eq!(lit.contains("let steps = 4u;"), hwrt); + } + } + + /// The uniform tables' buffers: each array fits a uniform binding of 64 KiB, the media's holds every medium a + /// scene may have, and the four are at their tables' bindings. A copy takes what the table and its uniform buffer + /// both hold, in whole words; a scene with more entries than an array is said, by name. + #[test] + fn the_uniform_tables_fit_their_bindings() { + use crate::merged::{TABLES, UNIFORM_LEN, UNIFORM_TABLES}; + let bindings = UNIFORM_TABLES.map(|k| MERGED_TABLE_BINDINGS[k]); + assert_eq!(bindings, [14, 15, 17, 18]); + for (k, &t) in UNIFORM_TABLES.iter().enumerate() { + let m = &TABLES[t]; + assert!(m.decl.starts_with(&format!("@group(0) @binding({}) var {}:", bindings[k], m.name)), "{}", m.decl); + let bytes = m.stride * UNIFORM_LEN[k] as u64; + assert!(bytes <= 65536 && bytes % 16 == 0, "{}: {bytes} bytes", m.name); + } + assert_eq!(UNIFORM_TABLES.map(|t| TABLES[t].name), ["tex_info", "objects", "instances", "media"]); + assert_eq!(UNIFORM_LEN[3] as usize, crate::scene::MAX_MEDIA); + assert_eq!(UNIFORM_LEN[2] as u64, 65536 / TABLES[5].stride); + // A table shorter than its array, one with the edits' room past it, one whose size is not the array's multiple. + assert_eq!([uniform_copy_bytes(512, 65536), uniform_copy_bytes(1 << 20, 65536), uniform_copy_bytes(65600, 65440), uniform_copy_bytes(0, 7168), uniform_copy_bytes(7, 64)], [512, 65536, 65440, 0, 4]); + assert!(uniform_overflow([4096, 4096, 409, 32]).is_empty()); + let notes = uniform_overflow([4096, 4099, 1024, 32]); + assert_eq!(notes.len(), 2); + assert!(notes[0].contains("4099 entries of `objects`") && notes[0].contains("holds 4096") && notes[1].contains("1024 entries of `instances`"), "{notes:?}"); + } + + /// The order of a layout's entries and of a bind group's is not the device's to see: wgpu puts both in binding + /// order before its backend gets them (wgpu-core 30: `bgl::EntryMap::from_entries` sorts a layout's, and + /// `create_bind_group` its entries), and takes two layouts of the same entries as one. So the kernel's layout, + /// whose acceleration structure (binding 13) is pushed last, needs no sorting here for a driver that minds the + /// order. Held on this machine's GPU: a bind group made of one order runs with a pipeline laid out in another, + /// without a validation error, and each binding reads its own buffer. (Skipped without a GPU.) + #[test] + fn wgpu_takes_bind_group_entries_in_any_order() { + let g = match Gpu::new_with(GpuOptions::default()) { + Ok(g) => g, + Err(e) => return eprintln!("skipped: {e}"), + }; + let d = &g.device; + let entry = |binding: u32, ty: wgpu::BufferBindingType| wgpu::BindGroupLayoutEntry { binding, visibility: wgpu::ShaderStages::COMPUTE, ty: wgpu::BindingType::Buffer { ty, has_dynamic_offset: false, min_binding_size: None }, count: None }; + let (uni, read, write) = (entry(0, wgpu::BufferBindingType::Uniform), entry(5, wgpu::BufferBindingType::Storage { read_only: true }), entry(9, wgpu::BufferBindingType::Storage { read_only: false })); + let scope = d.push_error_scope(wgpu::ErrorFilter::Validation); + let ordered = d.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: None, entries: &[uni, read, write] }); + let shuffled = d.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: None, entries: &[write, uni, read] }); + let module = d.create_shader_module(wgpu::ShaderModuleDescriptor { + label: None, + source: wgpu::ShaderSource::Wgsl("@group(0) @binding(0) var a: vec4;\n@group(0) @binding(5) var b: array;\n@group(0) @binding(9) var c: array;\n@compute @workgroup_size(1) fn main() { c[0] = a.x; c[1] = b[0]; }\n".into()), + }); + let layout = d.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: None, bind_group_layouts: &[Some(&ordered)], immediate_size: 0 }); + let pipeline = d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: None, layout: Some(&layout), module: &module, entry_point: Some("main"), compilation_options: Default::default(), cache: None }); + let buf = |words: [u32; 4], usage| d.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: None, contents: bytemuck::cast_slice(&words), usage }); + let (a, b, c) = (buf([11, 0, 0, 0], wgpu::BufferUsages::UNIFORM), buf([22, 0, 0, 0], wgpu::BufferUsages::STORAGE), buf([0; 4], wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC)); + fn e(binding: u32, b: &wgpu::Buffer) -> wgpu::BindGroupEntry<'_> { + wgpu::BindGroupEntry { binding, resource: b.as_entire_binding() } + } + let bg = d.create_bind_group(&wgpu::BindGroupDescriptor { label: None, layout: &shuffled, entries: &[e(9, &c), e(5, &b), e(0, &a)] }); + let out = d.create_buffer(&wgpu::BufferDescriptor { label: None, size: 16, usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false }); + let mut enc = d.create_command_encoder(&Default::default()); + { + let mut pass = enc.begin_compute_pass(&Default::default()); + pass.set_pipeline(&pipeline); + pass.set_bind_group(0, &bg, &[]); + pass.dispatch_workgroups(1, 1, 1); + } + enc.copy_buffer_to_buffer(&c, 0, &out, 0, 16); + g.queue.submit([enc.finish()]); + assert!(pollster::block_on(scope.pop()).is_none(), "a layout's entries in another order are another layout to wgpu"); + out.slice(..).map_async(wgpu::MapMode::Read, |_| {}); + d.poll(wgpu::PollType::wait_indefinitely()).unwrap(); + let words: Vec = bytemuck::cast_slice(&out.slice(..).get_mapped_range().unwrap()).to_vec(); + assert_eq!(words[..2], [11, 22]); + } + #[test] fn kernel_source_finds_its_anchors() { let base = include_str!("../shaders/pathtrace.wgsl"); diff --git a/renderer/src/lib.rs b/renderer/src/lib.rs index eda8fa8..58f3c78 100644 --- a/renderer/src/lib.rs +++ b/renderer/src/lib.rs @@ -30,6 +30,7 @@ pub mod texpack; #[cfg(not(target_arch = "wasm32"))] pub mod view; pub mod physics; +pub mod pipeline_cache; pub mod voxel; pub mod world; @@ -361,6 +362,18 @@ pub fn fog_params(fog: Option<&desc::Fog>) -> [[f32; 4]; 4] { ] } +/// Whether a baked environment map (an HDRI, a `SkyAtmosphere`'s bake) is light-sampled: THREERT_ENV_NEE=1|0, on +/// everywhere but Android. On Adreno 740 the kernel's reads of `env_cdf` for it went to the wrong memory (tens of +/// thousands of GPU page faults a second under a `SkyAtmosphere`, a trace up to twice as slow). Without it the sky +/// still lights through the paths' own rays, with their full weight; the sun is a directional light, sampled as ever. +pub fn env_nee_on() -> bool { + env_nee_of(std::env::var("THREERT_ENV_NEE").ok().as_deref(), cfg!(target_os = "android")) +} + +fn env_nee_of(var: Option<&str>, android: bool) -> bool { + var.map_or(!android, |v| v != "0") +} + /// `prepare` for an offline render: under software traversal every instance flattened unless THREERT_UNITS=1. pub fn prepare(g: &gpu::Gpu, sc: &mut scene::Scene, timings: &mut Vec<(String, f64)>) -> Prepared { prepare_as(g.hwrt, units_on(false), sc, timings) @@ -666,7 +679,8 @@ pub fn prepare_as(hwrt: bool, units: bool, sc: &mut scene::Scene, timings: &mut let env_max = env.texels.iter().map(|t| 0.2126 * t[0] + 0.7152 * t[1] + 0.0722 * t[2]).fold(0.0f32, f32::max).max(0.2126 * sky_color.x + 0.7152 * sky_color.y + 0.0722 * sky_color.z); let em_area: f32 = emitters[..emitter_tris.len()].iter().map(|e| e[3]).sum(); let env_nee_floor = if emitter_tris.is_empty() { 0.0 } else { 0.01 * all_power / em_area.max(1e-6) }; - let env_nee = env_max >= env_nee_floor; + // (A uniform sky is sampled regardless of `env_nee_on`: it reads no table.) + let env_nee = env_max >= env_nee_floor && (!has_env || env_nee_on()); // The kernel's light order (`light_order`). let class = light_class; let order = light_order(&sc.lights); diff --git a/renderer/src/live.rs b/renderer/src/live.rs index 1272f29..cefa8a3 100644 --- a/renderer/src/live.rs +++ b/renderer/src/live.rs @@ -261,9 +261,10 @@ impl Passes { "bloom_h5", "bloom_v5", "bloom_h4", "bloom_v4", "bloom_h3", "bloom_v3", "bloom_h2", "bloom_v2", "bloom_h1", "bloom_v1", "bloom_h0", "bloom_v0", "post"] .into_iter() .map(|name| { - (name, d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some(name), layout: Some(&pl), module: &module, entry_point: Some(name), compilation_options: Default::default(), cache: None })) + (name, d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some(name), layout: Some(&pl), module: &module, entry_point: Some(name), compilation_options: Default::default(), cache: g.cache() })) }) .collect(); + g.save_pipeline_cache(); Passes { layout, pipes } } } @@ -775,6 +776,9 @@ struct EditState { /// Times the room grew (frame log `grows`), and whether the scene bind group must be made again. grows: u32, rebind: bool, + /// The uniform tables (`gpu::Gpu::uniform_tables`): an edit placed entries past what their arrays hold, and it + /// was said (once a scene, as its upload says it of the scene's own). + uniform_said: bool, /// (triangles, records, objects, words, TLAS instances) the scene itself has: the room starts there. base: (usize, usize, usize, usize, usize), /// The first pool triangle (instanced meshes' triangles are counted from it). @@ -1052,6 +1056,7 @@ fn load_scene(g: &gpu::Gpu, world: &World, assets: &Path, presets: &Path, w: u32 dead_meshes: vec![], grows: 0, rebind: false, + uniform_said: false, base: room_base, pool_first: prep.groups.iter().sum(), blobs, @@ -1187,6 +1192,9 @@ pub struct LiveCore { /// Dynamic resolution: consecutive measured frames well over the frame time, and the first frame at the current /// scale. over_budget: (u32, u32), + /// When this view was made, and when its render scale was last checked (`govern_resolution` by the clock). + born: Instant, + governed: Instant, /// THREERT_PROFILE=1: GPU time per pass (timestamp queries), printed every 2 s. profile: Option, /// GPU time per frame without stalls (timestamp queries; None: measured by draining the queue twice a second). @@ -1373,8 +1381,9 @@ impl Raster { depth_stencil: Some(wgpu::DepthStencilState { format: wgpu::TextureFormat::Depth32Float, depth_write_enabled: Some(true), depth_compare: Some(wgpu::CompareFunction::Greater), stencil: Default::default(), bias: Default::default() }), multisample: Default::default(), multiview_mask: None, - cache: None, + cache: g.cache(), }); + g.save_pipeline_cache(); let tex = |format, usage| { d.create_texture(&wgpu::TextureDescriptor { label: Some("raster"), size: wgpu::Extent3d { width: w, height: h, depth_or_array_layers: 1 }, mip_level_count: 1, sample_count: 1, dimension: wgpu::TextureDimension::D2, format, usage, view_formats: &[] }) .create_view(&Default::default()) @@ -1433,6 +1442,13 @@ const COLD_FALLBACK_FRAMES: u32 = 30; /// pavement under the headlights 0.71 -> 0.77. 1024 adds nothing more. const SAMPLE_CEILING: f32 = 256.0; +/// THREERT_GOVERN_CLOCK=1|0 (default 1 on Android): dynamic resolution also checks by the clock, so a view whose +/// frames are far slower than 60 a second comes down to its frame rate in seconds (`LiveCore::govern_resolution`). +fn govern_by_clock() -> bool { + static F: std::sync::OnceLock = std::sync::OnceLock::new(); + *F.get_or_init(|| std::env::var("THREERT_GOVERN_CLOCK").map(|v| v == "1").unwrap_or(cfg!(target_os = "android"))) +} + fn cache_flags() -> u32 { static F: std::sync::OnceLock = std::sync::OnceLock::new(); *F.get_or_init(|| std::env::var("THREERT_CACHE_FLAGS").ok().and_then(|v| v.parse().ok()).unwrap_or(0)) @@ -1632,7 +1648,7 @@ impl LiveCore { // brings it, `gpu::loop_guard_of`.) let guard_read = (gpu::loop_guard_of(!g.hwrt).is_some() && targets.kernel.guard.is_some()) .then(|| (g.device.create_buffer(&wgpu::BufferDescriptor { label: Some("loop guard read"), size: 4, usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false }), Arc::new(AtomicU32::new(0)))); - let mut core = LiveCore { particle_spawns: Default::default(), grid_sources: Default::default(), g, passes, targets, sc, presets: presets.to_path_buf(), assets: assets.to_path_buf(), w, h, seq: 0, aspect: w as f32 / h as f32, cam, prev_cam: cam, frame: 0, history_len: 0.0, reset: true, bloom: ([0.0; 6], [0.0; 6]), vignette: 0.0, post: [0.0; 4], blur: (0.0, 0.0), inflight: Default::default(), gpu_ms: 0.0, gpu_frame: 0, gpu_trace_ms: 0.0, edit_ms: 0.0, instanced_ms: 0.0, edit_split: EditSplit::default(), rebuild_ms: 0.0, fallback: None, fallbacks: 0, bounces: 4, rays: true, view_mode: std::env::var("THREERT_VIEW_MODE").ok().and_then(|v| v.parse().ok()).unwrap_or(0), fixed_dt: None, motion_blur_off: false, triangles, scene_name, time: 0.0, frame_dt: 1.0 / 60.0, last_frame: Instant::now(), post_key: String::new(), render_scale: 1.0, dynamic_res: false, min_scale: MIN_SCALE, display_scale: 1.0, over_budget: (0, 0), profile, timer, raster, cache, restir, restir_prev: None, kernel_bg, prev_span: None, reposed: false, emitters_moved: false, cold_frames: 0, pose_span: None, pose_reposed: false, inst_moving: vec![], inst_moved: false, prev_isz: (0, 0), lit_changed: 0, spark_lights: false, shown_spp: 1, cpu_skin_ms: 0.0, cpu_upload_ms: 0.0, gpu_blas_ms: 0.0, cpu_wait_ms: 0.0, sky_cover: 0.0, guard_trips: guard_read.as_ref().map(|_| 0), guard_read }; + let mut core = LiveCore { particle_spawns: Default::default(), grid_sources: Default::default(), g, passes, targets, sc, presets: presets.to_path_buf(), assets: assets.to_path_buf(), w, h, seq: 0, aspect: w as f32 / h as f32, cam, prev_cam: cam, frame: 0, history_len: 0.0, reset: true, bloom: ([0.0; 6], [0.0; 6]), vignette: 0.0, post: [0.0; 4], blur: (0.0, 0.0), inflight: Default::default(), gpu_ms: 0.0, gpu_frame: 0, gpu_trace_ms: 0.0, edit_ms: 0.0, instanced_ms: 0.0, edit_split: EditSplit::default(), rebuild_ms: 0.0, fallback: None, fallbacks: 0, bounces: 4, rays: true, view_mode: std::env::var("THREERT_VIEW_MODE").ok().and_then(|v| v.parse().ok()).unwrap_or(0), fixed_dt: None, motion_blur_off: false, triangles, scene_name, time: 0.0, frame_dt: 1.0 / 60.0, last_frame: Instant::now(), post_key: String::new(), render_scale: 1.0, dynamic_res: false, min_scale: MIN_SCALE, display_scale: 1.0, over_budget: (0, 0), born: Instant::now(), governed: Instant::now(), profile, timer, raster, cache, restir, restir_prev: None, kernel_bg, prev_span: None, reposed: false, emitters_moved: false, cold_frames: 0, pose_span: None, pose_reposed: false, inst_moving: vec![], inst_moved: false, prev_isz: (0, 0), lit_changed: 0, spark_lights: false, shown_spp: 1, cpu_skin_ms: 0.0, cpu_upload_ms: 0.0, gpu_blas_ms: 0.0, cpu_wait_ms: 0.0, sky_cover: 0.0, guard_trips: guard_read.as_ref().map(|_| 0), guard_read }; core.spark_lights = core.sc.tables.emitter_tris.iter().any(|e| e.nominal); core.sky_cover = core.sc.tables.cloud_cover; core.apply_post(&pp); @@ -1991,8 +2007,9 @@ impl LiveCore { fn put_sky(&mut self, (env, cdf, max): (gpu::EnvMap, Vec, f32)) { let Some(sky) = self.sc.sky.as_ref() else { return }; self.g.write_env(&self.sc.up, &env, &cdf); - // As the load decides it: the sky is light-sampled unless it is negligible next to the emitters (night). - self.sc.base.counts2[1] = (max >= sky.env_nee_floor) as u32; + // As the load decides it: the sky is light-sampled unless it is negligible next to the emitters (night), or a + // baked sky is not light-sampled at all (`env_nee_on`). + self.sc.base.counts2[1] = (max >= sky.env_nee_floor && crate::env_nee_on()) as u32; self.lit_changed = self.frame.max(1); } @@ -3499,6 +3516,14 @@ impl LiveCore { let _ = g.device.poll(wgpu::PollType::wait_indefinitely()); self.gpu_blas_ms += t.elapsed().as_secs_f32() * 1000.0; } + // The uniform tables: entries the edits placed past what a uniform buffer holds read as its last one. + if g.uniform_tables && !edit.uniform_said { + let notes = gpu::uniform_overflow([0, (edit.objects.end * tables.stride) as u64, edit.records.end as u64, 0]); + for note in ¬es { + eprintln!("threert: {note}"); + } + edit.uniform_said = !notes.is_empty(); + } split.objects = ch.added.len() + ch.removed.len(); self.edit_split.add(&split); if emitters_changed { @@ -3638,6 +3663,12 @@ impl LiveCore { /// heavy one otherwise cost up to 30 frames at 35-48 ms (the 3 x 3 district's drive, its far end). Just over the /// frame time waits for the regular check, as before. After a change, only frames rendered at the new scale count: /// GPU time lags the scale by the frames in flight, and acting on the old scale's time overshot to the floor. + /// + /// "Twice a second" and "the first 2 s" are counted in frames (every 30th, the first 120), which is the same thing + /// at 60 frames a second. Where frames are much slower (`govern_by_clock`: a phone's GPU, on which a scene opens + /// at its full scale at 6 frames a second) they are also read off the clock: counted in frames alone the first + /// check that may act at once came after 20 s and the regular one every 5 s, and a game opened as a slide show + /// for half a minute (Adreno 740, OpenStrikeRT: 25 s to its frame rate; 3 s with the clock). pub fn govern_resolution(&mut self, fps: u32) { if !self.dynamic_res || self.gpu_ms <= 0.0 { return; @@ -3645,13 +3676,19 @@ impl LiveCore { if (self.gpu_frame.wrapping_sub(self.over_budget.1) as i32) < 0 { return; } + let clock = govern_by_clock(); // Not in the first 2 s: those frames run while caches fill (Kabukichō's first frames at 0.7 take ~47 ms), and // the regular check settles the start as before. - self.over_budget.0 = if self.frame >= 120 && self.gpu_ms > 1250.0 / fps as f32 && self.render_scale > self.min_scale { self.over_budget.0 + 1 } else { 0 }; + let warm = self.frame >= 120 || (clock && self.born.elapsed().as_secs_f32() >= 2.0); + self.over_budget.0 = if warm && self.gpu_ms > 1250.0 / fps as f32 && self.render_scale > self.min_scale { self.over_budget.0 + 1 } else { 0 }; let urgent = self.over_budget.0 >= 2; - if self.frame % 30 != 1 && !urgent { + // (By the clock only when 30 frames take longer than 0.75 s: at 40 frames a second and above the frame count + // comes first, as it always did.) + let due = self.frame % 30 == 1 || (clock && self.governed.elapsed().as_secs_f32() >= 0.75); + if !due && !urgent { return; } + self.governed = Instant::now(); self.over_budget.0 = 0; let before = self.render_scale; let budget = 1000.0 / fps as f32 * 0.9; @@ -3855,10 +3892,6 @@ impl LiveCore { p.counts[3] = p.counts[3].min(2); g.queue.write_buffer(&self.sc.up.params, 0, bytemuck::bytes_of(&p)); let mut e = g.device.create_command_encoder(&Default::default()); - // (The merged layout: the tables as this frame's edits left them, where the kernel reads them; merged.rs.) - if k == 0 { - g.sync_tables(&self.sc.up, &mut e); - } // Particles: placed for this frame's time and camera, then the dynamic BLASes rebuilt (before the raster and // the trace). // The grid programs (density grids changed on the GPU), before anything reads the grids. @@ -3894,6 +3927,11 @@ impl LiveCore { let blas_ts = self.profile.as_ref().map(|p| (&p.queries, n * 2 + 2)); g.run_particles(&self.sc.up, cam.pos.to_array(), cam.fwd.to_array(), shader_time(self.time), self.frame_dt, b.time[2], b.time[1], &mut e, pass_ts, blas_ts); } + // (The merged layout and the uniform tables: the tables as this frame's edits and the particle passes above + // left them, where the kernel reads them; merged.rs.) + if k == 0 { + g.sync_tables(&self.sc.up, &mut e); + } if let (Some(r), 0, false) = (&self.raster, k, self.sc.instanced) { let tw = self.timer.as_ref().zip(timed).map(|(t, slot)| wgpu::RenderPassTimestampWrites { query_set: &t.queries, beginning_of_pass_write_index: Some(slot as u32 * 8 + 4), end_of_pass_write_index: Some(slot as u32 * 8 + 5) }); r.record(g, &mut e, &self.sc.up.verts, self.sc.tris.len() as u32, &cam, self.aspect, (iw, ih), jitter, tw); @@ -4859,4 +4897,95 @@ mod instanced_tests { let r = core.sc.gpu_instances[first + 1]; assert_eq!((r.m[0][3], r.m[1][3], r.p[0][3], r.p[1][3]), (3.0, 1.0, 2.0, 0.0), "moved: last frame's rows kept"); } + + /// The uniform tables (`gpu::GpuOptions::uniform_tables`: what a device is opened with on Android), on this + /// machine's GPU: a lit scene with an instanced mesh renders offline to the default device's image, bit for bit; + /// live, the four uniform buffers hold what their tables hold when a frame is traced, and after an edit in place + /// what the tables hold then. (Skipped without a GPU.) + #[test] + fn the_uniform_tables_render_what_the_tables_do() { + let open = |uniform: bool| gpu::Gpu::new_with(gpu::GpuOptions { uniform_tables: Some(uniform), ..Default::default() }); + let (plain, uniform) = match (open(false), open(true)) { + (Ok(a), Ok(b)) => (a, b), + (Err(e), _) | (_, Err(e)) => return eprintln!("skipped: {e}"), + }; + if uniform.merged { + return eprintln!("skipped: the merged layout has no uniform tables"); + } + assert!(uniform.uniform_tables && !plain.uniform_tables); + let mut w = World::default(); + let mat = |color: f32, emissive: f32| serde_json::from_value(serde_json::json!({ "name": "m", "color": [color, color * 0.5, color], "opacity": 1, "metalness": 0, "roughness": 1, "emissive": [1, 1, 1], "emissiveIntensity": emissive, "doubleSided": true })).unwrap(); + w.set_material(1, mat(0.8, 0.0)); + w.set_material(5, mat(0.0, 40.0)); + let quad = |s: f32| Geometry { positions: vec![-s, -s, 0.0, s, -s, 0.0, s, s, 0.0, -s, s, 0.0], normals: [0.0, 0.0, 1.0].repeat(4), uvs: vec![0.0; 8], indices: vec![0, 1, 2, 0, 2, 3], groups: vec![], colors: vec![] }; + w.set_geometry(2, quad(1.0)).unwrap(); + w.set_geometry(3, quad(0.25)).unwrap(); + let mesh = |geometry, material, own_blas| Kind::Mesh { geometry, material, own_blas, materials: vec![] }; + // A wall, a small movable panel in front of it, and an emissive quad that lights both from the side. + w.set_object(10, "wall".into(), None, mesh(2, 1, false), true, false).unwrap(); + w.set_object(11, "panel".into(), None, mesh(3, 1, true), true, true).unwrap(); + w.set_transform(11, Vec3::new(-0.4, 0.0, 0.5), glam::Quat::IDENTITY, Vec3::ONE).unwrap(); + w.set_object(12, "lamp".into(), None, mesh(3, 5, false), true, false).unwrap(); + w.set_transform(12, Vec3::new(0.6, 0.3, 1.5), glam::Quat::from_rotation_y(2.6), Vec3::ONE).unwrap(); + let cam = serde_json::from_value(serde_json::json!({ "fov": 50, "near": 0.1, "far": 100, "exposure": { "mode": "manual", "ev100": 0 } })).unwrap(); + w.set_object(4, "camera".into(), None, Kind::Camera(cam), true, false).unwrap(); + w.set_transform(4, Vec3::new(0.0, 0.0, 5.0), glam::Quat::IDENTITY, Vec3::ONE).unwrap(); + w.set_settings(crate::world::Settings { camera: Some(4), ..Default::default() }); + let presets = crate::default_presets_dir(); + + // Offline: the same bytes. + let offline = |g: &gpu::Gpu| { + let mut sc = crate::scene::build(&w, Path::new("."), &presets, None).unwrap(); + (sc.width, sc.height) = (64, 36); + let prep = crate::prepare(g, &mut sc, &mut vec![]); + let frame = g.render(&sc, &prep.tri_verts, prep.groups, bytemuck::cast_slice(&prep.nodes), &prep.emitters, &prep.env, &prep.env_cdf, &prep.energy, prep.params, 16, &mut vec![]).unwrap(); + frame.color.iter().flatten().map(|v| v.to_bits()).collect::>() + }; + let (a, b) = (offline(&plain), offline(&uniform)); + let lit = a.iter().filter(|v| f32::from_bits(**v) > 0.0).count(); + assert!(lit > a.len() / 16, "the scene is lit ({lit} of {} numbers)", a.len()); + let differ = a.iter().zip(&b).filter(|(x, y)| x != y).count(); + assert_eq!(differ, 0, "{differ} of {} numbers differ between the storage tables and the uniform ones", a.len()); + drop(plain); + + // Live: the uniform buffers follow the tables. + let mut core = LiveCore::new(uniform, &w, Path::new("."), &presets, 64, 36).unwrap(); + w.take_changes(); + let read = |core: &LiveCore, buf: &wgpu::Buffer, bytes: u64| -> Vec { + let d = &core.g.device; + let to = d.create_buffer(&wgpu::BufferDescriptor { label: None, size: bytes, usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false }); + let mut e = d.create_command_encoder(&Default::default()); + e.copy_buffer_to_buffer(buf, 0, &to, 0, bytes); + core.g.queue.submit([e.finish()]); + to.slice(..).map_async(wgpu::MapMode::Read, |_| {}); + d.poll(wgpu::PollType::wait_indefinitely()).unwrap(); + let bytes = to.slice(..).get_mapped_range().unwrap().to_vec(); + bytes + }; + // Each table's bytes and its uniform copy's, over what both hold. + let pairs = |core: &LiveCore| -> Vec<(Vec, Vec)> { + let up = &core.sc.up; + let copies = up.uniform.as_ref().expect("the uniform tables' buffers"); + [&up.bg_parts.texinfo, &up.objects, &up.instances, &up.media].into_iter().zip(copies).map(|(t, c)| { + let n = t.size().min(c.size()) & !3; + (read(core, t, n), read(core, c, n)) + }).collect() + }; + core.render(1, |_, _| {}); + let first = pairs(&core); + for (k, (table, copy)) in first.iter().enumerate() { + assert!(!table.is_empty() && table == copy, "table {k} and its uniform buffer differ after a frame"); + } + assert!(first[1].0.iter().any(|b| *b != 0), "the objects' table has the scene's objects"); + // The panel moves in place: its object's entry and its instance's record change, and the next frame's + // uniform buffers have them. + w.set_transform(11, Vec3::new(0.3, 0.2, 0.75), glam::Quat::IDENTITY, Vec3::ONE).unwrap(); + assert!(!core.commit(&mut w).unwrap(), "moved in place"); + core.render(1, |_, _| {}); + let moved = pairs(&core); + for (k, (table, copy)) in moved.iter().enumerate() { + assert!(table == copy, "table {k} and its uniform buffer differ after an edit"); + } + assert!(moved[1].1 != first[1].1 || moved[2].1 != first[2].1, "the move reached the uniform buffers"); + } } diff --git a/renderer/src/merged.rs b/renderer/src/merged.rs index e5caaa9..19e8605 100644 --- a/renderer/src/merged.rs +++ b/renderer/src/merged.rs @@ -186,7 +186,7 @@ struct Rule { name: &'static str, /// What `name[` becomes, and what goes before the `]` that closes it. open: String, - close: &'static str, + close: String, /// What `arrayLength(&name)` becomes. len: String, } @@ -223,7 +223,7 @@ fn rewrite(src: &str, rules: &[Rule]) -> String { let mut out: Vec = Vec::with_capacity(b.len() + b.len() / 16); // Blocks and brackets open; the brackets a rewritten index closes with more than `]`. let (mut depth, mut brackets, mut parens) = (0usize, 0usize, 0usize); - let mut closes: Vec<(usize, &'static str)> = vec![]; + let mut closes: Vec<(usize, &str)> = vec![]; // Names shadowed, each with the block depth it lives in; those declared in the statement being read. let mut shadows: Vec<(usize, usize)> = vec![]; let mut pending: Vec<(usize, usize)> = vec![]; @@ -320,7 +320,7 @@ fn rewrite(src: &str, rules: &[Rule]) -> String { out.extend_from_slice(rules[r].open.as_bytes()); brackets += 1; if !rules[r].close.is_empty() { - closes.push((brackets, rules[r].close)); + closes.push((brackets, &rules[r].close)); } last = b'['; i = after + 1; @@ -389,11 +389,50 @@ pub fn kernel(src: &str, caps: &Caps) -> String { }; take(&mut s, m.decl, &with); } - let rule = |var: &str, m: &Member, cap: u32| Rule { name: m.name, open: format!("{var}.{}[", m.name), close: "", len: format!("{cap}u") }; + let rule = |var: &str, m: &Member, cap: u32| Rule { name: m.name, open: format!("{var}.{}[", m.name), close: String::new(), len: format!("{cap}u") }; let rules: Vec = TABLES.iter().zip(caps.tables).map(|(m, c)| rule("MT", m, c)).chain(KERNEL_LIVE.iter().zip(caps.live).map(|(m, c)| rule("ML", m, c))).collect(); rewrite(&s, &rules) } +/// The uniform tables (`gpu::GpuOptions::uniform_tables`; docs/RUNTIME.md "Android"): which of `TABLES` the kernel +/// reads from uniform buffers instead of storage buffers, in the layout as written. Qualcomm's Adreno 740 driver +/// (Vulkan, Android 13) gave the kernel other buffers' contents for these four: a model's textures came out as +/// another table's numbers, instance records sent hits to the wrong triangles, an emissive object's power read as +/// something else's and lit nothing. As uniform buffers they read right. They are the four of the kernel's read-only +/// tables that are bounded (the others: `uniform_tables`). +pub const UNIFORM_TABLES: [usize; 4] = [3, 4, 5, 6]; +/// Their arrays' lengths (a uniform buffer's array is sized): what 64 KiB hold, the size every device gives a uniform +/// binding (WebGPU's `maxUniformBufferBindingSize`, Vulkan's least `maxUniformBufferRange` on the phones this is +/// for), and for the media what a scene may have (`scene::MAX_MEDIA`). In a scene's terms: 2,048 textures, 1,365 +/// objects (1,024 with instance colours), 409 instance records, every medium. +pub const UNIFORM_LEN: [u32; 4] = [4096, 4096, 409, 32]; + +/// The kernel (as `kernel` takes it: a scene's variant, its functions spliced in) with the uniform tables: the four +/// declarations as sized uniform arrays at their own bindings, and every index into them held to the array +/// (`min(u32(i), last)`). A scene with more entries than an array holds reads the last one for the rest (the renderer +/// says so when it uploads one, `Gpu::upload`); nothing is read outside the buffer, which the kernel's unchecked +/// indexing would otherwise do. Everything else is the source as written. +/// +/// The kernel's other storage buffers stay storage buffers, each for its reason: +/// - `materials` (288 B each), `lights` (80 B): the device drew right with them as they are, and 64 KiB would hold +/// 227 materials; +/// - `emitters` (the emitter CDF and the light tree), `env_cdf` (the environment's CDF, then the procedural +/// geometry's data), `tri_verts`, `attrs`, `nodes`, `prev_verts` (and the vertex colours after them): megabytes, +/// which no uniform binding holds; +/// - `accum` and group 1 (the radiance cache, the reservoirs, the trace's outputs): written by the kernel. +pub fn uniform_tables(src: &str) -> String { + let mut s = src.to_string(); + let mut rules = vec![]; + for (&k, &len) in UNIFORM_TABLES.iter().zip(&UNIFORM_LEN) { + let m = &TABLES[k]; + let sized = m.decl.replace("var", "var").replace(&format!("array<{}>", m.ty), &format!("array<{}, {len}>", m.ty)); + assert!(sized.contains("var") && sized.ends_with(&format!(", {len}>;")), "uniform tables: the declaration {:?}", m.decl); + take(&mut s, m.decl, &sized); + rules.push(Rule { name: m.name, open: format!("{}[min(u32(", m.name), close: format!("), {}u)", len - 1), len: format!("{len}u") }); + } + rewrite(&s, &rules) +} + /// The live passes (live.wgsl) in the merged layout: sample_px, out_px and ae_bins as they are, the slabs, and the /// offsets `LO`. pub fn live(src: &str) -> String { @@ -414,7 +453,7 @@ pub fn live(src: &str) -> String { // After the uniforms' declaration (the file's first binding). let anchor = "@group(0) @binding(0) var L: LiveParams;\n"; take(&mut s, anchor, &format!("{anchor}{head}")); - let rules: Vec = LIVE_BUFS.iter().map(|b| Rule { name: b.name, open: format!("{}[LO.{}.x + (", SLABS[b.slab].name, b.name), close: ")", len: format!("LO.{}.y", b.name) }).collect(); + let rules: Vec = LIVE_BUFS.iter().map(|b| Rule { name: b.name, open: format!("{}[LO.{}.x + (", SLABS[b.slab].name, b.name), close: ")".into(), len: format!("LO.{}.y", b.name) }).collect(); rewrite(&s, &rules) } @@ -563,6 +602,62 @@ mod tests { } } + /// Every kernel variant with the uniform tables: valid WGSL; the four tables are uniform buffers at their own + /// bindings, sized as the Rust side makes their buffers; the other sixteen storage buffers are as written; and + /// every index into the four is held to its array, so nothing is read outside a uniform buffer. + #[test] + fn the_uniform_tables_are_four_sized_uniform_arrays() { + // A name's index sites in code (not comments): `name[` that is not part of a longer name or a member. + let sites = |src: &str, name: &str| { + let code: String = src.lines().map(|l| l.split("//").next().unwrap_or("")).collect::>().join("\n"); + let open = format!("{name}["); + code.match_indices(&open).filter(|(k, _)| !code[..*k].ends_with(|c: char| c.is_alphanumeric() || c == '_' || c == '.')).count() + }; + let globals = |m: &naga::Module, uniform: bool| -> Vec<(String, u32, u32)> { + let mut out: Vec<(String, u32, u32)> = m + .global_variables + .iter() + .filter(|(_, v)| if uniform { v.space == naga::AddressSpace::Uniform } else { matches!(v.space, naga::AddressSpace::Storage { .. }) }) + .map(|(_, v)| { + let b = v.binding.as_ref().expect("a bound global"); + let bytes = match &m.types[v.ty].inner { + naga::TypeInner::Array { size: naga::ArraySize::Constant(n), stride, .. } => n.get() * stride, + _ => 0, + }; + (v.name.clone().unwrap_or_default(), b.group * 100 + b.binding, bytes) + }) + .collect(); + out.sort_by_key(|g| g.1); + out + }; + let (written, _) = parse(KERNEL); + 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 (m, info) = parse(&src); + assert_eq!(storage_buffers(&m, &info).1, 16, "features {features:#x}"); + let tables: Vec<(String, u32, u32)> = globals(&m, true).into_iter().filter(|g| g.0 != "P" && g.0 != "energy_lut").collect(); + let want: Vec<(String, u32, u32)> = UNIFORM_TABLES.iter().zip(UNIFORM_LEN).zip([14, 15, 17, 18]).map(|((&k, len), binding)| (TABLES[k].name.to_string(), binding, (TABLES[k].stride * len as u64) as u32)).collect(); + assert_eq!(tables, want, "features {features:#x}"); + // The storage buffers that stay: the same names at the same bindings as the kernel as written. + let kept = |m: &naga::Module| globals(m, false).into_iter().map(|g| (g.0, g.1)).filter(|g| !want.iter().any(|w| w.0 == g.0)).collect::>(); + assert_eq!(kept(&m), kept(&written), "features {features:#x}"); + for (&k, len) in UNIFORM_TABLES.iter().zip(UNIFORM_LEN) { + let name = TABLES[k].name; + let held = format!("{name}[min(u32("); + assert!(sites(&base, name) >= 3, "{name}: the kernel indexes it"); + assert_eq!(src.matches(&held).count(), sites(&base, name), "{name}: every index is held (features {features:#x})"); + assert_eq!(sites(&src, name), src.matches(&held).count(), "{name}: no index is left as written"); + assert!(src.contains(&format!("), {}u)]", len - 1))); + } + } + // One index, as it is rewritten: an expression with brackets of its own, and a table inside another's index. + let src = "@group(0) @binding(14) var tex_info: array>;\n@group(0) @binding(15) var objects: array>;\n@group(0) @binding(17) var instances: array;\n@group(0) @binding(18) var media: array;\nfn f(a: array, k: i32) -> f32 { return objects[a[1] * 3u + 1u].w + tex_info[u32(objects[k].x)].y + media[0].b.x; }\n"; + let out = uniform_tables(src); + assert!(out.starts_with("@group(0) @binding(14) var tex_info: array, 4096>;\n@group(0) @binding(15) var objects: array, 4096>;\n@group(0) @binding(17) var instances: array;\n@group(0) @binding(18) var media: array;\n"), "{out}"); + assert!(out.ends_with("{ return objects[min(u32(a[1] * 3u + 1u), 4095u)].w + tex_info[min(u32(u32(objects[min(u32(k), 4095u)].x)), 4095u)].y + media[min(u32(0), 31u)].b.x; }\n"), "{out}"); + } + /// The same for the live passes (28 as written). #[test] fn the_merged_live_passes_fit_eight_storage_buffers() { @@ -595,7 +690,7 @@ mod tests { /// declaration to the end of its block; comments are not read. #[test] fn locals_that_shadow_a_buffer_are_left_alone() { - let rules = [Rule { name: "refl", open: "ML.refl[".into(), close: "", len: "7u".into() }, Rule { name: "fa", open: "MV4[LO.fa.x + (".into(), close: ")", len: "LO.fa.y".into() }]; + let rules = [Rule { name: "refl", open: "ML.refl[".into(), close: String::new(), len: "7u".into() }, Rule { name: "fa", open: "MV4[LO.fa.x + (".into(), close: ")".into(), len: "LO.fa.y".into() }]; let src = "struct S { refl: f32 };\nfn f(i: u32) -> f32 {\n let a = refl[i].x; // refl[0] }\n {\n let refl = refl[i + 1u].y;\n if (refl > 0.0) { return refl; }\n }\n var s: S;\n return s.refl + refl[arrayLength(&refl) - 1u].x + fa[fa[i].x].y;\n}\nfn g(refl: array) -> f32 { return refl[1]; }\nfn h() -> f32 { return refl[2].x; }\n"; let want = "struct S { refl: f32 };\nfn f(i: u32) -> f32 {\n let a = ML.refl[i].x; // refl[0] }\n {\n let refl = ML.refl[i + 1u].y;\n if (refl > 0.0) { return refl; }\n }\n var s: S;\n return s.refl + ML.refl[7u - 1u].x + MV4[LO.fa.x + (MV4[LO.fa.x + (i)].x)].y;\n}\nfn g(refl: array) -> f32 { return refl[1]; }\nfn h() -> f32 { return ML.refl[2].x; }\n"; assert_eq!(rewrite(src, &rules), want); @@ -629,7 +724,7 @@ mod tests { #[test] #[should_panic(expected = "`refl` is used in a way the generator does not know (line 2)")] fn a_use_the_generator_does_not_know_stops_it() { - rewrite("fn f() {\n let p = &refl;\n}\n", &[Rule { name: "refl", open: "ML.refl[".into(), close: "", len: "7u".into() }]); + rewrite("fn f() {\n let p = &refl;\n}\n", &[Rule { name: "refl", open: "ML.refl[".into(), close: String::new(), len: "7u".into() }]); } #[test] diff --git a/renderer/src/pipeline_cache.rs b/renderer/src/pipeline_cache.rs new file mode 100644 index 0000000..6c20f47 --- /dev/null +++ b/renderer/src/pipeline_cache.rs @@ -0,0 +1,239 @@ +//! Compiled pipelines kept on disk between runs (`gpu::GpuOptions::pipeline_cache`; THREERT_PIPELINE_CACHE=). +//! A phone's driver compiles a kernel for seconds (Adreno 740: ten), and every variant a game uses again each time it +//! is opened; with the cache the second opening reads what the first compiled (docs/RUNTIME.md "Android"). Off unless +//! a host names a directory; only where the backend has such a cache (Vulkan: wgpu's `PIPELINE_CACHE` feature). +//! +//! The file, one per adapter: +//! - **Its name** is wgpu's key for the adapter (`wgpu::util::pipeline_cache_key`: backend, vendor, device), so two +//! GPUs never read each other's. Inside, wgpu's own header holds the driver's cache identity: data from another +//! driver version is not used (`fallback`: an empty cache then, and the next save replaces the file). +//! - **Its bytes** are what `wgpu::PipelineCache::get_data` returned, sealed: a magic, the length and the data's +//! SHA-256 before them. A file that is cut short, damaged or not ours is not handed to the driver (which need not +//! check what it is given); it counts as no file. +//! - **Two processes may share it** (two games of one app, a game opened twice). A save writes the whole file under +//! a name of its own (process id, a count) beside it and renames it over the old one: a reader finds the old file +//! or the new one, whole, and a process stopped halfway leaves the old one. The last save wins; what the other +//! process compiled meanwhile is compiled once more, by whoever misses it, and saved then. +//! - **It may be deleted** at any time, and is not read past `MAX_BYTES` (a cache nobody cleared starts again). + +use sha2::{Digest, Sha256}; +use std::path::{Path, PathBuf}; +use std::sync::atomic::{AtomicBool, AtomicU64, Ordering}; + +const MAGIC: &[u8; 8] = b"3RTPIPE1"; +const HEADER: usize = 8 + 8 + 32; +/// The largest file read: past it the cache starts empty (every kernel variant of every game ever opened adds to it). +pub const MAX_BYTES: u64 = 256 << 20; +/// A save's own file that old was left by a process that stopped: removed when the cache is opened. +const STALE: std::time::Duration = std::time::Duration::from_secs(600); + +/// The file's bytes for a cache's data. +pub fn seal(data: &[u8]) -> Vec { + let mut out = Vec::with_capacity(HEADER + data.len()); + out.extend_from_slice(MAGIC); + out.extend_from_slice(&(data.len() as u64).to_le_bytes()); + out.extend_from_slice(&Sha256::digest(data)); + out.extend_from_slice(data); + out +} + +/// The cache's data in a file's bytes, if they are what `seal` wrote. +pub fn unseal(file: &[u8]) -> Option<&[u8]> { + let (head, data) = file.split_at_checked(HEADER)?; + let len = u64::from_le_bytes(head[8..16].try_into().ok()?); + (&head[..8] == MAGIC && len == data.len() as u64 && head[16..] == Sha256::digest(data)[..]).then_some(data) +} + +/// The cache file's name for an adapter (None: its backend has no pipeline cache). +pub fn file_name(info: &wgpu::AdapterInfo) -> Option { + wgpu::util::pipeline_cache_key(info).map(|key| format!("{key}.bin")) +} + +/// `bytes` as the whole of `path`: written beside it under a name no other save has, then renamed over it. +fn write_whole(path: &Path, bytes: &[u8]) -> std::io::Result<()> { + static SAVES: AtomicU64 = AtomicU64::new(0); + if let Some(dir) = path.parent() { + std::fs::create_dir_all(dir)?; + } + let mut name = path.file_name().unwrap_or_default().to_os_string(); + name.push(format!(".{}.{}.tmp", std::process::id(), SAVES.fetch_add(1, Ordering::Relaxed))); + let tmp = path.with_file_name(name); + let done = std::fs::write(&tmp, bytes).and_then(|_| std::fs::rename(&tmp, path)); + if done.is_err() { + let _ = std::fs::remove_file(&tmp); + } + done +} + +/// Removes the saves' own files beside `path` that a stopped process left (`STALE`). +fn sweep(path: &Path) { + let (Some(dir), Some(name)) = (path.parent(), path.file_name().and_then(|n| n.to_str())) else { return }; + let Ok(entries) = std::fs::read_dir(dir) else { return }; + for e in entries.flatten() { + let ours = e.file_name().to_str().is_some_and(|n| n.starts_with(name) && n.ends_with(".tmp")); + let old = e.metadata().and_then(|m| m.modified()).ok().and_then(|t| t.elapsed().ok()).is_some_and(|age| age > STALE); + if ours && old { + let _ = std::fs::remove_file(e.path()); + } + } +} + +/// The sealed data of the file at `path`, if there is one to use. +fn load(path: &Path) -> Option> { + if std::fs::metadata(path).ok()?.len() > MAX_BYTES { + return None; + } + let file = std::fs::read(path).ok()?; + unseal(&file).map(|d| d.to_vec()) +} + +/// A device's pipeline cache and its file. +pub struct PipelineCache { + cache: wgpu::PipelineCache, + path: PathBuf, + /// A pipeline was made with the cache since the last save. + dirty: AtomicBool, + /// The SHA-256 of the data last read or written: a save of the same data writes nothing. + saved: std::sync::Mutex<[u8; 32]>, +} + +impl PipelineCache { + /// The feature the device must be made with for `open`, if the adapter has a cache to keep in `dir`. + pub fn feature(adapter: &wgpu::Adapter, info: &wgpu::AdapterInfo) -> Option { + (adapter.features().contains(wgpu::Features::PIPELINE_CACHE) && file_name(info).is_some()).then_some(wgpu::Features::PIPELINE_CACHE) + } + + /// The cache of `device` (made with `feature`), filled from its file in `dir` when that holds this adapter's data. + pub fn open(device: &wgpu::Device, dir: &Path, info: &wgpu::AdapterInfo) -> Option { + let path = dir.join(file_name(info)?); + sweep(&path); + let data = load(&path); + // Safety: `data` is what `get_data` returned for an adapter of this key (the file's name), whole (`unseal`); + // wgpu checks its own header against this adapter and driver, and with `fallback` starts empty otherwise. + let cache = unsafe { device.create_pipeline_cache(&wgpu::PipelineCacheDescriptor { label: Some("threert"), data: data.as_deref(), fallback: true }) }; + let saved = data.map_or([0; 32], |d| Sha256::digest(&d).into()); + Some(PipelineCache { cache, path, dirty: AtomicBool::new(false), saved: std::sync::Mutex::new(saved) }) + } + + /// The cache, for the descriptor of a pipeline about to be made (the next `save` writes it). + pub fn for_pipeline(&self) -> &wgpu::PipelineCache { + self.dirty.store(true, Ordering::Relaxed); + &self.cache + } + + /// The cache to its file, if a pipeline was made since the last save and the data changed (a pipeline the cache + /// already had changes nothing). A failure to write is not an error: the next run compiles again. + pub fn save(&self) { + if !self.dirty.swap(false, Ordering::Relaxed) { + return; + } + let Some(data) = self.cache.get_data() else { return }; + let hash: [u8; 32] = Sha256::digest(&data).into(); + let mut saved = self.saved.lock().unwrap(); + if *saved != hash && write_whole(&self.path, &seal(&data)).is_ok() { + *saved = hash; + } + } +} + +#[cfg(test)] +mod tests { + use super::*; + + fn scratch(name: &str) -> PathBuf { + let dir = std::env::temp_dir().join(format!("threert-pipeline-cache-{name}-{}", std::process::id())); + let _ = std::fs::remove_dir_all(&dir); + dir + } + + /// A file is used only if it is whole and ours: cut short, a byte changed, another length, another magic, or the + /// bare data without the seal all count as no file. + #[test] + fn only_a_whole_file_is_used() { + let data: Vec = (0..1000u32).flat_map(|k| k.to_le_bytes()).collect(); + let file = seal(&data); + assert_eq!(unseal(&file), Some(&data[..])); + assert_eq!(unseal(&seal(&[])), Some(&[][..])); + assert_eq!(unseal(&file[..file.len() - 1]), None); + assert_eq!(unseal(&file[..HEADER - 1]), None); + assert_eq!(unseal(&[]), None); + assert_eq!(unseal(&data), None); + let mut flipped = file.clone(); + *flipped.last_mut().unwrap() ^= 1; + assert_eq!(unseal(&flipped), None); + let mut longer = file.clone(); + longer.push(0); + assert_eq!(unseal(&longer), None); + let mut other = file; + other[0] = b'4'; + assert_eq!(unseal(&other), None); + } + + /// The file is named for the adapter (wgpu's key: backend, vendor, device), so another GPU's is another file; a + /// backend without a pipeline cache has none. + #[test] + fn the_file_is_named_for_the_adapter() { + let mut adreno = wgpu::AdapterInfo::new(wgpu::DeviceType::IntegratedGpu, wgpu::Backend::Vulkan); + (adreno.vendor, adreno.device) = (0x5143, 0x43050a01); + let name = file_name(&adreno).unwrap(); + assert_eq!(name, format!("{}.bin", wgpu::util::pipeline_cache_key(&adreno).unwrap())); + assert!(name.contains(&0x5143.to_string()) && name.contains(&0x43050a01u32.to_string()), "{name}"); + let mut other = adreno.clone(); + other.device += 1; + assert_ne!(file_name(&other), Some(name)); + assert_eq!(file_name(&wgpu::AdapterInfo::new(wgpu::DeviceType::IntegratedGpu, wgpu::Backend::Metal)), None); + } + + /// Saves from several threads at once (as two processes would: each under a name of its own, renamed over the + /// file) leave one whole file that is one of theirs, and no file of a save behind; a reader meanwhile never sees + /// a file that is not whole. + #[test] + fn concurrent_saves_leave_a_whole_file() { + let dir = scratch("saves"); + let path = dir.join("cache.bin"); + let blobs: Vec> = (0..8u8).map(|k| vec![k; 200_000 + k as usize]).collect(); + write_whole(&path, &seal(&blobs[0])).unwrap(); + std::thread::scope(|s| { + for b in &blobs { + let path = &path; + s.spawn(move || { + for _ in 0..20 { + write_whole(path, &seal(b)).unwrap(); + } + }); + } + for _ in 0..200 { + let file = std::fs::read(&path).unwrap(); + assert!(unseal(&file).is_some(), "a reader saw a file that is not whole ({} bytes)", file.len()); + } + }); + let last = load(&path).expect("a whole file"); + assert!(blobs.contains(&last)); + assert_eq!(std::fs::read_dir(&dir).unwrap().count(), 1, "the saves' own files are gone"); + let _ = std::fs::remove_dir_all(&dir); + } + + /// What a stopped process left is swept when it is old, a file that is not whole is not loaded, and neither is one + /// past the size limit's check (a directory in its place: no file). + #[test] + fn a_damaged_file_and_a_stopped_save_are_ignored() { + let dir = scratch("stale"); + let path = dir.join("cache.bin"); + write_whole(&path, &seal(b"pipelines")).unwrap(); + assert_eq!(load(&path).as_deref(), Some(&b"pipelines"[..])); + // A save in progress (young) stays; one from long ago goes. + let (young, old) = (dir.join("cache.bin.1.0.tmp"), dir.join("cache.bin.2.0.tmp")); + std::fs::write(&young, b"half").unwrap(); + std::fs::write(&old, b"half").unwrap(); + let long_ago = std::time::SystemTime::now() - 2 * STALE; + std::fs::File::options().write(true).open(&old).unwrap().set_modified(long_ago).unwrap(); + sweep(&path); + assert!(young.exists() && !old.exists() && path.exists()); + // Cut short on disk (a copy that failed, a full disk): no file. + let whole = std::fs::read(&path).unwrap(); + std::fs::write(&path, &whole[..whole.len() - 3]).unwrap(); + assert_eq!(load(&path), None); + assert_eq!(load(&dir.join("missing.bin")), None); + let _ = std::fs::remove_dir_all(&dir); + } +}