Merge remote-tracking branch 'origin/main' into fix/environment-light-no-effect-cast-shadows

crates/gpu/src/context.rs

@@ -351,6 +351,7 @@ pub struct GpuContext {
     pub queue: wgpu::Queue,
     pub(crate) name: String,
     pub(crate) max_dim: u32,
+    pub(crate) median_kernel: crate::fx_noise::MedianKernel,
     bgl: wgpu::BindGroupLayout,
     bgl_ext: wgpu::BindGroupLayout,
     pipelines: HashMap<&'static str, wgpu::ComputePipeline>,
@@ -481,6 +482,9 @@ impl GpuContext {
     /// [`Self::device_failure_notifier`] for immediate readback cancellation.
     pub fn new(adapter: &wgpu::Adapter, device: wgpu::Device, queue: wgpu::Queue) -> Result<GpuContext, String> {
         let info = adapter.get_info();
+        let median_kernel = crate::fx_noise::MedianKernel::for_adapter(info.backend, info.device_type);
+        #[cfg(test)]
+        let median_kernel = median_kernel.test_override()?;
         if !adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) {
             return Err(format!("{} ({:?}): no compute shaders", info.name, info.backend));
         }
@@ -500,6 +504,7 @@ impl GpuContext {
             ("max_compute_workgroup_size_x", l.max_compute_workgroup_size_x, 64),
             ("max_compute_workgroup_size_y", l.max_compute_workgroup_size_y, 16),
             ("max_compute_invocations_per_workgroup", l.max_compute_invocations_per_workgroup, 256),
+            ("max_compute_workgroup_storage_size", l.max_compute_workgroup_storage_size, median_kernel.workgroup_storage()),
         ];
         if let Some((limit, have, need)) = needs.iter().find(|(_, have, need)| have < need) {
             return Err(format!("{} ({:?}): device limit {limit} is {have}, needs {need}", info.name, info.backend));
@@ -525,6 +530,7 @@ impl GpuContext {
             include_str!("shaders/fx_key.wgsl"),
             include_str!("shaders/fx_stylize.wgsl"),
             include_str!("shaders/fx_noise.wgsl"),
+            median_kernel.source(),
             include_str!("shaders/fx_tone.wgsl"),
             include_str!("shaders/fx_warp.wgsl"),
             include_str!("shaders/fx_extra.wgsl"),
@@ -706,6 +712,7 @@ impl GpuContext {
             queue,
             name,
             max_dim,
+            median_kernel,
             bgl,
             bgl_ext,
             pipelines,

crates/gpu/src/fx_depth.rs

@@ -280,7 +280,7 @@ fn cryptomatte(e: &mut Enc, ctx: &EffectCtx, b: GBuf, a: &AuxChannels) -> Option
     // cores with the rank planes resolved once.
     // The ID colours are only read by the Colors display; Matte, Preview and Matte Only need the
     // matte alone.
-    // `tests_fx_depth::cryptomatte_every_display_matches_the_cpu` guards this against the shader.
+    // `tests::fx_depth::cryptomatte_every_display_matches_the_cpu` guards this against the shader.
     let what = effectcraft_effects::CryptoAux::for_display(pr.e("display"), pr.b("matteOnly"));
     let img = effectcraft_effects::crypto_aux_image_with(a, &layer, &sel, what);
     let tex = e.g.upload_image(&img)?;

crates/gpu/src/fx_noise.rs

@@ -73,6 +73,53 @@ pub(crate) const IDS: &[&str] = &[
 /// not its square.
 pub(crate) const MEDIAN_BISECT_R: usize = 4;
 
+/// Physical GPUs retain the original parallel private histogram. Software Vulkan
+/// uses bounded workgroup storage to avoid its expensive first-use private-array
+/// lowering. This trades some warmed software throughput for bounded first use.
+#[derive(Clone, Copy, Debug, PartialEq, Eq)]
+pub(crate) enum MedianKernel {
+    Private64,
+    Software2,
+}
+
+impl MedianKernel {
+    pub(crate) fn for_adapter(backend: wgpu::Backend, device_type: wgpu::DeviceType) -> Self {
+        if backend == wgpu::Backend::Vulkan && device_type == wgpu::DeviceType::Cpu { Self::Software2 } else { Self::Private64 }
+    }
+
+    pub(crate) fn rows_per_workgroup(self) -> u32 {
+        match self {
+            Self::Private64 => 64,
+            Self::Software2 => 2,
+        }
+    }
+
+    pub(crate) fn workgroup_storage(self) -> u32 {
+        match self {
+            Self::Private64 => 0,
+            Self::Software2 => 16_384,
+        }
+    }
+
+    pub(crate) fn source(self) -> &'static str {
+        match self {
+            Self::Private64 => include_str!("shaders/median_private.wgsl"),
+            Self::Software2 => include_str!("shaders/median_software.wgsl"),
+        }
+    }
+
+    /// Separate test processes can qualify both implementations against existing oracles.
+    #[cfg(test)]
+    pub(crate) fn test_override(self) -> Result<Self, String> {
+        match std::env::var("EFFECTCRAFT_TEST_MEDIAN_KERNEL") {
+            Err(std::env::VarError::NotPresent) => Ok(self),
+            Ok(value) if value == "private64" => Ok(Self::Private64),
+            Ok(value) if value == "software2" => Ok(Self::Software2),
+            value => Err(format!("invalid EFFECTCRAFT_TEST_MEDIAN_KERNEL {value:?}; expected private64 or software2")),
+        }
+    }
+}
+
 /// Run effect `id` (one of [`IDS`]); `None` = this parameter combination runs on the CPU.
 pub(crate) fn apply(e: &mut Enc, id: &str, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
     match id {
@@ -166,7 +213,8 @@ pub(crate) fn median_image(e: &mut Enc, img: &GpuImage, r: usize, straight: bool
     // sliding along the rest.
     let seg = (4 * r as u32).clamp(32, 512).min(q.width);
     let out = e.scratch(q.width, q.height);
-    e.dispatch("fxn_median_huang", &with_u0([r as u32, seg, 0, 0]), &q, None, &out, Some(&levels), (q.height.div_ceil(64), q.width.div_ceil(seg)));
+    let groups = (q.height.div_ceil(e.g.median_kernel.rows_per_workgroup()), q.width.div_ceil(seg));
+    e.dispatch("fxn_median_huang", &with_u0([r as u32, seg, 0, 0]), &q, None, &out, Some(&levels), groups);
     out
 }
 

crates/gpu/src/lib.rs

@@ -407,59 +407,3 @@ impl effectcraft_effects::psim::ParticleSim for Gpu {
 
 #[cfg(test)]
 mod tests;
-#[cfg(test)]
-mod tests_3d;
-#[cfg(test)]
-mod tests_adjust;
-#[cfg(test)]
-mod tests_adv3d;
-#[cfg(test)]
-mod tests_deferred;
-#[cfg(test)]
-mod tests_deferred_jobs;
-#[cfg(test)]
-mod tests_fx_color;
-#[cfg(test)]
-mod tests_fx_depth;
-#[cfg(test)]
-mod tests_fx_distort;
-#[cfg(test)]
-mod tests_fx_extra;
-#[cfg(test)]
-mod tests_fx_gen2;
-#[cfg(test)]
-mod tests_fx_generate;
-#[cfg(test)]
-mod tests_fx_key;
-#[cfg(test)]
-mod tests_fx_light;
-#[cfg(test)]
-mod tests_fx_lut;
-#[cfg(test)]
-mod tests_fx_noise;
-#[cfg(test)]
-mod tests_fx_particles;
-#[cfg(test)]
-mod tests_fx_pixel2;
-#[cfg(test)]
-mod tests_fx_sim;
-#[cfg(test)]
-mod tests_fx_stylize;
-#[cfg(test)]
-mod tests_fx_text;
-#[cfg(test)]
-mod tests_fx_time;
-#[cfg(test)]
-mod tests_fx_tone;
-#[cfg(test)]
-mod tests_fx_transition;
-#[cfg(test)]
-mod tests_fx_vr;
-#[cfg(test)]
-mod tests_fx_warp;
-#[cfg(all(test, not(target_arch = "wasm32")))]
-mod tests_init;
-#[cfg(test)]
-mod tests_particles;
-#[cfg(test)]
-mod tests_wgsl_grammar;

crates/gpu/src/shaders/fx_noise.wgsl

@@ -82,69 +82,6 @@ fn fxn_median(@builtin(global_invocation_id) gid: vec3<u32>) {
     textureStore(out, p, vec4<f32>(data[level.x], data[level.y], data[level.z], data[level.w]));
 }
 
-// The same median for large radii, as the CPU computes it: a sliding 512-bin histogram per
-// channel (Huang). One invocation per row segment of u[0].y pixels: it fills the histogram for
-// the segment's first window, then slides right one column at a time (remove the column
-// leaving, add the one entering) and moves each channel's median by the counts below it.
-// u[0] = (r, segment length); dispatched over (rows, segments) with 64 × 1 workgroups.
-// data = the 512 CPU-rounded level / 511 values.
-@compute @workgroup_size(64, 1)
-fn fxn_median_huang(@builtin(global_invocation_id) gid: vec3<u32>) {
-    let dims = out_dims();
-    let y = i32(gid.x);
-    let seg = i32(P.u[0].y);
-    let x0 = i32(gid.y) * seg;
-    if (y >= dims.y || x0 >= dims.x) {
-        return;
-    }
-    let x1 = min(x0 + seg, dims.x);
-    let r = i32(P.u[0].x);
-    let half = u32((2 * r + 1) * (2 * r + 1) / 2);
-    var hist: array<u32, 2048>;
-    for (var dy = -r; dy <= r; dy++) {
-        for (var dx = -r; dx <= r; dx++) {
-            let v = vec4<u32>(tex_get_clamped(src, x0 + dx, y + dy));
-            hist[v.x] += 1u;
-            hist[512u + v.y] += 1u;
-            hist[1024u + v.z] += 1u;
-            hist[1536u + v.w] += 1u;
-        }
-    }
-    // Per channel: the median level and the count of values below it.
-    var med = vec4<u32>(0u);
-    var lt = vec4<u32>(0u);
-    for (var x = x0; x < x1; x++) {
-        if (x > x0) {
-            for (var dy = -r; dy <= r; dy++) {
-                let o = vec4<u32>(tex_get_clamped(src, x - 1 - r, y + dy));
-                let n = vec4<u32>(tex_get_clamped(src, x + r, y + dy));
-                for (var c = 0u; c < 4u; c++) {
-                    hist[c * 512u + o[c]] -= 1u;
-                    if (o[c] < med[c]) {
-                        lt[c] -= 1u;
-                    }
-                    hist[c * 512u + n[c]] += 1u;
-                    if (n[c] < med[c]) {
-                        lt[c] += 1u;
-                    }
-                }
-            }
-        }
-        for (var c = 0u; c < 4u; c++) {
-            let base = c * 512u;
-            while (lt[c] > half) {
-                med[c] -= 1u;
-                lt[c] -= hist[base + med[c]];
-            }
-            while (lt[c] + hist[base + med[c]] <= half) {
-                lt[c] += hist[base + med[c]];
-                med[c] += 1u;
-            }
-        }
-        textureStore(out, vec2<i32>(x, y), vec4<f32>(data[med.x], data[med.y], data[med.z], data[med.w]));
-    }
-}
-
 // noise::median_px.
 fn fxn_median_px(o: vec4<f32>, m: vec4<f32>, on_alpha: bool) -> vec4<f32> {
     if (on_alpha) {

crates/gpu/src/shaders/median_private.wgsl

@@ -0,0 +1,63 @@
+// The same median for large radii, as the CPU computes it: a sliding 512-bin histogram per
+// channel (Huang). One invocation per row segment of u[0].y pixels: it fills the histogram for
+// the segment's first window, then slides right one column at a time (remove the column
+// leaving, add the one entering) and moves each channel's median by the counts below it.
+// u[0] = (r, segment length); dispatched over (rows, segments) with 64 × 1 workgroups.
+// data = the 512 CPU-rounded level / 511 values.
+@compute @workgroup_size(64, 1)
+fn fxn_median_huang(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let dims = out_dims();
+    let y = i32(gid.x);
+    let seg = i32(P.u[0].y);
+    let x0 = i32(gid.y) * seg;
+    if (y >= dims.y || x0 >= dims.x) {
+        return;
+    }
+    let x1 = min(x0 + seg, dims.x);
+    let r = i32(P.u[0].x);
+    let half = u32((2 * r + 1) * (2 * r + 1) / 2);
+    var hist: array<u32, 2048>;
+    for (var dy = -r; dy <= r; dy++) {
+        for (var dx = -r; dx <= r; dx++) {
+            let v = vec4<u32>(tex_get_clamped(src, x0 + dx, y + dy));
+            hist[v.x] += 1u;
+            hist[512u + v.y] += 1u;
+            hist[1024u + v.z] += 1u;
+            hist[1536u + v.w] += 1u;
+        }
+    }
+    // Per channel: the median level and the count of values below it.
+    var med = vec4<u32>(0u);
+    var lt = vec4<u32>(0u);
+    for (var x = x0; x < x1; x++) {
+        if (x > x0) {
+            for (var dy = -r; dy <= r; dy++) {
+                let o = vec4<u32>(tex_get_clamped(src, x - 1 - r, y + dy));
+                let n = vec4<u32>(tex_get_clamped(src, x + r, y + dy));
+                for (var c = 0u; c < 4u; c++) {
+                    hist[c * 512u + o[c]] -= 1u;
+                    if (o[c] < med[c]) {
+                        lt[c] -= 1u;
+                    }
+                    hist[c * 512u + n[c]] += 1u;
+                    if (n[c] < med[c]) {
+                        lt[c] += 1u;
+                    }
+                }
+            }
+        }
+        for (var c = 0u; c < 4u; c++) {
+            let base = c * 512u;
+            while (lt[c] > half) {
+                med[c] -= 1u;
+                lt[c] -= hist[base + med[c]];
+            }
+            while (lt[c] + hist[base + med[c]] <= half) {
+                lt[c] += hist[base + med[c]];
+                med[c] += 1u;
+            }
+        }
+        textureStore(out, vec2<i32>(x, y), vec4<f32>(data[med.x], data[med.y], data[med.z], data[med.w]));
+    }
+}
+

crates/gpu/src/shaders/median_software.wgsl

@@ -0,0 +1,65 @@
+// The same median for large radii, as the CPU computes it: a sliding 512-bin histogram per
+// channel (Huang). One invocation per row segment of u[0].y pixels: it fills the histogram for
+// the segment's first window, then slides right one column at a time (remove the column
+// leaving, add the one entering) and moves each channel's median by the counts below it.
+// Each invocation owns one 8 KiB slice of the workgroup's zero-initialized histogram.
+// Two slices fit the guaranteed 16 KiB workgroup storage limit. No bins are shared.
+var<workgroup> median_hist: array<array<u32, 2048>, 2>;
+// u[0] = (r, segment length); dispatched over (rows, segments) with 2 × 1 workgroups.
+// data = the 512 CPU-rounded level / 511 values.
+@compute @workgroup_size(2, 1)
+fn fxn_median_huang(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(local_invocation_index) lane: u32) {
+    let dims = out_dims();
+    let y = i32(gid.x);
+    let seg = i32(P.u[0].y);
+    let x0 = i32(gid.y) * seg;
+    if (y >= dims.y || x0 >= dims.x) {
+        return;
+    }
+    let x1 = min(x0 + seg, dims.x);
+    let r = i32(P.u[0].x);
+    let half = u32((2 * r + 1) * (2 * r + 1) / 2);
+    for (var dy = -r; dy <= r; dy++) {
+        for (var dx = -r; dx <= r; dx++) {
+            let v = vec4<u32>(tex_get_clamped(src, x0 + dx, y + dy));
+            median_hist[lane][v.x] += 1u;
+            median_hist[lane][512u + v.y] += 1u;
+            median_hist[lane][1024u + v.z] += 1u;
+            median_hist[lane][1536u + v.w] += 1u;
+        }
+    }
+    // Per channel: the median level and the count of values below it.
+    var med = vec4<u32>(0u);
+    var lt = vec4<u32>(0u);
+    for (var x = x0; x < x1; x++) {
+        if (x > x0) {
+            for (var dy = -r; dy <= r; dy++) {
+                let o = vec4<u32>(tex_get_clamped(src, x - 1 - r, y + dy));
+                let n = vec4<u32>(tex_get_clamped(src, x + r, y + dy));
+                for (var c = 0u; c < 4u; c++) {
+                    median_hist[lane][c * 512u + o[c]] -= 1u;
+                    if (o[c] < med[c]) {
+                        lt[c] -= 1u;
+                    }
+                    median_hist[lane][c * 512u + n[c]] += 1u;
+                    if (n[c] < med[c]) {
+                        lt[c] += 1u;
+                    }
+                }
+            }
+        }
+        for (var c = 0u; c < 4u; c++) {
+            let base = c * 512u;
+            while (lt[c] > half) {
+                med[c] -= 1u;
+                lt[c] -= median_hist[lane][base + med[c]];
+            }
+            while (lt[c] + median_hist[lane][base + med[c]] <= half) {
+                lt[c] += median_hist[lane][base + med[c]];
+                med[c] += 1u;
+            }
+        }
+        textureStore(out, vec2<i32>(x, y), vec4<f32>(data[med.x], data[med.y], data[med.z], data[med.w]));
+    }
+}
+

crates/gpu/src/tests.rs

@@ -26,6 +26,36 @@ use effectcraft_time::{FrameRate, Tick};
 
 use crate::Gpu;
 
+mod adjust;
+mod adv3d;
+mod deferred;
+mod deferred_jobs;
+mod fx_color;
+mod fx_depth;
+mod fx_distort;
+mod fx_extra;
+mod fx_gen2;
+mod fx_generate;
+mod fx_key;
+mod fx_light;
+mod fx_lut;
+mod fx_noise;
+mod fx_particles;
+mod fx_pixel2;
+mod fx_sim;
+mod fx_stylize;
+mod fx_text;
+mod fx_time;
+mod fx_tone;
+mod fx_transition;
+mod fx_vr;
+mod fx_warp;
+#[cfg(all(test, not(target_arch = "wasm32")))]
+mod init;
+mod particles;
+mod three_d;
+mod wgsl_grammar;
+
 /// One GPU test at a time in this process: the calling test thread holds a process-wide lock
 /// until it exits (libtest runs every test on a thread of its own). wgpu's OpenGL backend (Mesa
 /// llvmpipe on FreeBSD, Linux without Vulkan) has one context lock per adapter and panics when a
@@ -52,9 +82,29 @@ pub(crate) fn gpu() -> Option<&'static Gpu> {
     static G: OnceLock<Result<Gpu, String>> = OnceLock::new();
     hold_gpu_lock();
     let g = G.get_or_init(|| {
+        // Readback mapping/allocation errors become None without necessarily retiring the
+        // context. Preserve their original messages alongside the failure's health snapshot.
+        struct GpuTestLog;
+        impl log::Log for GpuTestLog {
+            fn enabled(&self, metadata: &log::Metadata<'_>) -> bool {
+                metadata.level() <= log::Level::Warn && metadata.target().starts_with("effectcraft_gpu")
+            }
+
+            fn log(&self, record: &log::Record<'_>) {
+                if self.enabled(record.metadata()) {
+                    eprintln!("{} {}: {}", record.level(), record.target(), record.args());
+                }
+            }
+
+            fn flush(&self) {}
+        }
+        static LOGGER: GpuTestLog = GpuTestLog;
+        if log::set_logger(&LOGGER).is_ok() {
+            log::set_max_level(log::LevelFilter::Warn);
+        }
         let g = Gpu::try_headless();
         match &g {
-            Ok(g) => eprintln!("effectcraft-gpu tests: adapter {}", g.ctx.name),
+            Ok(g) => eprintln!("effectcraft-gpu tests: adapter {}, median {:?}", g.ctx.name, g.ctx.median_kernel),
             Err(e) => eprintln!("effectcraft-gpu tests: no GPU ({e}), skipping GPU comparisons"),
         }
         g
@@ -481,7 +531,16 @@ pub(crate) fn effect_direct(id: &str, vals: &[(&str, Value)], adjustment: bool)
         let buf = effectcraft_effects::Buf { img, offset: [0.0, 0.0], scale };
         let ctx = || effectcraft_effects::EffectCtx { params: &params, time: 0.25, layer_size: size, seed: 11, adjustment, env: Default::default() };
         let cpu = (spec.render)(&ctx(), buf.clone());
-        let out = effectcraft_render::Accelerator::effects(g, &[effectcraft_render::FxStep { spec, ctx: ctx() }], &buf, None).expect("the GPU runs the effect");
+        let started = std::time::Instant::now();
+        let out = effectcraft_render::Accelerator::effects(g, &[effectcraft_render::FxStep { spec, ctx: ctx() }], &buf, None).unwrap_or_else(|| {
+            panic!(
+                "the GPU runs {id} {vals:?} adj {adjustment} scale {scale}; elapsed={:?}, health={:?}, unavailable={:?}, transfers={:?}",
+                started.elapsed(),
+                g.ctx.check_health(),
+                g.ctx.unavailable_kernels(),
+                g.ctx.transfer_stats()
+            )
+        });
         if expect_empty {
             assert!(out.img.data.iter().all(|p| p[3] == 0.0), "{id} {vals:?} adj {adjustment} scale {scale}: Cut Away 100 produced non-zero alpha");
         }

crates/gpu/src/tests/deferred_jobs.rs

@@ -8,10 +8,10 @@ use effectcraft_render::three_d::adv;
 use effectcraft_render::{Accelerator, Backend, LayerCache, NoFootage, RenderOpts, Renderer};
 use effectcraft_time::Tick;
 
+use crate::tests::adv3d::{off_share, run_of, scene, scene_project};
+use crate::tests::deferred::{deferred_gpu, render_passes};
+use crate::tests::particles::{Host, agree, params};
 use crate::tests::{Pattern, diff, n, tolerance};
-use crate::tests_adv3d::{off_share, run_of, scene, scene_project};
-use crate::tests_deferred::{deferred_gpu, render_passes};
-use crate::tests_particles::{Host, agree, params};
 
 /// The device's readbacks have landed (the browser's event loop; natively a poll).
 fn settle(g: &crate::Gpu) {
@@ -87,7 +87,7 @@ fn advanced_3d_resolves_in_passes() {
 #[test]
 fn job_frames_render_in_passes() {
     let Some((g, _lock)) = deferred_gpu() else { return };
-    let s = crate::tests_deferred::scene();
+    let s = crate::tests::deferred::scene();
     let t = Tick::from_seconds_f64(0.5);
     let cache = LayerCache::default();
     cache.set_gate(g.miss_gate());

crates/gpu/src/tests/fx_gen2.rs

@@ -7,8 +7,8 @@
 use effectcraft_keyframe::Value;
 use effectcraft_time::Tick;
 
+use crate::tests::fx_pixel2::{e, env_case, env_case_allow, env_case_at, off, on, pt};
 use crate::tests::{c, n};
-use crate::tests_fx_pixel2::{e, env_case, env_case_allow, env_case_at, off, on, pt};
 
 fn layer(id: u64) -> Value {
     Value::Layer(Some(id))

crates/gpu/src/tests/fx_noise.rs

@@ -22,6 +22,24 @@ fn pt(x: f64, y: f64) -> Value {
     Value::Vec2([x, y])
 }
 
+#[test]
+fn median_kernel_selection_preserves_physical_gpu_path() {
+    use crate::fx_noise::MedianKernel;
+    for backend in [wgpu::Backend::Vulkan, wgpu::Backend::Metal, wgpu::Backend::Dx12, wgpu::Backend::Gl, wgpu::Backend::BrowserWebGpu] {
+        for device_type in
+            [wgpu::DeviceType::DiscreteGpu, wgpu::DeviceType::IntegratedGpu, wgpu::DeviceType::VirtualGpu, wgpu::DeviceType::Cpu, wgpu::DeviceType::Other]
+        {
+            let selected = MedianKernel::for_adapter(backend, device_type);
+            let expected =
+                if backend == wgpu::Backend::Vulkan && device_type == wgpu::DeviceType::Cpu { MedianKernel::Software2 } else { MedianKernel::Private64 };
+            assert_eq!(selected, expected, "{backend:?} {device_type:?}");
+        }
+    }
+    assert_eq!(MedianKernel::Private64.rows_per_workgroup(), 64);
+    assert_eq!(MedianKernel::Software2.rows_per_workgroup(), 2);
+    assert_eq!(MedianKernel::Software2.workgroup_storage(), 16_384);
+}
+
 #[test]
 fn medians() {
     effect_case("ec.noise.median", &[("radius", n(2.0))]);
@@ -237,8 +255,17 @@ fn median_levels_match_cpu_division() {
         // A monotone row's median is its centre, including repeated edges. Exercise both
         // the bisection and sliding-histogram kernels at every quantised level.
         let mut enc = crate::context::Enc::new(&g.ctx);
+        let started = std::time::Instant::now();
         let median = crate::fx_noise::median_image(&mut enc, &input, radius, false);
-        let actual = enc.download(&median).unwrap();
+        let actual = enc.download(&median).unwrap_or_else(|| {
+            panic!(
+                "median download radius {radius}; elapsed={:?}, health={:?}, unavailable={:?}, transfers={:?}",
+                started.elapsed(),
+                g.ctx.check_health(),
+                g.ctx.unavailable_kernels(),
+                g.ctx.transfer_stats()
+            )
+        });
         for (level, (want, got)) in image.data.iter().zip(&actual.data).enumerate() {
             assert_eq!(got, want, "radius {radius}, level {level}");
         }
@@ -273,7 +300,16 @@ fn dust_strict_threshold_boundaries() {
             let buf = effectcraft_effects::Buf { img, offset: [0.0; 2], scale: 1.0 };
             let cpu = (spec.render)(&ctx(), buf.clone());
             assert_eq!(cpu.img.data[24] != original, replaced, "CPU radius {radius}, threshold {threshold}");
-            let gpu = effectcraft_render::Accelerator::effects(g, &[effectcraft_render::FxStep { spec, ctx: ctx() }], &buf, None).unwrap();
+            let started = std::time::Instant::now();
+            let gpu = effectcraft_render::Accelerator::effects(g, &[effectcraft_render::FxStep { spec, ctx: ctx() }], &buf, None).unwrap_or_else(|| {
+                panic!(
+                    "dust radius {radius}, threshold {threshold}; elapsed={:?}, health={:?}, unavailable={:?}, transfers={:?}",
+                    started.elapsed(),
+                    g.ctx.check_health(),
+                    g.ctx.unavailable_kernels(),
+                    g.ctx.transfer_stats()
+                )
+            });
             assert_eq!(gpu.img.data[24], cpu.img.data[24], "radius {radius}, threshold {threshold}");
         }
     }

crates/gpu/src/tests/fx_particles.rs

@@ -7,8 +7,8 @@ use effectcraft_project::{BitDepth, LayerSource, build};
 use effectcraft_render::RenderOpts;
 use effectcraft_time::Tick;
 
+use crate::tests::fx_sim::{case, composited, e, off, on};
 use crate::tests::{Scene, check, compare_at, n, opts, set};
-use crate::tests_fx_sim::{case, composited, e, off, on};
 
 fn v2(x: f64, y: f64) -> Value {
     Value::Vec2([x, y])

crates/gpu/src/tests/wgsl_grammar.rs

@@ -22,9 +22,22 @@ enum Tok<'a> {
 
 const BITWISE: [&str; 3] = ["&", "|", "^"];
 
+const OPS: [&str; 18] = ["&&", "||", "==", "!=", "<=", ">=", "<<", ">>", "<", ">", "+", "-", "*", "/", "%", "&", "|", "^"];
+
 /// Drop comments and template argument lists (`vec3<f32>`, `array<vec4<u32>>`), whose
 /// angle brackets are not relational operators.
 fn strip(src: &str) -> String {
+    let mut out = drop_comments(src);
+    loop {
+        let (next, changed) = drop_template_lists_once(&out);
+        out = next;
+        if !changed {
+            return out;
+        }
+    }
+}
+
+fn drop_comments(src: &str) -> String {
     let mut out = String::with_capacity(src.len());
     let b = src.as_bytes();
     let mut i = 0;
@@ -45,31 +58,104 @@ fn strip(src: &str) -> String {
             i += 1;
         }
     }
-    loop {
-        let mut changed = false;
-        let mut next = String::with_capacity(out.len());
-        let c: Vec<char> = out.chars().collect();
-        let mut i = 0;
-        while i < c.len() {
-            if c[i] == '<' && i > 0 && (c[i - 1].is_alphanumeric() || c[i - 1] == '_') {
-                let mut j = i + 1;
-                while j < c.len() && (c[j].is_alphanumeric() || c[j] == '_' || c[j] == ',' || c[j] == ' ') {
-                    j += 1;
-                }
-                if j < c.len() && c[j] == '>' {
-                    i = j + 1;
-                    changed = true;
-                    continue;
-                }
+    out
+}
+
+/// One pass over `out` dropping every template argument list. Returns whether one was dropped.
+fn drop_template_lists_once(out: &str) -> (String, bool) {
+    let mut changed = false;
+    let mut next = String::with_capacity(out.len());
+    let c: Vec<char> = out.chars().collect();
+    let mut i = 0;
+    while i < c.len() {
+        if let Some(end) = template_list_end(&c, i) {
+            i = end;
+            changed = true;
+            continue;
+        }
+        next.push(c[i]);
+        i += 1;
+    }
+    (next, changed)
+}
+
+/// If `c[i]` opens a template argument list after an identifier, the index just past its `>`.
+fn template_list_end(c: &[char], i: usize) -> Option<usize> {
+    let prev = i.checked_sub(1).and_then(|p| c.get(p))?;
+    if c.get(i) != Some(&'<') || !(prev.is_alphanumeric() || *prev == '_') {
+        return None;
+    }
+    let mut j = i + 1;
+    while let Some(&ch) = c.get(j) {
+        if !(ch.is_alphanumeric() || ch == '_' || ch == ',' || ch == ' ') {
+            break;
+        }
+        j += 1;
+    }
+    (c.get(j) == Some(&'>')).then_some(j + 1)
+}
+
+/// Scans the token starting at `*at` (not whitespace) and advances `*at` past it.
+fn next_token<'a>(s: &'a str, at: &mut usize) -> Tok<'a> {
+    let b = s.as_bytes();
+    let Some(&ch) = b.get(*at) else {
+        return Tok::Sep;
+    };
+    let mut i = *at;
+    let tok = if ch.is_ascii_alphanumeric() || ch == b'_' || ch == b'.' {
+        // Identifier, member access or number (`1e-6`, `0x9E3779B9u`).
+        let start = i;
+        while i < b.len() && (b[i].is_ascii_alphanumeric() || b[i] == b'_' || b[i] == b'.') {
+            if (b[i] == b'e' || b[i] == b'E')
+                && b[start].is_ascii_digit()
+                && !s[start..i].starts_with("0x")
+                && i + 1 < b.len()
+                && (b[i + 1] == b'-' || b[i + 1] == b'+')
+            {
+                i += 1;
             }
-            next.push(c[i]);
             i += 1;
         }
-        out = next;
-        if !changed {
-            return out;
+        Tok::Operand
+    } else if ch == b'(' || ch == b'[' {
+        i += 1;
+        Tok::Open
+    } else if ch == b')' || ch == b']' {
+        i += 1;
+        Tok::Close
+    } else if b[i..].starts_with(b"->") || b[i..].starts_with(b"++") || b[i..].starts_with(b"--") {
+        i += 2;
+        Tok::Sep
+    } else if matches!(ch, b',' | b';' | b'{' | b'}' | b':' | b'@') {
+        i += 1;
+        Tok::Sep
+    } else if let Some(op) = OPS.iter().find(|op| b[i..].starts_with(op.as_bytes())) {
+        i += op.len();
+        if b.get(i) == Some(&b'=') && !matches!(*op, "==" | "!=" | "<=" | ">=") {
+            i += 1;
+            Tok::Sep // compound assignment
+        } else {
+            Tok::Op(op)
         }
+    } else if ch == b'=' {
+        i += 1;
+        Tok::Sep
+    } else {
+        i += 1;
+        Tok::Prefix
+    };
+    *at = i;
+    tok
+}
+
+/// Records a binary operator at the current parenthesis level, reporting it if it mixes with
+/// one already seen there.
+fn note_operator<'a>(level: &mut Vec<&'a str>, op: &'a str, line: usize, out: &mut Vec<(usize, String)>) {
+    let mixes = |a: &str, b: &str| a != b && (BITWISE.contains(&a) || BITWISE.contains(&b));
+    if let Some(other) = level.iter().find(|o| mixes(o, op)) {
+        out.push((line, format!("mixing '{other}' and '{op}' requires parentheses")));
     }
+    level.push(op);
 }
 
 /// `(line, message)` for every bitwise operator mixed with another binary operator at
@@ -82,7 +168,6 @@ fn violations(src: &str) -> Vec<(usize, String)> {
     let mut levels: Vec<Vec<&str>> = vec![vec![]];
     let mut prev = Tok::Sep;
     let (mut i, mut line) = (0, 1);
-    const OPS: [&str; 18] = ["&&", "||", "==", "!=", "<=", ">=", "<<", ">>", "<", ">", "+", "-", "*", "/", "%", "&", "|", "^"];
     while i < b.len() {
         let ch = b[i];
         if ch == b'\n' {
@@ -94,53 +179,14 @@ fn violations(src: &str) -> Vec<(usize, String)> {
             i += 1;
             continue;
         }
-        let tok = if ch.is_ascii_alphanumeric() || ch == b'_' || ch == b'.' {
-            // Identifier, member access or number (`1e-6`, `0x9E3779B9u`).
-            let start = i;
-            while i < b.len() && (b[i].is_ascii_alphanumeric() || b[i] == b'_' || b[i] == b'.') {
-                if (b[i] == b'e' || b[i] == b'E')
-                    && b[start].is_ascii_digit()
-                    && !s[start..i].starts_with("0x")
-                    && i + 1 < b.len()
-                    && (b[i + 1] == b'-' || b[i + 1] == b'+')
-                {
-                    i += 1;
-                }
-                i += 1;
-            }
-            Tok::Operand
-        } else if ch == b'(' || ch == b'[' {
-            i += 1;
-            levels.push(vec![]);
-            Tok::Open
-        } else if ch == b')' || ch == b']' {
-            i += 1;
-            if levels.len() > 1 {
-                levels.pop();
-            }
-            Tok::Close
-        } else if b[i..].starts_with(b"->") || b[i..].starts_with(b"++") || b[i..].starts_with(b"--") {
-            i += 2;
-            Tok::Sep
-        } else if matches!(ch, b',' | b';' | b'{' | b'}' | b':' | b'@') {
-            i += 1;
-            Tok::Sep
-        } else if let Some(op) = OPS.iter().find(|op| b[i..].starts_with(op.as_bytes())) {
-            i += op.len();
-            if b.get(i) == Some(&b'=') && !matches!(*op, "==" | "!=" | "<=" | ">=") {
-                i += 1;
-                Tok::Sep // compound assignment
-            } else {
-                Tok::Op(op)
-            }
-        } else if ch == b'=' {
-            i += 1;
-            Tok::Sep
-        } else {
-            i += 1;
-            Tok::Prefix
-        };
+        let tok = next_token(&s, &mut i);
         match tok {
+            Tok::Open => levels.push(vec![]),
+            Tok::Close => {
+                if levels.len() > 1 {
+                    levels.pop();
+                }
+            }
             Tok::Sep => {
                 if let Some(l) = levels.last_mut() {
                     l.clear();
@@ -148,11 +194,7 @@ fn violations(src: &str) -> Vec<(usize, String)> {
             }
             Tok::Op(op) if matches!(prev, Tok::Operand | Tok::Close) => {
                 let level = levels.last_mut().expect("level");
-                let mixes = |a: &str, b: &str| a != b && (BITWISE.contains(&a) || BITWISE.contains(&b));
-                if let Some(other) = level.iter().find(|o| mixes(o, op)) {
-                    out.push((line, format!("mixing '{other}' and '{op}' requires parentheses")));
-                }
-                level.push(op);
+                note_operator(level, op, line, &mut out);
             }
             _ => {}
         }