Merge remote-tracking branch 'upstream/main' into HEAD

Cargo.lock

@@ -1612,6 +1612,7 @@ dependencies = [
  "js-sys",
  "log",
  "pollster 1.0.1",
+ "rayon",
  "wasm-bindgen",
  "web-time",
  "wgpu",

crates/effects/src/lib.rs

@@ -104,12 +104,12 @@ pub use generate3::{
 };
 pub use noise::GrainLook;
 pub use textfx::{BoltPlan, TextLook, TextPass, bolt_plan, text_passes};
-pub use transition::place_layer;
+pub use transition::{GridWipe, place_layer};
 // effectcraft-gpu fx_noise.
 pub use keying2::{InnerOuterPlan, inner_outer_plan};
 pub use noise::remove_grain_level;
 pub use noise3::{FractalGpu, NOISE_CELL, fractal_gpu};
-pub use stylize3::{BumpLight, bump_layer_id};
+pub use stylize3::{BumpLight, HexTile, bump_layer_id};
 pub use textfx::{TextLayer, text_layer};
 pub use time_fx::{PixelMotion, TimewarpPlan, TwStep, pmb_plan, posterized_time, time_frames, timewarp_crop, timewarp_plan};
 // effectcraft-gpu fx_tone.

crates/effects/src/stylize3.rs

@@ -364,32 +364,54 @@ fn hex_cell(x: f64, y: f64, r: f64) -> ((f64, f64), (i64, i64)) {
     ((px, py), (rx as i64, rz as i64))
 }
 
+/// Stable f64 hex-cell selection and source geometry for CPU and GPU renders.
+pub struct HexTile {
+    r: f64,
+    c: (f64, f64),
+    sn: f64,
+    cs: f64,
+    smear: f64,
+    mode: u32,
+}
+impl HexTile {
+    /// No image pixels are read; only buffer geometry and parameters are used.
+    pub fn new(ctx: &EffectCtx, b: &Buf) -> Self {
+        let r = (ctx.params.f("radius") * b.scale).max(0.5);
+        let c = b.to_px(ctx.params.v2("center"));
+        let rot = ctx.params.f("rotate").to_radians();
+        let (sn, cs) = rot.sin_cos();
+        let smear = (ctx.params.f("smearing") / 100.0).clamp(0.0, 1.0);
+        let mode = ctx.params.e("render");
+        Self { r, c, sn, cs, smear, mode }
+    }
+    /// Bilinear source position at a buffer pixel centre.
+    pub fn source_at(&self, x: u32, y: u32) -> (f64, f64) {
+        let Self { r, c: (cx, cy), sn, cs, smear, mode } = *self;
+        let (dx, dy) = (x as f64 + 0.5 - cx, y as f64 + 0.5 - cy);
+        // Into the rotated tiling frame.
+        let (u, v) = (dx * cs + dy * sn, -dx * sn + dy * cs);
+        let ((hx, hy), (qi, ri)) = hex_cell(u, v, r);
+        let (mut lu, lv) = (u - hx, v - hy);
+        let odd = (qi + ri).rem_euclid(2) == 1;
+        match mode {
+            1 if odd => lu = -lu,
+            2 if odd => lu = -lu,
+            _ => {}
+        }
+        let lv = if mode == 2 && odd { -lv } else { lv };
+        // Every tile shows the hexagon around the centre point.
+        let (ox, oy) = (lu * cs - lv * sn, lu * sn + lv * cs);
+        let (tx, ty) = (cx + ox, cy + oy);
+        let (sx, sy) = (tx + (x as f64 + 0.5 - tx) * smear, ty + (y as f64 + 0.5 - ty) * smear);
+        (sx, sy)
+    }
+}
 fn hextile(ctx: &EffectCtx, mut b: Buf) -> Buf {
-    let r = (ctx.params.f("radius") * b.scale).max(0.5);
-    let (cx, cy) = b.to_px(ctx.params.v2("center"));
-    let rot = ctx.params.f("rotate").to_radians();
-    let (sn, cs) = rot.sin_cos();
-    let smear = (ctx.params.f("smearing") / 100.0).clamp(0.0, 1.0);
-    let mode = ctx.params.e("render");
+    let geometry = HexTile::new(ctx, &b);
     let src = b.img.clone();
     b.img.rows_mut().for_each(|(y, row)| {
         for (x, px) in row.iter_mut().enumerate() {
-            let (dx, dy) = (x as f64 + 0.5 - cx, y as f64 + 0.5 - cy);
-            // Into the rotated tiling frame.
-            let (u, v) = (dx * cs + dy * sn, -dx * sn + dy * cs);
-            let ((hx, hy), (qi, ri)) = hex_cell(u, v, r);
-            let (mut lu, lv) = (u - hx, v - hy);
-            let odd = (qi + ri).rem_euclid(2) == 1;
-            match mode {
-                1 if odd => lu = -lu,
-                2 if odd => lu = -lu,
-                _ => {}
-            }
-            let lv = if mode == 2 && odd { -lv } else { lv };
-            // Every tile shows the hexagon around the centre point.
-            let (ox, oy) = (lu * cs - lv * sn, lu * sn + lv * cs);
-            let (tx, ty) = (cx + ox, cy + oy);
-            let (sx, sy) = (tx + (x as f64 + 0.5 - tx) * smear, ty + (y as f64 + 0.5 - ty) * smear);
+            let (sx, sy) = geometry.source_at(x as u32, y as u32);
             *px = src.sample_bilinear_clamped(sx, sy);
         }
     });

crates/effects/src/transition.rs

@@ -314,53 +314,86 @@ pub(crate) fn card_wipe_plan(ctx: &EffectCtx, b: &Buf) -> Option<PiecePlan> {
     Some(PiecePlan { pieces, front: PieceTex::Layer, back })
 }
 
+/// Double-precision grid geometry, shared with the GPU mask upload. A tiny
+/// rounding error at a tile boundary can otherwise flip a fully opaque pixel.
+pub struct GridWipe {
+    c: (f64, f64),
+    s: f64,
+    co: f64,
+    cs: f64,
+    maxd: f64,
+    shape: u32,
+    reverse: bool,
+    done: f64,
+    border: f64,
+}
+
+impl GridWipe {
+    /// Construct from buffer geometry; no pixel data is accessed.
+    pub fn new(ctx: &EffectCtx, b: &Buf) -> Self {
+        let done = completion(ctx);
+        let c = b.to_px(ctx.params.v2("center"));
+        let rot = ctx.params.f("rotation").to_radians();
+        let border = ctx.params.f("border").max(0.0) * b.scale;
+        let tiles = ctx.params.f("tiles").round().clamp(1.0, 500.0);
+        let shape = ctx.params.e("shape");
+        let reverse = ctx.params.b("reverse");
+        let (_, _, w, _) = layer_rect(ctx, b);
+        let cs = (w / tiles).max(1.0);
+        let (s, co) = rot.sin_cos();
+        let metric = |dx: f64, dy: f64| match shape {
+            0 => dx.abs(),
+            2 => dx.abs().max(dy.abs()),
+            _ => dx.hypot(dy),
+        };
+        let maxd = max_corner_dist(b, c, |dx, dy| {
+            let (qx, qy) = (dx * co + dy * s, -dx * s + dy * co);
+            metric(qx, qy)
+        });
+        GridWipe { c, s, co, cs, maxd, shape, reverse, done, border }
+    }
+
+    /// Alpha multiplier at a buffer pixel centre.
+    pub fn coverage(&self, x: u32, y: u32) -> f32 {
+        let Self { c, s, co, cs, maxd, shape, reverse, done, border } = *self;
+        let metric = |dx: f64, dy: f64| match shape {
+            0 => dx.abs(),
+            2 => dx.abs().max(dy.abs()),
+            _ => dx.hypot(dy),
+        };
+        let spread = 0.25;
+        let (dx, dy) = (x as f64 + 0.5 - c.0, y as f64 + 0.5 - c.1);
+        let (qx, qy) = (dx * co + dy * s, -dx * s + dy * co);
+        let (cx, cy) = (((qx / cs).floor() + 0.5) * cs, ((qy / cs).floor() + 0.5) * cs);
+        let mut o = (metric(cx, cy) / maxd).clamp(0.0, 1.0);
+        if reverse {
+            o = 1.0 - o;
+        }
+        let p = ((done * (1.0 + spread) - o) / spread).clamp(0.0, 1.0);
+        if p <= 0.0 {
+            return 1.0;
+        }
+        let half = cs * 0.5 * (1.0 - p);
+        let m = (qx - cx).abs().max((qy - cy).abs());
+        let k = if border > 0.0 {
+            ((half - m) / border + 0.5).clamp(0.0, 1.0)
+        } else if m < half {
+            1.0
+        } else {
+            0.0
+        };
+        k as f32
+    }
+}
+
 fn grid_wipe(ctx: &EffectCtx, mut b: Buf) -> Buf {
-    let done = completion(ctx);
-    if done <= 0.0 {
+    if completion(ctx) <= 0.0 {
         return b;
     }
-    let c = b.to_px(ctx.params.v2("center"));
-    let rot = ctx.params.f("rotation").to_radians();
-    let border = ctx.params.f("border").max(0.0) * b.scale;
-    let tiles = ctx.params.f("tiles").round().clamp(1.0, 500.0);
-    let shape = ctx.params.e("shape");
-    let reverse = ctx.params.b("reverse");
-    let (_, _, w, _) = layer_rect(ctx, &b);
-    let cs = (w / tiles).max(1.0);
-    let (s, co) = rot.sin_cos();
-    let metric = |dx: f64, dy: f64| match shape {
-        0 => dx.abs(),
-        2 => dx.abs().max(dy.abs()),
-        _ => dx.hypot(dy),
-    };
-    let maxd = max_corner_dist(&b, c, |dx, dy| {
-        let (qx, qy) = (dx * co + dy * s, -dx * s + dy * co);
-        metric(qx, qy)
-    });
-    let spread = 0.25;
+    let grid = GridWipe::new(ctx, &b);
     b.img.rows_mut().for_each(|(y, row)| {
         for (x, px) in row.iter_mut().enumerate() {
-            let (dx, dy) = (x as f64 + 0.5 - c.0, y as f64 + 0.5 - c.1);
-            let (qx, qy) = (dx * co + dy * s, -dx * s + dy * co);
-            let (cx, cy) = (((qx / cs).floor() + 0.5) * cs, ((qy / cs).floor() + 0.5) * cs);
-            let mut o = (metric(cx, cy) / maxd).clamp(0.0, 1.0);
-            if reverse {
-                o = 1.0 - o;
-            }
-            let p = ((done * (1.0 + spread) - o) / spread).clamp(0.0, 1.0);
-            if p <= 0.0 {
-                continue;
-            }
-            let half = cs * 0.5 * (1.0 - p);
-            let m = (qx - cx).abs().max((qy - cy).abs());
-            let k = if border > 0.0 {
-                ((half - m) / border + 0.5).clamp(0.0, 1.0)
-            } else if m < half {
-                1.0
-            } else {
-                0.0
-            };
-            scale_px(px, k as f32);
+            scale_px(px, grid.coverage(x as u32, y as u32));
         }
     });
     b

crates/engine/src/background/offload.rs

@@ -558,6 +558,7 @@ async fn fill(s: &mut Session, plan: crate::commands::content_fill::FillPlan, po
         true
     };
     let accel = job_accel();
+    let fill = accel.as_deref().and_then(|a| a.inpaint());
     let files = crate::commands::content_fill::run_plan(
         &plan,
         &s.project,
@@ -566,6 +567,7 @@ async fn fill(s: &mut Session, plan: crate::commands::content_fill::FillPlan, po
         &s.layer_cache,
         s.services.as_ref(),
         accel.as_deref(),
+        fill,
         &mut progress,
         &mut |_| {},
     )

crates/engine/src/background/remote.rs

@@ -417,6 +417,10 @@ impl FrameServer {
                             auto = Some((key, gpu));
                         }
                     }
+                    // A frame still waiting for readbacks is replaced: end it, so every `frame_begin` has its `frame_end`.
+                    if self.pending.take().is_some() {
+                        a.frame_end();
+                    }
                     a.frame_begin();
                 }
                 self.pending = Some(Pending { id, project: p, comp, time, opts, t0: web_time::Instant::now(), passes: 0, auto });
@@ -457,6 +461,10 @@ impl FrameServer {
         {
             pick.record(key, chose, ms);
         }
+        // The frame is done (its `frame_begin` ran in `handle`): the GPU pool ages its sizes.
+        if let Some(a) = &self.accel {
+            a.frame_end();
+        }
         vec![(
             FrameReply::Frame { id: pr.id, width: img.width, height: img.height, ms, passes: pr.passes, gpu, layer_misses, layer_hits },
             Some(rgba8_premultiplied(&img)),

crates/engine/src/background/remote/tests.rs

@@ -192,6 +192,8 @@ struct Deferred {
     asked: std::sync::Mutex<std::collections::HashSet<u64>>,
     gate: Arc<std::sync::atomic::AtomicBool>,
     chains: std::sync::atomic::AtomicUsize,
+    begins: std::sync::atomic::AtomicUsize,
+    ends: std::sync::atomic::AtomicUsize,
     auto: effectcraft_render::AutoPick,
 }
 
@@ -245,8 +247,12 @@ impl Accelerator for Deferred {
         Some(&self.auto)
     }
     fn frame_begin(&self) {
+        self.begins.fetch_add(1, std::sync::atomic::Ordering::Relaxed);
         self.ready.lock().unwrap().clear();
     }
+    fn frame_end(&self) {
+        self.ends.fetch_add(1, std::sync::atomic::Ordering::Relaxed);
+    }
     fn pass_begin(&self) {
         self.gate.store(false, std::sync::atomic::Ordering::Relaxed);
     }
@@ -316,3 +322,56 @@ fn deferred_readbacks_render_in_passes() {
     let st = acc.auto.stats(effectcraft_render::AutoKey::new(comp, opts.scale, false)).unwrap();
     assert!(st.cpu_frames > 0 && st.gpu_frames > 0, "{st:?}");
 }
+
+/// `frame_begin` and `frame_end` stay paired: early fails (before `frame_begin`), a pending frame replaced by a new render, a finished multi-pass frame, and a CPU frame.
+#[test]
+fn every_frame_begin_has_a_frame_end() {
+    use std::sync::atomic::Ordering::Relaxed;
+    let mut s = Session::default();
+    s.execute("file.openDemoProject", json!({})).unwrap();
+    let comp = s.active_comp_id().unwrap();
+    s.execute("layer.newSolid", json!({"color": "#3388ff", "name": "Blue", "width": 200, "height": 120})).unwrap();
+    s.execute("effect.apply", json!({"effect": "ec.blur.gaussian"})).unwrap();
+    s.execute("effect.apply", json!({"effect": "ec.stylize.glow"})).unwrap();
+    s.execute("layer.newSolid", json!({"color": "#ffcc00", "name": "Yellow", "width": 80, "height": 80})).unwrap();
+    s.execute("effect.apply", json!({"effect": "ec.color.levels"})).unwrap();
+    let opts = RenderOpts { scale: 0.25, backend: Backend::Gpu, ..Default::default() };
+    let t = Tick::from_seconds_f64(1.0);
+    let mut w = server(&s);
+    let acc = Arc::new(Deferred::default());
+    w.set_accel(Some(acc.clone()));
+    let mut m = Mirror::default();
+    w.handle(round(&m.sync(s.revision, &s.project).unwrap()));
+    let req = |id, revision, comp, opts: RenderOpts| {
+        round(&FrameMsg::Render { id, revision, comp, time: t, opts: Box::new(opts), disk: None, layers: false, prefetch: vec![] })
+    };
+    let load = || (acc.begins.load(Relaxed), acc.ends.load(Relaxed));
+    // Early fails never touch the frame clock.
+    assert!(matches!(w.handle(req(1, s.revision + 1, comp, opts))[0].0, FrameReply::Failed { .. }));
+    assert!(matches!(w.handle(req(2, s.revision, ItemId(999_999), opts))[0].0, FrameReply::Failed { .. }));
+    assert_eq!(load(), (0, 0));
+    // First GPU pass misses and leaves the frame pending.
+    assert!(w.handle(req(3, s.revision, comp, opts)).is_empty() && w.waiting());
+    assert_eq!(load(), (1, 0));
+    // A fail while pending leaves the frame and the counters alone.
+    assert!(matches!(w.handle(req(4, s.revision + 1, comp, opts))[0].0, FrameReply::Failed { .. }));
+    assert!(w.waiting());
+    assert_eq!(load(), (1, 0));
+    // A new valid render replaces the pending frame (ends the old one).
+    assert!(w.handle(req(5, s.revision, comp, opts)).is_empty() && w.waiting());
+    assert_eq!(load(), (2, 1));
+    let mut out = vec![];
+    for _ in 0..10 {
+        acc.settle();
+        out = w.resume();
+        if !w.waiting() {
+            break;
+        }
+    }
+    assert!(matches!(out[0].0, FrameReply::Frame { id: 5, .. }));
+    assert_eq!(load(), (2, 2));
+    // CPU still pairs begin/end on the accelerator.
+    let cpu = RenderOpts { backend: Backend::Cpu, ..opts };
+    assert!(matches!(w.handle(req(6, s.revision, comp, cpu))[0].0, FrameReply::Frame { id: 6, .. }));
+    assert_eq!(load(), (3, 3));
+}

crates/engine/src/cmd/content_fill.rs

@@ -13,7 +13,7 @@ use std::sync::Arc;
 use effectcraft_color::Label;
 use effectcraft_project::{Footage, FootageKind, ItemId, ItemKind, LayerId, LayerSource, build};
 use effectcraft_raster::Image;
-use effectcraft_raster::inpaint::{FillInput, FillMethod, FillOpts, dilate, fill_sequence};
+use effectcraft_raster::inpaint::{FillInput, FillMethod, FillOpts, dilate};
 use effectcraft_render::{Accelerator, EvalCtx, ExprHost, FootageSource, LayerCache, Renderer};
 use effectcraft_time::Tick;
 use serde::{Deserialize, Serialize};
@@ -185,8 +185,21 @@ fn memory_shortfall(frames: usize, w: u32, h: u32, available: u64) -> Option<Str
     })
 }
 
+/// The status line shown while the frames are filled. Edge Blend calls only the membrane, which
+/// the device backend leaves to the CPU reference, so it is not a device fill even when a `fill`
+/// backend is present; Object and Surface run the flow and PatchMatch on the device.
+fn filling_status(on_device: bool, method: FillMethod) -> &'static str {
+    if on_device && method != FillMethod::EdgeBlend { "Filling (GPU)" } else { "Filling" }
+}
+
 /// Run a fill: the layer's frames (rendered in passes with a job worker's GPU), the fill, the
 /// PNG sequence written into `plan.dir` (`Services::store_file`). Returns the files.
+///
+/// `accel` renders the frames (a job worker's GPU, as for every analysis: `None` on the
+/// desktop, whose analysis threads render on the CPU by design); `fill` is the fill's own
+/// backend ([`effectcraft_raster::inpaint::FillBackend`]), which is what makes Content-Aware
+/// Fill itself run on the GPU. They are separate on purpose: a desktop fill can run on the GPU
+/// while its input frames are composited on the CPU.
 #[allow(clippy::too_many_arguments)]
 pub(crate) async fn run_plan(
     plan: &FillPlan,
@@ -196,6 +209,7 @@ pub(crate) async fn run_plan(
     cache: &LayerCache,
     services: &dyn crate::Services,
     accel: Option<&dyn Accelerator>,
+    fill: Option<&dyn effectcraft_raster::inpaint::FillBackend>,
     progress: &mut dyn FnMut(u64, u64) -> bool,
     message: &mut dyn FnMut(&str),
 ) -> std::result::Result<Vec<String>, String> {
@@ -233,10 +247,13 @@ pub(crate) async fn run_plan(
             references.push((k, img));
         }
     }
-    message("Filling");
     let opts = FillOpts { method: FillMethod::from_name(&plan.method).unwrap_or_default(), lighting: plan.lighting, ..Default::default() };
+    message(filling_status(fill.is_some(), opts.method));
     let input = FillInput { frames, holes, references };
-    let out = fill_sequence(&input, &opts, &mut |d, t| progress(n + (d as u64 * n / t.max(1) as u64), n * 3)).ok_or(crate::render_queue::CANCELLED)?;
+    let cpu = effectcraft_raster::inpaint::CpuFill;
+    let backend: &dyn effectcraft_raster::inpaint::FillBackend = fill.unwrap_or(&cpu);
+    let out = effectcraft_raster::inpaint::fill_sequence_with(&input, &opts, backend, &mut |d, t| progress(n + (d as u64 * n / t.max(1) as u64), n * 3))
+        .ok_or(crate::render_queue::CANCELLED)?;
     message("Writing the fill layer");
     let dir = &plan.dir;
     let mut paths = vec![];
@@ -324,7 +341,12 @@ fn generate(s: &mut Session, p: &Value) -> Result<Value> {
     let services = s.services.clone();
     let importer = s.importer.clone();
     let name = layer.name.clone();
+    // The fill's own GPU backend (Mercury GPU Acceleration; none under Software Only or
+    // without a compositor). The frames themselves still render on the CPU here, as for every
+    // desktop analysis (`offload::job_accel`): only the fill's own work moves to the device.
+    let accel = if s.project.settings.gpu_acceleration { s.accel.clone() } else { None };
     s.spawn_task("contentFill", format!("Content-Aware Fill: {name}"), wait, move |ctl| {
+        let fill = accel.as_deref().and_then(|a| a.inpaint());
         let paths = effectcraft_render::passes::block_on(run_plan(
             &plan,
             &lf.project,
@@ -333,6 +355,7 @@ fn generate(s: &mut Session, p: &Value) -> Result<Value> {
             &lf.cache,
             services.as_ref(),
             None,
+            fill,
             &mut |d, t| ctl.progress(d, t),
             &mut |m| ctl.message(m),
         ))?;
@@ -445,7 +468,18 @@ pub fn specs() -> Vec<CommandSpec> {
 
 #[cfg(test)]
 mod tests {
-    use super::memory_shortfall;
+    use super::{filling_status, memory_shortfall};
+
+    /// Edge Blend fills with the membrane, which the device backend leaves to the CPU reference:
+    /// it must not claim a GPU fill just because a backend is available.
+    #[test]
+    fn edge_blend_is_not_reported_as_a_device_fill() {
+        use effectcraft_raster::inpaint::FillMethod;
+        assert_eq!(filling_status(true, FillMethod::Object), "Filling (GPU)");
+        assert_eq!(filling_status(true, FillMethod::Surface), "Filling (GPU)");
+        assert_eq!(filling_status(true, FillMethod::EdgeBlend), "Filling");
+        assert_eq!(filling_status(false, FillMethod::Object), "Filling");
+    }
 
     /// #1789: 10 s of 1080p30 (300 frames) needs at least 20 GB, which 4 GB free can't hold.
     #[test]

crates/engine/src/settings/prefs/document.rs

@@ -76,7 +76,15 @@ impl Prefs {
             Err(_) => {
                 let mut base = serde_json::to_value(Prefs::default()).unwrap_or_default();
                 if let (Some(b), Some(o)) = (base.as_object_mut(), v.as_object()) {
+                    // Only the settings' own pages are tried one by one (each trial parses the
+                    // whole document): any other key is kept as it is, so a document of many
+                    // unknown keys costs one pass, not one parse per key.
+                    let pages: Vec<String> = b.keys().cloned().collect();
                     for (k, x) in o {
+                        if !pages.contains(k) {
+                            b.insert(k.clone(), x.clone());
+                            continue;
+                        }
                         let mut trial = b.clone();
                         trial.insert(k.clone(), x.clone());
                         if serde_json::from_value::<Prefs>(Value::Object(trial)).is_ok() {

crates/engine/src/tests/settings.rs

@@ -40,6 +40,22 @@ fn prefs_round_trip_through_json() {
     assert_eq!(Prefs::from_json(&Prefs::default().to_json()), Prefs::default());
 }
 
+/// A document with one bad page and many unknown keys recovers its good pages in one pass per
+/// page, not one full parse per key.
+#[test]
+fn prefs_recovery_tries_only_the_settings_pages() {
+    let mut doc = serde_json::Map::new();
+    doc.insert("general".into(), json!({"undoLevels": "not a number"}));
+    doc.insert("labels".into(), serde_json::to_value(Prefs::default()).unwrap()["labels"].clone());
+    for i in 0..20_000 {
+        doc.insert(format!("unknown{i}"), json!({"x": i}));
+    }
+    let t = std::time::Instant::now();
+    let p = Prefs::from_json(&serde_json::Value::Object(doc).to_string());
+    assert!(t.elapsed() < std::time::Duration::from_secs(10), "{:?}", t.elapsed());
+    assert_eq!(p.general, Prefs::default().general);
+}
+
 #[test]
 fn prefs_migrate_old_layout_and_keep_unknown_keys() {
     let old = r#"{

crates/export/src/lib.rs

@@ -176,7 +176,9 @@ pub fn export(job: &Job, progress: &mut dyn FnMut(&Progress) -> bool) -> Result<
 pub async fn export_async(job: &Job<'_>, progress: &mut dyn FnMut(&Progress) -> bool) -> Result<Report> {
     let t0 = Instant::now();
     let cx = Cx::new(job);
+    let render = RenderEnd(job.accel);
     let mut r = export_with(&cx, progress, t0).await;
+    drop(render);
     if matches!(r, Err(ExportError::Cancelled)) {
         cx.remove_cancelled_outputs()?;
         return r;
@@ -189,6 +191,18 @@ pub async fn export_async(job: &Job<'_>, progress: &mut dyn FnMut(&Progress) ->
     r
 }
 
+/// Calls [`effectcraft_render::Accelerator::render_end`] when dropped: after the export, and
+/// also when it panics or its future is dropped (a browser worker aborting it).
+struct RenderEnd<'a>(Option<&'a dyn effectcraft_render::Accelerator>);
+
+impl Drop for RenderEnd<'_> {
+    fn drop(&mut self) {
+        if let Some(a) = self.0 {
+            a.render_end();
+        }
+    }
+}
+
 async fn export_with(cx: &Cx<'_>, progress: &mut dyn FnMut(&Progress) -> bool, t0: Instant) -> Result<Report> {
     let comp = cx.comp().ok_or(ExportError::NoComp)?;
     let total = cx.settings.frame_count(comp);

crates/export/src/pipeline.rs

@@ -391,11 +391,14 @@ impl<'a> Cx<'a> {
         }
         let t0 = Instant::now();
         let r = self.renderer();
-        let planes = r.active_accel()?.comp_frame_yuv420(&r, self.job.comp, t, comp.background)?;
+        let a = r.active_accel()?;
+        let planes = a.comp_frame_yuv420(&r, self.job.comp, t, comp.background)?;
         let (cw, ch) = (w.div_ceil(2) as usize, h.div_ceil(2) as usize);
         if planes.len() != w as usize * h as usize + 2 * cw * ch {
             return None;
         }
+        // A frame rendered off `passes::comp_frame` ends here; on a `None` the fallback renders and ends it.
+        a.frame_end();
         if let Some(ft) = &self.frame_times {
             ft.lock().unwrap_or_else(|e| e.into_inner()).push((i, t0.elapsed().as_secs_f64()));
         }

crates/export/tests/it/h264_gpu_yuv.rs

@@ -24,11 +24,21 @@ struct Fake {
     /// Frames asking for planes right now, and the most at once after the stop.
     asking: AtomicUsize,
     most_after_stop: AtomicUsize,
+    /// `Accelerator::frame_end` calls: one per rendered comp frame, planes or not.
+    frame_ends: AtomicUsize,
 }
 
 impl Fake {
     fn new(yuv: bool, deferred: bool, stop_after: usize) -> Self {
-        Fake { yuv, yuv_frames: AtomicUsize::new(0), deferred, stop_after, asking: AtomicUsize::new(0), most_after_stop: AtomicUsize::new(0) }
+        Fake {
+            yuv,
+            yuv_frames: AtomicUsize::new(0),
+            deferred,
+            stop_after,
+            asking: AtomicUsize::new(0),
+            most_after_stop: AtomicUsize::new(0),
+            frame_ends: AtomicUsize::new(0),
+        }
     }
 }
 
@@ -54,6 +64,9 @@ impl Accelerator for Fake {
         }
         Some(effectcraft_raster::yuv::yuv420_over(&r.comp_frame_cpu(comp, t), bg))
     }
+    fn frame_end(&self) {
+        self.frame_ends.fetch_add(1, Ordering::SeqCst);
+    }
     fn available(&self) -> bool {
         self.yuv_frames.load(Ordering::SeqCst) < self.stop_after
     }
@@ -122,6 +135,8 @@ fn accelerator_planes_encode_like_the_cpu_conversion() {
     let b = run(&p, cid, Some(&rgba), &s, &d.join("rgba.mp4"));
     let c = run(&p, cid, None, &s, &d.join("cpu.mp4"));
     assert_eq!(yuv.yuv_frames.load(Ordering::SeqCst), 20, "every frame came from the accelerator");
+    assert_eq!(yuv.frame_ends.load(Ordering::SeqCst), 20, "every planes frame ends once");
+    assert_eq!(rgba.frame_ends.load(Ordering::SeqCst), 20, "every RGBA frame ends once");
     assert!(a == b && b == c, "the three files differ");
     let f = effectcraft_media::probe(&yuv_mp4).expect("probe yuv.mp4");
     assert_eq!((f.width, f.height), (64, 48));
@@ -144,6 +159,7 @@ fn frames_the_accelerator_cannot_take_fall_back_to_rgba() {
     let a = run(&p, cid, Some(&yuv), &odd, &d.join("yuv.mp4"));
     let c = run(&p, cid, None, &odd, &d.join("cpu.mp4"));
     assert_eq!(yuv.yuv_frames.load(Ordering::SeqCst), 0);
+    assert_eq!(yuv.frame_ends.load(Ordering::SeqCst), 6, "every fallback frame ends once");
     assert!(a == c);
 }
 

crates/gpu/Cargo.toml

@@ -21,6 +21,7 @@ effectcraft-keyframe = { workspace = true }
 wgpu = { workspace = true }
 bytemuck = { workspace = true }
 log = { workspace = true }
+rayon = { workspace = true }
 web-time = { workspace = true }
 
 [target.'cfg(not(target_arch = "wasm32"))'.dependencies]

crates/gpu/src/adv3d.rs

@@ -489,7 +489,7 @@ fn raster_into(e: &mut Enc, s: &Scene) -> Option<(wgpu::Texture, wgpu::Texture)>
     let zout = tex(DEPTH_OUT, rt);
     let depth = tex(DEPTH, wgpu::TextureUsages::RENDER_ATTACHMENT);
     let (cv, zv, dv) = (color.create_view(&Default::default()), zout.create_view(&Default::default()), depth.create_view(&Default::default()));
-    let enc = e.encoder();
+    let enc = e.record(&[]);
     {
         let mut pass = enc.begin_render_pass(&wgpu::RenderPassDescriptor {
             label: Some("adv3d"),
@@ -763,7 +763,7 @@ fn dispatch(e: &mut Enc, kernel: &str, u0: [u32; 4], f0: [f32; 4], b: Binds, gro
     }
     entries.push(wgpu::BindGroupEntry { binding: 10, resource: wgpu::BindingResource::TextureView(b.out.unwrap_or(&post.dummy_out)) });
     let bg = g.device.create_bind_group(&wgpu::BindGroupDescriptor { label: Some(kernel), layout: &post.bgl, entries: &entries });
-    let enc = e.encoder();
+    let enc = e.record(&[]);
     let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some(kernel), timestamp_writes: None });
     pass.set_pipeline(pipe);
     pass.set_bind_group(0, &bg, &[]);
@@ -799,7 +799,7 @@ fn read_buffer(e: &mut Enc, buf: &wgpu::Buffer, size: u64) -> Option<Vec<u8>> {
         usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
         mapped_at_creation: false,
     });
-    e.encoder().copy_buffer_to_buffer(buf, 0, &rb, 0, size);
+    e.record(&[]).copy_buffer_to_buffer(buf, 0, &rb, 0, size);
     e.submit();
     let out = e.g.read_buffer(&rb, row, row, 1).map_err(|e| log::error!("gpu 3d readback: {e}")).ok()?;
     Some(out)
@@ -922,6 +922,7 @@ fn depth_of_field(e: &mut Enc, acc: &wgpu::Buffer, depth: &wgpu::Buffer, (w, h):
 
 /// The resolved run as an image in the canvas's encoding, composited over `canvas` when given.
 pub(crate) fn finish(e: &mut Enc, r: &Resolved, encode: bool, canvas: Option<&GpuImage>) -> GpuImage {
+    canvas.inspect(|c| e.uses(c));
     let out = e.image(r.width, r.height);
     let ov = out.texture.create_view(&Default::default());
     let cv = canvas.map(|c| c.texture.create_view(&Default::default()));
@@ -1023,7 +1024,7 @@ fn render_prepared_deferred(g: &GpuContext, d: &std::sync::Arc<crate::deferred::
         usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
         mapped_at_creation: false,
     });
-    e.encoder().copy_buffer_to_buffer(&r.depth, 0, &rb, 0, size);
+    e.record(&[]).copy_buffer_to_buffer(&r.depth, 0, &rb, 0, size);
     e.submit();
     let done = d.start(kz, w, h, [0.0; 2], 1.0);
     g.map_readback(&rb, row, row, 1, move |r| done(r.map_err(|e| log::error!("gpu depth readback: {e}")).ok()));
@@ -1080,6 +1081,7 @@ pub(crate) fn sky(e: &mut Enc, sky: &effectcraft_render::three_d::SkyDraw, canva
 /// Wireframe-quality layer outlines: `canvas` with `pixels` set to opaque white
 /// ([`effectcraft_render::Renderer::wireframe_pixels`]).
 pub(crate) fn wireframe(e: &mut Enc, canvas: &GpuImage, pixels: &[(u32, u32)]) -> GpuImage {
+    e.uses(canvas);
     let out = e.image(canvas.width, canvas.height);
     e.copy_into(canvas, &out, 0, 0);
     if pixels.is_empty() {

crates/gpu/src/context.rs

@@ -42,7 +42,17 @@ pub const FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Rgba32Float;
 const UPLOAD_BUDGET: usize = 1536 << 20;
 
 /// Working textures kept for reuse once their last image is dropped (bytes).
-const POOL_BUDGET: usize = 768 << 20;
+const POOL_BUDGET: usize = 256 << 20;
+
+/// Frames after which a size the pool was not asked for is stale: only stale sizes are
+/// evicted to make room (see [`Pool::make_free`]).
+///
+/// The clock is coarse: it advances on every finished frame ([`GpuContext::pool_frame_end`])
+/// from any thread, so with many frames in flight (an export renders up to 16 at once) this
+/// means "not asked for during the last 2 completions", not "during 2 frames' rendering". A
+/// size a slower in-flight frame still needs can go stale and be evicted. That costs a pool
+/// miss (a new allocation), never a wrong pixel.
+const STALE_FRAMES: u64 = 2;
 
 /// Readback staging buffers kept for reuse across sizes (bytes).
 const STAGING_BUDGET: u64 = 256 << 20;
@@ -66,6 +76,20 @@ impl GpuImage {
     pub fn new(texture: wgpu::Texture, width: u32, height: u32) -> GpuImage {
         GpuImage { texture, width, height, _lease: None, levels: None }
     }
+
+    /// This image leaves the [`Enc`] that made it (read by an encoder that may stay open, see
+    /// [`Enc::uses`]): when its last clone drops, its texture waits for every open encoder
+    /// ([`Pool::release`]) instead of going back to that `Enc` for immediate reuse.
+    pub(crate) fn escape(&self) {
+        if let Some(l) = &self._lease {
+            l.escaped.store(true, std::sync::atomic::Ordering::Relaxed);
+        }
+    }
+
+    /// A pooled image of another `Enc` than `owner` (see [`GpuImage::escape`]).
+    fn foreign_to(&self, owner: u64) -> bool {
+        self._lease.as_ref().is_some_and(|l| l.owner != owner)
+    }
 }
 
 /// Hands a pooled texture back to its [`Pool`] when the last [`GpuImage`] holding it drops.
@@ -74,6 +98,10 @@ struct Lease {
     pool: Weak<Mutex<Pool>>,
     /// The [`Enc`] that took it (see [`Pool::owned`]).
     owner: u64,
+    /// Another [`Enc`] recorded work on it ([`Enc::uses`], [`GpuImage::escape`]):
+    /// its owner can no longer reuse it at once, since that other encoder may still be open, so
+    /// it is released through the guarded [`Pool::release`].
+    escaped: std::sync::atomic::AtomicBool,
     /// The pool's [`Pool::generation`] when it was taken.
     generation: u64,
 }
@@ -89,7 +117,15 @@ impl Drop for Lease {
         if let Some(pool) = self.pool.upgrade()
             && let Ok(mut p) = pool.lock()
         {
-            p.release_owned(self.texture.clone(), self.owner, self.generation);
+            // Always undo `pooled`'s leased count first: `Pool::discard` does not reset
+            // `leased`, and a stale-generation lease must still clear its contribution.
+            p.leased = p.leased.saturating_sub(tex_bytes(&self.texture));
+            if !self.escaped.load(std::sync::atomic::Ordering::Relaxed) {
+                p.release_owned(self.texture.clone(), self.owner, self.generation);
+            } else if self.generation == p.generation {
+                // Like `release_owned`, nothing taken before a failed frame comes back (#1319).
+                p.release(self.texture.clone());
+            }
         }
     }
 }
@@ -103,16 +139,42 @@ impl Drop for Lease {
 /// it was released has been submitted (or dropped).
 #[derive(Default)]
 pub(crate) struct Pool {
-    free: HashMap<(u32, u32), Vec<wgpu::Texture>>,
+    /// Free textures by size, each with the [`Pool::stamp`] at which it was freed (oldest first).
+    free: HashMap<(u32, u32), std::collections::VecDeque<(wgpu::Texture, u64)>>,
+    /// Counts frees, so the free list evicts the oldest texture first.
+    stamp: u64,
+    /// Bytes the free list may hold (0 = [`POOL_BUDGET`]); tests lower it per context.
+    budget: usize,
     /// (texture, newest encoder open at its release).
     released: Vec<(wgpu::Texture, u64)>,
     /// Encoders with recorded, unsubmitted work.
     open: BTreeSet<u64>,
     next: u64,
     bytes: usize,
+    /// Bytes currently held by live [`Lease`]s.
+    leased: usize,
     /// Working textures reused / newly allocated (diagnostics, tests).
     hits: u64,
     misses: u64,
+    /// Free textures dropped to make room for newer ones (diagnostics, tests).
+    evicted: u64,
+    /// Output frames done so far ([`GpuContext::pool_frame_end`]).
+    frame: u64,
+    /// Per size, the [`Pool::frame`] in which it was last asked for (hit or miss).
+    asked: HashMap<(u32, u32), u64>,
+    /// Bytes of evicted textures whose GPU memory may not be freed yet: wgpu frees a dropped
+    /// texture only once the submissions that used it are done. Counted against the budget
+    /// for evicting, so evicting and reallocating cannot stack up freed-but-held memory.
+    dying: Arc<std::sync::atomic::AtomicUsize>,
+    /// The [`Pool::dying`] bytes left after a [`GpuContext::pool_frame_end`] wait that cleared
+    /// none of them (`0` = not stalled). While `dying` does not drop below it, later frames only
+    /// poll instead of waiting again (each wait could take up to [`crate::readback::TIMEOUT`]
+    /// when the queue stopped completing work).
+    stalled: usize,
+    /// Number of empty-submission waits from [`GpuContext::pool_frame_end`] (diagnostics, tests).
+    waits: u64,
+    /// Reports when the submissions that may use an evicted texture are done.
+    queue: Option<wgpu::Queue>,
     /// Textures released while the [`Enc`] that took them is alive, kept for that `Enc` only.
     /// Its later commands run after everything it recorded or submitted before (one queue,
     /// recording order), so it can overwrite them at once instead of allocating a new texture
@@ -234,13 +296,100 @@ impl Pool {
         }
     }
 
+    fn budget(&self) -> usize {
+        if self.budget == 0 { POOL_BUDGET } else { self.budget }
+    }
+
+    /// Keep `t` on the free list if it fits. Otherwise make room by evicting stale sizes, least
+    /// recently freed first: sizes not asked for in the last [`STALE_FRAMES`] output frames (a
+    /// blur bound or padded layer size that no longer occurs), so they age out instead of
+    /// filling the budget for good. Sizes still in use are never evicted for a newcomer: a
+    /// frame that frees more than the budget in a repeating order would otherwise evict, on
+    /// every free, the texture the next call takes first (LRU on a cycle), and miss on every
+    /// take. If nothing stale makes room, `t` is dropped (the kept set keeps hitting).
+    ///
+    /// Evicted textures count against the budget ([`Pool::dying`]) until the GPU work
+    /// submitted so far is done. That throttles eviction to roughly one budget ahead of the
+    /// GPU; it is not a hard bound (wgpu frees a dropped texture only when the device is
+    /// maintained, which can be later than the work-done callback).
     fn make_free(&mut self, t: wgpu::Texture) {
+        use std::sync::atomic::Ordering::Relaxed;
         let b = tex_bytes(&t);
-        if self.bytes + b > POOL_BUDGET {
+        let budget = self.budget();
+        if b > budget {
+            return;
+        }
+        let mut evicted = 0;
+        while self.bytes + b > budget {
+            let frame = self.frame;
+            let asked = &self.asked;
+            let stale = |k: &(u32, u32)| asked.get(k).is_none_or(|f| frame.saturating_sub(*f) >= STALE_FRAMES);
+            let Some(oldest) = self.free.iter().filter(|(k, _)| stale(k)).filter_map(|(k, q)| q.front().map(|(_, s)| (*s, *k))).min().map(|(_, k)| k) else {
+                break;
+            };
+            let ob = oldest.0 as usize * oldest.1 as usize * 16;
+            if self.dying.load(Relaxed) + evicted + ob > budget {
+                break;
+            }
+            let Some(q) = self.free.get_mut(&oldest) else { break };
+            if let Some((old, _)) = q.pop_front() {
+                self.bytes -= tex_bytes(&old);
+                evicted += ob;
+                self.evicted += 1;
+            }
+            if q.is_empty() {
+                self.free.remove(&oldest);
+            }
+        }
+        if evicted > 0 {
+            self.dying.fetch_add(evicted, Relaxed);
+            let dying = self.dying.clone();
+            match &self.queue {
+                Some(q) => q.on_submitted_work_done(move || {
+                    let _ = dying.try_update(Relaxed, Relaxed, |d| Some(d.saturating_sub(evicted)));
+                }),
+                None => {
+                    let _ = dying.try_update(Relaxed, Relaxed, |d| Some(d.saturating_sub(evicted)));
+                }
+            }
+        }
+        if self.bytes + b > budget {
             return;
         }
         self.bytes += b;
-        self.free.entry((t.width(), t.height())).or_default().push(t);
+        self.stamp += 1;
+        self.free.entry((t.width(), t.height())).or_default().push_back((t, self.stamp));
+    }
+
+    /// An output frame is done: sizes not asked for since age towards [`STALE_FRAMES`].
+    /// Returns the evicted bytes the caller should wait for: 0 when none still count, or when a
+    /// previous wait cleared none and none has been freed since.
+    fn frame_end(&mut self) -> usize {
+        self.frame += 1;
+        let (frame, free) = (self.frame, &self.free);
+        self.asked.retain(|k, f| frame - *f < STALE_FRAMES || free.contains_key(k));
+        let d = self.dying.load(std::sync::atomic::Ordering::Relaxed);
+        if d == 0 {
+            self.stalled = 0;
+            return 0;
+        }
+        if self.stalled != 0 && d >= self.stalled {
+            // No progress since the stalled wait: do not wait again.
+            return 0;
+        }
+        self.stalled = 0;
+        d
+    }
+
+    /// A [`GpuContext::pool_frame_end`] wait finished with `before` dying bytes still counted
+    /// beforehand. A wait that freed nothing (timed out, or callbacks not run synchronously,
+    /// e.g. on WebGPU where poll does not block) marks the pool stalled.
+    fn waited(&mut self, before: usize) {
+        self.waits += 1;
+        let d = self.dying.load(std::sync::atomic::Ordering::Relaxed);
+        if d >= before {
+            self.stalled = d;
+        }
     }
 
     fn open(&mut self) -> u64 {
@@ -291,14 +440,92 @@ impl Pool {
         self.release(t);
     }
 
+    /// An owner for [`Enc::lent`]. It is a fresh id each time, so a lease of an earlier lent
+    /// `Enc` that drops later (an image that outlived its call) finds no owner and goes
+    /// through [`Pool::release`], which waits for every open encoder.
+    fn lend(&mut self) -> u64 {
+        self.register()
+    }
+
+    /// [`Enc::lent`] is done with `owner` (its work submitted or discarded): its kept textures
+    /// go to the shared free list at once. The free list has one [`POOL_BUDGET`] (only stale
+    /// sizes are evicted), so finished calls no longer keep a budget each (16 frame threads
+    /// kept up to 16 × `POOL_BUDGET`). The budget caps the free list only: textures waiting
+    /// for open encoders ([`Pool::released`]) and those held by live `Enc`s are outside it.
+    ///
+    /// They skip [`Pool::release`]'s wait for open encoders: a kept texture's lease has
+    /// dropped and only this `Enc` recorded work on it, which is already submitted or
+    /// discarded, so later work cannot overtake it on the queue. That is enforced, not assumed:
+    /// a lease read by another `Enc` is marked escaped
+    /// ([`GpuImage::escape`]) and released through [`Pool::release`] whenever it drops, and a
+    /// lease that drops after this call finds no owner and goes the same way.
+    fn give_back(&mut self, owner: u64) {
+        if let Some(m) = self.owned.remove(&owner) {
+            for t in m.into_values().flatten() {
+                self.make_free(t);
+            }
+        }
+    }
+
+    /// Moves `owner`'s returned working textures onto the shared free list (dropping any over
+    /// budget) right after `owner`'s encoder has been submitted, while keeping `owner` registered.
+    fn flush_owned(&mut self, owner: u64) {
+        let Some(m) = self.owned.get_mut(&owner).map(std::mem::take) else { return };
+        for t in m.into_values().flatten() {
+            self.make_free(t);
+        }
+    }
+
+    /// Bytes the pool will still keep alive (live `Lease`s, other owners' in-flight working sets,
+    /// released textures still guarded by other open encoders, and up to `budget` on the shared
+    /// free list) once `owner` flushes its returned textures and `closing_enc` submits.
+    fn kept_after_flush(&self, owner: u64, closing_enc: Option<u64>) -> usize {
+        let oldest = self.open.iter().copied().find(|&o| Some(o) != closing_enc).unwrap_or(u64::MAX);
+        let mut ready = 0usize;
+        let mut still_released = 0usize;
+        for (t, w) in &self.released {
+            let b = tex_bytes(t);
+            if *w < oldest {
+                ready += b;
+            } else {
+                still_released += b;
+            }
+        }
+        let mut self_owned = 0usize;
+        let mut other_owned = 0usize;
+        for (&k, m) in &self.owned {
+            let b: usize = m.values().flatten().map(tex_bytes).sum();
+            if k == owner {
+                self_owned += b;
+            } else {
+                other_owned += b;
+            }
+        }
+        let free_after = self.bytes.max((self.bytes + ready + self_owned).min(self.budget()));
+        self.leased + other_owned + still_released + free_after
+    }
+
+    /// Bytes of working textures kept for reuse: the free list, those waiting for open
+    /// encoders and those kept by live `Enc`s.
+    fn retained(&self) -> usize {
+        let owned = self.owned.values().flat_map(|m| m.values().flatten());
+        self.released.iter().map(|(t, _)| t).chain(owned).map(tex_bytes).fold(self.bytes, usize::saturating_add)
+    }
+
     fn take_owned(&mut self, owner: u64, w: u32, h: u32) -> Option<wgpu::Texture> {
         let t = self.owned.get_mut(&owner)?.get_mut(&(w, h))?.pop()?;
+        self.asked.insert((w, h), self.frame);
         self.hits += 1;
         Some(t)
     }
 
     fn take(&mut self, w: u32, h: u32) -> Option<wgpu::Texture> {
-        let t = self.free.get_mut(&(w, h)).and_then(Vec::pop);
+        self.asked.insert((w, h), self.frame);
+        let q = self.free.get_mut(&(w, h));
+        let t = q.and_then(|q| q.pop_back()).map(|(t, _)| t);
+        if self.free.get(&(w, h)).is_some_and(|q| q.is_empty()) {
+            self.free.remove(&(w, h));
+        }
         match &t {
             Some(t) => {
                 self.bytes -= tex_bytes(t);
@@ -411,6 +638,11 @@ pub struct GpuContext {
     /// Export frame → YUV 4:2:0 kernel (built on first use).
     pub(crate) yuv: std::sync::OnceLock<Result<crate::yuv::YuvPipe, String>>,
     readbacks: Arc<crate::readback::Readbacks>,
+    /// Serializes `device.poll` across threads (`wgpu-core 30.0.1` can panic on concurrent
+    /// `Device::maintain` when two callers poll the same device at once). Shared by every
+    /// context on the same device.
+    poll_lock: Arc<Mutex<()>>,
+    memcap: crate::memcap::MemCap,
 }
 
 fn tex_entry(binding: u32) -> wgpu::BindGroupLayoutEntry {
@@ -571,9 +803,12 @@ impl GpuContext {
             include_str!("shaders/fx_time.wgsl"),
             include_str!("shaders/fx_pixel2.wgsl"),
             include_str!("shaders/fx_gen2.wgsl"),
+            include_str!("shaders/inpaint.wgsl"),
             include_str!("shaders/sky.wgsl"),
         ]
         .concat();
+        #[cfg(test)]
+        let src = src + crate::tests::MATH_PROBE_SHADER;
         let module = init_resource(&device, "kernels shader module", || {
             device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("effectcraft kernels"), source: wgpu::ShaderSource::Wgsl(src.into()) })
         })?;
@@ -657,6 +892,8 @@ impl GpuContext {
                 }
             }
         }
+        #[cfg(test)]
+        pipelines.insert("math_probe", build("math_probe", &layout)?);
         for e in EFFECT_ENTRIES
             .iter()
             .chain(crate::fx_color::KERNELS)
@@ -679,6 +916,7 @@ impl GpuContext {
             .chain(crate::fx_time::KERNELS)
             .chain(crate::fx_pixel2::KERNELS)
             .chain(crate::fx_gen2::KERNELS)
+            .chain(crate::inpaint::KERNELS)
         {
             match build(e, &layout) {
                 Ok(p) => {
@@ -739,6 +977,9 @@ impl GpuContext {
             init_resource(&device, "dummy storage buffer", || buffer_init(&device, &queue, Some("dummy"), &[0u8; 16], wgpu::BufferUsages::STORAGE))?;
         let dummy_tex = init_resource(&device, "dummy texture view", || dummy.create_view(&Default::default()))?;
         let max_dim = device.limits().max_texture_dimension_2d;
+        let memcap = crate::memcap::MemCap::new(&info);
+        let pool_budget = if memcap.cap() > 0 { POOL_BUDGET.min((memcap.cap() / 4) as usize) } else { 0 };
+        let pool = Arc::new(Mutex::new(Pool { queue: Some(queue.clone()), budget: pool_budget, ..Default::default() }));
         Ok(GpuContext {
             device,
             queue,
@@ -758,7 +999,7 @@ impl GpuContext {
             uploads: Mutex::new(Uploads::default()),
             upload_budget: std::sync::atomic::AtomicUsize::new(UPLOAD_BUDGET),
             upload_full_warned: Default::default(),
-            pool: Default::default(),
+            pool,
             staging: Default::default(),
             transfers: Default::default(),
             zeros: Default::default(),
@@ -770,6 +1011,8 @@ impl GpuContext {
             particle_states: Default::default(),
             deferred: None,
             readbacks,
+            poll_lock: Default::default(),
+            memcap,
         })
     }
 
@@ -803,10 +1046,25 @@ impl GpuContext {
         crate::readback::failure_notifier(&self.readbacks)
     }
 
+    /// Maximum time a blocking poll holds `poll_lock` in one slice so concurrent readbacks are not
+    /// starved behind a long `mem_gate` or readback wait.
+    const POLL_SLICE: std::time::Duration = std::time::Duration::from_millis(20);
+
     /// All device polling routes through this failure bridge.
     pub(crate) fn poll_readbacks(&self, kind: wgpu::PollType) -> Result<wgpu::PollStatus, String> {
         self.readbacks.check()?;
-        self.device.poll(kind).map_err(|e| {
+        let res = {
+            let _guard = match kind {
+                wgpu::PollType::Poll => match self.poll_lock.try_lock() {
+                    Ok(g) => g,
+                    Err(std::sync::TryLockError::Poisoned(e)) => e.into_inner(),
+                    Err(std::sync::TryLockError::WouldBlock) => return Ok(wgpu::PollStatus::Poll),
+                },
+                _ => self.poll_lock.lock().unwrap_or_else(|e| e.into_inner()),
+            };
+            self.device.poll(kind)
+        };
+        res.map_err(|e| {
             let error = format!("GPU readback poll: {e}");
             self.readbacks.retire(error.clone());
             log::error!("{error}");
@@ -814,6 +1072,73 @@ impl GpuContext {
         })
     }
 
+    /// Waits on `submission_index` up to `timeout` in [`Self::POLL_SLICE`] increments so wgpu
+    /// frees resources dropped before that submission without holding `poll_lock` for seconds at a
+    /// time. A timeout simply moves on without retiring the context or its readbacks; any other
+    /// poll error retires the readbacks through the failure path.
+    fn poll_wait(&self, submission_index: wgpu::SubmissionIndex, timeout: std::time::Duration) {
+        if self.readbacks.check().is_err() {
+            return;
+        }
+        let deadline = web_time::Instant::now() + timeout;
+        loop {
+            let remaining = deadline.saturating_duration_since(web_time::Instant::now());
+            let slice = remaining.min(Self::POLL_SLICE);
+            let res = {
+                let _guard = self.poll_lock.lock().unwrap_or_else(|e| e.into_inner());
+                self.device.poll(wgpu::PollType::Wait { submission_index: Some(submission_index.clone()), timeout: Some(slice) })
+            };
+            let done = match res {
+                Ok(_) => true,
+                Err(wgpu::PollError::Timeout) => false,
+                Err(e) => {
+                    let error = format!("GPU poll wait: {e}");
+                    self.readbacks.retire(error.clone());
+                    log::error!("{error}");
+                    true
+                }
+            };
+            if done || remaining <= slice {
+                break;
+            }
+        }
+    }
+
+    /// Bytes held by wgpu's device allocator (`0` on backends that do not report).
+    pub(crate) fn allocated_bytes(&self) -> u64 {
+        self.device.generate_allocator_report().map_or(0, |r| r.total_allocated_bytes)
+    }
+
+    /// Bytes held in long-lived or still-referenced GPU caches/leases that submitting and waiting
+    /// on `owner` (`closing_enc`) cannot free.
+    fn kept_bytes(&self, owner: u64, closing_enc: Option<u64>) -> u64 {
+        let uploads = self.uploads.lock().map(|u| u.bytes as u64).unwrap_or(0);
+        let pool = self.pool.lock().map(|p| p.kept_after_flush(owner, closing_enc) as u64).unwrap_or(0);
+        let staging: u64 = self.staging.lock().map(|s| s.iter().map(|(k, v)| k * v.len() as u64).sum()).unwrap_or(0);
+        let zeros: u64 = self.zeros.lock().map(|z| z.keys().map(|&(w, h)| w as u64 * h as u64 * 16).sum()).unwrap_or(0);
+        uploads + pool + staging + zeros
+    }
+
+    #[cfg(test)]
+    pub(crate) fn memcap_waits(&self) -> u64 {
+        self.memcap.waits()
+    }
+
+    #[cfg(test)]
+    pub(crate) fn memcap_warned(&self) -> bool {
+        self.memcap.warned()
+    }
+
+    #[cfg(test)]
+    pub(crate) fn set_memcap_wait_timeout(&self, timeout: std::time::Duration) {
+        self.memcap.set_wait_timeout(timeout);
+    }
+
+    #[cfg(test)]
+    pub(crate) fn lock_poll_for_test(&self) -> std::sync::MutexGuard<'_, ()> {
+        self.poll_lock.lock().unwrap_or_else(|e| e.into_inner())
+    }
+
     /// Start a bounded buffer mapping. Retired contexts reject new requests immediately.
     pub(crate) fn map_readback(
         &self,
@@ -848,25 +1173,52 @@ impl GpuContext {
         });
     }
 
-    /// A bounded native readback: polling and receiving share one deadline.
+    /// A bounded native readback: polling and receiving share one deadline, started once
+    /// `poll_lock` is first acquired and polled in [`Self::POLL_SLICE`] slices so lock contention
+    /// never consumes the readback deadline or holds the mutex across the whole timeout.
     #[cfg(not(target_arch = "wasm32"))]
     pub(crate) fn read_buffer(&self, buffer: &wgpu::Buffer, row: usize, stride: usize, height: usize) -> Result<Vec<u8>, String> {
-        let deadline = std::time::Instant::now() + crate::readback::TIMEOUT;
         let (tx, rx) = std::sync::mpsc::channel();
         self.map_readback(buffer, row, stride, height, move |r| {
             let _ = tx.send(r);
         });
+        let mut deadline = None;
         loop {
             self.readbacks.check()?;
-            let remaining = deadline.saturating_duration_since(std::time::Instant::now());
+            if let Ok(result) = rx.try_recv() {
+                return result;
+            }
+            let (remaining, slice, poll_res) = {
+                let _guard = self.poll_lock.lock().unwrap_or_else(|e| e.into_inner());
+                self.readbacks.check()?;
+                let d = *deadline.get_or_insert_with(|| std::time::Instant::now() + crate::readback::TIMEOUT);
+                let remaining = d.saturating_duration_since(std::time::Instant::now());
+                if remaining.is_zero() {
+                    (remaining, std::time::Duration::ZERO, Ok(wgpu::PollStatus::Poll))
+                } else {
+                    let slice = remaining.min(Self::POLL_SLICE);
+                    let res = self.device.poll(wgpu::PollType::Wait { submission_index: None, timeout: Some(slice) });
+                    (remaining, slice, res)
+                }
+            };
+            if let Ok(result) = rx.try_recv() {
+                return result;
+            }
             if remaining.is_zero() {
                 let error = "GPU readback receive timed out; context retired".to_string();
                 self.readbacks.retire(error.clone());
                 return Err(error);
             }
-            // Block until the GPU is done (bounded by the deadline: a timeout retires the
-            // context) rather than polling and sleeping, which adds latency to every readback.
-            self.poll_readbacks(wgpu::PollType::Wait { submission_index: None, timeout: Some(remaining) })?;
+            match poll_res {
+                Ok(_) => {}
+                Err(wgpu::PollError::Timeout) if remaining > slice => {}
+                Err(e) => {
+                    let error = format!("GPU readback poll: {e}");
+                    self.readbacks.retire(error.clone());
+                    log::error!("{error}");
+                    return Err(error);
+                }
+            }
             // The callback may be running on another thread that polled the same device.
             match rx.recv_timeout(remaining.min(std::time::Duration::from_millis(1))) {
                 Ok(result) => return result,
@@ -1005,6 +1357,8 @@ impl GpuContext {
             readbacks: self.readbacks.clone(),
             software: self.software,
             yuv: std::sync::OnceLock::new(),
+            poll_lock: self.poll_lock.clone(),
+            memcap: self.memcap.fresh(),
         }
     }
 
@@ -1051,10 +1405,19 @@ impl GpuContext {
     /// one is zeroed.
     fn pooled(&self, w: u32, h: u32, owner: u64) -> (GpuImage, bool) {
         let (w, h) = (w.max(1), h.max(1));
-        let (reused, generation) = self.pool.lock().map(|mut p| (p.take_owned(owner, w, h).or_else(|| p.take(w, h)), p.generation)).unwrap_or((None, u64::MAX));
+        let bytes = w as usize * h as usize * 16;
+        let (reused, generation) = self
+            .pool
+            .lock()
+            .map(|mut p| {
+                let t = p.take_owned(owner, w, h).or_else(|| p.take(w, h));
+                p.leased += bytes;
+                (t, p.generation)
+            })
+            .unwrap_or((None, u64::MAX));
         let was_reused = reused.is_some();
         let texture = reused.unwrap_or_else(|| self.texture(w, h));
-        let lease = Arc::new(Lease { texture: texture.clone(), pool: Arc::downgrade(&self.pool), owner, generation });
+        let lease = Arc::new(Lease { texture: texture.clone(), pool: Arc::downgrade(&self.pool), owner, escaped: Default::default(), generation });
         (GpuImage { texture, width: w, height: h, _lease: Some(lease), levels: None }, was_reused)
     }
 
@@ -1081,6 +1444,73 @@ impl GpuContext {
         }
     }
 
+    /// Bytes of working textures the pool keeps for reuse (see [`GpuContext::trim_pool`]).
+    pub fn pool_retained_bytes(&self) -> usize {
+        self.pool.lock().map(|p| p.retained()).unwrap_or_default()
+    }
+
+    /// Drop the working textures kept for reuse (a render is done). Textures still waiting for
+    /// an open encoder and those of live `Enc`s are unaffected.
+    pub fn trim_pool(&self) {
+        if let Ok(mut p) = self.pool.lock() {
+            p.clear();
+        }
+    }
+
+    /// An output frame is done: the pool ages the sizes not asked for since (only sizes unused
+    /// for [`STALE_FRAMES`] frames are evicted to make room for others). When evicted textures
+    /// still count ([`Pool::dying`]), waits once so their work-done callbacks can clear that
+    /// count; if the wait frees none of them, later frames only poll until `dying` drops (a
+    /// stopped queue must not re-wait [`crate::readback::TIMEOUT`] on every frame).
+    pub fn pool_frame_end(&self) {
+        let before = self.pool.lock().map(|mut p| p.frame_end()).unwrap_or_default();
+        if before > 0 {
+            // A work-done callback waits for a submission; an empty one closes the wait.
+            let index = self.queue.submit([]);
+            self.poll_wait(index, crate::readback::TIMEOUT);
+            if let Ok(mut p) = self.pool.lock() {
+                p.waited(before);
+            }
+        } else {
+            let _ = self.poll_readbacks(wgpu::PollType::Poll);
+        }
+    }
+
+    /// Bytes the pool's free list may keep (tests lower it on a context of their own).
+    #[cfg(test)]
+    pub(crate) fn set_pool_budget(&self, bytes: usize) {
+        if let Ok(mut p) = self.pool.lock() {
+            p.budget = bytes;
+        }
+    }
+
+    /// Free textures evicted to make room for newer ones.
+    #[cfg(test)]
+    pub(crate) fn pool_evictions(&self) -> u64 {
+        self.pool.lock().map(|p| p.evicted).unwrap_or_default()
+    }
+
+    /// Empty-submission waits from [`GpuContext::pool_frame_end`].
+    #[cfg(test)]
+    pub(crate) fn pool_waits(&self) -> u64 {
+        self.pool.lock().map(|p| p.waits).unwrap_or_default()
+    }
+
+    /// Add or subtract phantom [`Pool::dying`] bytes without a work-done callback (tests that
+    /// simulate a queue which stopped completing work).
+    #[cfg(test)]
+    pub(crate) fn pool_charge_dying(&self, bytes: isize) {
+        use std::sync::atomic::Ordering::Relaxed;
+        if let Ok(p) = self.pool.lock() {
+            if bytes >= 0 {
+                p.dying.fetch_add(bytes as usize, Relaxed);
+            } else {
+                let sub = (-bytes) as usize;
+                let _ = p.dying.try_update(Relaxed, Relaxed, |d| Some(d.saturating_sub(sub)));
+            }
+        }
+    }
+
     /// Texture pool counters: (working textures reused, newly allocated).
     pub fn pool_stats(&self) -> (u64, u64) {
         self.pool.lock().map(|p| (p.hits, p.misses)).unwrap_or_default()
@@ -1192,11 +1622,14 @@ pub(crate) struct Enc<'g> {
     pub g: &'g GpuContext,
     /// The open command encoder and its id in the texture pool.
     enc: Option<(wgpu::CommandEncoder, u64)>,
+    last_submit: Option<wgpu::SubmissionIndex>,
     pending: usize,
     /// A dispatch named a kernel the device doesn't have (see [`Enc::take_missing`]).
     missing: bool,
     /// This `Enc`'s id in the texture pool (textures it may reuse at once).
     owner: u64,
+    /// `owner` is lent ([`Enc::lent`]) and outlives this `Enc`.
+    lent: bool,
 }
 
 /// Images bound to a compute kernel's source, optional auxiliary and output slots.
@@ -1213,7 +1646,11 @@ impl Drop for Enc<'_> {
             if let Some((_, id)) = self.enc.take() {
                 p.close(id);
             }
-            p.unregister(self.owner);
+            if self.lent {
+                p.give_back(self.owner);
+            } else {
+                p.unregister(self.owner);
+            }
         }
     }
 }
@@ -1221,7 +1658,27 @@ impl Drop for Enc<'_> {
 impl<'g> Enc<'g> {
     pub fn new(g: &'g GpuContext) -> Enc<'g> {
         let owner = g.pool.lock().map(|mut p| p.register()).unwrap_or(0);
-        Enc { g, enc: None, pending: 0, missing: false, owner }
+        Enc { g, enc: None, last_submit: None, pending: 0, missing: false, owner, lent: false }
+    }
+
+    /// An `Enc` for self-contained work (an effect chain): once it is dropped, its working
+    /// textures can be reused at once ([`Pool::give_back`]). Its images stay with it, except a
+    /// result another `Enc` records work on ([`Enc::uses`]), whose lease then releases through
+    /// the guarded [`Pool::release`]. A plain [`Enc::new`] hands its
+    /// textures back through the shared pool, where they wait for every open encoder, and a
+    /// GPU frame keeps its encoder open while the CPU renders layers with GPU effects, so every
+    /// effect's intermediates stayed allocated until the frame ended.
+    pub(crate) fn lent(g: &'g GpuContext) -> Enc<'g> {
+        let owner = g.pool.lock().map(|mut p| p.lend()).unwrap_or(0);
+        Enc { g, enc: None, last_submit: None, pending: 0, missing: false, owner, lent: true }
+    }
+
+    /// `img` is used by this `Enc`: if another `Enc` made it, that one must not reuse it at once
+    /// (this encoder may still be open when its lease drops), so mark it escaped.
+    pub(crate) fn uses(&self, img: &GpuImage) {
+        if img.foreign_to(self.owner) {
+            img.escape();
+        }
     }
 
     /// Whether a dispatch since the last call named a kernel the device doesn't have (one
@@ -1230,7 +1687,21 @@ impl<'g> Enc<'g> {
         std::mem::take(&mut self.missing)
     }
 
-    pub(crate) fn encoder(&mut self) -> &mut wgpu::CommandEncoder {
+    /// The command encoder, to record work that touches `imgs` (every pooled image the commands
+    /// read or write; buffers and textures of the caller's own need not be named). The only
+    /// way to record outside this module: naming the images marks foreign leases escaped
+    /// ([`Enc::uses`]). Callers must name every pooled image the encoder's work uses, here or
+    /// with separate [`Enc::uses`] calls (`record(&[])` plus `uses()` is allowed); this is a
+    /// convention, not enforced, and skipping it reintroduces the reuse race it guards (#1441).
+    pub(crate) fn record(&mut self, imgs: &[&GpuImage]) -> &mut wgpu::CommandEncoder {
+        for img in imgs {
+            self.uses(img);
+        }
+        self.encoder()
+    }
+
+    /// The raw encoder: within this module only, after [`Enc::uses`] on every pooled image.
+    fn encoder(&mut self) -> &mut wgpu::CommandEncoder {
         let g = self.g;
         &mut self
             .enc
@@ -1329,7 +1800,8 @@ impl<'g> Enc<'g> {
                 ],
             })
         });
-        let enc = self.encoder();
+        let imgs: Vec<&GpuImage> = [Some(src), aux, Some(out)].into_iter().flatten().collect();
+        let enc = self.record(&imgs);
         {
             let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some(entry), timestamp_writes: None });
             pass.set_pipeline(pipe);
@@ -1354,7 +1826,7 @@ impl<'g> Enc<'g> {
 
     /// Copy `src` into a larger zeroed image at (`x`, `y`) (Buf::pad).
     pub fn copy_into(&mut self, src: &GpuImage, dst: &GpuImage, x: u32, y: u32) {
-        let enc = self.encoder();
+        let enc = self.record(&[src, dst]);
         enc.copy_texture_to_texture(
             wgpu::TexelCopyTextureInfo { texture: &src.texture, mip_level: 0, origin: wgpu::Origin3d::ZERO, aspect: wgpu::TextureAspect::All },
             wgpu::TexelCopyTextureInfo { texture: &dst.texture, mip_level: 0, origin: wgpu::Origin3d { x, y, z: 0 }, aspect: wgpu::TextureAspect::All },
@@ -1371,7 +1843,7 @@ impl<'g> Enc<'g> {
             (b > a).then(|| (a as u32, (a - o) as u32, (b - a) as u32))
         };
         let (Some((sx, dx, cw)), Some((sy, dy, ch))) = (span(x, w, src.width), span(y, h, src.height)) else { return dst };
-        let enc = self.encoder();
+        let enc = self.record(&[src, &dst]);
         enc.copy_texture_to_texture(
             wgpu::TexelCopyTextureInfo { texture: &src.texture, mip_level: 0, origin: wgpu::Origin3d { x: sx, y: sy, z: 0 }, aspect: wgpu::TextureAspect::All },
             wgpu::TexelCopyTextureInfo { texture: &dst.texture, mip_level: 0, origin: wgpu::Origin3d { x: dx, y: dy, z: 0 }, aspect: wgpu::TextureAspect::All },
@@ -1383,7 +1855,7 @@ impl<'g> Enc<'g> {
     /// Submit recorded work.
     pub fn submit(&mut self) {
         if let Some((enc, id)) = self.enc.take() {
-            self.g.queue.submit([enc.finish()]);
+            self.last_submit = Some(self.g.queue.submit([enc.finish()]));
             if let Ok(mut p) = self.g.pool.lock() {
                 p.close(id);
             }
@@ -1391,8 +1863,34 @@ impl<'g> Enc<'g> {
         self.pending = 0;
     }
 
+    /// Submits and waits on this encoder's own submission when the device allocator is over the
+    /// memory cap and enough of that memory is reclaimable, so wgpu frees textures dropped so far.
+    pub fn mem_gate(&mut self) {
+        if !self.g.memcap.due() {
+            return;
+        }
+        let allocated = self.g.allocated_bytes();
+        if allocated <= self.g.memcap.cap() {
+            let _ = self.g.memcap.should_wait(allocated, 0);
+            return;
+        }
+        let closing = self.enc.as_ref().map(|(_, id)| *id);
+        let kept = self.g.kept_bytes(self.owner, closing);
+        if !self.g.memcap.should_wait(allocated, kept) {
+            return;
+        }
+        self.submit();
+        if let Ok(mut p) = self.g.pool.lock() {
+            p.flush_owned(self.owner);
+        }
+        let idx = self.last_submit.take().unwrap_or_else(|| self.g.queue.submit([]));
+        self.g.poll_wait(idx, self.g.memcap.wait_timeout());
+        self.g.memcap.record_wait(self.g.allocated_bytes(), self.g.kept_bytes(self.owner, None));
+    }
+
     /// Read an image back to the CPU (submits and waits). `None` where readback is impossible.
     pub fn download(&mut self, img: &GpuImage) -> Option<Image> {
+        self.uses(img);
         let layout = crate::readback::Layout::new(img.width, img.height, 16, self.g.device.limits().max_buffer_size)
             .map_err(|e| log::error!("gpu image readback: {e}"))
             .ok()?;
@@ -1401,23 +1899,31 @@ impl<'g> Enc<'g> {
         data.resize(layout.len / 16, [0.0; 4]);
         let mut out = Image { width: img.width, height: img.height, data };
         let dst: &mut [u8] = bytemuck::cast_slice_mut(&mut out.data);
-        self.read_texture_into(&img.texture, img.width, img.height, 16, dst)?;
+        self.read_texture_into(&img.texture, [0, 0, img.width, img.height], 16, dst)?;
         Some(out)
     }
 
     /// Tightly packed texel bytes of a texture (`bpp` bytes per pixel).
     pub fn read_texture(&mut self, texture: &wgpu::Texture, w: u32, h: u32, bpp: u32) -> Option<Vec<u8>> {
+        self.read_rect(texture, 0, 0, w, h, bpp)
+    }
+
+    /// [`Enc::read_texture`] of a sub-rectangle at (`x`, `y`): a fill only changes its hole, so
+    /// only the hole's bounds need to cross back (`bytes_per_row` is still row-aligned inside).
+    pub fn read_rect(&mut self, texture: &wgpu::Texture, x: u32, y: u32, w: u32, h: u32, bpp: u32) -> Option<Vec<u8>> {
         let layout =
             crate::readback::Layout::new(w, h, bpp, self.g.device.limits().max_buffer_size).map_err(|e| log::error!("gpu texture readback: {e}")).ok()?;
         let mut out = Vec::new();
         out.try_reserve_exact(layout.len).map_err(|e| log::error!("gpu readback allocation: {e}")).ok()?;
         out.resize(layout.len, 0);
-        self.read_texture_into(texture, w, h, bpp, &mut out)?;
+        self.read_texture_into(texture, [x, y, w, h], bpp, &mut out)?;
         Some(out)
     }
 
-    /// [`Enc::read_texture`] into `out` (`w × bpp × h` bytes), through a reused staging buffer.
-    fn read_texture_into(&mut self, texture: &wgpu::Texture, w: u32, h: u32, bpp: u32, out: &mut [u8]) -> Option<()> {
+    /// [`Enc::read_rect`] into `out` (`w × bpp × h` bytes), through a reused staging buffer.
+    /// `rect` is `[x, y, w, h]`.
+    fn read_texture_into(&mut self, texture: &wgpu::Texture, rect: [u32; 4], bpp: u32, out: &mut [u8]) -> Option<()> {
+        let [x, y, w, h] = rect;
         if !self.g.can_wait() {
             return None;
         }
@@ -1440,7 +1946,7 @@ impl<'g> Enc<'g> {
             })
         });
         self.encoder().copy_texture_to_buffer(
-            wgpu::TexelCopyTextureInfo { texture, mip_level: 0, origin: wgpu::Origin3d::ZERO, aspect: wgpu::TextureAspect::All },
+            wgpu::TexelCopyTextureInfo { texture, mip_level: 0, origin: wgpu::Origin3d { x, y, z: 0 }, aspect: wgpu::TextureAspect::All },
             wgpu::TexelCopyBufferInfo { buffer: &buf, layout: wgpu::TexelCopyBufferLayout { offset: 0, bytes_per_row: Some(padded), rows_per_image: Some(h) } },
             wgpu::Extent3d { width: w, height: h, depth_or_array_layers: 1 },
         );
@@ -1509,6 +2015,7 @@ impl<'g> Enc<'g> {
     /// once they arrived (`Some(None)`: the readback failed), else `None` after starting the
     /// copy if needed and marking the pass missed.
     pub(crate) fn read_keyed(&mut self, img: &GpuImage, key: u128, offset: [f64; 2], scale: f64) -> Option<Option<crate::deferred::Readback>> {
+        self.uses(img);
         use crate::deferred::Lookup;
         let d = self.g.deferred.clone()?;
         match d.lookup(key) {
@@ -1536,6 +2043,7 @@ impl<'g> Enc<'g> {
     /// Copy an image into a storage buffer (RGBA f32 rows) for kernels that read a third image
     /// through `data`: (buffer, row length in pixels).
     pub fn image_rows(&mut self, img: &GpuImage) -> (wgpu::Buffer, u32) {
+        self.uses(img);
         let row = (img.width * 16).div_ceil(wgpu::COPY_BYTES_PER_ROW_ALIGNMENT) * wgpu::COPY_BYTES_PER_ROW_ALIGNMENT;
         let buf = self.g.device.create_buffer(&wgpu::BufferDescriptor {
             label: Some("image rows"),
@@ -1565,6 +2073,7 @@ impl<'g> Enc<'g> {
     /// Convert to the viewer's premultiplied RGBA8 texture (egui-wgpu's native texture format),
     /// with a full mip chain so the viewer minifies it without aliasing (#417).
     pub fn display(&mut self, img: &GpuImage) -> wgpu::Texture {
+        self.uses(img);
         let g = self.g;
         let levels = mip_levels(img.width, img.height);
         let tex = g.device.create_texture(&wgpu::TextureDescriptor {
@@ -1607,6 +2116,30 @@ pub(crate) fn mip_levels(w: u32, h: u32) -> u32 {
     u32::BITS - w.max(h).max(1).leading_zeros()
 }
 
+#[cfg(test)]
+mod pool_stall_tests {
+    use super::Pool;
+    use std::sync::atomic::Ordering::Relaxed;
+
+    #[test]
+    fn frame_end_skips_wait_while_stalled_until_dying_drops() {
+        let mut p = Pool::default();
+        p.dying.store(100, Relaxed);
+        assert_eq!(p.frame_end(), 100);
+        p.waited(100);
+        assert_eq!(p.stalled, 100);
+        assert_eq!(p.waits, 1);
+        assert_eq!(p.frame_end(), 0);
+        p.dying.store(150, Relaxed);
+        assert_eq!(p.frame_end(), 0);
+        p.dying.store(50, Relaxed);
+        assert_eq!(p.frame_end(), 50);
+        assert_eq!(p.stalled, 0);
+        p.dying.store(0, Relaxed);
+        assert_eq!(p.frame_end(), 0);
+    }
+}
+
 #[cfg(test)]
 mod adapter_name_tests {
     use super::{match_adapter_name, wanted_adapter_name};

crates/gpu/src/effects.rs

@@ -101,6 +101,7 @@ pub(crate) fn run_chain(e: &mut Enc, chain: &[FxStep], buf: &Buf, levels: Option
     e.take_missing();
     for step in chain {
         b = apply(e, step.spec.id, &step.ctx, b)?;
+        e.mem_gate();
         if e.take_missing() {
             return None;
         }
@@ -158,6 +159,7 @@ impl effectcraft_render::FxTarget for GpuFx<'_, '_> {
         self.e.take_missing();
         for step in steps {
             let Some(next) = apply(self.e, step.spec.id, &step.ctx, cur) else { return false };
+            self.e.mem_gate();
             // A kernel the device couldn't build: the effects render on the CPU (#613).
             if self.e.take_missing() {
                 return false;

crates/gpu/src/fx_pixel2.rs

@@ -12,9 +12,11 @@
 //! Layer is chosen; the colouring runs here.
 
 use effectcraft_color::BlendMode;
-use effectcraft_effects::EffectCtx;
+use effectcraft_effects::{Buf, EffectCtx};
 use effectcraft_raster::Image;
 
+use rayon::prelude::*;
+
 use crate::context::{Enc, GpuImage, Params};
 use crate::effects::{GBuf, box_passes, gaussian_blur};
 use crate::ops::mode_id;
@@ -380,14 +382,24 @@ fn clamp_plane(e: &mut Enc, pl: &GpuImage, lo: f32, hi: f32) -> GpuImage {
 // ---------------------------------------------------------------- CC HexTile / CC Mr. Smoothie (stylize3)
 
 fn hextile(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
-    let r = (ctx.params.f("radius") * b.scale).max(0.5);
-    let c = b.to_px(ctx.params.v2("center"));
-    let (sn, cs) = ctx.params.f("rotate").to_radians().sin_cos();
-    let mut p = Params::default();
-    p.u[0][0] = ctx.params.e("render");
-    p.f[0] = [c.0 as f32, c.1 as f32, r as f32, (ctx.params.f("smearing") / 100.0).clamp(0.0, 1.0) as f32];
-    p.f[1] = [sn as f32, cs as f32, 0.0, 0.0];
-    run(e, "fp2_hextile", &p, b, None, None)
+    let count = (b.img.width as usize).checked_mul(b.img.height as usize)?;
+    if count.checked_mul(8)? > e.g.device.limits().max_storage_buffer_binding_size as usize {
+        return None;
+    }
+    let geometry = effectcraft_effects::HexTile::new(
+        ctx,
+        &Buf { img: Image { width: b.img.width, height: b.img.height, data: vec![] }, offset: b.offset, scale: b.scale },
+    );
+    let stride = (b.img.width as usize).checked_mul(2)?.max(2);
+    let mut coords = vec![0.0; count.checked_mul(2)?];
+    coords.par_chunks_mut(stride).enumerate().for_each(|(y, row)| {
+        for (x, xy) in row.as_chunks_mut::<2>().0.iter_mut().enumerate() {
+            let (sx, sy) = geometry.source_at(x as u32, y as u32);
+            *xy = [sx as f32, sy as f32];
+        }
+    });
+    let data = e.data(&coords);
+    run(e, "fp2_hextile", &Params::default(), b, None, Some(&data))
 }
 
 fn smoothie(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {

crates/gpu/src/fx_tone.rs

@@ -174,6 +174,62 @@ fn parse_map(s: &str) -> Vec<f32> {
     vals.iter().map(|v| (v * k).clamp(0.0, 1.0)).collect()
 }
 
+/// Piecewise-linear lookup in the CPU effect's f64 phase domain. Store the
+/// first representable f32 at each boundary, its value and the interval slope.
+/// Computing floor/wrap in shader f32 can select the end of the map instead of
+/// its beginning at a wrap (a full-scale colour error).
+fn arbitrary_map_table(map: &[f32], phase: f64) -> Option<Vec<f32>> {
+    let n = map.len();
+    if !(2..=4096).contains(&n) || !phase.is_finite() {
+        return None;
+    }
+    let scale = (n - 1) as f64;
+    let offset = phase * scale / 255.0;
+    if !offset.is_finite() || offset.abs() >= (1u64 << 52) as f64 {
+        return None;
+    }
+    let at = |v: f32| {
+        let x = (v as f64 * scale + offset).rem_euclid(n as f64);
+        let i = x.floor() as usize;
+        let a = *map.get(i).unwrap_or(&0.0);
+        let b = *map.get((i + 1).min(n - 1)).unwrap_or(&a);
+        (a + (b - a) * (x - x.floor()) as f32, (b - a) * scale as f32)
+    };
+    let mut points = vec![0.0f32];
+    // Search ordered positive f32 encodings using the CPU's actual f64 sum.
+    // Merely rounding a mathematical boundary upward misses the half-ulp
+    // region where f64 addition itself rounds across a wrap.
+    let first = offset.floor() + 1.0;
+    for i in 0..=n {
+        let boundary = first + i as f64;
+        if boundary > scale + offset {
+            break;
+        }
+        let (mut lo, mut hi) = (0u32, 1.0f32.to_bits());
+        while lo < hi {
+            let mid = lo + (hi - lo) / 2;
+            if f32::from_bits(mid) as f64 * scale + offset < boundary {
+                lo = mid + 1;
+            } else {
+                hi = mid;
+            }
+        }
+        let start = f32::from_bits(lo);
+        if points.last().is_none_or(|last| start > *last) {
+            points.push(start);
+        }
+    }
+    if points.last() != Some(&1.0) {
+        points.push(1.0);
+    }
+    let mut table = Vec::with_capacity(points.len() * 3);
+    for v in points {
+        let (value, slope) = at(v);
+        table.extend([v, value, slope]);
+    }
+    Some(table)
+}
+
 fn point(e: &mut Enc, id: &str, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
     let pr = ctx.params;
     let f = |k: &str| pr.f(k) as f32;
@@ -308,9 +364,9 @@ fn point(e: &mut Enc, id: &str, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
         }
         "ec.color.psarbitrarymap" => {
             let map = parse_map(pr.s("map"));
-            p.u[0] = [12, pr.b("applyPhaseMapToAlpha") as u32, map.len() as u32, 0];
-            p.f[0][0] = (pr.f("phase") * (map.len() - 1) as f64 / 255.0) as f32;
-            data = Some(map);
+            let table = arbitrary_map_table(&map, pr.f("phase"))?;
+            p.u[0] = [12, pr.b("applyPhaseMapToAlpha") as u32, (table.len() / 3) as u32, 0];
+            data = Some(table);
         }
         "ec.color.videolimiter" => {
             const CLIP_LEVELS: [f32; 5] = [0.9, 0.95, 1.0, 1.05, 1.09];

crates/gpu/src/fx_transition.rs

@@ -15,6 +15,8 @@
 use effectcraft_effects::{Buf, EffectCtx};
 use effectcraft_raster::Image;
 
+use rayon::prelude::*;
+
 use crate::context::{Enc, GpuImage, Params};
 use crate::effects::{GBuf, gaussian_blur};
 
@@ -203,25 +205,22 @@ fn grid_wipe(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
     if done <= 0.0 {
         return Some(b);
     }
-    let pr = ctx.params;
-    let c = b.to_px(pr.v2("center"));
-    let rot = pr.f("rotation").to_radians();
-    let border = pr.f("border").max(0.0) * b.scale;
-    let tiles = pr.f("tiles").round().clamp(1.0, 500.0);
-    let shape = pr.e("shape");
-    let cs = (ctx.layer_size[0] * b.scale / tiles).max(1.0);
-    let (s, co) = rot.sin_cos();
-    let metric = |dx: f64, dy: f64| match shape {
-        0 => dx.abs(),
-        2 => dx.abs().max(dy.abs()),
-        _ => dx.hypot(dy),
-    };
-    let maxd = max_corner_dist(&b, c, |dx, dy| metric(dx * co + dy * s, -dx * s + dy * co));
-    let mut p = Params::default();
-    p.u[0] = [shape, pr.b("reverse") as u32, 0, 0];
-    p.f[0] = f4([c.0, c.1, s, co]);
-    p.f[1] = f4([cs, maxd, done, border]);
-    run(e, "ftr_grid", &p, b, None, None)
+    // Decide discontinuous tile coverage in the CPU effect's f64 geometry.
+    // Only a scalar mask is uploaded; source pixels and multiplication stay resident.
+    let count = (b.img.width as usize).checked_mul(b.img.height as usize)?;
+    if count.checked_mul(4)? > e.g.device.limits().max_storage_buffer_binding_size as usize {
+        return None;
+    }
+    let geometry = shape(&b);
+    let grid = effectcraft_effects::GridWipe::new(ctx, &geometry);
+    let mut mask = vec![0.0; count];
+    mask.par_chunks_mut((b.img.width as usize).max(1)).enumerate().for_each(|(y, row)| {
+        for (x, value) in row.iter_mut().enumerate() {
+            *value = grid.coverage(x as u32, y as u32);
+        }
+    });
+    let data = e.data(&mask);
+    run(e, "ftr_grid", &Params::default(), b, None, Some(&data))
 }
 
 fn radial_scale_wipe(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {

crates/gpu/src/inpaint.rs

@@ -0,0 +1,729 @@
+//! Content-Aware Fill's per-frame synthesis and flow on the GPU: PatchMatch's multi-scale EM
+//! loop, propagation, random search and vote of `effectcraft_raster::inpaint`, and the block
+//! matching of `effectcraft_raster::flow::block_flow`, as compute kernels following the CPU's
+//! steps (`shaders/inpaint.wgsl`).
+//!
+//! The fill's input frames are composited on the CPU by design (analyses do not use the job
+//! worker's GPU: `effectcraft_engine::background::offload::job_accel`), so this module is about the fill
+//! itself — the part that runs for minutes on a large mask, and whose parallel search is what a
+//! device wins on. The membrane (harmonic) solve, the flow's hole completion, the orchestration
+//! ([`effectcraft_raster::inpaint::fill_sequence_with`]), the flow-guided propagation and the
+//! per-frame flow control stay on the CPU: the membrane is bandwidth-bound (see
+//! [`FillBackend::patch`]), and the propagation's coupling is sequential, needing every frame of
+//! the range resident at once.
+//!
+//! Parity with the CPU reference is a quality contract, not bit-exactness: the CPU's
+//! scan-order propagation becomes a two-pass checkerboard, and its in-place relaxation becomes
+//! Jacobi reads. Both solve the same problem to the same quality (`tests/inpaint.rs`).
+
+use effectcraft_raster::Image;
+use effectcraft_raster::flow::{Flow, block_flow};
+use effectcraft_raster::inpaint::{FillBackend, complete_flow, dilate};
+
+use crate::Gpu;
+use crate::context::{Enc, GpuContext, GpuImage, KernelImages, Params};
+
+/// The kernels this module adds to the shared pipeline registry (`context.rs`).
+pub(crate) const KERNELS: &[&str] = &[
+    "imp_down",
+    "imp_up",
+    "imp_jacobi",
+    "imp_pm_down",
+    "imp_pm_up",
+    "imp_nnf_up",
+    "imp_pm_init",
+    "imp_pm_prop",
+    "imp_pm_rand",
+    "imp_nnf_store",
+    "imp_pm_vote",
+    "imp_hist",
+    "imp_fl_half",
+    "imp_fl_search",
+    "imp_fl_subpix",
+    "imp_fl_median",
+];
+
+/// Distance histogram bins (`imp_hist`); bin 0's `.g` holds the largest finite distance the
+/// dispatch saw, for the HDR re-bin.
+const BINS: usize = 1024;
+
+/// The distance range `imp_hist`'s bins cover for SDR input. An image whose channels are all ≤ 1
+/// has `Σ_c (a − b)² ≤ 4`, so no patch distance exceeds it and the bins need no rescaling.
+const SDR_DISTANCE: f32 = 4.0;
+
+/// `u`/`f` blocks of a kernel's parameter buffer (see `inpaint.wgsl`'s header).
+fn params(u0: [u32; 4], u1: [u32; 4], u2: [u32; 4], f0: [f32; 4]) -> Params {
+    Params { u: [u0, u1, u2, [0; 4]], f: [f0, [0.0; 4], [0.0; 4], [0.0; 4], [0.0; 4], [0.0; 4], [0.0; 4], [0.0; 4], [0.0; 4], [0.0; 4], [0.0; 4], [0.0; 4]] }
+}
+
+/// `p` with this level's mask offset in `u[3].x` (`inpaint.wgsl`'s `data`).
+fn with_off(p: Params, off: u32) -> Params {
+    Params { u: [p.u[0], p.u[1], p.u[2], [off, 0, 0, 0]], f: p.f }
+}
+
+/// A relaxation sweep count of `membrane_core`. The CPU uses 200 for the small levels and 40 for
+/// the large ones; this uses 64 and 40. The small levels are where the difference lies: Jacobi
+/// needs O(n²) sweeps to converge on an n-wide grid, so at the levels the coarse-to-fine schedule
+/// actually reaches (n ≤ 16, where n² ≤ 256) 64 sweeps is already converged, and finer levels
+/// start from a good guess. Matching the CPU's 200 measured ≈ 6% slower on the whole fill with no
+/// change in the result to two decimals (67.0 dB either way), so the schedule is deliberately
+/// not the CPU's — Jacobi and Gauss–Seidel do not converge alike, and the parity contract is
+/// quality, not bit-exactness.
+fn sweep_count(w: usize, h: usize) -> usize {
+    if w * h <= 4096 { 64 } else { 40 }
+}
+
+/// PatchMatch's patch radius at a level (`pm_level`).
+fn radius(patch: usize, w: u32, h: u32) -> u32 {
+    (patch / 2).min(w.min(h).saturating_sub(1) as usize / 2).max(1) as u32
+}
+
+/// The bounding box of `mask`'s set pixels, `[x0, y0, x1, y1)`, or `None` when it is empty.
+fn bbox(mask: &[bool], w: usize, h: usize) -> Option<[u32; 4]> {
+    let (mut x0, mut y0, mut x1, mut y1) = (usize::MAX, usize::MAX, 0usize, 0usize);
+    for y in 0..h {
+        for x in 0..w {
+            if mask.get(y * w + x).copied().unwrap_or(false) {
+                x0 = x0.min(x);
+                y0 = y0.min(y);
+                x1 = x1.max(x + 1);
+                y1 = y1.max(y + 1);
+            }
+        }
+    }
+    if x0 == usize::MAX { None } else { Some([x0 as u32, y0 as u32, x1 as u32, y1 as u32]) }
+}
+
+/// [`bbox`] grown by `n` pixels, clipped to the image.
+fn grow(b: [u32; 4], n: u32, w: u32, h: u32) -> [u32; 4] {
+    [b[0].saturating_sub(n), b[1].saturating_sub(n), b[2].saturating_add(n).min(w), b[3].saturating_add(n).min(h)]
+}
+
+/// The bounding box of the pixels a membrane level still has to solve (`known` false), or
+/// `None` when every pixel is known.
+fn unknown_bbox(known: &[f32], w: usize, h: usize) -> Option<[u32; 4]> {
+    let (mut x0, mut y0, mut x1, mut y1) = (usize::MAX, usize::MAX, 0usize, 0usize);
+    for y in 0..h {
+        for x in 0..w {
+            if known.get(y * w + x).copied().unwrap_or(1.0) <= 0.5 {
+                x0 = x0.min(x);
+                y0 = y0.min(y);
+                x1 = x1.max(x + 1);
+                y1 = y1.max(y + 1);
+            }
+        }
+    }
+    if x0 == usize::MAX { None } else { Some([x0 as u32, y0 as u32, x1 as u32, y1 as u32]) }
+}
+
+/// One dispatch's images, its per-pixel `data` table and the region it covers. Grouped into a
+/// struct because a kernel's bindings plus its bounds exceed clippy's argument limit, and the
+/// workspace denies `too_many_arguments`.
+struct Pass<'a> {
+    src: &'a GpuImage,
+    aux: Option<&'a GpuImage>,
+    out: &'a GpuImage,
+    data: Option<&'a wgpu::Buffer>,
+    bb: [u32; 4],
+}
+
+/// Dispatch a kernel over `pass.bb`. The box's origin is written into `P.u[1].xy` (the kernels
+/// place their other `u[1]` fields in `.zw`): deriving it from the box the dispatch is sized from
+/// keeps a dispatch's covered region and the origin its shader reads from drifting apart, which
+/// is the subtle way an edge of the region goes unwritten.
+fn over(e: &mut Enc, entry: &str, p: &Params, pass: Pass<'_>) {
+    let bb = pass.bb;
+    let p = Params { u: [p.u[0], [bb[0], bb[1], p.u[1][2], p.u[1][3]], p.u[2], p.u[3]], f: p.f };
+    let groups = ((bb[2] - bb[0]).div_ceil(16), (bb[3] - bb[1]).div_ceil(16));
+    e.dispatch(entry, &p, KernelImages { src: pass.src, aux: pass.aux, out: pass.out }, pass.data, groups);
+}
+
+// ------------------------------------------------------------------ membrane
+
+/// Relax `fine`'s unknown pixels — the harmonic interpolation of their border — with the CPU's
+/// coarse-to-fine schedule (`membrane_core`). `spare` is a free texture of the same size; the
+/// result is returned (it may be `fine`, `spare` or one of the coarse levels).
+fn membrane_tex(e: &mut Enc, fine: &GpuImage, spare: &GpuImage, hole: &[bool], w: u32, h: u32) -> Option<GpuImage> {
+    let (uw, uh) = (w as usize, h as usize);
+    let n = uw * uh;
+    if hole.len() < n || !hole.iter().take(n).any(|b| *b) || !hole.iter().take(n).any(|b| !*b) {
+        return Some(fine.clone());
+    }
+    // The recursion's levels: halve while the level is wider than 8.
+    let mut dims = vec![(uw, uh)];
+    while let Some(&(cw, ch)) = dims.last() {
+        if !(cw > 8 && ch > 8) || dims.len() >= 32 {
+            break;
+        }
+        dims.push((cw.div_ceil(2), ch.div_ceil(2)));
+    }
+    // The known mask of every level: a coarse pixel is known where any of its 2×2 is.
+    let mut knowns: Vec<Vec<f32>> = Vec::with_capacity(dims.len());
+    knowns.push(hole.iter().take(n).map(|b| if *b { 0.0 } else { 1.0 }).collect());
+    for k in 1..dims.len() {
+        let (pw, ph) = dims[k - 1];
+        let (cw, ch) = dims[k];
+        let prev = &knowns[k - 1];
+        let mut cur = vec![0.0f32; cw * ch];
+        for y in 0..ch {
+            for x in 0..cw {
+                for (dx, dy) in [(0usize, 0usize), (1, 0), (0, 1), (1, 1)] {
+                    let (fx, fy) = (2 * x + dx, 2 * y + dy);
+                    if fx < pw && fy < ph && prev[fy * pw + fx] > 0.5 {
+                        cur[y * cw + x] = 1.0;
+                        break;
+                    }
+                }
+            }
+        }
+        knowns.push(cur);
+    }
+    // One buffer for every level's mask (a buffer per level costs a device allocation each).
+    let mut packed: Vec<f32> = Vec::with_capacity(knowns.iter().map(|k| k.len()).sum());
+    let mut offs: Vec<u32> = Vec::with_capacity(knowns.len());
+    for k in &knowns {
+        offs.push(packed.len() as u32);
+        packed.extend_from_slice(k);
+    }
+    let datas = e.data(&packed);
+    let moff = |k: usize| offs.get(k).copied().unwrap_or(0);
+    // The bounds of each level's unknown pixels: the relaxation only touches them (known pixels
+    // never move), which keeps a fill's cost proportional to the hole, not the frame.
+    let mut uns: Vec<Option<[u32; 4]>> = Vec::with_capacity(dims.len());
+    for (k, &(cw, ch)) in dims.iter().enumerate() {
+        uns.push(unknown_bbox(knowns.get(k)?, cw, ch));
+    }
+    // Two textures per coarse level, for the ping-pong sweeps.
+    let mut lv: Vec<GpuImage> = Vec::with_capacity(2 * dims.len());
+    for k in 1..dims.len() {
+        let (cw, ch) = dims[k];
+        lv.push(e.scratch(cw as u32, ch as u32));
+        lv.push(e.scratch(cw as u32, ch as u32));
+    }
+    // Fine to coarse.
+    for k in 1..dims.len() {
+        let (pw, ph) = dims[k - 1];
+        let (cw, ch) = dims[k];
+        let p = with_off(params([pw as u32, ph as u32, cw as u32, ch as u32], [0; 4], [0; 4], [0.0; 4]), moff(k - 1));
+        let src = if k == 1 { fine } else { lv.get(2 * (k - 2))? };
+        e.pixels("imp_down", &p, src, None, lv.get(2 * (k - 1))?, Some(&datas));
+    }
+    // Relax, coarse to fine. `first[k]` says whether level k's image is the first of its pair
+    // (level 0's pair is `fine`/`spare`).
+    let top = dims.len() - 1;
+    let mut first = vec![true; dims.len()];
+    if top >= 1 {
+        let (a, b) = (lv.get(2 * (top - 1))?, lv.get(2 * (top - 1) + 1)?);
+        // Both textures must hold the level before the sweeps: they now write only the hole.
+        e.copy_into(a, b, 0, 0);
+        first[top] = sweeps(e, &datas, moff(top), a, b, Region { w: dims[top].0 as u32, h: dims[top].1 as u32, bb: uns.get(top).copied().flatten() });
+    }
+    for k in (1..top).rev() {
+        let (cw, ch) = dims[k];
+        let (pw, ph) = dims[k + 1];
+        let coarse = lv.get(2 * k + usize::from(!first[k + 1]))?;
+        let (fa, fb) = (lv.get(2 * (k - 1))?, lv.get(2 * (k - 1) + 1)?);
+        e.copy_into(fa, fb, 0, 0);
+        if let Some(bb) = uns.get(k).copied().flatten() {
+            let p = with_off(params([cw as u32, ch as u32, pw as u32, ph as u32], [0; 4], [0; 4], [0.0; 4]), moff(k));
+            over(e, "imp_up", &p, Pass { src: coarse, aux: None, out: fb, data: Some(&datas), bb });
+        }
+        // `imp_up` wrote into `fb`, the pair's second.
+        first[k] = !sweeps(e, &datas, moff(k), fb, fa, Region { w: cw as u32, h: ch as u32, bb: uns.get(k).copied().flatten() });
+    }
+    if top >= 1 {
+        let coarse = lv.get(usize::from(!first[1]))?;
+        e.copy_into(fine, spare, 0, 0);
+        let bb = uns.first().copied().flatten()?;
+        let p = with_off(params([w, h, dims[1].0 as u32, dims[1].1 as u32], [0; 4], [0; 4], [0.0; 4]), moff(0));
+        over(e, "imp_up", &p, Pass { src: coarse, aux: None, out: spare, data: Some(&datas), bb });
+        return Some(if sweeps(e, &datas, moff(0), spare, fine, Region { w, h, bb: Some(bb) }) { spare.clone() } else { fine.clone() });
+    }
+    let bb = uns.first().copied().flatten()?;
+    e.copy_into(fine, spare, 0, 0);
+    Some(if sweeps(e, &datas, moff(0), fine, spare, Region { w, h, bb: Some(bb) }) { fine.clone() } else { spare.clone() })
+}
+
+/// A texture's dimensions and the region a dispatch covers (`None`: nothing to do).
+struct Region {
+    w: u32,
+    h: u32,
+    bb: Option<[u32; 4]>,
+}
+
+/// Jacobi relaxation sweeps of a texture pair, over the unknown pixels' bounds `region.bb` (the
+/// state starts in `a`); returns whether the result is in `a`.
+fn sweeps(e: &mut Enc, data: &wgpu::Buffer, off: u32, a: &GpuImage, b: &GpuImage, region: Region) -> bool {
+    let Some(bb) = region.bb else { return true };
+    let p = with_off(params([region.w, region.h, 0, 0], [0; 4], [0; 4], [0.0; 4]), off);
+    let mut in_a = true;
+    for _ in 0..sweep_count(region.w as usize, region.h as usize) {
+        let (src, out) = if in_a { (a, b) } else { (b, a) };
+        over(e, "imp_jacobi", &p, Pass { src, aux: None, out, data: Some(data), bb });
+        in_a = !in_a;
+    }
+    in_a
+}
+
+/// `img` with the pixels `mask` marks inside `bb` replaced by `rect` — the tightly packed RGBA f32
+/// rows of the rectangle that was read back. Pixels outside the mask stay the input's, bit for bit.
+fn patch_back(img: &Image, mask: &[bool], bb: [u32; 4], rect: &[u8]) -> Image {
+    let (w, rw) = (img.width as usize, (bb[2] - bb[0]) as usize);
+    let mut out = img.clone();
+    for y in bb[1]..bb[3] {
+        for x in bb[0]..bb[2] {
+            let i = y as usize * w + x as usize;
+            if !mask.get(i).copied().unwrap_or(false) {
+                continue;
+            }
+            let off = ((y - bb[1]) as usize * rw + (x - bb[0]) as usize) * 16;
+            let Some(b) = rect.get(off..off + 16) else { continue };
+            let mut px = [0.0f32; 4];
+            for (c, v) in px.iter_mut().enumerate() {
+                *v = f32::from_le_bytes([b[c * 4], b[c * 4 + 1], b[c * 4 + 2], b[c * 4 + 3]]);
+            }
+            out.data[i] = px;
+        }
+    }
+    out
+}
+
+impl FillBackend for Gpu {
+    /// PatchMatch synthesis — the fill's dominant cost, and the stage a device wins on: two to
+    /// three times faster than the CPU's 18-rayon-thread loop at 1920×1080 with a 300×200 hole
+    /// on an Intel Arc iGPU, and far more on a discrete one.
+    ///
+    /// The membrane ([`FillBackend::membrane`]) deliberately stays on the CPU: a harmonic solve
+    /// over a hole is only a few tens of milliseconds, and on the same measurement it is *slower*
+    /// on the device (≈ 90 ms against ≈ 55 ms). Most of that is the fill's full-frame data
+    /// movement — uploading the frame, copying it and building the known pyramid — which does not
+    /// shrink with the hole, while the relaxation itself does. The CPU's in-place Gauss–Seidel
+    /// over many cores is already competitive there, so the device only takes the stage where its
+    /// parallel search actually pays. (`membrane_tex` still runs on the device *inside* this
+    /// stage, where it is the coarsest level's start and never crosses back to the CPU.)
+    fn patch(&self, img: &Image, hole: &[bool], init: Option<&Image>, patch: usize, seed: u32) -> Image {
+        // The error scopes make an allocation or validation failure local to this frame: the
+        // CPU fills it instead, rather than an uncaptured error retiring the shared context.
+        if self.ctx.can_wait()
+            && let Ok(Some(out)) = self.within_memory(|| patch_image(&self.ctx, img, hole, init, patch, seed))
+        {
+            return out;
+        }
+        effectcraft_raster::inpaint::patch_fill(img, hole, init, patch, seed)
+    }
+
+    /// The block matching that the sequence's flow-guided propagation starts from, which is the
+    /// flow's dominant cost. The field is a small `cols × rows` grid (8 px tiles, so 240 × 135 at
+    /// 1080p), and the search is texture-load-bound: at 1920×1080 the CPU's `block_flow` is
+    /// ≈ 65–70 ms against ≈ 55 ms here, most of it the coarsest level's window (`radius 6` over
+    /// up to six seeds) and the finest level's 8×8 blocks. That is a smaller margin than
+    /// [`patch`](Self::patch), but the flow runs twice per frame pair, so on a four-frame 1080p
+    /// sequence it is worth ≈ 0.15× of the fill's total (1.42× against 1.27× with the CPU flow).
+    ///
+    /// The field is exhaustive, not heuristic: it matches the CPU's to a float rounding, so the
+    /// propagation sees the same motion. Both A's block per candidate (hoisted into registers)
+    /// and the CPU's `Plane::luma` (computed here by [`effectcraft_raster::flow::luma_pair`] and
+    /// uploaded, because a device may contract `a·b + c` into a multiply-add and round the luma
+    /// differently) are why it is that close.
+    ///
+    /// The completion — marking the tiles that touch either hole and interpolating the surrounding
+    /// motion over them — stays on the CPU ([`complete_flow`]). It is a membrane solve over that
+    /// same tiny grid whose cost is the mask, not the solve, so there is little for the device to
+    /// take; keeping it here also keeps the completed field the CPU's to within the search's
+    /// rounding.
+    fn flow(&self, a: &Image, b: &Image, hole_a: &[bool], hole_b: &[bool]) -> Flow {
+        // Frame-local error scopes, as for the patch: a failed GPU search falls back to the CPU.
+        if self.ctx.can_wait()
+            && let Ok(Some(f)) = self.within_memory(|| flow_image(&self.ctx, a, b))
+        {
+            return complete_flow(f, hole_a, hole_b, a.width as usize, a.height as usize);
+        }
+        complete_flow(block_flow(a, b, 8, 6), hole_a, hole_b, a.width as usize, a.height as usize)
+    }
+}
+
+// ------------------------------------------------------------------ block flow
+
+/// `effectcraft_raster::flow::block_flow(a, b, 8, 6)` on the device: the two frames' luma
+/// pyramids, the coarse-to-fine search, the sub-pixel fit and the 3×3 median, returning the raw
+/// field ([`complete_flow`] does the hole completion). `None` where the device cannot run it and
+/// the caller falls back to the CPU.
+pub(crate) fn flow_image(g: &GpuContext, a: &Image, b: &Image) -> Option<Flow> {
+    let (w, h) = (a.width, a.height);
+    if w == 0 || h == 0 || w != b.width || h != b.height || !g.fits(w, h) {
+        return None;
+    }
+    let (tile, win, radius) = (8u32, 8u32, 6u32);
+    let cols = (w as usize).div_ceil(tile as usize).max(1);
+    let rows = (h as usize).div_ceil(tile as usize).max(1);
+    // block_flow_window's early out: a frame too small to hold one block and window.
+    if w < tile || h < tile || w < win || h < win {
+        return Some(Flow { width: w, height: h, block: tile, cols, rows, v: vec![[0.0; 2]; cols * rows] });
+    }
+    let mut e = Enc::new(g);
+    // The luma pyramid: `Plane::luma` (packed .x/.y), halved while both sides stay worth matching
+    // (`Plane::half`), at most four levels.
+    let mut dims = vec![(w, h)];
+    while dims.len() < 4 {
+        let (cw, ch) = *dims.last()?;
+        if cw / 2 < 16 || ch / 2 < 16 {
+            break;
+        }
+        dims.push(((cw / 2).max(1), (ch / 2).max(1)));
+    }
+    // The luma is the CPU's (`luma_pair`), so the field matches `block_flow`'s bit for bit; the
+    // halves below are exact additions and a multiply by a power of two.
+    let pair = g.upload_image(&effectcraft_raster::flow::luma_pair(a, b))?;
+    let mut luma: Vec<GpuImage> = Vec::with_capacity(dims.len());
+    luma.push(pair);
+    for k in 1..dims.len() {
+        let (pw, ph) = dims[k - 1];
+        let (cw, ch) = dims[k];
+        let dst = e.scratch(cw, ch);
+        e.pixels("imp_fl_half", &params([pw, ph, cw, ch], [0; 4], [0; 4], [0.0; 4]), luma.get(k - 1)?, None, &dst, None);
+        luma.push(dst);
+    }
+    // The coarse-to-fine search: the coarsest level from the zero vector, each finer level from
+    // the doubled vectors of its tile and its four neighbours.
+    let top = dims.len() - 1;
+    let groups = (cols.div_ceil(16) as u32, rows.div_ceil(16) as u32);
+    let mut field: Vec<GpuImage> = Vec::with_capacity(dims.len());
+    for _ in 0..dims.len() {
+        field.push(e.scratch(cols as u32, rows as u32));
+    }
+    for k in (0..dims.len()).rev() {
+        let (pw, ph) = dims[k];
+        let coarser = (k != top).then(|| field.get(k + 1)).flatten();
+        let p = params([pw, ph, cols as u32, rows as u32], [tile, win, 1u32 << k, if k == top { radius } else { 1 }], [u32::from(k != top), 0, 0, 0], [0.0; 4]);
+        e.dispatch("imp_fl_search", &p, KernelImages { src: luma.get(k)?, aux: coarser, out: field.get(k)? }, None, groups);
+    }
+    // The full-resolution sub-pixel fit, then the vector median.
+    let fine = e.scratch(cols as u32, rows as u32);
+    e.dispatch(
+        "imp_fl_subpix",
+        &params([w, h, cols as u32, rows as u32], [tile, win, 0, 0], [0; 4], [0.0; 4]),
+        KernelImages { src: luma.first()?, aux: Some(field.first()?), out: &fine },
+        None,
+        groups,
+    );
+    let out = e.scratch(cols as u32, rows as u32);
+    e.dispatch(
+        "imp_fl_median",
+        &params([w, h, cols as u32, rows as u32], [0; 4], [0; 4], [0.0; 4]),
+        KernelImages { src: &fine, aux: None, out: &out },
+        None,
+        groups,
+    );
+    if e.take_missing() {
+        return None;
+    }
+    // Only the field crosses back to the CPU: `cols × rows` entries.
+    let rect = e.read_rect(&out.texture, 0, 0, cols as u32, rows as u32, 16)?;
+    let mut v = Vec::with_capacity(cols * rows);
+    for i in 0..cols * rows {
+        let b = rect.get(i * 16..i * 16 + 16)?;
+        let f = |o: usize| b.get(o).copied().unwrap_or(0);
+        v.push([f32::from_le_bytes([f(0), f(1), f(2), f(3)]), f32::from_le_bytes([f(4), f(5), f(6), f(7)])]);
+    }
+    Some(Flow { width: w, height: h, block: tile, cols, rows, v })
+}
+
+// ------------------------------------------------------------------ PatchMatch
+
+/// One PatchMatch level's plan.
+struct Lv {
+    w: u32,
+    h: u32,
+    r: u32,
+    /// The target centres: pixels whose patch overlaps the hole (`near`).
+    near: [u32; 4],
+    /// The hole's own bounding box (the vote's dispatch).
+    holes: [u32; 4],
+    /// Where the nearest-neighbour field must be defined: the vote reads it `r` pixels out.
+    field: [u32; 4],
+    flags: wgpu::Buffer,
+}
+
+/// The multi-scale PatchMatch synthesis of `patch_fill`: `init` is the starting guess (the
+/// previous frame's fill warped onto this one), which refines the finest level only.
+pub(crate) fn patch_image(g: &GpuContext, img: &Image, hole: &[bool], init: Option<&Image>, patch: usize, seed: u32) -> Option<Image> {
+    let (w, h) = (img.width, img.height);
+    if w == 0 || h == 0 || !g.fits(w, h) {
+        return None;
+    }
+    let (uw, uh) = (w as usize, h as usize);
+    let n = uw * uh;
+    if hole.len() < n || !hole.iter().take(n).any(|b| *b) {
+        return Some(img.clone());
+    }
+    let hole: Vec<bool> = hole.iter().take(n).copied().collect();
+    let patch = patch.max(3) | 1;
+    let mut e = Enc::new(g);
+
+    // A refinement of one frame (`patch_fill`'s `init`): the previous frame's fill seeds the
+    // hole and only the finest level runs, from 2 EM rounds of 4 iterations.
+    if let Some(init) = init {
+        if init.width != w || init.height != h {
+            return None;
+        }
+        let mut start = img.clone();
+        for (i, v) in start.data.iter_mut().enumerate() {
+            if hole.get(i).copied().unwrap_or(false) {
+                *v = *init.data.get(i)?;
+            }
+        }
+        let lv = plan_level(&mut e, &hole, w, h, patch)?;
+        let image = g.upload_image(&start)?;
+        let spare = e.scratch(w, h);
+        let nnf = e.image(w, h);
+        let (out, _) = pm_level(&mut e, &lv, PmRun { seed, em: 2, iters: 4 }, &image, &spare, &nnf)?;
+        if e.take_missing() {
+            return None;
+        }
+        let bb = bbox(&hole, uw, uh)?;
+        let rect = e.read_rect(&out.texture, bb[0], bb[1], bb[2] - bb[0], bb[3] - bb[1], 16)?;
+        return Some(patch_back(img, &hole, bb, &rect));
+    }
+
+    // The pyramid: halve while the level is still a few patches across.
+    let mut dims = vec![(w, h)];
+    let mut holes = vec![hole.clone()];
+    while dims.len() < 6 {
+        let &(lw, lh) = dims.last()?;
+        let hl = holes.last()?;
+        if lw.min(lh) / 2 < patch as u32 * 4 || !hl.iter().any(|b| *b) {
+            break;
+        }
+        let (cw, ch) = (lw.div_ceil(2) as usize, lh.div_ceil(2) as usize);
+        let mut nh = vec![false; cw * ch];
+        for y in 0..ch {
+            for x in 0..cw {
+                for (dx, dy) in [(0usize, 0usize), (1, 0), (0, 1), (1, 1)] {
+                    let (fx, fy) = ((2 * x + dx).min(lw as usize - 1), (2 * y + dy).min(lh as usize - 1));
+                    if hl[fy * lw as usize + fx] {
+                        nh[y * cw + x] = true;
+                        break;
+                    }
+                }
+            }
+        }
+        if !nh.iter().any(|b| !*b) {
+            break;
+        }
+        dims.push((cw as u32, ch as u32));
+        holes.push(nh);
+    }
+    let levels = dims.len();
+    let top = levels - 1;
+
+    let mut plan = Vec::with_capacity(levels);
+    for k in 0..levels {
+        plan.push(plan_level(&mut e, holes.get(k)?, dims[k].0, dims[k].1, patch)?);
+    }
+
+    // One image per level (the coarse-to-fine steps and the vote ping-pong in fresh textures).
+    let mut ima: Vec<GpuImage> = Vec::with_capacity(levels);
+    ima.push(g.upload_image(img)?);
+    for k in 1..levels {
+        let (cw, ch) = dims[k];
+        ima.push(e.scratch(cw, ch));
+    }
+    for k in 1..levels {
+        let (pw, ph) = dims[k - 1];
+        let (cw, ch) = dims[k];
+        e.pixels("imp_pm_down", &params([pw, ph, cw, ch], [0; 4], [0; 4], [0.0; 4]), &ima[k - 1], None, &ima[k], Some(&plan[k - 1].flags));
+    }
+    // The coarsest level starts as its membrane fill.
+    let (cw, ch) = dims[top];
+    let spare = e.scratch(cw, ch);
+    ima[top] = membrane_tex(&mut e, &ima[top], &spare, &holes[top], cw, ch)?;
+
+    // Coarse to fine: upsample, then the EM rounds of propagation and voting.
+    let mut nnf: Option<GpuImage> = None;
+    for k in (0..levels).rev() {
+        let lv = &plan[k];
+        nnf = Some(match nnf {
+            Some(coarse) => {
+                let (cw, ch) = dims[k + 1];
+                let fine = e.scratch(lv.w, lv.h);
+                e.pixels("imp_nnf_up", &params([lv.w, lv.h, cw, ch], [0; 4], [0; 4], [0.0; 4]), &coarse, None, &fine, None);
+                fine
+            }
+            None => e.image(lv.w, lv.h),
+        });
+        if k < top {
+            let (cw, ch) = dims[k + 1];
+            let dst = e.scratch(lv.w, lv.h);
+            e.copy_into(&ima[k], &dst, 0, 0);
+            e.pixels("imp_pm_up", &params([lv.w, lv.h, cw, ch], [0; 4], [0; 4], [0.0; 4]), &ima[k + 1], None, &dst, Some(&lv.flags));
+            ima[k] = dst;
+        }
+        let (em, iters) = if k == top { (8, 5) } else { (4, 5) };
+        let spare = e.scratch(lv.w, lv.h);
+        let run = PmRun { seed: seed.wrapping_add(k as u32 * 101), em, iters };
+        let (out, out_nnf) = pm_level(&mut e, lv, run, &ima[k], &spare, nnf.as_ref()?)?;
+        ima[k] = out;
+        nnf = Some(out_nnf);
+    }
+    if e.take_missing() {
+        return None;
+    }
+    // Only the hole changes, so only the hole's bounds cross back to the CPU.
+    let bb = bbox(&hole, uw, uh)?;
+    let rect = e.read_rect(&ima[0].texture, bb[0], bb[1], bb[2] - bb[0], bb[3] - bb[1], 16)?;
+    Some(patch_back(img, &hole, bb, &rect))
+}
+
+/// The plan of one level: its flags (`bit 0` hole, `bit 1` a target centre) and the dispatch
+/// bounds of its targets, holes and field.
+fn plan_level(e: &mut Enc, hole: &[bool], w: u32, h: u32, patch: usize) -> Option<Lv> {
+    let (iw, ih) = (w as usize, h as usize);
+    let r = radius(patch, w, h);
+    let near = dilate(hole, iw, ih, r as usize);
+    let near_bb = bbox(&near, iw, ih)?;
+    let flags: Vec<f32> = (0..iw * ih).map(|i| (u32::from(hole[i]) + 2 * u32::from(near[i])) as f32).collect();
+    Some(Lv { w, h, r, near: near_bb, holes: bbox(hole, iw, ih)?, field: grow(near_bb, r + 1, w, h), flags: e.data(&flags) })
+}
+
+/// The PatchMatch run at one level: the random stream's seed and the EM round and iteration
+/// counts.
+struct PmRun {
+    seed: u32,
+    em: usize,
+    iters: usize,
+}
+
+/// `pm_level`: `run.em` EM rounds of propagation and voting at one level. `img` holds the level's
+/// image and `img_spare` is a free texture of the same size; the nearest-neighbour field starts
+/// at `nnf`. Returns the level's filled image and its field.
+fn pm_level(e: &mut Enc, lv: &Lv, run: PmRun, img: &GpuImage, img_spare: &GpuImage, nnf: &GpuImage) -> Option<(GpuImage, GpuImage)> {
+    let (seed, em, iters) = (run.seed, run.em, run.iters);
+    let (r, w, h) = (lv.r, lv.w, lv.h);
+    let field_a = e.scratch(w, h);
+    let field_b = e.scratch(w, h);
+    let hist = e.scratch(BINS as u32, 1);
+    let (img_a, img_b) = (img.clone(), img_spare.clone());
+    let (mut field_nnf, mut nnf_spare) = (nnf.clone(), e.scratch(w, h));
+    // The propagation and random search write the targets and the one-pixel ring around them, so
+    // every neighbour they read has a defined alpha (`grow(near, 1)`). The dispatch origin must
+    // be the ring's, not the targets': `over` sizes the dispatch from the box it is given.
+    let ring = grow(lv.near, 1, w, h);
+    let mut in_a = true;
+    for round in 0..em {
+        // The round's initial field: the inherited source where it is still valid, else random.
+        // Every pixel of the field's bounds is written, so the previous round's field is free.
+        let img = if in_a { &img_a } else { &img_b };
+        let p = params([w, h, r, 0], [0; 4], [seed, 0, 0, round as u32], [0.0; 4]);
+        over(e, "imp_pm_init", &p, Pass { src: &field_nnf, aux: Some(img), out: &field_a, data: Some(&lv.flags), bb: lv.field });
+        let mut f_in_a = true;
+        for it in 0..iters {
+            for parity in 0..2u32 {
+                let (s, o) = if f_in_a { (&field_a, &field_b) } else { (&field_b, &field_a) };
+                let p = params([w, h, r, parity], [0; 4], [seed, it as u32, 0, round as u32], [0.0; 4]);
+                over(e, "imp_pm_prop", &p, Pass { src: s, aux: Some(img), out: o, data: Some(&lv.flags), bb: ring });
+                f_in_a = !f_in_a;
+            }
+            let (s, o) = if f_in_a { (&field_a, &field_b) } else { (&field_b, &field_a) };
+            let p = params([w, h, r, 0], [0; 4], [seed, it as u32, 0, round as u32], [0.0; 4]);
+            over(e, "imp_pm_rand", &p, Pass { src: s, aux: Some(img), out: o, data: Some(&lv.flags), bb: ring });
+            f_in_a = !f_in_a;
+        }
+        let field = if f_in_a { &field_a } else { &field_b };
+        // The vote's sigma²: the 75th percentile of the targets' match distances (`pm_level`).
+        let sigma2 = sigma2(e, field, lv, &hist)?;
+        let next_img = if in_a { &img_b } else { &img_a };
+        e.copy_into(img, next_img, 0, 0);
+        let p = params([w, h, r, 0], [0; 4], [0; 4], [sigma2, 0.0, 0.0, 0.0]);
+        over(e, "imp_pm_vote", &p, Pass { src: field, aux: Some(img), out: next_img, data: Some(&lv.flags), bb: lv.holes });
+        // Carry the round's sources into the nearest-neighbour field (targets only).
+        e.pixels("imp_nnf_store", &params([w, h, 0, 0], [0; 4], [0; 4], [0.0; 4]), field, Some(&field_nnf), &nnf_spare, Some(&lv.flags));
+        std::mem::swap(&mut field_nnf, &mut nnf_spare);
+        in_a = !in_a;
+    }
+    Some((if in_a { img_a } else { img_b }, field_nnf))
+}
+
+/// The 75th percentile of a 1024-bin distance histogram (`imp_hist`'s readback) covering
+/// `0..range`, and the histogram's total. The bin's centre is the estimate: within half a bin
+/// (`range/2048`) of the CPU's exact percentile.
+fn percentile(bytes: &[u8], range: f32) -> (f32, f64) {
+    let bin = |i: usize| -> f64 {
+        match bytes.get(i * 16..i * 16 + 4) {
+            Some(b) => f32::from_le_bytes([b[0], b[1], b[2], b[3]]) as f64,
+            None => 0.0,
+        }
+    };
+    let total: f64 = (0..BINS).map(bin).sum();
+    if total < 1.0 {
+        return (1.0, total);
+    }
+    let rank = total * 3.0 / 4.0;
+    let mut cum = 0.0f64;
+    let mut at = BINS - 1;
+    for i in 0..BINS {
+        cum += bin(i);
+        if cum > rank {
+            at = i;
+            break;
+        }
+    }
+    let centre = ((at as f32) + 0.5) * range / BINS as f32;
+    (centre.max(1e-5), total)
+}
+
+/// The largest finite distance `imp_hist` saw, from bin 0's `.g` channel.
+fn max_distance(bytes: &[u8]) -> f32 {
+    match bytes.get(4..8) {
+        Some(b) => f32::from_le_bytes([b[0], b[1], b[2], b[3]]),
+        None => 0.0,
+    }
+}
+
+/// `imp_hist` over `bb`, the targets' bounds: the 75th percentile of the finite distances and the
+/// sample count. SDR input is binned over the fixed [`SDR_DISTANCE`] range; a frame with values
+/// above 1 (32-bpc HDR) has larger distances, so the second dispatch bins over the observed
+/// maximum instead. Clamping those to the SDR range would report a sigma² of ≈ 4 for a distance
+/// of 1600, and `imp_pm_vote`'s weights would round to zero and keep the prior image.
+///
+/// One 16 KB readback per EM round (≈ 28 per 1080p patch) on the SDR path: a few percent of the
+/// stage, and the obvious next thing to remove if the patch has to be faster (a device-side
+/// percentile would need a reduction instead). HDR frames pay one more dispatch and readback.
+fn histogram(e: &mut Enc, field: &GpuImage, flags: &wgpu::Buffer, w: u32, h: u32, bb: [u32; 4], hist: &GpuImage) -> Option<(f32, f64)> {
+    let (bw, bh) = (bb[2] - bb[0], bb[3] - bb[1]);
+    let mut p = params([w, h, 0, 0], [bb[0], bb[1], bw, bh], [0; 4], [0.0; 4]);
+    e.dispatch("imp_hist", &p, KernelImages { src: field, aux: None, out: hist }, Some(flags), (1, 1));
+    let bytes = e.read_texture(&hist.texture, BINS as u32, 1, 16)?;
+    let max = max_distance(&bytes);
+    if max <= SDR_DISTANCE {
+        return Some(percentile(&bytes, SDR_DISTANCE));
+    }
+    // The observed maximum itself has to be in the last bin, so widen the range by a hair rather
+    // than binning it to 1024.
+    let range = max * (1.0 + 1e-6);
+    p.f[0] = [range, 0.0, 0.0, 0.0];
+    e.dispatch("imp_hist", &p, KernelImages { src: field, aux: None, out: hist }, Some(flags), (1, 1));
+    let bytes = e.read_texture(&hist.texture, BINS as u32, 1, 16)?;
+    Some(percentile(&bytes, range))
+}
+
+/// The vote's sigma² at a level (`pm_level`): the 75th percentile of the targets' distances.
+fn sigma2(e: &mut Enc, field: &GpuImage, lv: &Lv, hist: &GpuImage) -> Option<f32> {
+    Some(histogram(e, field, &lv.flags, lv.w, lv.h, lv.near, hist)?.0)
+}
+
+/// [`imp_hist`] on its own: `field`'s match distance (`.b`, `.a` > 0.5 where set) over the
+/// pixels `near` marks, returning the 75th percentile and the sample count. Used by the test that
+/// checks the histogram counts the whole target population.
+#[cfg(test)]
+pub(crate) fn distance_histogram(g: &GpuContext, field: &Image, near: &[bool], w: u32, h: u32) -> Option<(f32, f64)> {
+    let flags: Vec<f32> = near.iter().map(|b| if *b { 2.0 } else { 0.0 }).collect();
+    let mut e = Enc::new(g);
+    let data = e.data(&flags);
+    let img = g.upload_image(field)?;
+    let hist = e.scratch(BINS as u32, 1);
+    let out = histogram(&mut e, &img, &data, w, h, [0, 0, w, h], &hist);
+    if e.take_missing() {
+        return None;
+    }
+    out
+}

crates/gpu/src/lib.rs

@@ -79,6 +79,8 @@ mod fx_tone;
 mod fx_transition;
 mod fx_vr;
 mod fx_warp;
+mod inpaint;
+mod memcap;
 mod ops;
 mod particles;
 mod readback;
@@ -397,7 +399,7 @@ impl Accelerator for Gpu {
 
     fn effects(&self, chain: &[FxStep], buf: &Buf, levels: Option<f32>) -> Option<Buf> {
         self.ctx.check_health().ok()?;
-        effects::run_chain(&mut Enc::new(&self.ctx), chain, buf, levels)
+        effects::run_chain(&mut Enc::lent(&self.ctx), chain, buf, levels)
     }
 
     fn raster_3d(&self, scene: &effectcraft_render::three_d::adv::Scene) -> Option<effectcraft_render::three_d::adv::Target> {
@@ -415,16 +417,41 @@ impl Accelerator for Gpu {
         self.ctx.can_readback().then_some(self as &dyn effectcraft_effects::psim::ParticleSim)
     }
 
+    fn inpaint(&self) -> Option<&dyn effectcraft_raster::inpaint::FillBackend> {
+        self.ctx.check_health().ok()?;
+        // Readback is synchronous-only, and a kernel that failed to build would make every stage
+        // fall back to the CPU *after* paying the upload and a failed dispatch (and the fill
+        // would still report the GPU path). Neither is a backend at all.
+        if !self.ctx.can_wait() || self.ctx.unavailable_kernels().iter().any(|(e, _)| crate::inpaint::KERNELS.contains(e)) {
+            return None;
+        }
+        Some(self as &dyn effectcraft_raster::inpaint::FillBackend)
+    }
+
     fn auto_pick(&self) -> Option<&effectcraft_render::AutoPick> {
         Some(&self.auto)
     }
 
+    fn render_end(&self) {
+        let (hits, misses) = self.ctx.pool_stats();
+        let pct = hits.saturating_mul(100) / hits.saturating_add(misses).max(1);
+        log::debug!(
+            target: "effectcraft_gpu",
+            "gpu pool: {hits} working textures reused, {misses} allocated ({pct}% reuse) so far"
+        );
+        self.ctx.trim_pool();
+    }
+
     fn frame_begin(&self) {
         if let Some(d) = self.ctx.deferred() {
             d.frame_begin();
         }
     }
 
+    fn frame_end(&self) {
+        self.ctx.pool_frame_end();
+    }
+
     fn pass_begin(&self) {
         if let Some(d) = self.ctx.deferred() {
             d.pass_begin();

crates/gpu/src/memcap.rs

@@ -0,0 +1,236 @@
+//! Memory back-pressure for frames in flight.
+//!
+//! wgpu destroys a texture or buffer we drop only once the submissions that used it have
+//! completed, and pooled textures released while other frames' encoders are open stay pending
+//! until then. With several 4K frames in flight, the frames record work faster than the GPU
+//! retires it, and freed-but-not-destroyed memory runs the device out of memory (#879). Once the
+//! allocator holds more than the cap and enough of that total is reclaimable (not held by the
+//! upload cache, the texture pool, or live images), a frame submits its encoder and waits for its
+//! own submission so wgpu frees what was dropped. Frames are not held back for memory that cannot
+//! be reclaimed; a warning is logged once.
+//!
+//! wgpu has no safe query for the adapter's memory size (`max_buffer_size` is capped at `i32::MAX`
+//! on Direct3D 12 and Linux/Mesa Vulkan), so the default cap is conservative per device type;
+//! `EFFECTCRAFT_GPU_MEMCAP_MB` overrides it (0 = off). Backends without an allocator report
+//! (Metal, WebGPU, GL) are never gated.
+
+use std::sync::atomic::{AtomicBool, AtomicU64, Ordering::Relaxed};
+use std::time::Duration;
+
+/// Check the allocator only every this many gate calls (effect calls and layers): the report
+/// lists every allocation, and between checks a frame adds at most a few textures.
+const EVERY: u64 = 8;
+
+const GIB: u64 = 1 << 30;
+const MIB: u64 = 1 << 20;
+
+/// Overrides the default cap: whole MiB, 0 turns the gate off.
+const MEMCAP_ENV: &str = "EFFECTCRAFT_GPU_MEMCAP_MB";
+
+pub(crate) struct MemCap {
+    cap: u64,
+    hysteresis: u64,
+    calls: AtomicU64,
+    waits: AtomicU64,
+    /// Reclaimable bytes remaining after the last wait: another wait fires only once reclaimable
+    /// memory grows by `hysteresis` above this floor (or drops below it first).
+    floor: AtomicU64,
+    /// Set after the first over-cap sample where non-reclaimable memory alone exceeds the cap.
+    warned: AtomicBool,
+    #[cfg(test)]
+    timeout_ns: AtomicU64,
+}
+
+impl MemCap {
+    pub(crate) fn new(info: &wgpu::AdapterInfo) -> MemCap {
+        let env = std::env::var(MEMCAP_ENV).ok();
+        Self::from_env(env.as_deref(), info.device_type)
+    }
+
+    fn from_env(var: Option<&str>, device_type: wgpu::DeviceType) -> MemCap {
+        let cap = cap_from_env(var, device_type);
+        MemCap {
+            cap,
+            hysteresis: default_hysteresis(cap),
+            calls: AtomicU64::new(0),
+            waits: AtomicU64::new(0),
+            floor: AtomicU64::new(0),
+            warned: AtomicBool::new(false),
+            #[cfg(test)]
+            timeout_ns: AtomicU64::new(crate::readback::TIMEOUT.as_nanos() as u64),
+        }
+    }
+
+    /// The same cap with fresh counters (a test context sharing another's device).
+    #[cfg(test)]
+    pub(crate) fn fresh(&self) -> MemCap {
+        MemCap {
+            cap: self.cap,
+            hysteresis: self.hysteresis,
+            calls: AtomicU64::new(0),
+            waits: AtomicU64::new(0),
+            floor: AtomicU64::new(0),
+            warned: AtomicBool::new(false),
+            timeout_ns: AtomicU64::new(self.timeout_ns.load(Relaxed)),
+        }
+    }
+
+    pub(crate) fn cap(&self) -> u64 {
+        self.cap
+    }
+
+    /// Whether to sample the allocator on this call.
+    pub(crate) fn due(&self) -> bool {
+        self.cap != 0 && self.calls.fetch_add(1, Relaxed).is_multiple_of(EVERY)
+    }
+
+    /// Whether `allocated` bytes in the allocator with `kept` non-reclaimable bytes (upload
+    /// cache, pool free list, live frames) should trigger a submit and wait.
+    pub(crate) fn should_wait(&self, allocated: u64, kept: u64) -> bool {
+        if self.cap == 0 || allocated <= self.cap {
+            self.floor.store(0, Relaxed);
+            return false;
+        }
+        if kept > self.cap && !self.warned.swap(true, Relaxed) {
+            log::warn!(
+                target: "effectcraft_gpu",
+                "gpu: {} MiB held by live frames, the upload cache and the texture pool is over the {} MiB memory cap: waiting cannot free it, so frames are not held back for it ({MEMCAP_ENV} sets the cap)",
+                kept >> 20,
+                self.cap >> 20,
+            );
+        }
+        let reclaimable = allocated.saturating_sub(kept);
+        let floor = self.floor.load(Relaxed);
+        if reclaimable < floor {
+            self.floor.store(reclaimable, Relaxed);
+            return false;
+        }
+        reclaimable >= floor.saturating_add(self.hysteresis)
+    }
+
+    pub(crate) fn record_wait(&self, allocated_after: u64, kept_after: u64) {
+        self.floor.store(allocated_after.saturating_sub(kept_after), Relaxed);
+        let n = self.waits.fetch_add(1, Relaxed);
+        if n == 0 {
+            log::info!("gpu: allocator over the {} MiB memory cap, frames wait for the GPU to catch up", self.cap >> 20);
+        }
+    }
+
+    pub(crate) fn wait_timeout(&self) -> Duration {
+        #[cfg(test)]
+        {
+            Duration::from_nanos(self.timeout_ns.load(Relaxed))
+        }
+        #[cfg(not(test))]
+        {
+            crate::readback::TIMEOUT
+        }
+    }
+
+    #[cfg(test)]
+    pub(crate) fn waits(&self) -> u64 {
+        self.waits.load(Relaxed)
+    }
+
+    #[cfg(test)]
+    pub(crate) fn warned(&self) -> bool {
+        self.warned.load(Relaxed)
+    }
+
+    #[cfg(test)]
+    pub(crate) fn set_wait_timeout(&self, timeout: Duration) {
+        self.timeout_ns.store(timeout.as_nanos() as u64, Relaxed);
+    }
+}
+
+/// The cap from the `EFFECTCRAFT_GPU_MEMCAP_MB` value: whole MiB (0 turns the gate off), or the
+/// default for `device_type` when unset, or (with a warning) when not a number or too large.
+fn cap_from_env(var: Option<&str>, device_type: wgpu::DeviceType) -> u64 {
+    let Some(v) = var else { return default_cap(device_type) };
+    match v.trim().parse::<u64>().ok().and_then(|mb| mb.checked_mul(MIB)) {
+        Some(cap) => cap,
+        None => {
+            let cap = default_cap(device_type);
+            log::warn!(target: "effectcraft_gpu", "{MEMCAP_ENV}={v:?} is not a whole number of MiB that fits in 64 bits: using {} MiB", cap / MIB);
+            cap
+        }
+    }
+}
+
+/// Discrete cards: 3 GiB, low enough to fire on 4–8 GB cards with headroom for the driver and
+/// desktop while leaving ~1 GiB above the default 1.5 GiB upload cache + 256 MB pool (larger
+/// cards can raise it with `EFFECTCRAFT_GPU_MEMCAP_MB`). Integrated and software adapters share
+/// system memory: 2 GiB.
+fn default_cap(t: wgpu::DeviceType) -> u64 {
+    match t {
+        wgpu::DeviceType::DiscreteGpu => 3 * GIB,
+        _ => 2 * GIB,
+    }
+}
+
+/// Minimum reclaimable bytes before a wait is worth draining an encoder (1/8 of the cap,
+/// clamped between 1 MiB and 256 MiB, or 1/4 of caps smaller than 4 MiB).
+fn default_hysteresis(cap: u64) -> u64 {
+    if cap == 0 { 0 } else { (cap / 8).clamp(MIB.min(cap / 4), 256 * MIB) }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn defaults_and_cadence() {
+        assert_eq!(default_cap(wgpu::DeviceType::DiscreteGpu), 3 * GIB);
+        assert_eq!(default_cap(wgpu::DeviceType::IntegratedGpu), 2 * GIB);
+        let m = MemCap::from_env(None, wgpu::DeviceType::IntegratedGpu);
+        assert_eq!(m.cap(), 2 * GIB);
+        let due = (0..2 * EVERY).filter(|_| m.due()).count();
+        assert_eq!(due, 2);
+        let off = MemCap::from_env(Some("0"), wgpu::DeviceType::DiscreteGpu);
+        assert_eq!(off.cap(), 0);
+        assert!((0..EVERY).all(|_| !off.due()));
+    }
+
+    #[test]
+    fn override_is_whole_mib_or_the_default() {
+        let cap = |v| MemCap::from_env(Some(v), wgpu::DeviceType::DiscreteGpu).cap();
+        assert_eq!(cap("2048"), 2 * GIB);
+        assert_eq!(cap(" 8 "), 8 * MIB);
+        assert_eq!(cap("17592186044415"), 17592186044415 * MIB, "the largest value that fits");
+        for bad in ["", "lots", "-1", "1.5", "17592186044416", "99999999999999999999"] {
+            assert_eq!(cap(bad), 3 * GIB, "{bad:?} falls back to the default");
+        }
+    }
+
+    #[test]
+    fn hysteresis_skips_non_reclaimable_and_unfreed_memory() {
+        let m = MemCap::from_env(Some("2048"), wgpu::DeviceType::IntegratedGpu);
+        // Upload cache + pool + live frames (2.15 GiB) sit over the 2 GiB cap with only 50 MiB
+        // reclaimable: waiting would free almost nothing, so the gate must stay quiet.
+        let kept = 2150 * MIB;
+        assert!(!m.should_wait(kept + 50 * MIB, kept));
+        assert!(!m.should_wait(kept + 200 * MIB, kept));
+        // Once 256 MiB of dropped textures pile up, the gate fires.
+        assert!(m.should_wait(kept + 256 * MIB, kept));
+        // If a wait freed nothing (e.g. held by another thread's open encoder or timed out),
+        // the floor prevents firing again until another 256 MiB of reclaimable memory accumulates.
+        m.record_wait(kept + 256 * MIB, kept);
+        assert!(!m.should_wait(kept + 256 * MIB, kept));
+        assert!(!m.should_wait(kept + 400 * MIB, kept));
+        assert!(m.should_wait(kept + 512 * MIB, kept));
+        // Once the other thread finishes and reclaimable memory drops, the floor drops with it.
+        assert!(!m.should_wait(kept + 16 * MIB, kept));
+        assert!(m.should_wait(kept + (16 + 256) * MIB, kept));
+    }
+
+    #[test]
+    fn kept_memory_over_the_cap_warns_once_and_never_waits() {
+        let m = MemCap::from_env(Some("8"), wgpu::DeviceType::IntegratedGpu);
+        // Cap 8 MiB: kept alone over the cap with reclaimable under the hysteresis → no wait, warn once.
+        assert!(!m.should_wait(20 * MIB, 20 * MIB));
+        assert!(m.warned());
+        assert!(!m.should_wait(20 * MIB, 20 * MIB));
+        assert!(!m.should_wait(20 * MIB, 19 * MIB + MIB / 2));
+        assert!(m.should_wait(20 * MIB, 4 * MIB), "16 MiB reclaimable is over the hysteresis");
+    }
+}

crates/gpu/src/shaders/common.wgsl

@@ -50,6 +50,91 @@ fn acos_p(x: f32) -> f32 {
     return 1.5707963267948966 - asin_p(x);
 }
 
+// Some hardware's native trig instructions saturate cos near a full turn and
+// have errors large enough to move particles and sphere samples visibly. Reduce
+// to [-pi/4, pi/4] with a split pi/2, then evaluate the Taylor polynomials with
+// explicit fused operations. The effect angles use this bounded range; retain
+// the builtins for non-finite or very large arguments rather than cast them.
+fn sincos_p(x: f32) -> vec2<f32> {
+    if (!(abs(x) <= 8192.0)) {
+        return vec2<f32>(sin(x), cos(x));
+    }
+    let k = round(x * 0.6366197723675814);
+    // Keep the coarse reduction in integers. Some drivers reassociate even
+    // explicit fma chains into x - k*pi/2, losing low bits for large angles.
+    // 3216/2048 is exact; the bounded scaled input and product fit in i32.
+    let scaled = x * 2048.0;
+    let whole = floor(scaled);
+    var r = f32(i32(whole) - i32(k) * 3216) / 2048.0;
+    r += (scaled - whole) / 2048.0;
+    r = fma(-k, 0.0004838267893064022, r);
+    r = fma(-k, 5.564431554416771e-12, r);
+    let z = r * r;
+    var sp = fma(z, -1.0 / 39916800.0, 1.0 / 362880.0);
+    sp = fma(z, sp, -1.0 / 5040.0);
+    sp = fma(z, sp, 1.0 / 120.0);
+    sp = fma(z, sp, -1.0 / 6.0);
+    let sn = fma(r * z, sp, r);
+    var cp = fma(z, -1.0 / 3628800.0, 1.0 / 40320.0);
+    cp = fma(z, cp, -1.0 / 720.0);
+    cp = fma(z, cp, 1.0 / 24.0);
+    cp = fma(z, cp, -0.5);
+    let cs = fma(z, cp, 1.0);
+    switch imod(i32(k), 4) {
+        case 1: { return vec2<f32>(cs, -sn); }
+        case 2: { return vec2<f32>(-sn, -cs); }
+        case 3: { return vec2<f32>(-cs, sn); }
+        default: { return vec2<f32>(sn, cs); }
+    }
+}
+
+fn sin_p(x: f32) -> f32 { return sincos_p(x).x; }
+fn cos_p(x: f32) -> f32 { return sincos_p(x).y; }
+
+// Correctly-rounded positive x/a, clamped to [0,1]. Native division and fma
+// need not have IEEE precision on every shader backend. Discontinuous colour
+// effects need the exact CPU f32 value, not just an approximation within an ulp.
+// Long division uses only bounded 24-bit significands (at most 24 iterations).
+fn unit_div(x: f32, a: f32) -> f32 {
+    if (x <= 0.0 || a <= 0.0) { return 0.0; }
+    if (x >= a) { return 1.0; }
+    if (a == 1.0) { return x; }
+    let xb = bitcast<u32>(x);
+    let ab = bitcast<u32>(a);
+    var nx = xb & 0x7fffffu;
+    var na = ab & 0x7fffffu;
+    var ex = i32((xb >> 23u) & 255u) - 127;
+    var ea = i32((ab >> 23u) & 255u) - 127;
+    if (ex == -127) {
+        let shift = 23u - firstLeadingBit(nx);
+        nx <<= shift;
+        ex = -126 - i32(shift);
+    } else { nx |= 0x800000u; }
+    if (ea == -127) {
+        let shift = 23u - firstLeadingBit(na);
+        na <<= shift;
+        ea = -126 - i32(shift);
+    } else { na |= 0x800000u; }
+    var exponent = ex - ea;
+    if (nx < na) { nx <<= 1u; exponent -= 1; }
+    var rem = nx - na;
+    if (exponent < -150) { return 0.0; }
+    if (exponent == -150) { return bitcast<f32>(select(0u, 1u, rem != 0u)); }
+    let fraction = u32(min(23, exponent + 149));
+    var q = 1u;
+    for (var i = 0u; i <= fraction; i++) {
+        rem <<= 1u;
+        q <<= 1u;
+        if (rem >= na) { rem -= na; q |= 1u; }
+    }
+    let guard = q & 1u;
+    q >>= 1u;
+    if (guard != 0u && (rem != 0u || (q & 1u) != 0u)) { q += 1u; }
+    if (exponent < -126) { return bitcast<f32>(q); }
+    if (q == 0x1000000u) { q >>= 1u; exponent += 1; }
+    return bitcast<f32>((u32(exponent + 127) << 23u) | (q & 0x7fffffu));
+}
+
 // ---------------------------------------------------------------- sampling (raster::Image)
 
 fn tex_get(t: texture_2d<f32>, x: i32, y: i32) -> vec4<f32> {
@@ -157,13 +242,16 @@ fn powz(x: f32, y: f32) -> f32 {
     return pow(x, y);
 }
 
-// Rust's f32::round (half away from zero; WGSL's round() is half to even).
+// Rust's f32::round (half away from zero; WGSL's round() is half to even). Decided by comparing
+// `x` with `t ± 0.5` (exact for integer `t`) rather than subtracting: drivers may fuse a product
+// that feeds `x - t` into one fma (AMD does), skipping the product's rounding, so `5/6 · 3`,
+// which rounds to exactly 2.5 on the CPU, came out 0.49999994 above 2 and rounded down (#1605).
 fn round_away(x: f32) -> f32 {
     let t = trunc(x);
-    if (abs(x - t) >= 0.5) {
-        return t + sign(x);
+    if (x >= 0.0) {
+        return select(t, t + 1.0, x >= t + 0.5);
     }
-    return t;
+    return select(t, t - 1.0, x <= t - 0.5);
 }
 
 // f32::rem_euclid(1.0).

crates/gpu/src/shaders/fx_pixel2.wgsl

@@ -633,52 +633,9 @@ fn fp2_hextile(@builtin(global_invocation_id) gid: vec3<u32>) {
     if (!fp2_inside(p)) {
         return;
     }
-    let cx = P.f[0].x;
-    let cy = P.f[0].y;
-    let r = P.f[0].z;
-    let smear = P.f[0].w;
-    let sn = P.f[1].x;
-    let cs = P.f[1].y;
-    let mode = P.u[0].x;
-    let dx = f32(p.x) + 0.5 - cx;
-    let dy = f32(p.y) + 0.5 - cy;
-    let u = dx * cs + dy * sn;
-    let v = -dx * sn + dy * cs;
-    // stylize3::hex_cell (cube rounding).
-    let q = (0.5773502691896257 * u - v / 3.0) / r;
-    let rr = (2.0 / 3.0 * v) / r;
-    // Evaluate the third cube coordinate symmetrically with q so the centerline
-    // tie is not broken by subtracting separately rounded coordinates.
-    let ccy = (-0.5773502691896257 * u - v / 3.0) / r;
-    var rx = round_away(q);
-    let ry = round_away(ccy);
-    var rz = round_away(rr);
-    let ddx = abs(rx - q);
-    let ddy = abs(ry - ccy);
-    let ddz = abs(rz - rr);
-    if (ddx > ddy && ddx > ddz) {
-        rx = -ry - rz;
-    } else if (ddy <= ddz) {
-        rz = -rx - ry;
-    }
-    let hx = r * 1.7320508075688772 * (rx + rz / 2.0);
-    let hy = r * 1.5 * rz;
-    var lu = u - hx;
-    var lv = v - hy;
-    let odd = imod(i32(rx) + i32(rz), 2) == 1;
-    if ((mode == 1u || mode == 2u) && odd) {
-        lu = -lu;
-    }
-    if (mode == 2u && odd) {
-        lv = -lv;
-    }
-    let ox = lu * cs - lv * sn;
-    let oy = lu * sn + lv * cs;
-    let tx = cx + ox;
-    let ty = cy + oy;
-    let sx = tx + (f32(p.x) + 0.5 - tx) * smear;
-    let sy = ty + (f32(p.y) + 0.5 - ty) * smear;
-    textureStore(out, p, sample_bilinear_clamped(src, sx, sy));
+    // Cell membership and reflection were decided in the CPU effect's f64 geometry.
+    let i = (gid.y * u32(out_dims().x) + gid.x) * 2u;
+    textureStore(out, p, sample_bilinear_clamped(src, data[i], data[i + 1u]));
 }
 
 // ---------------------------------------------------------------- CC Mr. Smoothie (stylize3::smoothie)

crates/gpu/src/shaders/fx_sim.wgsl

@@ -114,8 +114,8 @@ fn fxm_sprite(at: u32, x: f32, y: f32) -> FxmShade {
             cov = min(0.12 * select(0.0, 1.0, d <= 1.0) + ring * 0.8 + hl, 1.0);
         }
         case 5u: {
-            let sn = sin(data[at + 8u]);
-            let cs = cos(data[at + 8u]);
+            let sn = sin_p(data[at + 8u]);
+            let cs = cos_p(data[at + 8u]);
             let ux = (dx * cs + dy * sn) / r;
             let uy = (-dx * sn + dy * cs) / r;
             let d = sqrt(ux * ux + uy * uy);
@@ -125,23 +125,23 @@ fn fxm_sprite(at: u32, x: f32, y: f32) -> FxmShade {
             cov = min(core + stx + sty, 1.0);
         }
         case 6u: {
-            let sn = sin(data[at + 8u]);
-            let cs = cos(data[at + 8u]);
+            let sn = sin_p(data[at + 8u]);
+            let cs = cos_p(data[at + 8u]);
             let ux = dx * cs + dy * sn;
             let uy = -dx * sn + dy * cs;
             let e = r * 0.75;
             cov = fxm_aa(e - abs(ux)) * fxm_aa(e - abs(uy));
         }
         case 7u: {
-            let sn = sin(data[at + 8u]);
-            let cs = cos(data[at + 8u]);
+            let sn = sin_p(data[at + 8u]);
+            let cs = cos_p(data[at + 8u]);
             let ux = dx * cs + dy * sn;
             let uy = -dx * sn + dy * cs;
             let e = r * 0.5;
             var m = 3.4e38;
             for (var k = 0; k < 3; k++) {
                 let a = 1.5707964 + f32(k) * 6.2831855 / 3.0;
-                m = min(m, e - (ux * cos(a) + uy * sin(a)));
+                m = min(m, e - (ux * cos_p(a) + uy * sin_p(a)));
             }
             cov = fxm_aa(m);
         }
@@ -431,8 +431,8 @@ fn fxm_ripples(x: f32, y: f32) -> f32 {
         let r = data[i * 4u + 2u];
         let u = (d - r) / P.f[0].y;
         if (abs(u) < 1.0) {
-            let env = 0.5 + 0.5 * cos(u * 3.1415927);
-            v += data[i * 4u + 3u] * env * cos((d - r) / P.f[0].x * 6.2831855);
+            let env = 0.5 + 0.5 * cos_p(u * 3.1415927);
+            v += data[i * 4u + 3u] * env * cos_p((d - r) / P.f[0].x * 6.2831855);
         }
     }
     return v;

crates/gpu/src/shaders/fx_stylize.wgsl

@@ -42,7 +42,11 @@ fn fxs_qdiv(x: vec3<f32>, a: f32) -> vec3<f32> {
     let r = fma(-q, vec3<f32>(a), x);
     // `q + r / a`, written as an fma so drivers (NVIDIA Vulkan) can't contract it another way
     // and round differently from the CPU.
-    return fma(r / a, vec3<f32>(1.0), q);
+    let corrected = fma(r / a, vec3<f32>(1.0), q);
+    return vec3<f32>(
+        select(corrected.x, unit_div(x.x, a), x.x >= 0.0 && x.x <= a),
+        select(corrected.y, unit_div(x.y, a), x.y >= 0.0 && x.y <= a),
+        select(corrected.z, unit_div(x.z, a), x.z >= 0.0 && x.z <= a));
 }
 
 fn fxs_rem2(i: i32) -> i32 {
@@ -63,6 +67,32 @@ fn fxs_bin(v: f32, t: f32, inv: bool) -> f32 {
     return select(0.0, 1.0, (v >= t) != inv);
 }
 
+// Preserve the CPU's separately-rounded f32 multiply before rounding to a bin.
+// A shader compiler may reassociate division and multiplication across a half
+// threshold. The 24-bit mantissa times at most 254 fits in u32 exactly.
+fn fxs_posterize_bin(v: f32, levels: u32) -> f32 {
+    let bits = bitcast<u32>(clamp(v, 0.0, 1.0));
+    let exponent = (bits >> 23u) & 255u;
+    if (exponent == 0u) { return 0.0; }
+    var product = ((bits & 0x7fffffu) | 0x800000u) * levels;
+    let high = firstLeadingBit(product);
+    if (high > 23u) {
+        let lost = high - 23u;
+        let mask = (1u << lost) - 1u;
+        let remainder = product & mask;
+        let half = 1u << (lost - 1u);
+        var kept = product >> lost;
+        if (remainder > half || (remainder == half && (kept & 1u) != 0u)) { kept += 1u; }
+        product = kept << lost;
+    }
+    let shift = 150u - exponent;
+    if (shift > 32u) { return 0.0; }
+    if (shift == 32u) { return select(0.0, 1.0, product >= 0x80000000u); }
+    let mask = (1u << shift) - 1u;
+    let integer = (product >> shift) + select(0u, 1u, (product & mask) >= (1u << (shift - 1u)));
+    return f32(integer);
+}
+
 fn fxs_strobe_op(op: u32, v: f32, k: f32) -> f32 {
     switch op {
         case 1u: { return v + k; }
@@ -96,8 +126,8 @@ fn fxs_point(@builtin(global_invocation_id) gid: vec3<u32>) {
     switch mode {
         case 1u: {
             let lv = P.f[0].x;
-            let c = clamp(fxs_qdiv(px.xyz, a), vec3<f32>(0.0), vec3<f32>(1.0)) * lv;
-            let o = vec3<f32>(round_away(c.x), round_away(c.y), round_away(c.z)) / lv;
+            let c = vec3<f32>(unit_div(px.x, a), unit_div(px.y, a), unit_div(px.z, a));
+            let o = vec3<f32>(fxs_posterize_bin(c.x, u32(lv)), fxs_posterize_bin(c.y, u32(lv)), fxs_posterize_bin(c.z, u32(lv))) / lv;
             textureStore(out, p, vec4<f32>(o * a, a));
         }
         case 2u: {

crates/gpu/src/shaders/fx_tone.wgsl

@@ -129,16 +129,17 @@ fn fxt_knee(v: f32, start: f32, limit: f32) -> f32 {
     return start + r * (1.0 - exp(-(v - start) / r));
 }
 
-// PS Arbitrary Map lookup (map in data, n entries; f[0].x = phase offset in entries).
+// PS Arbitrary Map: CPU-rounded interval starts, values and slopes in data.
 fn fxt_lookup(v: f32, n: u32) -> f32 {
-    let nf = f32(n);
-    var x = clamp(v, 0.0, 1.0) * (nf - 1.0) + P.f[0].x;
-    x = x - nf * floor(x / nf);
-    let fl = floor(x);
-    let i0 = u32(fl) % n;
-    let i1 = select(i0, i0 + 1u, i0 + 1u < n);
-    let t = x - fl;
-    return data[i0] + (data[i1] - data[i0]) * t;
+    let x = clamp(v, 0.0, 1.0);
+    var lo = 0u;
+    var hi = n;
+    while (lo + 1u < hi) {
+        let mid = lo + (hi - lo) / 2u;
+        if (x < data[mid * 3u]) { hi = mid; } else { lo = mid; }
+    }
+    let i = lo * 3u;
+    return data[i + 1u] + (x - data[i]) * data[i + 2u];
 }
 
 fn fxt_eq_look(table: u32, n: u32, v: f32) -> f32 {
@@ -447,7 +448,11 @@ fn fxt_point(@builtin(global_invocation_id) gid: vec3<u32>) {
         case 12u: {
             let n = u0.z;
             let aa = max(a, 0.0);
-            let r = vec3<f32>(fxt_lookup(c.x, n), fxt_lookup(c.y, n), fxt_lookup(c.z, n));
+            var straight = vec3<f32>(0.0);
+            if (a > 1e-6) {
+                straight = vec3<f32>(unit_div(px.x, a), unit_div(px.y, a), unit_div(px.z, a));
+            }
+            let r = vec3<f32>(fxt_lookup(straight.x, n), fxt_lookup(straight.y, n), fxt_lookup(straight.z, n));
             var na = aa;
             if (u0.y != 0u) {
                 na = fxt_lookup(aa, n);

crates/gpu/src/shaders/fx_transition.wgsl

@@ -224,28 +224,8 @@ fn ftr_grid(@builtin(global_invocation_id) gid: vec3<u32>) {
         return;
     }
     let px = textureLoad(src, p, 0);
-    let d = ftr_centre(p) - P.f[0].xy;
-    let s = P.f[0].z;
-    let co = P.f[0].w;
-    let q = vec2<f32>(d.x * co + d.y * s, -d.x * s + d.y * co);
-    let cs = P.f[1].x;
-    let cc = (floor(q / cs) + 0.5) * cs;
-    var o = clamp(ftr_metric(P.u[0].x, cc) / P.f[1].y, 0.0, 1.0);
-    if (P.u[0].y != 0u) {
-        o = 1.0 - o;
-    }
-    let spread = 0.25;
-    let pr = clamp((P.f[1].z * (1.0 + spread) - o) / spread, 0.0, 1.0);
-    if (pr <= 0.0) {
-        textureStore(out, p, px);
-        return;
-    }
-    let half = cs * 0.5 * (1.0 - pr);
-    let m = max(abs(q.x - cc.x), abs(q.y - cc.y));
-    var k = select(0.0, 1.0, m < half);
-    if (P.f[1].w > 0.0) {
-        k = clamp((half - m) / P.f[1].w + 0.5, 0.0, 1.0);
-    }
+    // CPU-generated coverage preserves exact f64 tile and hard-edge decisions.
+    let k = data[gid.y * u32(out_dims().x) + gid.x];
     textureStore(out, p, px * k);
 }
 

crates/gpu/src/shaders/fx_vr.wgsl

@@ -16,7 +16,7 @@ fn fxv_unpremul(p: vec4<f32>) -> vec4<f32> {
 fn fxv_equi_dir(x: f32, y: f32, w: f32, h: f32) -> vec3<f32> {
     let lon = (x / w - 0.5) * 2.0 * FXV_PI;
     let lat = (0.5 - y / h) * FXV_PI;
-    return vec3<f32>(cos(lat) * sin(lon), sin(lat), cos(lat) * cos(lon));
+    return vec3<f32>(cos_p(lat) * sin_p(lon), sin_p(lat), cos_p(lat) * cos_p(lon));
 }
 
 // vr::equi_px.
@@ -168,7 +168,7 @@ fn fxv_dir(kind: u32, fov: f32, x: f32, y: f32, w: f32, h: f32) -> vec4<f32> {
             cu = u / r;
             cv = v / r;
         }
-        return vec4<f32>(sin(th) * cu, cos(th), sin(th) * cv, 1.0);
+        return vec4<f32>(sin_p(th) * cu, cos_p(th), sin_p(th) * cv, 1.0);
     }
     if (kind == 4u) {
         let half = radians(clamp(fov, 1.0, 360.0) / 2.0);
@@ -186,7 +186,7 @@ fn fxv_dir(kind: u32, fov: f32, x: f32, y: f32, w: f32, h: f32) -> vec4<f32> {
             cu = u / r;
             cv = v / r;
         }
-        return vec4<f32>(sin(th) * cu, sin(th) * cv, cos(th), 1.0);
+        return vec4<f32>(sin_p(th) * cu, sin_p(th) * cv, cos_p(th), 1.0);
     }
     let t = tan(radians(clamp(fov, 1.0, 179.0) / 2.0));
     let u = (x / w * 2.0 - 1.0) * t;
@@ -519,8 +519,8 @@ fn fxv_gen(@builtin(global_invocation_id) gid: vec3<u32>) {
                 var s = px;
                 if (al >= 1e-9) {
                     let new_th = th * (1.0 - P.f[0][c] * weight);
-                    let perp = normalize(d - center * cos(th));
-                    let nd = center * cos(new_th) + perp * sin(new_th);
+                    let perp = normalize(d - center * cos_p(th));
+                    let nd = center * cos_p(new_th) + perp * sin_p(new_th);
                     let q = fxv_equi_px(nd, w, h);
                     s = fxv_sample_equi(q.x, q.y);
                 }

crates/gpu/src/shaders/inpaint.wgsl

@@ -0,0 +1,796 @@
+// GPU Content-Aware Fill: the membrane (harmonic) fill and PatchMatch synthesis of
+// effectcraft_raster::inpaint, operation for operation (the CPU reference is the oracle; see
+// crates/gpu/src/tests/inpaint.rs).
+//
+// Bindings are the shared ones (common.wgsl): `src` / `aux` are the input textures (an image, a
+// coarser level, or the nearest-neighbour field), `out` is the storage texture and `data` is a
+// per-pixel float table — the level's flags:
+//   * membrane kernels: data[i] = 1 where the pixel is known (the hole's complement);
+//   * PatchMatch kernels: data[i] = bit 0 hole, bit 1 near (a pixel whose patch touches the hole).
+//
+// Parameter block (P.u = u32, P.f = f32):
+//   membrane kernels:  u[0] = (w, h, coarseW, coarseH), u[1] = (x0, y0, 0, 0) (dispatch
+//                      origin), u[3].x = this level's offset into `data`
+//   down kernels:      u[0] = (srcW, srcH, dstW, dstH)
+//   PatchMatch kernels: u[0] = (w, h, r, parity), u[1] = (x0, y0, 0, 0) (dispatch origin),
+//                       u[2] = (seed, iteration, 0, EM round), f[0] = (sigma², 0, 0, 0)
+//   histogram:          u[0] = (w, h, 0, 0), u[1] = (x0, y0, w, h) (target bounds),
+//                       f[0].x = the distance range (0: the SDR range 0…4)
+// The nearest-neighbour field texture holds one entry per pixel: r, g = source centre ("px"),
+// b = SSD distance (inf where unset), a = 1 where the entry is set.
+
+const IMP_INF: f32 = 1e30;
+
+fn imp_flags(i: i32) -> u32 {
+    return u32(max(data[i], 0.0));
+}
+
+fn imp_hole(i: i32) -> bool {
+    return (imp_flags(i) & 1u) != 0u;
+}
+
+fn imp_near(i: i32) -> bool {
+    return (imp_flags(i) & 2u) != 0u;
+}
+
+// A source centre whose patch is fully known and inside the image (Pm::valid_source).
+fn imp_valid_source(x: i32, y: i32, w: i32, h: i32, r: i32) -> bool {
+    if (x < r || y < r || x + r >= w || y + r >= h) {
+        return false;
+    }
+    return !imp_near(y * w + x);
+}
+
+// PatchMatch's hashing rnd (raster::inpaint::rnd), bit for bit: the CPU's random stream is a
+// contract of the algorithm, not an implementation detail.
+fn imp_rnd(seed: u32, a: u32, b: u32) -> u32 {
+    var h = (seed * 0x9E3779B9u) ^ (a * 0x85EBCA6Bu) ^ (b * 0xC2B2AE35u);
+    h = h ^ (h >> 16u);
+    h = h * 0x7FEB352Du;
+    h = h ^ (h >> 15u);
+    h = h * 0x846CA68Bu;
+    return h ^ (h >> 16u);
+}
+
+// SSD of the patch at target (tx, ty) (clipped to the image) against source (sx, sy)
+// (Pm::dist). `cap` short-circuits a match that is already worse.
+fn imp_dist(tx: i32, ty: i32, sx: i32, sy: i32, w: i32, h: i32, r: i32, cap: f32) -> f32 {
+    var d = 0.0;
+    var n = 0.0;
+    for (var dy = -r; dy <= r; dy = dy + 1) {
+        let ty2 = ty + dy;
+        if (ty2 < 0 || ty2 >= h) {
+            continue;
+        }
+        let sy2 = sy + dy;
+        for (var dx = -r; dx <= r; dx = dx + 1) {
+            let tx2 = tx + dx;
+            if (tx2 < 0 || tx2 >= w) {
+                continue;
+            }
+            let a = textureLoad(aux, vec2<i32>(tx2, ty2), 0);
+            let b = textureLoad(aux, vec2<i32>(sx + dx, sy2), 0);
+            let e = a - b;
+            d = d + e.x * e.x + e.y * e.y + e.z * e.z + e.w * e.w;
+            n = n + 1.0;
+        }
+        if (d > cap * max(n, 1.0)) {
+            return IMP_INF;
+        }
+    }
+    return d / max(n, 1.0);
+}
+
+// ---------------------------------------------------------------- membrane
+
+// membrane_core's coarse level: the known-weighted 2×2 average.
+@compute @workgroup_size(16, 16)
+fn imp_down(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let cw = i32(P.u[0].z);
+    let ch = i32(P.u[0].w);
+    let cx = i32(gid.x);
+    let cy = i32(gid.y);
+    if (cx >= cw || cy >= ch) {
+        return;
+    }
+    var s = vec4<f32>(0.0);
+    var n = 0.0;
+    for (var dy = 0; dy < 2; dy = dy + 1) {
+        for (var dx = 0; dx < 2; dx = dx + 1) {
+            let fx = 2 * cx + dx;
+            let fy = 2 * cy + dy;
+            if (fx < w && fy < h && data[i32(P.u[3].x) + fy * w + fx] > 0.5) {
+                s = s + textureLoad(src, vec2<i32>(fx, fy), 0);
+                n = n + 1.0;
+            }
+        }
+    }
+    var v = vec4<f32>(0.0);
+    if (n > 0.0) {
+        v = s / n;
+    }
+    textureStore(out, vec2<i32>(cx, cy), v);
+}
+
+// membrane_core's coarse-to-fine step: unknown pixels take the nearest coarse value. The fine
+// image is copied into `out` first, so writing only the unknown pixels keeps the known ones.
+// Dispatched over the unknown pixels' bounds (`P.u[1]`); the coarse level is valid everywhere.
+@compute @workgroup_size(16, 16)
+fn imp_up(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let cw = i32(P.u[0].z);
+    let ch = i32(P.u[0].w);
+    let x = i32(P.u[1].x) + i32(gid.x);
+    let y = i32(P.u[1].y) + i32(gid.y);
+    if (x >= w || y >= h || data[i32(P.u[3].x) + y * w + x] > 0.5) {
+        return;
+    }
+    let cx = min(x / 2, cw - 1);
+    let cy = min(y / 2, ch - 1);
+    textureStore(out, vec2<i32>(x, y), textureLoad(src, vec2<i32>(cx, cy), 0));
+}
+
+// One relaxation of the unknown pixels (membrane_core's Jacobi sweep). `src` is the current
+// state and `aux` is unused. Dispatched over the unknown pixels' bounds (`P.u[1]`): known
+// pixels never move, so nothing outside them needs writing. Both `src` and `out` are valid
+// across the image (the fine level's, a copy of it, or the full downsample).
+@compute @workgroup_size(16, 16)
+fn imp_jacobi(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let x = i32(P.u[1].x) + i32(gid.x);
+    let y = i32(P.u[1].y) + i32(gid.y);
+    if (x >= w || y >= h) {
+        return;
+    }
+    var v = textureLoad(src, vec2<i32>(x, y), 0);
+    if (data[i32(P.u[3].x) + y * w + x] <= 0.5) {
+        var s = vec4<f32>(0.0);
+        var n = 0.0;
+        if (x > 0) {
+            s = s + textureLoad(src, vec2<i32>(x - 1, y), 0);
+            n = n + 1.0;
+        }
+        if (x + 1 < w) {
+            s = s + textureLoad(src, vec2<i32>(x + 1, y), 0);
+            n = n + 1.0;
+        }
+        if (y > 0) {
+            s = s + textureLoad(src, vec2<i32>(x, y - 1), 0);
+            n = n + 1.0;
+        }
+        if (y + 1 < h) {
+            s = s + textureLoad(src, vec2<i32>(x, y + 1), 0);
+            n = n + 1.0;
+        }
+        if (n > 0.0) {
+            v = s / n;
+        }
+    }
+    textureStore(out, vec2<i32>(x, y), v);
+}
+
+// ---------------------------------------------------------------- PatchMatch
+
+// PatchMatch's pyramid downsample (downsample): the level's image, hole pixels excluded from the
+// average, over `out`'s size.
+@compute @workgroup_size(16, 16)
+fn imp_pm_down(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let cw = i32(P.u[0].z);
+    let ch = i32(P.u[0].w);
+    let cx = i32(gid.x);
+    let cy = i32(gid.y);
+    if (cx >= cw || cy >= ch) {
+        return;
+    }
+    var s = vec4<f32>(0.0);
+    var n = 0.0;
+    for (var dy = 0; dy < 2; dy = dy + 1) {
+        for (var dx = 0; dx < 2; dx = dx + 1) {
+            let fx = min(2 * cx + dx, w - 1);
+            let fy = min(2 * cy + dy, h - 1);
+            if (!imp_hole(fy * w + fx)) {
+                s = s + textureLoad(src, vec2<i32>(fx, fy), 0);
+                n = n + 1.0;
+            }
+        }
+    }
+    var v = vec4<f32>(0.0);
+    if (n > 0.0) {
+        v = s / n;
+    }
+    textureStore(out, vec2<i32>(cx, cy), v);
+}
+
+// PatchMatch's coarse-to-fine step: bilinear upsample of the coarse image into this level's hole
+// (patch_fill's upsample loop). The fine image is copied into `out` first.
+@compute @workgroup_size(16, 16)
+fn imp_pm_up(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let cw = i32(P.u[0].z);
+    let ch = i32(P.u[0].w);
+    let x = i32(P.u[1].x) + i32(gid.x);
+    let y = i32(P.u[1].y) + i32(gid.y);
+    if (x >= w || y >= h || !imp_hole(y * w + x)) {
+        return;
+    }
+    let fx = clamp((f32(x) + 0.5) / 2.0 - 0.5, 0.0, f32(cw - 1));
+    let fy = clamp((f32(y) + 0.5) / 2.0 - 0.5, 0.0, f32(ch - 1));
+    let x0 = i32(floor(fx));
+    let y0 = i32(floor(fy));
+    let x1 = min(x0 + 1, cw - 1);
+    let y1 = min(y0 + 1, ch - 1);
+    let tx = fx - f32(x0);
+    let ty = fy - f32(y0);
+    let a = textureLoad(src, vec2<i32>(x0, y0), 0);
+    let b = textureLoad(src, vec2<i32>(x1, y0), 0);
+    let c = textureLoad(src, vec2<i32>(x0, y1), 0);
+    let d = textureLoad(src, vec2<i32>(x1, y1), 0);
+    let top = a + (b - a) * tx;
+    let bot = c + (d - c) * tx;
+    textureStore(out, vec2<i32>(x, y), top + (bot - top) * ty);
+}
+
+// The NNF's coarse-to-fine step: ×2 (patch_fill's nnf upsample).
+@compute @workgroup_size(16, 16)
+fn imp_nnf_up(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let cw = i32(P.u[0].z);
+    let ch = i32(P.u[0].w);
+    let x = i32(gid.x);
+    let y = i32(gid.y);
+    if (x >= w || y >= h) {
+        return;
+    }
+    let cx = min(x / 2, cw - 1);
+    let cy = min(y / 2, ch - 1);
+    let c = textureLoad(src, vec2<i32>(cx, cy), 0);
+    var o = vec4<f32>(0.0);
+    if (c.a > 0.5) {
+        o = vec4<f32>(c.r * 2.0 + f32(x % 2), c.g * 2.0 + f32(y % 2), 0.0, 1.0);
+    }
+    textureStore(out, vec2<i32>(x, y), o);
+}
+
+// A random valid source centre, by rejection (the CPU draws from its compacted source list; the
+// distribution over the valid centres is the same, the stream differs from the list's order).
+fn imp_any_source(seed: u32, i: u32, k: u32, w: i32, h: i32, r: i32) -> vec2<i32> {
+    for (var n = 0u; n < 32u; n = n + 1u) {
+        let a = imp_rnd(seed, i, k + n);
+        let x = i32(a % u32(w));
+        let y = i32((a >> 16u) % u32(h));
+        if (imp_valid_source(x, y, w, h, r)) {
+            return vec2<i32>(x, y);
+        }
+    }
+    return vec2<i32>(-1, -1);
+}
+
+// A round's initial field: the inherited source where it is still valid, a random one otherwise
+// (pm_level's initialise/validate step).
+@compute @workgroup_size(16, 16)
+fn imp_pm_init(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let r = i32(P.u[0].z);
+    let x = i32(P.u[1].x) + i32(gid.x);
+    let y = i32(P.u[1].y) + i32(gid.y);
+    if (x >= w || y >= h) {
+        return;
+    }
+    let i = y * w + x;
+    var o = vec4<f32>(0.0, 0.0, IMP_INF, 0.0);
+    if (imp_near(i)) {
+        let prev = textureLoad(src, vec2<i32>(x, y), 0);
+        var s = vec2<i32>(i32(prev.r), i32(prev.g));
+        if (prev.a <= 0.5 || !imp_valid_source(s.x, s.y, w, h, r)) {
+            // rnd(seed ^ e, i, 0), the round's random source (see imp_any_source).
+            s = imp_any_source(P.u[2].x ^ P.u[2].w, u32(i), 0u, w, h, r);
+        }
+        if (s.x >= 0) {
+            let d = imp_dist(x, y, s.x, s.y, w, h, r, IMP_INF);
+            o = vec4<f32>(f32(s.x), f32(s.y), d, 1.0);
+        }
+    }
+    textureStore(out, vec2<i32>(x, y), o);
+}
+
+// One propagation pass: every target pixel tries its four neighbours' sources shifted by the
+// step that would move that neighbour's match onto this pixel. The checkerboard parity in P.u[0].w
+// makes the pass race-free (a pixel of one parity only reads the other's).
+@compute @workgroup_size(16, 16)
+fn imp_pm_prop(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let r = i32(P.u[0].z);
+    let x = i32(P.u[1].x) + i32(gid.x);
+    let y = i32(P.u[1].y) + i32(gid.y);
+    if (x >= w || y >= h) {
+        return;
+    }
+    let cur = textureLoad(src, vec2<i32>(x, y), 0);
+    let i = y * w + x;
+    if (!imp_near(i) || ((x + y) & 1) != i32(P.u[0].w)) {
+        textureStore(out, vec2<i32>(x, y), cur);
+        return;
+    }
+    var best = vec2<i32>(i32(cur.r), i32(cur.g));
+    var bd = cur.b;
+    if (cur.a <= 0.5) {
+        bd = IMP_INF;
+    }
+    for (var k = 0; k < 4; k = k + 1) {
+        var nx = x;
+        var ny = y;
+        var sx = 0;
+        var sy = 0;
+        if (k == 0) {
+            nx = x - 1;
+            sx = 1;
+        } else if (k == 1) {
+            nx = x + 1;
+            sx = -1;
+        } else if (k == 2) {
+            ny = y - 1;
+            sy = 1;
+        } else {
+            ny = y + 1;
+            sy = -1;
+        }
+        if (nx < 0 || ny < 0 || nx >= w || ny >= h) {
+            continue;
+        }
+        let n = textureLoad(src, vec2<i32>(nx, ny), 0);
+        if (n.a <= 0.5) {
+            continue;
+        }
+        let c = vec2<i32>(i32(n.r) + sx, i32(n.g) + sy);
+        if (c.x == best.x && c.y == best.y) {
+            continue;
+        }
+        if (!imp_valid_source(c.x, c.y, w, h, r)) {
+            continue;
+        }
+        let d = imp_dist(x, y, c.x, c.y, w, h, r, bd);
+        if (d < bd) {
+            bd = d;
+            best = c;
+        }
+    }
+    // alpha = the entry is set: `bd` stays INF only if no match was found.
+    textureStore(out, vec2<i32>(x, y), vec4<f32>(f32(best.x), f32(best.y), bd, select(0.0, 1.0, bd < IMP_INF)));
+}
+
+// The random search (halving radius) and one random source anywhere, after propagation.
+@compute @workgroup_size(16, 16)
+fn imp_pm_rand(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let r = i32(P.u[0].z);
+    let x = i32(P.u[1].x) + i32(gid.x);
+    let y = i32(P.u[1].y) + i32(gid.y);
+    if (x >= w || y >= h) {
+        return;
+    }
+    let cur = textureLoad(src, vec2<i32>(x, y), 0);
+    let i = y * w + x;
+    if (!imp_near(i) || cur.a <= 0.5) {
+        textureStore(out, vec2<i32>(x, y), cur);
+        return;
+    }
+    var best = vec2<i32>(i32(cur.r), i32(cur.g));
+    var bd = cur.b;
+    let seed = P.u[2].x;
+    let it = P.u[2].y;
+    let em = P.u[2].w;
+    let iu = u32(i);
+    var rad = max(w, h);
+    var k = 0u;
+    while (rad >= 1) {
+        let a = imp_rnd(seed + it * 977u + em * 31u, iu, k);
+        let ox = i32(a % u32(2 * rad + 1)) - rad;
+        let oy = i32((a >> 16u) % u32(2 * rad + 1)) - rad;
+        let c = vec2<i32>(best.x + ox, best.y + oy);
+        if (imp_valid_source(c.x, c.y, w, h, r)) {
+            let d = imp_dist(x, y, c.x, c.y, w, h, r, bd);
+            if (d < bd) {
+                bd = d;
+                best = c;
+            }
+        }
+        rad = rad / 2;
+        k = k + 1u;
+    }
+    let any = imp_any_source(seed ^ 0xABCDu, iu, it + em * 64u, w, h, r);
+    if (any.x >= 0 && (any.x != best.x || any.y != best.y)) {
+        let d = imp_dist(x, y, any.x, any.y, w, h, r, bd);
+        if (d < bd) {
+            bd = d;
+            best = any;
+        }
+    }
+    textureStore(out, vec2<i32>(x, y), vec4<f32>(f32(best.x), f32(best.y), bd, select(0.0, 1.0, bd < IMP_INF)));
+}
+
+// Carry a round's field back into the nearest-neighbour field (only targets, as on the CPU).
+@compute @workgroup_size(16, 16)
+fn imp_nnf_store(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let x = i32(gid.x);
+    let y = i32(gid.y);
+    if (x >= w || y >= h) {
+        return;
+    }
+    let i = y * w + x;
+    var o = textureLoad(aux, vec2<i32>(x, y), 0);
+    if (imp_near(i)) {
+        let f = textureLoad(src, vec2<i32>(x, y), 0);
+        o = vec4<f32>(f.r, f.g, 0.0, 1.0);
+    }
+    textureStore(out, vec2<i32>(x, y), o);
+}
+
+// The vote: every hole pixel becomes the distance-weighted mean of the samples its overlapping
+// patches carry (pm_level's vote step). `aux` is the current image; the image is copied into
+// `out` first, so only hole pixels are written.
+@compute @workgroup_size(16, 16)
+fn imp_pm_vote(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let r = i32(P.u[0].z);
+    let x = i32(P.u[1].x) + i32(gid.x);
+    let y = i32(P.u[1].y) + i32(gid.y);
+    if (x >= w || y >= h) {
+        return;
+    }
+    let i = y * w + x;
+    if (!imp_hole(i)) {
+        return;
+    }
+    let sigma2 = P.f[0].x;
+    var acc = vec4<f32>(0.0);
+    var wsum = 0.0;
+    for (var dy = -r; dy <= r; dy = dy + 1) {
+        let py = y + dy;
+        if (py < 0 || py >= h) {
+            continue;
+        }
+        for (var dx = -r; dx <= r; dx = dx + 1) {
+            let px = x + dx;
+            if (px < 0 || px >= w) {
+                continue;
+            }
+            let f = textureLoad(src, vec2<i32>(px, py), 0);
+            if (f.a <= 0.5 || f.b >= IMP_INF) {
+                continue;
+            }
+            let s = vec2<i32>(i32(f.r) - dx, i32(f.g) - dy);
+            let wt = exp(-f.b / (2.0 * sigma2));
+            acc = acc + textureLoad(aux, vec2<i32>(s.x, s.y), 0) * wt;
+            wsum = wsum + wt;
+        }
+    }
+    if (wsum > 0.0) {
+        textureStore(out, vec2<i32>(x, y), acc / wsum);
+    } else {
+        textureStore(out, vec2<i32>(x, y), textureLoad(aux, vec2<i32>(x, y), 0));
+    }
+}
+
+// The workgroup's 1024 distance bins (`imp_hist`); the name is unique across the concatenated
+// shader sources because `var<workgroup>` is module scope.
+var<workgroup> imp_hist_bins: array<atomic<u32>, 1024>;
+// The largest finite distance this dispatch saw (as its f32 bit pattern: positive floats are
+// ordered by their bits). `sigma2` uses it to learn whether the frame needs a wider range.
+var<workgroup> imp_hist_max: atomic<u32>;
+
+// One workgroup's histogram of the targets' match distances, for the vote's sigma² (the CPU takes
+// the 75th percentile of the finite distances). The single workgroup accumulates the whole
+// population through the shared atomic bins above.
+//
+// `P.f[0].x` is the range the bins cover; 0 selects the SDR range 0…4, where an image whose
+// channels are all ≤ 1 lives: each channel's squared difference is ≤ 4, so the mean SSD is too.
+// Frames with values above 1 (32-bpc HDR) have larger distances, so `sigma2` re-dispatches with
+// the observed maximum as the range instead of clamping them to bin 1023, which would report a
+// sigma² of ≈ 4 for a distance of 1600 and round the vote's weights to zero. The observed maximum
+// rides in bin 0's unused `.g` channel on every dispatch (its `.x` stays the count).
+@compute @workgroup_size(256, 1, 1)
+fn imp_hist(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let x0 = i32(P.u[1].x);
+    let y0 = i32(P.u[1].y);
+    let bw = i32(P.u[1].z);
+    let bh = i32(P.u[1].w);
+    let t = gid.x;
+    let range = select(4.0, P.f[0].x, P.f[0].x > 0.0);
+    let scale = 1024.0 / range;
+    if (t == 0u) {
+        atomicStore(&imp_hist_max, 0u);
+    }
+    // One workgroup, so the whole population is counted through shared atomics: a per-thread
+    // strided pass would keep only the pixels whose bin selected that thread (a biased ~1/256
+    // subsample, correlated with the pixel index).
+    for (var i = t; i < 1024u; i = i + 256u) {
+        atomicStore(&imp_hist_bins[i], 0u);
+    }
+    workgroupBarrier();
+    let total = u32(max(bw, 0) * max(bh, 0));
+    for (var idx = t; idx < total; idx = idx + 256u) {
+        let x = x0 + i32(idx) % bw;
+        let y = y0 + i32(idx) / bw;
+        let i = y * w + x;
+        if (!imp_near(i)) {
+            continue;
+        }
+        let f = textureLoad(src, vec2<i32>(x, y), 0);
+        if (f.a <= 0.5 || f.b >= IMP_INF) {
+            continue;
+        }
+        atomicMax(&imp_hist_max, bitcast<u32>(f.b));
+        // Clamp in f32 before the conversion: a distance far outside the range (a mismatched patch
+        // in an HDR frame) would overflow the i32 conversion, which is indeterminate in WGSL.
+        let b = u32(clamp(f.b * scale, 0.0, 1023.0));
+        atomicAdd(&imp_hist_bins[b], 1u);
+    }
+    workgroupBarrier();
+    // 256 threads × 4 bins cover the histogram.
+    for (var j = 0u; j < 4u; j = j + 1u) {
+        let bin = t * 4u + j;
+        textureStore(out, vec2<i32>(i32(bin), 0), vec4<f32>(f32(atomicLoad(&imp_hist_bins[bin])), 0.0, 0.0, 0.0));
+    }
+    // Thread 0 is bin 0's only writer, so this store follows the one above. `.g` carries the
+    // observed maximum: host readbacks keep the row 1024 texels wide (a 16 KB, 256-aligned copy).
+    if (t == 0u) {
+        textureStore(out, vec2<i32>(0, 0), vec4<f32>(f32(atomicLoad(&imp_hist_bins[0])), bitcast<f32>(atomicLoad(&imp_hist_max)), 0.0, 0.0));
+    }
+}
+
+// ---------------------------------------------------------------- block flow (flow completion)
+
+// The block matcher of effectcraft_raster::flow::block_flow(a, b, 8, 6), which the fill uses to
+// compute the motion between neighbouring membrane-filled frames before completing it inside the
+// holes (`FillBackend::flow`). `src` packs the two frames' luma planes at the current level:
+// .x is frame A's, .y is frame B's (`effectcraft_raster::flow::Plane::luma`). The kernels mirror
+// the CPU operation for operation, so the field matches the CPU's (crates/gpu/src/tests/inpaint.rs).
+
+// Plane::half of the packed pair: the clamped 2×2 average of each channel.
+@compute @workgroup_size(16, 16)
+fn imp_fl_half(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let cx = i32(gid.x);
+    let cy = i32(gid.y);
+    if (cx >= i32(P.u[0].z) || cy >= i32(P.u[0].w)) {
+        return;
+    }
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    var sx = 0.0;
+    var sy = 0.0;
+    for (var dy = 0; dy < 2; dy = dy + 1) {
+        for (var dx = 0; dx < 2; dx = dx + 1) {
+            let v = textureLoad(src, vec2<i32>(min(2 * cx + dx, w - 1), min(2 * cy + dy, h - 1)), 0);
+            sx = sx + v.x;
+            sy = sy + v.y;
+        }
+    }
+    textureStore(out, vec2<i32>(cx, cy), vec4<f32>(sx * 0.25, sy * 0.25, 0.0, 0.0));
+}
+
+// Plane::get: the packed luma at (x, y), clamped to the plane.
+fn imp_fl_tex(x: i32, y: i32, w: i32, h: i32) -> vec2<f32> {
+    let v = textureLoad(src, clamp(vec2<i32>(x, y), vec2<i32>(0, 0), vec2<i32>(w - 1, h - 1)), 0);
+    return vec2<f32>(v.x, v.y);
+}
+
+// The block match's SAD with A's block already loaded (`la`, row-major): B's luma is the packed
+// .y. The sub-pixel kernel's five fits share the window's block, so it is read once; the search
+// hoists its own 4×4/8×8 block the same way. The loops' bounds are literals so the backend can
+// unroll them and keep the loads in flight — a runtime-bound loop runs an order of magnitude
+// slower.
+fn imp_fl_sad8h(bx: i32, by: i32, dx: i32, dy: i32, w: i32, h: i32, la: array<f32, 64>) -> f32 {
+    var s = 0.0;
+    for (var yy = 0; yy < 8; yy = yy + 1) {
+        for (var xx = 0; xx < 8; xx = xx + 1) {
+            s = s + abs(la[yy * 8 + xx] - imp_fl_tex(bx + xx + dx, by + yy + dy, w, h).y);
+        }
+    }
+    return s;
+}
+
+// block_flow_window's coarse-to-fine search at one level: one integer vector per tile, seeded by
+// the previous (coarser) level's vector at the tile and its four neighbours. `aux` is that field
+// (`P.u[2].x` = 0 at the coarsest level, where the only seed is the zero vector).
+@compute @workgroup_size(16, 16)
+fn imp_fl_search(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let cols = i32(P.u[0].z);
+    let rows = i32(P.u[0].w);
+    let c = i32(gid.x);
+    let r = i32(gid.y);
+    if (c >= cols || r >= rows) {
+        return;
+    }
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let bs = i32(P.u[1].x);
+    let k = i32(P.u[1].z);
+    let rad = i32(P.u[1].w);
+    let lbs = max(i32(P.u[1].y) / k, 4);
+    let bx = c * bs / k;
+    let by = r * bs / k;
+    // The zero vector plus up to five neighbour seeds (the tile and its four neighbours).
+    var seeds: array<vec2<i32>, 6>;
+    var ns = 0;
+    seeds[ns] = vec2<i32>(0, 0);
+    ns = ns + 1;
+    if (P.u[2].x != 0u) {
+        var nbr = array<vec2<i32>, 5>(vec2<i32>(0, 0), vec2<i32>(-1, 0), vec2<i32>(1, 0), vec2<i32>(0, -1), vec2<i32>(0, 1));
+        for (var j = 0; j < 5; j = j + 1) {
+            let nc = c + nbr[j].x;
+            let nr = r + nbr[j].y;
+            if (nc >= 0 && nr >= 0 && nc < cols && nr < rows) {
+                let pv = textureLoad(aux, vec2<i32>(nc, nr), 0);
+                seeds[ns] = vec2<i32>(i32(round(pv.x)) * 2, i32(round(pv.y)) * 2);
+                ns = ns + 1;
+            }
+        }
+    }
+    // sort_unstable + dedup: distinct seeds in lexicographic (x, then y) order.
+    for (var i = 1; i < ns; i = i + 1) {
+        let key = seeds[i];
+        var j = i - 1;
+        while (j >= 0 && (seeds[j].x > key.x || (seeds[j].x == key.x && seeds[j].y > key.y))) {
+            seeds[j + 1] = seeds[j];
+            j = j - 1;
+        }
+        seeds[j + 1] = key;
+    }
+    var m = 0;
+    for (var i = 0; i < ns; i = i + 1) {
+        if (m == 0 || seeds[i].x != seeds[m - 1].x || seeds[i].y != seeds[m - 1].y) {
+            seeds[m] = seeds[i];
+            m = m + 1;
+        }
+    }
+    ns = m;
+    // A's block is the same for every candidate, so load it once (the unrolled loops keep it in
+    // registers) and the search only reads B's shifted block.
+    var la4: array<f32, 16>;
+    var la8: array<f32, 64>;
+    if (lbs >= 8) {
+        for (var yy = 0; yy < 8; yy = yy + 1) {
+            for (var xx = 0; xx < 8; xx = xx + 1) {
+                la8[yy * 8 + xx] = imp_fl_tex(bx + xx, by + yy, w, h).x;
+            }
+        }
+    } else {
+        for (var yy = 0; yy < 4; yy = yy + 1) {
+            for (var xx = 0; xx < 4; xx = xx + 1) {
+                la4[yy * 4 + xx] = imp_fl_tex(bx + xx, by + yy, w, h).x;
+            }
+        }
+    }
+    var best = vec2<i32>(0, 0);
+    var best_e = 1e30;
+    for (var si = 0; si < ns; si = si + 1) {
+        let s = seeds[si];
+        for (var dy = -rad; dy <= rad; dy = dy + 1) {
+            for (var dx = -rad; dx <= rad; dx = dx + 1) {
+                let d0 = s.x + dx;
+                let d1 = s.y + dy;
+                var acc = 0.0;
+                if (lbs >= 8) {
+                    for (var yy = 0; yy < 8; yy = yy + 1) {
+                        for (var xx = 0; xx < 8; xx = xx + 1) {
+                            acc = acc + abs(la8[yy * 8 + xx] - imp_fl_tex(bx + xx + d0, by + yy + d1, w, h).y);
+                        }
+                    }
+                } else {
+                    for (var yy = 0; yy < 4; yy = yy + 1) {
+                        for (var xx = 0; xx < 4; xx = xx + 1) {
+                            acc = acc + abs(la4[yy * 4 + xx] - imp_fl_tex(bx + xx + d0, by + yy + d1, w, h).y);
+                        }
+                    }
+                }
+                let e = acc + 1e-4 * f32(abs(d0) + abs(d1));
+                if (e < best_e) {
+                    best_e = e;
+                    best = vec2<i32>(d0, d1);
+                }
+            }
+        }
+    }
+    textureStore(out, vec2<i32>(c, r), vec4<f32>(f32(best.x), f32(best.y), 0.0, 0.0));
+}
+
+// The parabola fit of the SAD around the best vector (block_flow_window's sub-pixel step).
+fn imp_fl_fit(em: f32, ep: f32, e0: f32) -> f32 {
+    let den = em - 2.0 * e0 + ep;
+    if (den > 1e-6) {
+        return clamp(0.5 * (em - ep) / den, -0.5, 0.5);
+    }
+    return 0.0;
+}
+
+// Sub-pixel refinement at full resolution: `aux` is the integer field of level 0, `src` the
+// level-0 packed luma (`P.u[1]` = (tile, window, 0, 0)).
+@compute @workgroup_size(16, 16)
+fn imp_fl_subpix(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let cols = i32(P.u[0].z);
+    let rows = i32(P.u[0].w);
+    let c = i32(gid.x);
+    let r = i32(gid.y);
+    if (c >= cols || r >= rows) {
+        return;
+    }
+    let w = i32(P.u[0].x);
+    let h = i32(P.u[0].y);
+    let bs = i32(P.u[1].x);
+    let win = i32(P.u[1].y);
+    let bx = c * bs;
+    let by = r * bs;
+    let f = textureLoad(aux, vec2<i32>(c, r), 0);
+    let d = vec2<i32>(i32(round(f.x)), i32(round(f.y)));
+    // The window's A block is the same for the five SADs: load it once (the loop is unrolled).
+    var la: array<f32, 64>;
+    for (var yy = 0; yy < 8; yy = yy + 1) {
+        for (var xx = 0; xx < 8; xx = xx + 1) {
+            la[yy * 8 + xx] = imp_fl_tex(bx + xx, by + yy, w, h).x;
+        }
+    }
+    let e0 = imp_fl_sad8h(bx, by, d.x, d.y, w, h, la);
+    let sx = imp_fl_fit(imp_fl_sad8h(bx, by, d.x - 1, d.y, w, h, la), imp_fl_sad8h(bx, by, d.x + 1, d.y, w, h, la), e0);
+    let sy = imp_fl_fit(imp_fl_sad8h(bx, by, d.x, d.y - 1, w, h, la), imp_fl_sad8h(bx, by, d.x, d.y + 1, w, h, la), e0);
+    textureStore(out, vec2<i32>(c, r), vec4<f32>(f32(d.x) + sx, f32(d.y) + sy, 0.0, 0.0));
+}
+
+// The 3×3 component-wise vector median against outliers (each component sorted on its own).
+@compute @workgroup_size(16, 16)
+fn imp_fl_median(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let cols = i32(P.u[0].z);
+    let rows = i32(P.u[0].w);
+    let c = i32(gid.x);
+    let r = i32(gid.y);
+    if (c >= cols || r >= rows) {
+        return;
+    }
+    var xs: array<f32, 9>;
+    var ys: array<f32, 9>;
+    var n = 0;
+    for (var dr = -1; dr <= 1; dr = dr + 1) {
+        for (var dc = -1; dc <= 1; dc = dc + 1) {
+            let nc = c + dc;
+            let nr = r + dr;
+            if (nc >= 0 && nr >= 0 && nc < cols && nr < rows) {
+                let v = textureLoad(src, vec2<i32>(nc, nr), 0);
+                xs[n] = v.x;
+                ys[n] = v.y;
+                n = n + 1;
+            }
+        }
+    }
+    for (var i = 1; i < n; i = i + 1) {
+        let kx = xs[i];
+        var j = i - 1;
+        while (j >= 0 && xs[j] > kx) {
+            xs[j + 1] = xs[j];
+            j = j - 1;
+        }
+        xs[j + 1] = kx;
+        let ky = ys[i];
+        j = i - 1;
+        while (j >= 0 && ys[j] > ky) {
+            ys[j + 1] = ys[j];
+            j = j - 1;
+        }
+        ys[j + 1] = ky;
+    }
+    textureStore(out, vec2<i32>(c, r), vec4<f32>(xs[n / 2], ys[n / 2], 0.0, 0.0));
+}

crates/gpu/src/shaders/particles.wgsl

@@ -53,12 +53,12 @@ fn ps_norm3(v: vec3<f32>) -> vec3<f32> {
 fn ps_rand_dir(i: u32, k: u32, flat: bool) -> vec3<f32> {
     if (flat) {
         let a = ps_h(i, k) * PS_TAU;
-        return vec3<f32>(cos(a), sin(a), 0.0);
+        return vec3<f32>(cos_p(a), sin_p(a), 0.0);
     }
     let z = ps_hs(i, k);
     let a = ps_h(i, k + 1u) * PS_TAU;
     let r = sqrt(max(1.0 - z * z, 0.0));
-    return vec3<f32>(r * cos(a), r * sin(a), z);
+    return vec3<f32>(r * cos_p(a), r * sin_p(a), z);
 }
 
 struct Part {
@@ -189,7 +189,7 @@ fn ps_spawn_cannon(id: u32) -> Part {
     let vel = S.f[1].y;
     let vel_spread = S.f[1].z;
     let a = radians(angle + ps_hs(id, 1u) * spread * 0.5);
-    let d = vec2<f32>(sin(a), -cos(a));
+    let d = vec2<f32>(sin_p(a), -cos_p(a));
     let speed = max(vel + ps_hs(id, 3u) * vel_spread, 0.0);
     var off: vec2<f32>;
     if (r >= 0.0) {
@@ -197,7 +197,7 @@ fn ps_spawn_cannon(id: u32) -> Part {
     } else {
         let aa = ps_h(id, 5u) * PS_TAU;
         let rr = sqrt(ps_h(id, 6u)) * -r;
-        off = vec2<f32>(cos(aa) * rr, sin(aa) * rr);
+        off = vec2<f32>(cos_p(aa) * rr, sin_p(aa) * rr);
     }
     let v = d * speed;
     let dt = S.f[2].z;

crates/gpu/src/tests.rs

@@ -52,10 +52,108 @@ mod fx_vr;
 mod fx_warp;
 #[cfg(all(test, not(target_arch = "wasm32")))]
 mod init;
+mod inpaint;
 mod particles;
 mod three_d;
 mod wgsl_grammar;
 
+pub(crate) const MATH_PROBE_SHADER: &str = r#"
+@compute @workgroup_size(16, 16)
+fn math_probe(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let p = vec2<i32>(gid.xy);
+    if (p.x >= out_dims().x || p.y >= out_dims().y) { return; }
+    let v = textureLoad(src, p, 0).x;
+    if (P.u[0].x == 1u) {
+        let a = textureLoad(src, p, 0).w;
+        let q = unit_div(v, a);
+        textureStore(out, p, vec4<f32>(q, fxs_posterize_bin(q, P.u[0].y), 0.0, 1.0));
+        return;
+    }
+    textureStore(out, p, vec4<f32>(sin_p(v), cos_p(v), 0.0, 1.0));
+}
+"#;
+
+/// Direction errors accumulate into visible particle offsets and spherical
+/// remapping errors. Check the numerical primitive before rendering a scene.
+#[test]
+fn angular_math_preserves_particle_and_vr_directions() {
+    let Some(g) = gpu() else { return };
+    for extent in [std::f32::consts::TAU, 100.0, 8192.0] {
+        let mut img = Image::new(4097, 1);
+        for (i, px) in img.data.iter_mut().enumerate() {
+            px[0] = (i as f32 - 2048.0) * extent / 2048.0;
+        }
+        let input = g.ctx.upload_image(&img).unwrap();
+        let mut e = crate::context::Enc::new(&g.ctx);
+        let out = e.scratch(img.width, img.height);
+        e.pixels("math_probe", &Default::default(), &input, None, &out, None);
+        let result = e.download(&out).unwrap();
+        for (i, (v, got)) in img.data.iter().zip(result.data).enumerate() {
+            let want = [(v[0] as f64).sin() as f32, (v[0] as f64).cos() as f32];
+            for c in 0..2 {
+                assert!((want[c] - got[c]).abs() <= 2e-7, "sample {i}, angle {}: CPU {want:?}, GPU {got:?}", v[0]);
+            }
+        }
+    }
+}
+
+#[test]
+fn alpha_division_and_posterization_match_f32_rounding() {
+    let Some(g) = gpu() else { return };
+    let mut values = vec![];
+    for levels in [1u32, 3, 8, 254] {
+        values.clear();
+        for a in [1.0f32, 0.1, 0.3, 1.0 / 255.0, 6.0 / 255.0] {
+            for i in 0..levels {
+                let half = (i as f32 + 0.5) / levels as f32;
+                for v in [0.0, 1e-6, half.next_down(), half, half.next_up(), 1.0] {
+                    values.push([v * a, 0.0, 0.0, a]);
+                }
+            }
+        }
+        let img = Image { width: values.len() as u32, height: 1, data: values.clone() };
+        let input = g.ctx.upload_image(&img).unwrap();
+        let mut e = crate::context::Enc::new(&g.ctx);
+        let out = e.scratch(img.width, 1);
+        let mut params = crate::context::Params::default();
+        params.u[0] = [1, levels, 0, 0];
+        e.pixels("math_probe", &params, &input, None, &out, None);
+        let result = e.download(&out).unwrap();
+        for (i, (v, got)) in values.iter().zip(result.data).enumerate() {
+            let q = (v[0] / v[3]).clamp(0.0, 1.0);
+            assert_eq!(q.to_bits(), got[0].to_bits(), "levels {levels}, pixel {i}, input {v:?}");
+            assert_eq!((q * levels as f32).round(), got[1], "levels {levels}, pixel {i}, input {v:?}");
+        }
+    }
+}
+
+/// Prove that the numerical fixes run on the device, without per-effect image
+/// downloads or CPU fallback. Compiler-limited DX12 devices retain their tested
+/// fallback contract; these kernels must compile on the Vulkan test adapters.
+#[test]
+fn numerical_effects_remain_resident_on_vulkan() {
+    let Some(g) = gpu() else { return };
+    if !g.ctx.name.ends_with("(Vulkan)") {
+        return;
+    }
+    for entry in ["fxs_point", "fxt_point", "ftr_grid", "fp2_hextile"] {
+        assert!(!g.ctx.unavailable_kernels().iter().any(|(name, _)| *name == entry), "{entry}: {:?}", g.ctx.unavailable_kernels());
+    }
+    for (id, vals) in [
+        ("ec.stylize.posterize", vec![("level", n(4.0))]),
+        ("ec.color.psarbitrarymap", vec![("phase", n(40.0))]),
+        ("ec.transition.ccgridwipe", vec![("completion", n(40.0))]),
+        ("ec.stylize.cchextile", vec![("radius", n(9.3))]),
+    ] {
+        let before = g.ctx.transfer_stats();
+        effect_direct(id, &vals, false);
+        let after = g.ctx.transfer_stats();
+        // Full- and half-resolution runs each upload once and download once.
+        assert_eq!(after.uploads - before.uploads, 2, "{id}: uploads");
+        assert_eq!(after.readbacks - before.readbacks, 2, "{id}: readbacks");
+    }
+}
+
 /// 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
@@ -784,7 +882,7 @@ fn display_frames_are_mipmapped() {
         view_formats: &[],
     });
     let mut e = crate::context::Enc::new(&g.ctx);
-    e.encoder().copy_texture_to_texture(
+    e.record(&[]).copy_texture_to_texture(
         wgpu::TexelCopyTextureInfo { texture: &f.texture, mip_level: 1, origin: wgpu::Origin3d::ZERO, aspect: wgpu::TextureAspect::All },
         wgpu::TexelCopyTextureInfo { texture: &tex, mip_level: 0, origin: wgpu::Origin3d::ZERO, aspect: wgpu::TextureAspect::All },
         wgpu::Extent3d { width: w1, height: h1, depth_or_array_layers: 1 },
@@ -920,7 +1018,7 @@ fn a_failed_frame_returns_no_texture_to_the_pool() {
     let rejected = rejected_texture(&g);
     let earlier = g.within_memory(|| crate::context::Enc::new(&g.ctx).image(48, 16)).expect("an earlier frame");
     let mut other = crate::context::Enc::new(&g.ctx);
-    other.encoder();
+    other.record(&[]);
     g.within_memory(|| (crate::context::Enc::new(&g.ctx).image(64, 32), rejected.create_view(&Default::default()))).expect_err("the frame fails");
     drop(earlier);
     drop(other);
@@ -1315,3 +1413,485 @@ fn owned_devices_favour_memory_usage() {
     assert_eq!(d.label, Some("effectcraft gpu"));
     assert_eq!(d.required_limits.max_texture_dimension_2d, wgpu::Limits::default().max_texture_dimension_2d);
 }
+
+/// A GPU context of its own with a small pool budget, so other tests (sharing [`gpu`]) neither
+/// see the budget nor add to what the pool keeps.
+fn private_gpu(budget: usize) -> Option<Gpu> {
+    gpu()?;
+    let g = Gpu::headless()?;
+    g.context().set_pool_budget(budget);
+    Some(g)
+}
+
+/// Three effect calls while a frame encoder is open with recorded work reuse one set of
+/// textures: lent calls give them back at once, without waiting for that frame to submit.
+#[test]
+fn three_effect_calls_in_a_frame_reuse_one_set_of_textures() {
+    let Some(g) = private_gpu(8 << 20) else { return };
+    let ctx = g.context();
+    let mut frame = crate::context::Enc::new(ctx);
+    // Open the frame encoder with real work so the pool has an open id (an empty Enc::new does not).
+    let src = frame.image(64, 64);
+    let dst = frame.image(64, 64);
+    frame.copy_into(&src, &dst, 0, 0);
+    let (hits0, misses0) = ctx.pool_stats();
+    for _ in 0..3 {
+        let mut e = crate::context::Enc::lent(ctx);
+        let a = e.image(256, 256);
+        let b = e.image(256, 256);
+        drop((a, b));
+        drop(e);
+    }
+    let (hits, misses) = (ctx.pool_stats().0 - hits0, ctx.pool_stats().1 - misses0);
+    assert_eq!(misses, 2, "only the first effect call allocates: {hits} hits, {misses} misses");
+    assert_eq!(hits, 4, "calls 2 and 3 reuse both textures: {hits} hits, {misses} misses");
+    frame.submit();
+    drop((src, dst, frame));
+}
+
+/// Review of #1523: after the first frame fills the free list, later identical frames hit
+/// instead of allocating again.
+#[test]
+fn a_repeated_frame_reuses_after_the_first() {
+    let Some(g) = private_gpu(8 << 20) else { return };
+    let ctx = g.context();
+    let frame = || {
+        let mut e = crate::context::Enc::lent(ctx);
+        drop((e.image(256, 256), e.image(128, 64)));
+        drop(e);
+        ctx.pool_frame_end();
+    };
+    let (hits0, misses0) = ctx.pool_stats();
+    frame();
+    let (hits1, misses1) = (ctx.pool_stats().0 - hits0, ctx.pool_stats().1 - misses0);
+    assert_eq!(misses1, 2, "first frame allocates both sizes: {hits1} hits, {misses1} misses");
+    let (hits_a, misses_a) = ctx.pool_stats();
+    for _ in 0..3 {
+        frame();
+    }
+    let (hits, misses) = (ctx.pool_stats().0 - hits_a, ctx.pool_stats().1 - misses_a);
+    assert_eq!(misses, 0, "frames 2..4 allocate nothing: {hits} hits, {misses} misses");
+    assert!(hits >= 6, "frames 2..4 reuse both sizes: {hits} hits");
+}
+
+/// Review of #1523: every idle effect-call owner kept up to the pool budget of its own,
+/// outside the shared accounting, and there are as many as calls ever ran at once (16 frame
+/// threads: ~12 GB after a 4K render). Finished calls share the free list's one budget, and
+/// the end of a render drops it.
+#[test]
+fn idle_effect_calls_share_one_pool_budget() {
+    let budget = 8 << 20;
+    let Some(g) = private_gpu(budget) else { return };
+    let ctx = g.context();
+    // 16 effect calls in flight at once, each with three 1 MB working textures.
+    let mut encs: Vec<_> = (0..16).map(|_| crate::context::Enc::lent(ctx)).collect();
+    for e in &mut encs {
+        let imgs: Vec<_> = (0..3).map(|_| e.image(256, 256)).collect();
+        drop(imgs);
+    }
+    drop(encs);
+    let kept = ctx.pool_retained_bytes();
+    assert!(kept <= budget, "idle effect calls keep {} MB, budget {} MB", kept >> 20, budget >> 20);
+    // A later call reuses them instead of allocating.
+    let (hits, _) = ctx.pool_stats();
+    drop(crate::context::Enc::lent(ctx).image(256, 256));
+    assert_eq!(ctx.pool_stats().0, hits + 1, "a later effect call reuses a kept texture");
+    effectcraft_render::Accelerator::render_end(&g);
+    assert_eq!(ctx.pool_retained_bytes(), 0, "the end of a render drops the kept textures");
+}
+
+/// Review of #1523: a full free list rejected every new texture instead of evicting, so once
+/// sizes that no longer occur (per-frame blur bounds, padded layer sizes) filled it, a steady
+/// size never got a hit again: ~0.75 GB of dead weight plus an allocation per working texture.
+#[test]
+fn stale_sizes_do_not_starve_a_steady_size() {
+    let budget = 8 << 20;
+    let Some(g) = private_gpu(budget) else { return };
+    let ctx = g.context();
+    for k in 0..16 {
+        let mut e = crate::context::Enc::lent(ctx);
+        drop((e.image(256 + k, 256), e.image(256 + k, 256)));
+        drop(e);
+        ctx.pool_frame_end();
+    }
+    let (hits0, misses0) = ctx.pool_stats();
+    for _ in 0..10 {
+        let mut e = crate::context::Enc::lent(ctx);
+        drop((e.image(300, 300), e.image(300, 300)));
+        drop(e);
+        ctx.pool_frame_end();
+    }
+    let (hits, misses) = (ctx.pool_stats().0 - hits0, ctx.pool_stats().1 - misses0);
+    assert!(ctx.pool_retained_bytes() <= budget);
+    assert!(ctx.pool_evictions() > 0, "the stale sizes were evicted");
+    assert!(hits >= 18 && misses <= 2, "steady 300x300 calls after stale sizes: {hits} hits, {misses} misses");
+}
+
+/// Review of #1523: a `pool_frame_end` wait that clears none of the evicted bytes (timeout, or
+/// callbacks that never run) must not re-wait on every later frame while `dying` stays put.
+#[test]
+fn a_wait_without_progress_is_not_repeated() {
+    let Some(g) = private_gpu(8 << 20) else { return };
+    let ctx = g.context();
+    let phantom = 1 << 20;
+    ctx.pool_charge_dying(phantom as isize);
+    ctx.pool_frame_end();
+    assert_eq!(ctx.pool_waits(), 1, "first frame with dying bytes waits once");
+    for _ in 0..3 {
+        ctx.pool_frame_end();
+    }
+    assert_eq!(ctx.pool_waits(), 1, "no progress: later frames must not wait again");
+    ctx.pool_charge_dying(-(phantom as isize));
+    ctx.pool_frame_end();
+    assert_eq!(ctx.pool_waits(), 1, "dying cleared: poll only, stall cleared");
+    ctx.pool_charge_dying(phantom as isize);
+    ctx.pool_frame_end();
+    assert_eq!(ctx.pool_waits(), 2, "a new stall waits once");
+    // Partial progress: more dying while stalled → still no wait; drop below stalled → wait.
+    let extra = phantom / 2;
+    ctx.pool_charge_dying(extra as isize);
+    ctx.pool_frame_end();
+    assert_eq!(ctx.pool_waits(), 2, "dying above stalled: poll only");
+    ctx.pool_charge_dying(-(extra as isize + 1));
+    ctx.pool_frame_end();
+    assert_eq!(ctx.pool_waits(), 3, "dying dropped below stalled: wait again");
+    ctx.pool_charge_dying(-((phantom - 1) as isize));
+}
+
+/// RTX 3090 4K runs of #1523: evicting the least recently freed texture to fit every newcomer
+/// thrashes when a frame frees more than the budget in a repeating order (blur/glow
+/// intermediates): each free evicts the size the next call takes first, so every take
+/// missed (d_blurglow misses back at the old level, and evict + reallocate raised the peak).
+/// Sizes still in use stay; a kept subset of the cycle must keep hitting.
+#[test]
+fn repeating_sizes_over_the_budget_keep_hitting() {
+    let budget = 8 << 20;
+    let Some(g) = private_gpu(budget) else { return };
+    let ctx = g.context();
+    // Five ~2.1 MB sizes a frame (10.5 MB > 8 MB), one effect call each, same order every frame.
+    let sizes: Vec<(u32, u32)> = (0..5).map(|k| (256 + 8 * k, 512)).collect();
+    let frame = || {
+        for &(w, h) in &sizes {
+            drop(crate::context::Enc::lent(ctx).image(w, h));
+        }
+        ctx.pool_frame_end();
+    };
+    for _ in 0..3 {
+        frame();
+    }
+    let (hits0, misses0) = ctx.pool_stats();
+    let evicted0 = ctx.pool_evictions();
+    for _ in 0..10 {
+        frame();
+    }
+    let (hits, misses) = (ctx.pool_stats().0 - hits0, ctx.pool_stats().1 - misses0);
+    assert!(ctx.pool_retained_bytes() <= budget);
+    assert!(hits >= 30, "10 frames of a 5-size cycle over the budget: {hits} hits, {misses} misses");
+    assert_eq!(ctx.pool_evictions(), evicted0, "sizes in use were evicted");
+}
+
+/// Review of #1523: a lent call's image read by another, still open encoder must not go back to
+/// the lent call for reuse when its lease drops first: that encoder's read is not submitted yet.
+#[test]
+fn image_read_by_an_open_encoder_is_not_reused_by_its_lent_call() {
+    let Some(g) = private_gpu(8 << 20) else { return };
+    let ctx = g.context();
+    let mut frame = crate::context::Enc::new(ctx);
+    let mut lent = crate::context::Enc::lent(ctx);
+    let img = lent.image(64, 64);
+    let dst = frame.image(64, 64);
+    frame.copy_into(&img, &dst, 0, 0);
+    drop(img);
+    let (hits, _) = ctx.pool_stats();
+    drop(lent.image(64, 64));
+    assert_eq!(ctx.pool_stats().0, hits, "the lent call reused a texture an open encoder still reads");
+    drop(lent);
+    let (hits, _) = ctx.pool_stats();
+    drop(crate::context::Enc::lent(ctx).image(64, 64));
+    // The other 64x64 (the lent call's second image) is free; the read one still waits.
+    let (hits2, _) = ctx.pool_stats();
+    assert_eq!(hits2, hits + 1);
+    let (h, _) = ctx.pool_stats();
+    let mut e = crate::context::Enc::lent(ctx);
+    let a = e.image(64, 64);
+    let b = e.image(64, 64);
+    assert_eq!(ctx.pool_stats().0, h + 1, "the read texture waits for the open encoder");
+    drop((a, b, e));
+    frame.submit();
+    drop(dst);
+}
+
+/// The same through [`Gpu::effects`] from 16 threads at once, as an export's frame threads.
+#[test]
+fn effect_calls_on_16_threads_keep_one_pool_budget() {
+    if on_gl() {
+        return; // one context lock per adapter on wgpu's GL backend (see hold_gpu_lock)
+    }
+    let budget = 8 << 20;
+    let Some(g) = private_gpu(budget) else { return };
+    let g = &g;
+    let ctx = g.context();
+    let spec = effectcraft_effects::find("ec.blur.gaussian").unwrap();
+    let size = [240.0, 240.0];
+    let mut params =
+        effectcraft_effects::Params { values: spec.params.iter().map(|ps| (ps.id.to_string(), effectcraft_effects::default_value(ps, size))).collect() };
+    params.values.insert("blurriness".into(), Value::Scalar(6.0));
+    let buf = effectcraft_effects::Buf { img: pattern(3, 240, 240), offset: [0.0, 0.0], scale: 1.0 };
+    let go = std::sync::Barrier::new(16);
+    std::thread::scope(|s| {
+        for i in 0..16 {
+            let (params, buf, go) = (&params, &buf, &go);
+            s.spawn(move || {
+                go.wait();
+                for _ in 0..3 {
+                    let fx = effectcraft_effects::EffectCtx { params, time: 0.0, layer_size: size, seed: i, adjustment: false, env: Default::default() };
+                    effectcraft_render::Accelerator::effects(g, &[effectcraft_render::FxStep { spec, ctx: fx }], buf, None).expect("the GPU runs the effect");
+                }
+            });
+        }
+    });
+    let kept = ctx.pool_retained_bytes();
+    assert!(kept <= budget, "16 threads of effect calls keep {} MB, budget {} MB", kept >> 20, budget >> 20);
+    effectcraft_render::Accelerator::render_end(g);
+    assert_eq!(ctx.pool_retained_bytes(), 0, "the end of a render drops the kept textures");
+}
+
+/// Sceptic r2 of #1523: only the export pipeline ended frames, so off export (preview, remote
+/// frames) the pool's clock stood still, no size ever went stale and a full free list dropped
+/// every newcomer again. Frames rendered through [`effectcraft_render::passes::comp_frame`]
+/// (no explicit `pool_frame_end`) must age the sizes they did not ask for.
+#[test]
+fn frames_outside_the_export_pipeline_age_the_pool() {
+    let budget = 8 << 20;
+    let Some(g) = private_gpu(budget) else { return };
+    let ctx = g.context();
+    // Fill the free list with sizes no frame asks for again.
+    for k in 0..16 {
+        let mut e = crate::context::Enc::lent(ctx);
+        drop((e.image(256 + k, 256), e.image(256 + k, 256)));
+    }
+    let evicted0 = ctx.pool_evictions();
+    let s = blend_scene(BitDepth::Bpc8, BlendMode::Screen);
+    let mut r = Renderer::new(&s.p, &Pattern, opts());
+    r.accel = Some(&g);
+    for _ in 0..3 {
+        effectcraft_render::passes::block_on(effectcraft_render::passes::comp_frame(&r, s.cid, Tick::ZERO));
+    }
+    for _ in 0..4 {
+        let mut e = crate::context::Enc::lent(ctx);
+        drop((e.image(300, 300), e.image(300, 300)));
+    }
+    assert!(ctx.pool_retained_bytes() <= budget);
+    assert!(ctx.pool_evictions() > evicted0, "rendered frames did not age the pool: stale sizes were never evicted");
+}
+
+/// #1550: with a small `EFFECTCRAFT_GPU_MEMCAP_MB` on a backend that reports allocator usage
+/// (Vulkan, DX12), `Enc::mem_gate` stays quiet when the bytes over the cap are non-reclaimable
+/// (live uploads), submits and waits when dropped working textures push reclaimable bytes over
+/// the cap, carries on without retiring the context if a wait times out, warns once when a
+/// frame's kept memory alone exceeds the cap, waits between two frames that fit only one at a
+/// time, and recovers budget when a frame is dropped unfinished.
+#[test]
+fn mem_gate_reclaims_over_cap_and_tolerates_timeout() {
+    const ENV: &str = "EFFECTCRAFT_GPU_MEMCAP_TEST_CHILD";
+    const SKIP: &str = "mem_gate_reclaims_over_cap_and_tolerates_timeout: skipped";
+    if std::env::var_os(ENV).is_none() {
+        let exe = std::env::current_exe().expect("test binary path");
+        let out = std::process::Command::new(exe)
+            .env(ENV, "1")
+            .env("EFFECTCRAFT_GPU_MEMCAP_MB", "8")
+            .args(["--exact", "tests::mem_gate_reclaims_over_cap_and_tolerates_timeout", "--nocapture"])
+            .output()
+            .expect("spawn child test process");
+        assert!(
+            out.status.success(),
+            "child test failed (status {:?}):\nstdout:\n{}\nstderr:\n{}",
+            out.status.code(),
+            String::from_utf8_lossy(&out.stdout),
+            String::from_utf8_lossy(&out.stderr),
+        );
+        // The child's output is captured: pass its skip line on.
+        String::from_utf8_lossy(&out.stderr).lines().filter(|l| l.starts_with(SKIP)).for_each(|l| eprintln!("{l}"));
+        return;
+    }
+
+    // EFFECTCRAFT_REQUIRE_GPU=1 (inherited from the parent) fails on a GPU that didn't start.
+    let g = Gpu::try_headless();
+    let Some(g) = require_gpu("mem_gate child", g.as_ref()) else {
+        eprintln!("{SKIP}: no GPU ({})", g.as_ref().err().map_or("", String::as_str));
+        return;
+    };
+    let ctx = g.context();
+    if ctx.allocated_bytes() == 0 {
+        // The backend has no allocator report (Metal, GL, WebGPU): the gate never engages.
+        eprintln!("{SKIP}: adapter {} reports no allocator usage", ctx.name);
+        return;
+    }
+
+    // 1. Non-reclaimable memory (upload cache) over the 8 MiB cap must not trigger waits.
+    let img = pattern(7, 512, 512); // 4 MiB in RGBA f32
+    let held: Vec<_> = (0..3)
+        .map(|i| {
+            let b = std::sync::Arc::new(effectcraft_effects::Buf { img: img.clone(), offset: [i as f64, 0.0], scale: 1.0 });
+            ctx.upload_buf(&b).expect("upload");
+            b
+        })
+        .collect();
+    let idx = ctx.queue.submit([]);
+    let _ = ctx.poll_readbacks(wgpu::PollType::Wait { submission_index: Some(idx), timeout: Some(std::time::Duration::from_secs(1)) });
+    assert!(ctx.allocated_bytes() > (8 << 20), "uploads should put the allocator over the 8 MiB cap (got {} B)", ctx.allocated_bytes());
+    {
+        let mut e = crate::context::Enc::new(ctx);
+        for _ in 0..16 {
+            e.mem_gate();
+        }
+    }
+    assert_eq!(ctx.memcap_waits(), 0, "mem_gate must not wait when over-cap memory is held by live uploads");
+    assert!(ctx.memcap_warned(), "uploads alone over the cap log the one-time warning");
+    drop(held);
+    ctx.clear_uploads();
+    ctx.pool_frame_end();
+
+    // 2. Dropped working textures over the 2 MiB pool budget and 8 MiB cap trigger a wait and
+    // free the dropped textures in wgpu's allocator.
+    let mut e = crate::context::Enc::new(ctx);
+    let mut peak_before_wait = 0u64;
+    for k in 0..16u32 {
+        let a = e.image(512 + k * 16, 256); // ~2 MiB each, distinct sizes
+        let b = e.image(512 + k * 16, 256);
+        e.copy_into(&a, &b, 0, 0);
+        drop((a, b));
+        peak_before_wait = peak_before_wait.max(ctx.allocated_bytes());
+        e.mem_gate();
+    }
+    drop(e);
+    assert!(peak_before_wait > (8 << 20), "working textures should push allocator over 8 MiB before a wait (peak {} B)", peak_before_wait);
+    let waits = ctx.memcap_waits();
+    assert!(waits > 0, "mem_gate should have waited once reclaimable bytes exceeded the cap");
+    assert!(
+        ctx.allocated_bytes() < peak_before_wait,
+        "waiting should reclaim dropped textures (peak {} B, after {} B)",
+        peak_before_wait,
+        ctx.allocated_bytes()
+    );
+    assert!(ctx.check_health().is_ok());
+    ctx.pool_frame_end();
+
+    // 3. A frame bigger than the cap: live GpuImages count in Pool::leased (kept for every owner),
+    // so three live 512×512 RGBA f32 images (12 MiB) make kept alone exceed the 8 MiB cap.
+    // Waiting cannot free them; mem_gate must not add waits (warning already asserted in section 1).
+    let waits_before_big = ctx.memcap_waits();
+    {
+        let mut e = crate::context::Enc::new(ctx);
+        let live = [e.image(512, 512), e.image(512, 512), e.image(512, 512)];
+        for _ in 0..16 {
+            e.mem_gate();
+        }
+        assert_eq!(ctx.memcap_waits(), waits_before_big, "a frame whose live images alone exceed the cap must not wait");
+        assert!(ctx.check_health().is_ok());
+        drop(live);
+    }
+    ctx.pool_frame_end();
+
+    // 4. Two frames that fit only one at a time: frame 1 creates ~6 MiB, copies, submits its
+    // work and drops its images, so wgpu holds them until the submission is done and polled.
+    // Frame 2 keeps ~6 MiB of new live images and waits for frame 1's reclaimable memory.
+    {
+        let mut e1 = crate::context::Enc::new(ctx);
+        let hold = [e1.image(512, 384), e1.image(512, 384)]; // 6 MiB
+        e1.copy_into(&hold[0], &hold[1], 0, 0);
+        e1.submit();
+        drop((hold, e1));
+    }
+    let allocated_after_f1 = ctx.allocated_bytes();
+    let waits_before_pair = ctx.memcap_waits();
+    let allocated_before_gate;
+    let downloaded;
+    {
+        let mut e2 = crate::context::Enc::new(ctx);
+        let a = e2.image(520, 384);
+        let b = e2.image(504, 384); // ~6 MiB together, sizes distinct from each other and frame 1
+        e2.copy_into(&b, &a, 0, 0);
+        allocated_before_gate = ctx.allocated_bytes();
+        assert!(
+            allocated_before_gate > (8 << 20),
+            "frame 1's unreclaimed memory ({allocated_after_f1} B) plus frame 2's live images should exceed the 8 MiB cap (got {allocated_before_gate} B)"
+        );
+        for _ in 0..16 {
+            e2.mem_gate();
+        }
+        assert!(ctx.memcap_waits() > waits_before_pair, "frame 2 should wait for frame 1's reclaimable memory (its own images stay live)");
+        assert!(
+            ctx.allocated_bytes() < allocated_before_gate,
+            "waiting should reclaim frame 1's dropped textures (before {} B, after {} B)",
+            allocated_before_gate,
+            ctx.allocated_bytes()
+        );
+        downloaded = e2.download(&a);
+    }
+    assert!(downloaded.is_some(), "frame 2 must finish with a correct download after waiting");
+    ctx.pool_frame_end();
+
+    // 5. A failed / unfinished frame releases its budget: drop an Enc mid-flight, then a
+    // following frame must not wait once the pool has drained.
+    let baseline = ctx.allocated_bytes();
+    {
+        let mut e = crate::context::Enc::new(ctx);
+        let _hold = [e.image(512, 512), e.image(512, 512), e.image(480, 480)];
+        // Drop without finishing: same path as a frame that errors before submit.
+    }
+    ctx.pool_frame_end();
+    let _ = ctx.poll_readbacks(wgpu::PollType::Poll);
+    assert!(
+        ctx.allocated_bytes() <= baseline + (2 << 20) + (1 << 20),
+        "dropped frame must release to baseline + pool budget (2 MiB) with slack (got {} B, baseline {} B)",
+        ctx.allocated_bytes(),
+        baseline
+    );
+    let waits_after_drop = ctx.memcap_waits();
+    {
+        let mut e = crate::context::Enc::lent(ctx);
+        for _ in 0..16 {
+            e.mem_gate();
+        }
+    }
+    assert_eq!(ctx.memcap_waits(), waits_after_drop, "a following frame must not wait after a dropped frame released its memory");
+
+    // 6. A zero-timeout wait must not retire the context or break subsequent renders.
+    ctx.set_memcap_wait_timeout(std::time::Duration::ZERO);
+    {
+        let mut e = crate::context::Enc::lent(ctx);
+        for k in 0..16u32 {
+            let a = e.image(384 + k * 16, 256);
+            let b = e.image(384 + k * 16, 256);
+            e.copy_into(&a, &b, 0, 0);
+            drop((a, b));
+            e.mem_gate();
+        }
+    }
+    assert!(ctx.memcap_waits() > waits_after_drop, "mem_gate should have fired again under zero timeout");
+    assert!(ctx.check_health().is_ok(), "a mem_gate wait timeout must carry on without retiring the context");
+    let s = blend_scene(BitDepth::Bpc8, BlendMode::Screen);
+    let r = Renderer::new(&s.p, &Pattern, opts());
+    assert!(g.render(&r, s.cid, Tick::ZERO).is_some(), "context renders normally after a zero-timeout mem_gate wait");
+
+    // 7. Non-blocking polls (`Gpu::poll`, `pool_frame_end`) skip when `poll_lock` is held by
+    // another thread, and `read_buffer` waits for `poll_lock` without consuming its deadline.
+    // Drain any dying pool textures first so the timed path takes Poll (try_lock) rather than Wait.
+    ctx.set_memcap_wait_timeout(crate::readback::TIMEOUT);
+    ctx.pool_frame_end();
+    let (locked_tx, locked_rx) = std::sync::mpsc::channel();
+    std::thread::scope(|sc| {
+        sc.spawn(|| {
+            let _guard = ctx.lock_poll_for_test();
+            locked_tx.send(()).unwrap();
+            std::thread::sleep(std::time::Duration::from_millis(60));
+        });
+        locked_rx.recv().unwrap();
+        let t0 = std::time::Instant::now();
+        g.poll();
+        ctx.pool_frame_end();
+        assert!(t0.elapsed() < std::time::Duration::from_millis(30), "non-blocking polls must not wait on a contended poll_lock (took {:?})", t0.elapsed());
+        assert!(g.render(&r, s.cid, Tick::ZERO).is_some(), "readback waits for poll_lock and succeeds without retiring the context");
+    });
+}

crates/gpu/src/tests/fx_stylize.rs

@@ -46,6 +46,66 @@ fn posterize_thresholds_strobe_vignette() {
     effect_case("ec.stylize.ccvignette", &[("amount", n(-60.0)), ("pinHighlights", n(50.0)), ("angleOfView", n(120.0)), ("center", pt(20.0, 15.0))]);
 }
 
+/// Premultiplied, quantized footage often lands exactly on a posterization
+/// threshold after division by alpha. Exercise that input before any placement
+/// or compositing can hide a wrong bin behind the normal image tolerance.
+#[test]
+fn posterize_quantized_premultiplied_boundaries() {
+    let Some(g) = crate::tests::gpu() else { return };
+    let spec = effectcraft_effects::find("ec.stylize.posterize").unwrap();
+    for level in [4.0, 9.0] {
+        let params = effectcraft_effects::Params { values: [("level".into(), n(level))].into() };
+        let ctx =
+            || effectcraft_effects::EffectCtx { params: &params, time: 0.0, layer_size: [70.0, 44.0], seed: 11, adjustment: false, env: Default::default() };
+        for seed in 1..=7 {
+            let mut img = crate::tests::pattern(seed, 70, 44);
+            effectcraft_render::color::quantize(&mut img, 255.0);
+            let input = effectcraft_effects::Buf { img, offset: [0.0; 2], scale: 1.0 };
+            let cpu = (spec.render)(&ctx(), input.clone());
+            let gpu = effectcraft_render::Accelerator::effects(g, &[effectcraft_render::FxStep { spec, ctx: ctx() }], &input, None).unwrap();
+            for (i, (want, got)) in cpu.img.data.iter().zip(&gpu.img.data).enumerate() {
+                for c in 0..4 {
+                    assert!((want[c] - got[c]).abs() <= 1e-6, "level {level} seed {seed} pixel {i}: input {:?}, CPU {want:?}, GPU {got:?}", input.img.data[i]);
+                }
+            }
+        }
+    }
+}
+
+/// Wrapped ramps have a large discontinuity at their end. Check signed and
+/// fractional phases, including a wrap a few double-precision ulps above zero.
+#[test]
+fn arbitrary_map_signed_fractional_and_near_zero_wraps() {
+    let Some(g) = crate::tests::gpu() else { return };
+    let spec = effectcraft_effects::find("ec.color.psarbitrarymap").unwrap();
+    for phase in [-1000.0, -20.5, 0.0, 40.0, 1000.0, 255.99999999999997] {
+        let params = effectcraft_effects::Params {
+            values: [
+                ("map".into(), Value::Str((0..256).map(|n| n.to_string()).collect::<Vec<_>>().join(","))),
+                ("phase".into(), n(phase)),
+                ("applyPhaseMapToAlpha".into(), on()),
+            ]
+            .into(),
+        };
+        let ctx =
+            || effectcraft_effects::EffectCtx { params: &params, time: 0.0, layer_size: [1024.0, 1.0], seed: 7, adjustment: false, env: Default::default() };
+        let mut img = effectcraft_raster::Image::new(1024, 1);
+        for (i, px) in img.data.iter_mut().enumerate() {
+            let v = if i < 512 { i as f32 * 1e-18 } else { (i - 512) as f32 / 511.0 };
+            let a = [1.0f32, 0.1, 0.3, 6.0 / 255.0][i % 4];
+            *px = [v * a, v * a, v * a, a];
+        }
+        let input = effectcraft_effects::Buf { img, offset: [0.0; 2], scale: 1.0 };
+        let cpu = (spec.render)(&ctx(), input.clone());
+        let gpu = effectcraft_render::Accelerator::effects(g, &[effectcraft_render::FxStep { spec, ctx: ctx() }], &input, None).unwrap();
+        for (i, (want, got)) in cpu.img.data.iter().zip(&gpu.img.data).enumerate() {
+            for c in 0..4 {
+                assert!((want[c] - got[c]).abs() <= 1e-6, "phase {phase}, pixel {i}: CPU {want:?}, GPU {got:?}");
+            }
+        }
+    }
+}
+
 #[test]
 fn scatter_brush_strokes() {
     effect_case("ec.stylize.scatter", &[("amount", n(5.0))]);
@@ -228,6 +288,36 @@ fn transform_motion_blur() {
     }
 }
 
+/// #1605: Posterize matches the CPU exactly on every 8 bpc premultiplied colour and alpha. A
+/// straight value of 5/6 at level 4 rounds to exactly 2.5 on the CPU; drivers that fused the
+/// product into `round_away`'s subtraction (AMD) rounded it down a level.
+#[test]
+fn posterize_matches_the_cpu_on_every_8bpc_pair() {
+    let Some(g) = crate::tests::gpu() else { return };
+    let spec = effectcraft_effects::find("ec.stylize.posterize").unwrap();
+    let size = [256.0, 255.0];
+    let mut params =
+        effectcraft_effects::Params { values: spec.params.iter().map(|ps| (ps.id.to_string(), effectcraft_effects::default_value(ps, size))).collect() };
+    for level in [4.0, 7.0, 9.0] {
+        params.values.insert("level".into(), n(level));
+        // Pixel (x, a - 1): premultiplied colour min(x, a)/255 on alpha a/255.
+        let mut img = effectcraft_raster::Image::new(256, 255);
+        for a in 1..=255u32 {
+            for x in 0..256u32 {
+                let v = x.min(a) as f32 / 255.0;
+                img.data[((a - 1) * 256 + x) as usize] = [v, v, v, a as f32 / 255.0];
+            }
+        }
+        let buf = effectcraft_effects::Buf { img, offset: [0.0, 0.0], scale: 1.0 };
+        let ctx = || effectcraft_effects::EffectCtx { params: &params, time: 0.0, layer_size: size, seed: 1, adjustment: false, 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 bad: Vec<(usize, usize)> =
+            cpu.img.data.iter().zip(&out.img.data).enumerate().filter(|(_, (c, o))| c != o).map(|(i, _)| ((i % 256).min(i / 256 + 1), i / 256 + 1)).collect();
+        assert!(bad.is_empty(), "level {level}: {} (x, a) pairs differ from the CPU, first {:?}", bad.len(), &bad[..bad.len().min(8)]);
+    }
+}
+
 /// Motion Tile's Output size on an adjustment layer with a Transform after it that scales the
 /// tiles into view, and cropping (Output under 100 %): the stack's grown or cropped frame goes
 /// back onto the canvas as on the CPU (#1727).

crates/gpu/src/tests/inpaint.rs

@@ -0,0 +1,310 @@
+//! GPU Content-Aware Fill vs the CPU reference: the CPU's fill is the oracle.
+//!
+//! Parity here is a quality contract, not bit-exactness (see `inpaint.rs`): PatchMatch's
+//! scan-order propagation becomes a two-pass checkerboard. So the tests compare both
+//! syntheses against the ground truth they are reconstructing, and check the GPU is not the
+//! worse of the two.
+
+use effectcraft_raster::Image;
+use effectcraft_raster::inpaint::{FillBackend, FillInput, FillOpts, fill_sequence_with, membrane_fill, patch_fill, psnr};
+
+use crate::tests::gpu;
+
+/// A textured background (periodic stripes + blocky noise), shifted by `dx` (the CPU tests'
+/// scene, so the numbers here are comparable with theirs).
+fn background(w: u32, h: u32, dx: f32) -> Image {
+    let mut img = Image::new(w, h);
+    for y in 0..h {
+        for x in 0..w {
+            let fx = x as f32 + dx;
+            let s = 0.5 + 0.3 * (fx * std::f32::consts::TAU / 10.0).sin();
+            let n = effectcraft_raster::hash_noise(((fx.floor() as i64).rem_euclid(20) / 4) as u32, (y % 20) / 4, 3);
+            img.set(x, y, [s, 0.3 + 0.4 * n, 0.6 - 0.3 * s * n, 1.0]);
+        }
+    }
+    img
+}
+
+fn cut(img: &Image, hole: &[bool]) -> Image {
+    let mut o = img.clone();
+    for (p, h) in o.data.iter_mut().zip(hole) {
+        if *h {
+            *p = [0.0; 4];
+        }
+    }
+    o
+}
+
+fn rect_hole(w: u32, h: u32, x0: u32, y0: u32, s: u32) -> Vec<bool> {
+    (0..w * h).map(|i| (x0..x0 + s).contains(&(i % w)) && (y0..y0 + s).contains(&(i / w))).collect()
+}
+
+/// The kernels must exist: a shader that fails to build falls back to the CPU silently, which
+/// would make every quality test below pass while testing nothing.
+#[test]
+fn the_inpaint_kernels_are_available() {
+    let Some(g) = gpu() else { return };
+    let bad: Vec<&str> = g.ctx.unavailable_kernels().iter().filter(|(e, _)| crate::inpaint::KERNELS.contains(e)).map(|(e, _)| *e).collect();
+    assert!(bad.is_empty(), "inpaint kernels did not build on {}: {bad:?}", g.ctx.name);
+}
+
+/// The device path (not the CPU fallback) is what the quality tests measure.
+#[test]
+fn the_gpu_runs_the_fill_itself() {
+    let Some(g) = gpu() else { return };
+    let (w, h) = (96u32, 72u32);
+    let gt = background(w, h, 0.0);
+    let hole = rect_hole(w, h, 40, 28, 16);
+    let src = cut(&gt, &hole);
+    assert!(crate::inpaint::patch_image(&g.ctx, &src, &hole, None, 7, 1).is_some(), "PatchMatch fell back to the CPU");
+}
+
+#[test]
+fn patchmatch_beats_the_membrane_on_texture() {
+    let Some(g) = gpu() else { return };
+    let (w, h) = (96u32, 72u32);
+    let gt = background(w, h, 0.0);
+    let hole = rect_hole(w, h, 40, 28, 16);
+    let src = cut(&gt, &hole);
+    let mem = membrane_fill(&src, &hole);
+    let pm = g.patch(&src, &hole, None, 7, 1);
+    let (pm_db, mem_db) = (psnr(&pm, &gt, Some(&hole)), psnr(&mem, &gt, Some(&hole)));
+    assert!(pm_db > mem_db + 3.0 && pm_db > 20.0, "PatchMatch {pm_db:.1} dB vs membrane {mem_db:.1} dB");
+    for i in 0..hole.len() {
+        if !hole[i] {
+            assert_eq!(pm.data[i], gt.data[i]);
+        }
+    }
+}
+
+#[test]
+fn patchmatch_agrees_with_the_cpu() {
+    let Some(g) = gpu() else { return };
+    let (w, h) = (96u32, 72u32);
+    let gt = background(w, h, 0.0);
+    let hole = rect_hole(w, h, 40, 28, 16);
+    let src = cut(&gt, &hole);
+    let (cpu, gpu_pm) = (patch_fill(&src, &hole, None, 7, 1), g.patch(&src, &hole, None, 7, 1));
+    // Both are syntheses of the same hole: neither is the other's ground truth, so they are
+    // compared as two reconstructions of the same thing. On this periodic background both are
+    // near-exact, so a relative dB check at machine precision says nothing; the device must be a
+    // good reconstruction *and* agree with the CPU (a wrong fill that still looks plausible would
+    // pass the first and fail the second).
+    let gpu_db = psnr(&gpu_pm, &gt, Some(&hole));
+    assert!(gpu_db > 25.0, "GPU {gpu_db:.1} dB");
+    assert!(psnr(&gpu_pm, &cpu, Some(&hole)) > 30.0, "the two syntheses disagree");
+}
+
+/// The `init` path (a frame refined from the previous frame's warped fill) is the finest level
+/// only, from two rounds of four iterations (`patch_fill`).
+#[test]
+fn patchmatch_refines_an_init_frame() {
+    let Some(g) = gpu() else { return };
+    let (w, h) = (96u32, 72u32);
+    let gt = background(w, h, 0.0);
+    let hole = rect_hole(w, h, 40, 28, 16);
+    let src = cut(&gt, &hole);
+    let init = gt.clone();
+    let (cpu, gpu_pm) = (patch_fill(&src, &hole, Some(&init), 7, 1), g.patch(&src, &hole, Some(&init), 7, 1));
+    assert!(psnr(&gpu_pm, &gt, Some(&hole)) > 25.0, "{:.1} dB", psnr(&gpu_pm, &gt, Some(&hole)));
+    assert!(psnr(&gpu_pm, &cpu, Some(&hole)) > 20.0, "the two syntheses disagree");
+}
+
+/// A hole across the whole multi-scale pyramid (the coarse levels' membrane start included):
+/// every level must come back finite and the pyramid must beat a single-level solve.
+#[test]
+fn patchmatch_survives_a_large_hole_across_levels() {
+    let Some(g) = gpu() else { return };
+    let (w, h) = (256u32, 192u32);
+    let gt = background(w, h, 0.0);
+    let hole = rect_hole(w, h, 90, 70, 80);
+    let src = cut(&gt, &hole);
+    let pm = g.patch(&src, &hole, None, 7, 1);
+    assert!(pm.data.iter().flatten().all(|v| v.is_finite()), "the synthesis left a non-finite value");
+    let db = psnr(&pm, &gt, Some(&hole));
+    assert!(db > 20.0, "{db:.1} dB");
+    for i in 0..hole.len() {
+        if !hole[i] {
+            assert_eq!(pm.data[i], gt.data[i], "known pixel {i}");
+        }
+    }
+}
+
+/// A moving object removed from a (static or panning) textured background, filled frame by
+/// frame through the device backend the sequence driver uses.
+#[test]
+fn object_fill_recovers_background_and_is_stable() {
+    let Some(g) = gpu() else { return };
+    let (w, h) = (96u32, 64u32);
+    for pan in [0.0f32, 1.0] {
+        let mut frames = vec![];
+        let mut holes = vec![];
+        let mut gts = vec![];
+        for t in 0..8u32 {
+            let gt = background(w, h, t as f32 * pan);
+            let hole = rect_hole(w, h, 10 + 8 * t, 22, 18);
+            frames.push(cut(&gt, &hole));
+            holes.push(hole);
+            gts.push(gt);
+        }
+        let input = FillInput { frames, holes, references: vec![] };
+        let out = fill_sequence_with(&input, &FillOpts::default(), g, &mut |_, _| true).expect("fill");
+        for t in 0..out.len() {
+            let db = psnr(&out[t], &gts[t], Some(&input.holes[t]));
+            assert!(db > 26.0, "frame {t} pan {pan}: {db:.1} dB");
+            assert_eq!(out[t].data[0][3], 1.0);
+        }
+        if pan == 0.0 {
+            for t in 1..out.len() {
+                let both: Vec<bool> = input.holes[t].iter().zip(&input.holes[t - 1]).map(|(a, b)| *a || *b).collect();
+                let db = psnr(&out[t], &out[t - 1], Some(&both));
+                assert!(db > 30.0, "frames {}–{t} differ: {db:.1} dB", t - 1);
+            }
+        }
+    }
+}
+
+/// The membrane stays on the CPU (see [`crate::inpaint`]): Edge Blend must still get the CPU's
+/// result through the backend, not a device one.
+#[test]
+fn edge_blend_takes_the_cpu_membrane() {
+    let Some(g) = gpu() else { return };
+    let (w, h) = (80u32, 60u32);
+    let gt = background(w, h, 0.0);
+    let hole = rect_hole(w, h, 30, 20, 14);
+    let input = FillInput { frames: vec![cut(&gt, &hole); 4], holes: vec![hole.clone(); 4], references: vec![] };
+    let opts = FillOpts { method: effectcraft_raster::inpaint::FillMethod::EdgeBlend, ..Default::default() };
+    let edge = fill_sequence_with(&input, &opts, g, &mut |_, _| true).expect("fill");
+    for t in 1..4 {
+        assert!(psnr(&edge[t], &edge[0], Some(&hole)) > 40.0, "frame {t}");
+    }
+    assert!(psnr(&edge[0], &membrane_fill(&input.frames[0], &hole), Some(&hole)) > 60.0);
+}
+
+/// The dilation the level plans are built from (kept honest: the flags' target bit is `near`).
+#[test]
+fn the_target_band_is_the_dilated_hole() {
+    let (w, h) = (24u32, 18u32);
+    let hole = rect_hole(w, h, 8, 6, 4);
+    let near = effectcraft_raster::inpaint::dilate(&hole, w as usize, h as usize, 3);
+    assert!(near.iter().filter(|b| **b).count() > hole.iter().filter(|b| **b).count());
+    assert!(near.iter().zip(&hole).all(|(n, h)| !*h || *n));
+}
+
+/// The flow's block matching runs on the device and finds the same motion the CPU does: the two
+/// frames are the membrane-filled neighbours the sequence driver builds, with a hole in each.
+#[test]
+fn flow_agrees_with_the_cpu() {
+    let Some(g) = gpu() else { return };
+    let (w, h) = (192u32, 128u32);
+    let a_gt = background(w, h, 0.0);
+    let b_gt = background(w, h, 5.0);
+    let hole_a = rect_hole(w, h, 30, 40, 40);
+    let mut hole_b = vec![false; (w * h) as usize];
+    for y in 50..90 {
+        for x in 120..160 {
+            hole_b[(y * w + x) as usize] = true;
+        }
+    }
+    let a = membrane_fill(&cut(&a_gt, &hole_a), &hole_a);
+    let b = membrane_fill(&cut(&b_gt, &hole_b), &hole_b);
+    assert!(crate::inpaint::flow_image(&g.ctx, &a, &b).is_some(), "the flow fell back to the CPU");
+    let cpu = effectcraft_raster::inpaint::CpuFill.flow(&a, &b, &hole_a, &hole_b);
+    let dev = g.flow(&a, &b, &hole_a, &hole_b);
+    assert_eq!((dev.cols, dev.rows, dev.block), (cpu.cols, cpu.rows, cpu.block));
+    // The field is the CPU's to a float rounding: the luma is the CPU's (`luma_pair`), the
+    // pyramid is exact additions, and the search enumerates candidates in the CPU's order.
+    let diff = cpu.v.iter().zip(&dev.v).map(|(c, d)| (c[0] - d[0]).abs().max((c[1] - d[1]).abs())).fold(0.0f32, f32::max);
+    assert!(diff <= 1e-6, "the flow fields differ by {diff}");
+    // The motion is real: the panning background's 5 px shift survives the completion. `b`'s
+    // content sits 5 px left of `a`'s (`background` adds `dx` inside the sample), so a → b is −5.
+    let m = dev.at(96.0, 64.0);
+    assert!((m[0] + 5.0).abs() < 0.5 && m[1].abs() < 0.5, "flow at centre: {m:?}");
+}
+
+/// The vote's σ² comes from a histogram that must count *every* target: the old per-thread
+/// strided pass kept only the pixels whose bin selected that thread (a biased ~1/256 subsample),
+/// which left small holes with a `total < 1` fallback and near-uniform vote weights.
+#[test]
+fn the_distance_histogram_counts_every_target() {
+    let Some(g) = gpu() else { return };
+    let (w, h) = (48u32, 32u32);
+    // Every target has the same distance, so the histogram is one spike. 0.5 is the SDR range
+    // (bin 128); 1600 is above it, which is reachable from a 32-bpc frame (the reviewer's example:
+    // clamping it to bin 1023 reported σ² ≈ 4, and the vote's weights then rounded to zero).
+    for (d, want) in [(0.5f32, 128.5f32 / 256.0), (1600.0, 1600.0)] {
+        let (mut field, near) = (Image::new(w, h), vec![true; (w * h) as usize]);
+        for px in field.data.iter_mut() {
+            *px = [0.0, 0.0, d, 1.0];
+        }
+        let (sigma2, total) = crate::inpaint::distance_histogram(&g.ctx, &field, &near, w, h).expect("histogram");
+        assert_eq!(total as usize, (w * h) as usize, "distance {d}: the histogram kept {total} of {} samples", w * h);
+        let tol = (want * 1e-3).max(1e-3);
+        assert!((sigma2 - want).abs() <= tol, "distance {d}: sigma2 {sigma2}, want {want}");
+    }
+}
+
+/// HDR frames (channels above 1, as a 32-bpc image has) must still fill like the CPU. Their patch
+/// distances are far past the SDR 0…4 range the histogram used to clamp at, so this covers the
+/// rescaled binning end to end; `the_distance_histogram_counts_every_target` pins the estimator
+/// itself (the old clamp read σ² ≈ 4 for a distance of 1600).
+#[test]
+fn hdr_frames_fill_like_the_cpu() {
+    let Some(g) = gpu() else { return };
+    let (w, h) = (96u32, 72u32);
+    let mut gt = background(w, h, 0.0);
+    for px in &mut gt.data {
+        // A 32-bpc frame's values are unbounded: at 24× the patch distances run into the thousands.
+        // The old fixed 0…4 histogram reported σ² ≈ 4 for those, `exp(-d/(2σ²))` rounded to zero,
+        // and the vote kept the prior image instead of synthesising.
+        for c in 0..3 {
+            px[c] *= 24.0;
+        }
+    }
+    let hole = rect_hole(w, h, 40, 28, 16);
+    let src = cut(&gt, &hole);
+    let (cpu, gpu_pm) = (patch_fill(&src, &hole, None, 7, 1), g.patch(&src, &hole, None, 7, 1));
+    let gpu_db = psnr(&gpu_pm, &gt, Some(&hole));
+    assert!(gpu_db > 25.0, "GPU {gpu_db:.1} dB");
+    assert!(psnr(&gpu_pm, &cpu, Some(&hole)) > 30.0, "the two syntheses disagree");
+}
+
+/// A frame too small for one block is the CPU's zero field, on the device path too.
+#[test]
+fn small_frames_get_a_zero_flow() {
+    let Some(g) = gpu() else { return };
+    let (w, h) = (6u32, 6u32);
+    let a = background(w, h, 0.0);
+    let b = background(w, h, 2.0);
+    let none = vec![false; (w * h) as usize];
+    let f = g.flow(&a, &b, &none, &none);
+    assert_eq!((f.cols, f.rows), (1, 1));
+    assert_eq!(f.v, vec![[0.0, 0.0]]);
+}
+
+/// The flow the sequence driver uses (through the `Gpu` backend) keeps the fill stable: the same
+/// frames filled with the CPU backend and the device backend agree, so the propagation sees the
+/// same motion.
+#[test]
+fn sequence_fill_with_the_device_flow_matches_the_cpu() {
+    let Some(g) = gpu() else { return };
+    let (w, h) = (96u32, 64u32);
+    let cpu = effectcraft_raster::inpaint::CpuFill;
+    for pan in [0.0f32, 2.0] {
+        let mut frames = vec![];
+        let mut holes = vec![];
+        for t in 0..6u32 {
+            let gt = background(w, h, t as f32 * pan);
+            let hole = rect_hole(w, h, 12 + 6 * t, 22, 18);
+            frames.push(cut(&gt, &hole));
+            holes.push(hole);
+        }
+        let input = FillInput { frames, holes, references: vec![] };
+        let dev = fill_sequence_with(&input, &FillOpts::default(), g, &mut |_, _| true).expect("gpu fill");
+        let ref_cpu = fill_sequence_with(&input, &FillOpts::default(), &cpu, &mut |_, _| true).expect("cpu fill");
+        for t in 0..dev.len() {
+            let db = psnr(&dev[t], &ref_cpu[t], Some(&input.holes[t]));
+            assert!(db > 40.0, "frame {t} pan {pan}: device vs CPU {db:.1} dB");
+        }
+    }
+}

crates/gpu/src/walk.rs

@@ -213,6 +213,12 @@ fn draw_collapsed<'a>(
 }
 
 fn draw_layer(e: &mut Enc, r: &Renderer, ctx: &EvalCtx, layer: &Layer, canvas: &mut GpuImage) -> Option<()> {
+    let out = draw_layer_on(e, r, ctx, layer, canvas);
+    e.mem_gate();
+    out
+}
+
+fn draw_layer_on(e: &mut Enc, r: &Renderer, ctx: &EvalCtx, layer: &Layer, canvas: &mut GpuImage) -> Option<()> {
     if layer.switches.adjustment {
         if draw_adjustment(e, r, ctx, layer, canvas).is_some() {
             return Some(());

crates/gpu/src/yuv.rs

@@ -130,7 +130,7 @@ pub(crate) fn record(e: &mut Enc, yp: &YuvPipe, img: &GpuImage, bg: [f32; 3]) ->
         ],
     });
     {
-        let enc = e.encoder();
+        let enc = e.record(&[img]);
         {
             let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("yuv420"), timestamp_writes: None });
             pass.set_pipeline(&yp.pipe);

crates/host/tests/it/plugins.rs

@@ -69,14 +69,18 @@ fn rust_plugins_apply_by_id() {
 /// Build the example plug-in (skipped when the wasm32 target isn't installed).
 fn build_example() -> Option<std::path::PathBuf> {
     let root = std::path::Path::new(env!("CARGO_MANIFEST_DIR")).join("../../examples/plugins/posterize-bands");
-    let target = std::env::var_os("CARGO_TARGET_DIR").map(std::path::PathBuf::from).unwrap_or_else(|| root.join("../../../target")).join("plugin-example");
-    let target = if target.is_absolute() { target } else { root.join("../../..").join(target) };
-    let st = std::process::Command::new(env!("CARGO"))
-        .current_dir(&root)
-        .args(["build", "--release", "--target", "wasm32-unknown-unknown", "--target-dir"])
-        .arg(&target)
-        .status()
-        .ok()?;
+    let workspace = std::path::Path::new(env!("CARGO_MANIFEST_DIR")).join("../..");
+    let assigned = std::env::var_os("CARGO_TARGET_DIR");
+    let target = assigned.as_ref().map(std::path::PathBuf::from).unwrap_or_else(|| workspace.join("target"));
+    let target = if target.is_absolute() { target } else { workspace.join(target) };
+    // The plug-in is part of this test's build. Inherit the assigned target and
+    // compiler guards rather than creating a second target at its root.
+    let mut command = std::process::Command::new(env!("CARGO"));
+    command.current_dir(&workspace).args(["build", "--release", "--target", "wasm32-unknown-unknown", "--manifest-path"]).arg(root.join("Cargo.toml"));
+    if assigned.is_none() {
+        command.env("CARGO_TARGET_DIR", &target);
+    }
+    let st = command.status().ok()?;
     if !st.success() {
         eprintln!("skipping: the example plug-in didn't build (rustup target add wasm32-unknown-unknown)");
         return None;

crates/raster/src/flow.rs

@@ -51,6 +51,21 @@ impl Flow {
     }
 }
 
+/// The `Plane::luma` of `a` and `b` packed into one image (`a` in the red channel, `b` in the
+/// green; the rest zero). A device block matcher takes this as its input rather than computing
+/// the luma itself: the weighted sum is not reproducible on a device, which may contract
+/// `a·b + c` into a fused multiply-add, and the search's near-tie choices would then differ.
+pub fn luma_pair(a: &Image, b: &Image) -> Image {
+    let luma = |p: &[f32; 4]| 0.2126 * p[0] + 0.7152 * p[1] + 0.0722 * p[2] + 0.25 * p[3];
+    let mut out = Image::new(a.width, a.height);
+    out.data.par_iter_mut().enumerate().for_each(|(i, q)| {
+        let la = a.data.get(i).map(luma).unwrap_or(0.0);
+        let lb = b.data.get(i).map(luma).unwrap_or(0.0);
+        *q = [la, lb, 0.0, 0.0];
+    });
+    out
+}
+
 /// A single-channel plane.
 struct Plane {
     w: usize,

crates/raster/src/inpaint.rs

@@ -562,23 +562,31 @@ pub fn patch_fill(img: &Image, hole: &[bool], init: Option<&Image>, patch: usize
 
 // ------------------------------------------------------------------ flow-guided propagation
 
-/// Block flow from `a` to `b` with the tiles that touch either hole replaced by the membrane
-/// interpolation of the surrounding motion (flow completion).
+/// Block flow `a` → `b` (tile 8, search radius 6) with the tiles that touch either hole replaced
+/// by the membrane interpolation of the surrounding motion (flow completion).
 fn completed_flow(a: &Image, b: &Image, hole_a: &[bool], hole_b: &[bool]) -> Flow {
-    let mut f = block_flow(a, b, 8, 6);
+    let f = block_flow(a, b, 8, 6);
+    complete_flow(f, hole_a, hole_b, a.width as usize, a.height as usize)
+}
+
+/// The completion half of `completed_flow`: the tiles of `f` (a `block_flow` of `w × h`) that
+/// touch either hole become the harmonic interpolation of the surrounding motion. Split out so a
+/// hardware backend can compute the block matching itself and hand the raw field here; the
+/// completion is cheap (the field is a `cols × rows` grid) and stays on the CPU.
+pub fn complete_flow(mut f: Flow, hole_a: &[bool], hole_b: &[bool], w: usize, h: usize) -> Flow {
     if f.v.is_empty() {
         return f;
     }
-    let (w, bs) = (a.width as usize, f.block as usize);
+    let bs = f.block as usize;
     let mut known = vec![true; f.v.len()];
     for (i, k) in known.iter_mut().enumerate() {
         let (c, r) = (i % f.cols, i / f.cols);
         // A tile is unreliable when any pixel within it (and a block's margin) is a hole.
         let (x0, y0) = ((c * bs).saturating_sub(bs / 2), (r * bs).saturating_sub(bs / 2));
-        let (x1, y1) = (((c + 1) * bs + bs / 2).min(w), ((r + 1) * bs + bs / 2).min(a.height as usize));
+        let (x1, y1) = (((c + 1) * bs + bs / 2).min(w), ((r + 1) * bs + bs / 2).min(h));
         'tile: for y in y0..y1 {
             for x in x0..x1 {
-                if hole_a[y * w + x] || hole_b[y * w + x] {
+                if hole_a.get(y * w + x).copied().unwrap_or(true) || hole_b.get(y * w + x).copied().unwrap_or(true) {
                     *k = false;
                     break 'tile;
                 }
@@ -660,8 +668,54 @@ fn trace(context: &TraceContext<'_>, start: TraceStart) -> Option<([f32; 4], usi
     None
 }
 
-/// Fill a range of frames. `progress(done, total)` returns `false` to cancel (→ `None`).
+/// A hardware backend for the fill's stages (the GPU compositor implements it; see
+/// `effectcraft_render::Accelerator::inpaint`). Every method has the CPU implementation as its
+/// default, so a backend that cannot take a stage returns the CPU's work and the fill stays
+/// correct — only slower.
+///
+/// The orchestration ([`fill_sequence_with`]) stays on the CPU: the sequence's frame-to-frame
+/// coupling is sequential, and only the per-frame and per-pair work is worth offloading.
+///
+/// **A backend must be safe to call concurrently.** [`fill_sequence_with`] runs the stages of
+/// independent frames and frame pairs in parallel (`rayon`), so `membrane`, `flow` and `patch`
+/// may all be entered from several threads at once — a backend that owns a device must serialise
+/// its use of it (the GPU compositor's per-call encoder does; a backend whose device cannot be
+/// used from several threads must lock).
+pub trait FillBackend: Send + Sync {
+    /// [`membrane_fill`]: the `hole` pixels take the harmonic interpolation of their border.
+    fn membrane(&self, img: &Image, hole: &[bool]) -> Image {
+        membrane_fill(img, hole)
+    }
+    /// `membrane_core`: harmonic interpolation of `data`'s unknown pixels in place (`known` is
+    /// true where the pixel is known).
+    fn membrane_known(&self, data: &mut [[f32; 4]], known: &[bool], w: usize, h: usize) {
+        membrane_core(data, known, w, h)
+    }
+    /// `completed_flow`: block flow `a` → `b` with the tiles touching either hole replaced by
+    /// the membrane interpolation of the surrounding motion.
+    fn flow(&self, a: &Image, b: &Image, hole_a: &[bool], hole_b: &[bool]) -> Flow {
+        completed_flow(a, b, hole_a, hole_b)
+    }
+    /// [`patch_fill`]: PatchMatch synthesis of `img`'s `hole` (`init`: a starting guess, e.g. the
+    /// previous frame's fill warped here).
+    fn patch(&self, img: &Image, hole: &[bool], init: Option<&Image>, patch: usize, seed: u32) -> Image {
+        patch_fill(img, hole, init, patch, seed)
+    }
+}
+
+/// The CPU backend: every stage is the reference implementation.
+pub struct CpuFill;
+
+impl FillBackend for CpuFill {}
+
+/// Fill a range of frames on the CPU. `progress(done, total)` returns `false` to cancel
+/// (→ `None`).
 pub fn fill_sequence(input: &FillInput, o: &FillOpts, progress: &mut dyn FnMut(usize, usize) -> bool) -> Option<Vec<Image>> {
+    fill_sequence_with(input, o, &CpuFill, progress)
+}
+
+/// [`fill_sequence`] with a hardware backend for the stages it can take (see [`FillBackend`]).
+pub fn fill_sequence_with(input: &FillInput, o: &FillOpts, backend: &dyn FillBackend, progress: &mut dyn FnMut(usize, usize) -> bool) -> Option<Vec<Image>> {
     let n = input.frames.len();
     if n == 0 {
         return Some(vec![]);
@@ -672,7 +726,7 @@ pub fn fill_sequence(input: &FillInput, o: &FillOpts, progress: &mut dyn FnMut(u
     if o.method == FillMethod::EdgeBlend {
         let mut out = Vec::with_capacity(n);
         for (i, (f, hl)) in input.frames.iter().zip(&input.holes).enumerate() {
-            out.push(opaque_fill(membrane_fill(f, hl), hl));
+            out.push(opaque_fill(backend.membrane(f, hl), hl));
             if !progress(i + 1, n) {
                 return None;
             }
@@ -691,10 +745,10 @@ pub fn fill_sequence(input: &FillInput, o: &FillOpts, progress: &mut dyn FnMut(u
         }
     }
     // 1. Flows between membrane-filled neighbours, completed inside the holes.
-    let pre: Vec<Image> = frames.par_iter().zip(&holes).map(|(f, hl)| membrane_fill(f, hl)).collect();
+    let pre: Vec<Image> = frames.par_iter().zip(&holes).map(|(f, hl)| backend.membrane(f, hl)).collect();
     let pairs: Vec<(Flow, Flow)> = (0..n.saturating_sub(1))
         .into_par_iter()
-        .map(|i| (completed_flow(&pre[i], &pre[i + 1], &holes[i], &holes[i + 1]), completed_flow(&pre[i + 1], &pre[i], &holes[i + 1], &holes[i])))
+        .map(|i| (backend.flow(&pre[i], &pre[i + 1], &holes[i], &holes[i + 1]), backend.flow(&pre[i + 1], &pre[i], &holes[i + 1], &holes[i])))
         .collect();
     let (fwd, bwd): (Vec<Flow>, Vec<Flow>) = pairs.into_iter().unzip();
     if !progress(1, total) {
@@ -746,7 +800,7 @@ pub fn fill_sequence(input: &FillInput, o: &FillOpts, progress: &mut dyn FnMut(u
                         diff[i] = std::array::from_fn(|c| if c < 3 { a[c] - v[c] } else { 0.0 });
                     }
                 }
-                membrane_core(&mut diff, &dk, w, h);
+                backend.membrane_known(&mut diff, &dk, w, h);
                 for (i, _) in &filled {
                     for c in 0..3 {
                         img.data[*i][c] += o.lighting * diff[*i][c];
@@ -770,7 +824,7 @@ pub fn fill_sequence(input: &FillInput, o: &FillOpts, progress: &mut dyn FnMut(u
         *slot = Some(propagated[t].0.clone());
     }
     let mut done = 2;
-    out[t0] = Some(patch_fill(&propagated[t0].0, &propagated[t0].1, None, o.patch, o.seed));
+    out[t0] = Some(backend.patch(&propagated[t0].0, &propagated[t0].1, None, o.patch, o.seed));
     done += 1;
     if !progress(done, total) {
         return None;
@@ -783,7 +837,7 @@ pub fn fill_sequence(input: &FillInput, o: &FillOpts, progress: &mut dyn FnMut(u
             // Warp the previous result here along the backward flow (t → t − 1).
             // Frames after `t0` are filled in order, so the previous one is set.
             let Some(prev) = out[t - 1].as_ref() else {
-                out[t] = Some(patch_fill(img, left, None, o.patch, o.seed));
+                out[t] = Some(backend.patch(img, left, None, o.patch, o.seed));
                 continue;
             };
             let mut warped = img.clone();
@@ -796,7 +850,7 @@ pub fn fill_sequence(input: &FillInput, o: &FillOpts, progress: &mut dyn FnMut(u
             }
             out[t] = Some(match o.method {
                 FillMethod::Surface => warped,
-                _ => patch_fill(img, left, Some(&warped), o.patch, o.seed),
+                _ => backend.patch(img, left, Some(&warped), o.patch, o.seed),
             });
         }
         done += 1;

crates/render/src/accel.rs

@@ -76,6 +76,16 @@ pub trait Accelerator: Send + Sync {
     fn particles(&self) -> Option<&dyn effectcraft_effects::psim::ParticleSim> {
         None
     }
+    /// A hardware backend for Content-Aware Fill's stages (the GPU compositor's kernels), with
+    /// the CPU fill's semantics. `None` = the fill runs on the CPU.
+    ///
+    /// The fill's input frames are read back to the CPU before it runs (analyses render on the
+    /// CPU by design, see `effectcraft_engine::background::offload::job_accel`), so a backend
+    /// here is about the fill's own work — PatchMatch and the flow's block matching today, one
+    /// stage at a time as each pays (`effectcraft_gpu::inpaint`) — not about compositing.
+    fn inpaint(&self) -> Option<&dyn effectcraft_raster::inpaint::FillBackend> {
+        None
+    }
     /// Timing history [`Backend::Auto`] uses to send each comp's top-level frames to the
     /// faster compositor. `None` = Auto always tries the accelerator first.
     fn auto_pick(&self) -> Option<&AutoPick> {
@@ -96,6 +106,19 @@ pub trait Accelerator: Send + Sync {
     /// results, `pass_begin` clears the miss. Synchronous accelerators never miss.
     fn frame_begin(&self) {}
     fn pass_begin(&self) {}
+    /// A rendered comp frame is done (the GPU pool ages sizes no frame asks for, and evicts
+    /// only those once full). The multi-frame render paths call it once per comp frame:
+    /// [`crate::passes::comp_frame`] (exports; a field-rendered output frame is two), the
+    /// export's 4:2:0 frames ([`Accelerator::comp_frame_yuv420`]), remote frames and the
+    /// preview workers. One-off renders (`Session::render`, single-frame export,
+    /// thumbnails) don't, so they don't advance the pool's clock. Without it the pool never
+    /// finds a stale size and keeps its first sizes, dropping newcomers.
+    fn frame_end(&self) {}
+    /// A render of many frames (an export) is done: drop what is kept only to speed up the
+    /// next frame (pooled working textures). Exports call it, even when they fail, panic or
+    /// are dropped; other multi-frame paths (preview, remote frames) don't, so there the pool
+    /// keeps up to its budget until the next export ends.
+    fn render_end(&self) {}
     fn pass_missed(&self) -> bool {
         false
     }

crates/render/src/passes.rs

@@ -59,11 +59,25 @@ pub async fn in_passes<'a, T>(accel: Option<&'a dyn Accelerator>, mut render: im
     Passed { value: render(None), passes: MAX_PASSES + 1, accelerated: false }
 }
 
+/// Ends the frame for the accelerator ([`Accelerator::frame_end`]) when dropped: also when the
+/// render panics or the future is dropped between passes.
+struct FrameEnd<'a>(&'a dyn Accelerator);
+impl Drop for FrameEnd<'_> {
+    fn drop(&mut self) {
+        self.0.frame_end();
+    }
+}
+
 /// A top-level comp frame through [`in_passes`] ([`Renderer::comp_frame`] in passes). Backend
 /// Auto picks CPU or GPU for the whole frame (all its passes) and is timed as a whole, as the
-/// frame workers do.
+/// frame workers do. Ends the frame for the accelerator ([`Accelerator::frame_end`]) on every
+/// exit, including a panic.
 pub async fn comp_frame(r: &Renderer<'_>, comp: ItemId, t: Tick) -> Image {
-    let Some(a) = r.accel.filter(|a| a.miss_gate().is_some()) else { return r.comp_frame(comp, t) };
+    let Some(a) = r.accel else { return r.comp_frame(comp, t) };
+    let _end = FrameEnd(a);
+    if a.miss_gate().is_none() {
+        return r.comp_frame(comp, t);
+    }
     let mut opts = r.opts;
     let mut auto = None;
     if opts.backend == Backend::Auto {

crates/render/src/tests.rs

@@ -838,6 +838,50 @@ impl crate::Accelerator for ProbeAccel {
     }
 }
 
+/// Counts [`Accelerator::frame_end`] and optionally panics inside [`Accelerator::comp_frame`].
+struct EndAccel {
+    ends: std::sync::atomic::AtomicUsize,
+    panic: std::sync::atomic::AtomicBool,
+}
+impl crate::Accelerator for EndAccel {
+    fn name(&self) -> String {
+        "end".into()
+    }
+    fn comp_frame(&self, _r: &crate::Renderer, _comp: ItemId, _t: Tick) -> Option<crate::Image> {
+        if self.panic.load(std::sync::atomic::Ordering::SeqCst) {
+            panic!("comp_frame");
+        }
+        Some(crate::Image::new(1, 1))
+    }
+    fn supports_effect(&self, _id: &str) -> bool {
+        false
+    }
+    fn effects(&self, _chain: &[crate::FxStep], _buf: &crate::Buf, _levels: Option<f32>) -> Option<crate::Buf> {
+        None
+    }
+    fn frame_end(&self) {
+        self.ends.fetch_add(1, std::sync::atomic::Ordering::SeqCst);
+    }
+}
+
+#[test]
+fn comp_frame_ends_the_frame_when_the_render_panics() {
+    use std::panic::{AssertUnwindSafe, catch_unwind};
+    use std::sync::atomic::Ordering;
+    let (mut p, cid, comp) = setup();
+    let l = solid(&mut p, &comp, [1.0, 0.0, 0.0], 50, 50);
+    p.comp_mut(cid).unwrap().layers.push(l);
+    let a = EndAccel { ends: Default::default(), panic: true.into() };
+    let mut r = crate::Renderer::new(&p, &crate::NoFootage, crate::RenderOpts { backend: crate::Backend::Gpu, ..Default::default() });
+    r.accel = Some(&a);
+    assert!(catch_unwind(AssertUnwindSafe(|| crate::passes::block_on(crate::passes::comp_frame(&r, cid, Tick::ZERO)))).is_err());
+    assert_eq!(a.ends.load(Ordering::SeqCst), 1);
+    a.panic.store(false, Ordering::SeqCst);
+    a.ends.store(0, Ordering::SeqCst);
+    let _ = crate::passes::block_on(crate::passes::comp_frame(&r, cid, Tick::ZERO));
+    assert_eq!(a.ends.load(Ordering::SeqCst), 1);
+}
+
 #[test]
 fn auto_backend_sends_2d_and_3d_comps_to_the_gpu() {
     let (mut p, cid, comp) = setup();

crates/ui-egui/src/frames.rs

@@ -39,10 +39,13 @@ use rayon::prelude::*;
 const MAX_KEEPERS: usize = 64;
 
 /// GPU frames kept for RAM preview (bytes of video memory) at first. wgpu can't tell how much
-/// video memory there is: each time the device runs out, the budget halves (to
-/// [`MIN_GPU_BUDGET`]) and the frame renders on the CPU (#106).
+/// video memory there is: allocation pressure halves the budget (to [`MIN_GPU_BUDGET`])
+/// and the frame renders on the CPU (#106). One burst is counted once; successful GPU
+/// renders gradually restore the budget after pressure subsides.
 const GPU_BUDGET: usize = 1 << 30;
 const MIN_GPU_BUDGET: usize = 64 << 20;
+const GPU_PRESSURE_COOLDOWN: std::time::Duration = std::time::Duration::from_secs(5);
+const GPU_RECOVERY_INTERVAL: std::time::Duration = std::time::Duration::from_secs(30);
 
 /// A rendered viewer frame: CPU pixels or a GPU texture.
 #[derive(Clone)]
@@ -151,6 +154,8 @@ struct Cache {
     budget: usize,
     /// Bytes of GPU frames to keep at most ([`GPU_BUDGET`]).
     gpu_budget: usize,
+    gpu_pressure: Option<web_time::Instant>,
+    gpu_adjusted: Option<web_time::Instant>,
     /// The playing preview, if any: evictions spare the frames it shows next.
     playhead: Option<Playhead>,
     /// Scale of the frame on screen (the newest urgent request): frames of other scales go first.
@@ -207,11 +212,28 @@ impl Cache {
             }
         }
     }
-    /// The GPU ran out of memory: keep half as many GPU frames from now on.
+    /// Several workers failing in one burst must not halve the shared budget repeatedly.
     fn gpu_out_of_memory(&mut self) {
-        self.gpu_budget = (self.gpu_budget / 2).max(MIN_GPU_BUDGET);
+        let now = web_time::Instant::now();
+        if self.gpu_pressure.is_none_or(|last| now.duration_since(last) >= GPU_PRESSURE_COOLDOWN) {
+            self.gpu_budget = (self.gpu_budget / 2).max(MIN_GPU_BUDGET);
+            self.gpu_adjusted = Some(now);
+        }
+        self.gpu_pressure = Some(now);
         self.evict_gpu();
     }
+    fn gpu_succeeded(&mut self) {
+        let now = web_time::Instant::now();
+        if self.gpu_pressure.is_some_and(|last| now.duration_since(last) >= GPU_RECOVERY_INTERVAL)
+            && self.gpu_adjusted.is_some_and(|last| now.duration_since(last) >= GPU_RECOVERY_INTERVAL)
+        {
+            self.gpu_budget = self.gpu_budget.saturating_add(MIN_GPU_BUDGET).min(GPU_BUDGET);
+            self.gpu_adjusted = Some(now);
+        }
+    }
+    fn gpu_ready(&self) -> bool {
+        self.gpu_pressure.is_none_or(|last| last.elapsed() >= GPU_PRESSURE_COOLDOWN)
+    }
     fn forget(&mut self, img: &FrameImage) {
         self.bytes -= img.bytes();
         if img.is_gpu() {
@@ -505,7 +527,7 @@ impl Default for Frames {
         #[cfg(target_arch = "wasm32")]
         let queue = Queue::default();
         Frames {
-            cache: Arc::new(Mutex::new(Cache { map: HashMap::new(), order: VecDeque::new(), bytes: 0, gpu_bytes: 0, budget: 3 << 30, gpu_budget: GPU_BUDGET, playhead: None, scale: None })),
+            cache: Arc::new(Mutex::new(Cache { map: HashMap::new(), order: VecDeque::new(), bytes: 0, gpu_bytes: 0, budget: 3 << 30, gpu_budget: GPU_BUDGET, gpu_pressure: None, gpu_adjusted: None, playhead: None, scale: None })),
             inflight: Arc::new(Mutex::new(HashSet::new())),
             queue: Arc::new(Mutex::new(queue)),
             #[cfg(not(target_arch = "wasm32"))]
@@ -914,6 +936,9 @@ impl Frames {
             c.order.clear();
             c.bytes = 0;
             c.gpu_bytes = 0;
+            c.gpu_budget = GPU_BUDGET;
+            c.gpu_pressure = None;
+            c.gpu_adjusted = None;
         }
         // Disk keys hash the footage files' size and time, but are remembered per revision; a
         // file reloaded from disk without a new revision (a re-rendered image sequence) must be
@@ -1382,10 +1407,15 @@ impl Worker {
     /// `None` if that panicked too.
     fn render_guarded(&self, job: &Job) -> Option<FrameImage> {
         let attempt = |gpu: bool| std::panic::catch_unwind(std::panic::AssertUnwindSafe(|| self.render_job(job, gpu))).ok();
-        attempt(true).or_else(|| {
+        let img = attempt(true).or_else(|| {
             log::error!("a frame render panicked (frame {})", job.key.frame);
             job.src.gpu.as_ref().and_then(|_| attempt(false))
-        })
+        });
+        // A preview frame is done: the GPU pool ages the sizes it did not ask for.
+        if let Some(g) = &job.src.gpu {
+            effectcraft_engine::render::Accelerator::frame_end(g);
+        }
+        img
     }
 
     /// A failed scoped attempt retries on CPU. Retired contexts cannot recover by reducing
@@ -1406,9 +1436,19 @@ impl Worker {
         }
     }
 
+    fn gpu_succeeded(&self) {
+        if let Ok(mut c) = self.cache.lock() {
+            c.gpu_succeeded();
+        }
+    }
+
+    fn gpu_ready(&self) -> bool {
+        !self.gpu_retired.load(Ordering::Acquire) && self.cache.lock().is_ok_and(|c| c.gpu_ready())
+    }
+
     /// The frame on the GPU, kept there for the viewer, if the compositor can do it alone.
     fn render_gpu(&self, job: &Job) -> Option<FrameImage> {
-        if self.gpu_retired.load(Ordering::Acquire) {
+        if !self.gpu_ready() {
             return None;
         }
         let g = job.src.gpu.as_ref()?;
@@ -1418,8 +1458,15 @@ impl Worker {
         r.cache = Some(&job.src.layer_cache);
         r.accel = Some(g as &dyn effectcraft_engine::render::Accelerator);
         r.active_accel()?;
-        match g.within_memory(|| g.render_display(&r, job.comp, job.t)) {
-            Ok(f) => f.map(|f| FrameImage::Gpu(Arc::new(f))),
+        let f = g.within_memory(|| g.render_display(&r, job.comp, job.t));
+        effectcraft_engine::render::Accelerator::frame_end(g);
+        match f {
+            Ok(f) => {
+                if f.is_some() {
+                    self.gpu_succeeded();
+                }
+                f.map(|f| FrameImage::Gpu(Arc::new(f)))
+            }
             Err(e) => {
                 self.gpu_attempt_failed(g, &e);
                 None
@@ -1432,7 +1479,7 @@ impl Worker {
         let footage = job.footage();
         let mut r = Renderer::new(&job.src.project, &footage, job.opts);
         r.expr = job.src.expr.as_deref();
-        let gpu_allowed = gpu && !self.gpu_retired.load(Ordering::Acquire);
+        let gpu_allowed = gpu && self.gpu_ready();
         // A failed GPU attempt may have populated intermediate entries before its scope
         // reported failure. The CPU retry must compute its own pixels from the snapshot.
         r.cache = gpu_allowed.then_some(job.src.layer_cache.as_ref());
@@ -1452,6 +1499,7 @@ impl Worker {
                 let t0 = web_time::Instant::now();
                 match g.within_memory(|| g.render_display(&r, job.comp, job.t)) {
                     Ok(Some(f)) => {
+                        self.gpu_succeeded();
                         if let Some((p, key)) = auto {
                             // Time the GPU's work, not just its recording.
                             g.wait();
@@ -1483,7 +1531,10 @@ impl Worker {
         if let Some(g) = gpu {
             // The GPU composites and the frame is read back.
             return match g.within_memory(|| r.comp_frame(job.comp, job.t)) {
-                Ok(img) => FrameImage::Cpu(Arc::new(to_color_image(&img))),
+                Ok(img) => {
+                    self.gpu_succeeded();
+                    FrameImage::Cpu(Arc::new(to_color_image(&img)))
+                }
                 Err(e) => {
                     self.gpu_attempt_failed(g, &e);
                     self.render_job(job, false)

crates/ui-egui/src/frames/tests.rs

@@ -469,18 +469,36 @@ fn a_newer_viewer_frame_drops_the_ones_still_waiting() {
     assert_eq!(f.inflight(), 3, "dropped frames aren't in flight: the viewer can ask again");
 }
 
-/// Each time the GPU runs out of memory, half as many GPU frames are kept, down to a floor.
+/// One OOM burst halves once; separate episodes back off to a floor and healthy GPU work recovers.
 #[test]
 fn running_out_of_video_memory_halves_the_gpu_frame_budget() {
     let f = Frames::default();
     let mut c = f.cache.lock().unwrap();
     assert_eq!(c.gpu_budget, GPU_BUDGET);
+    assert!(c.gpu_ready());
     c.gpu_out_of_memory();
+    assert!(!c.gpu_ready(), "allocation failures must not be retried for every frame");
     assert_eq!(c.gpu_budget, GPU_BUDGET / 2);
     for _ in 0..20 {
         c.gpu_out_of_memory();
     }
+    assert_eq!(c.gpu_budget, GPU_BUDGET / 2, "one burst must not collapse the shared budget");
+    for _ in 0..20 {
+        c.gpu_pressure = None;
+        c.gpu_out_of_memory();
+    }
     assert_eq!(c.gpu_budget, MIN_GPU_BUDGET);
+    let before = web_time::Instant::now() - GPU_RECOVERY_INTERVAL;
+    c.gpu_pressure = Some(before);
+    assert!(c.gpu_ready(), "GPU attempts resume after the cooldown");
+    c.gpu_adjusted = Some(before);
+    c.gpu_succeeded();
+    assert_eq!(c.gpu_budget, 2 * MIN_GPU_BUDGET);
+    c.gpu_succeeded();
+    assert_eq!(c.gpu_budget, 2 * MIN_GPU_BUDGET, "recovery is paced");
+    drop(c);
+    f.clear();
+    assert_eq!(f.cache.lock().unwrap().gpu_budget, GPU_BUDGET);
 }
 
 #[test]

crates/ui-egui/src/theme.rs

@@ -352,6 +352,7 @@ static NOTO_SYMBOLS2: &[u8] = include_bytes!("../../../assets/fonts/NotoSansSymb
 /// frame loop), so the fonts do not need a restart.
 pub fn install(ctx: &egui::Context, t: &Tokens, language: &str) {
     let mut fonts = FontDefinitions::default();
+    let proportional_fallbacks = fonts.families.get(&FontFamily::Proportional).cloned().unwrap_or_default();
     fonts.font_data.insert("inter".into(), Arc::new(FontData::from_static(INTER_REGULAR)));
     fonts.font_data.insert("inter-medium".into(), Arc::new(FontData::from_static(INTER_MEDIUM)));
     fonts.font_data.insert("inter-semibold".into(), Arc::new(FontData::from_static(INTER_SEMIBOLD)));
@@ -359,8 +360,13 @@ pub fn install(ctx: &egui::Context, t: &Tokens, language: &str) {
     fonts.font_data.insert("noto-symbols2".into(), Arc::new(FontData::from_static(NOTO_SYMBOLS2)));
     fonts.families.entry(FontFamily::Proportional).or_default().insert(0, "inter".into());
     fonts.families.entry(FontFamily::Monospace).or_default().insert(0, "jbmono".into());
-    fonts.families.insert(FontFamily::Name("semibold".into()), vec!["inter-semibold".into(), "inter".into()]);
-    fonts.families.insert(FontFamily::Name("medium".into()), vec!["inter-medium".into(), "inter".into()]);
+    // Named weights need the same symbol fallback chain as ordinary labels.
+    // The bundled Inter subset and JetBrains Mono do not cover every UI symbol.
+    for (family, face) in [("semibold", "inter-semibold"), ("medium", "inter-medium")] {
+        let mut stack = vec![face.into(), "inter".into()];
+        stack.extend(proportional_fallbacks.iter().cloned());
+        fonts.families.insert(FontFamily::Name(family.into()), stack);
+    }
     // Inter and egui's own fallbacks have no small triangles (the `▸` of "File ▸ Import"), which
     // would draw as missing-glyph boxes: JetBrains Mono fills in the symbols the others lack.
     // Noto Sans Symbols 2 follows it for the check/box glyphs (`✔`, `▣`) that JetBrains Mono

docs/hardware-parity.md

@@ -1,6 +1,6 @@
 # Hardware parity
 
-> **Last reviewed:** 2026-10-12 · **Last updated:** 2026-10-12 · **Change:** minor (closed issue references re-audited against `upstream/main` 470ea7a3) · **Target:** Adobe After Effects 2026
+> **Last reviewed:** 2026-10-29 · **Last updated:** 2026-10-29 · **Change:** minor (Content-Aware Fill's PatchMatch and flow search on the GPU) · **Target:** Adobe After Effects 2026
 
 Each piece of hardware After Effects uses, per platform, against EffectCraft. Ours comes from the
 code at `origin/main` 1446e27 (`crates/gpu`, `crates/render`, `crates/ui-egui`) and the issue
@@ -14,6 +14,7 @@ monitors are missing, and two GPU start-up failures are open.
 | Feature | After Effects 2026 | EffectCraft macOS | Windows | Linux | Web | Notes / evidence |
 |---|---|---|---|---|---|---|
 | GPU compositing and effects | Mercury GPU (Metal; CUDA / OpenCL / DirectX on Windows) | yes, Metal via wgpu | Vulkan yes; **Direct3D 12 fails** (FXC can't compile the `pointwise` kernel, #613) | Vulkan | WebGPU, in frame workers | 280 of 306 effects on the GPU, checked against the CPU renderer; Fractal and the analysis/tool effects stay on the CPU ([effects-parity.md](effects-parity.md)) |
+| GPU analysis | GPU-accelerated Content-Aware Fill | Content-Aware Fill's PatchMatch **and** flow block matching on the GPU (`Gpu::inpaint`); every other stage on the CPU | same | same | no (needs synchronous readback) | On an Intel Arc iGPU a four-frame 1920 × 1080 fill is ≈ 1.4× faster than the CPU's (PatchMatch is the dominant win, the flow search a smaller one), and even a 16 × 16 hole at 1920 × 1080 is ≈ 1.9× faster — the CPU's own stages are O(frame), not O(hole). The device's flow field matches `block_flow` to a float rounding. The membrane and the flow's hole completion stay on the CPU: both are bandwidth-bound, and the device's full-frame transfers cost more than they save — see `effectcraft_gpu::inpaint` |
 | CPU fallback and device-loss recovery | yes | yes | yes; Windows ARM64 Adreno panic after device loss (#591, closed) | yes | yes (#225 fixed) | Out of video memory falls back to the CPU (#106) |
 | GPU detection and choice | GPU info, renderer choice | Backend Auto per comp | "GPU isn't being detected" (#596, closed); Intel HD 5500 fix unverified (#198) | yes | yes | No settings page listing adapters yet |
 | Multi-core rendering | Multi-Frame Rendering | rayon pools; preview scheduler sized to the render pool | same | same | Web Workers, one engine per worker | No benchmark against Multi-Frame Rendering |
@@ -39,5 +40,6 @@ the Adreno and Intel cases.
 
 | Date | Change | Summary |
 |---|---|---|
+| 2026-10-29 | minor | Content-Aware Fill's PatchMatch and flow block matching run on the GPU (`FillBackend`/`Accelerator::inpaint`, `effectcraft_gpu::inpaint`) |
 | 2026-10-12 | minor | #591, #596 and #554 marked closed; #498 answer and open PR #819 |
 | 2026-10-10 | major | First hardware parity document |