diff --git a/Cargo.lock b/Cargo.lock index 288cce1ad..232fa6b18 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1612,6 +1612,7 @@ dependencies = [ "js-sys", "log", "pollster 1.0.1", + "rayon", "wasm-bindgen", "web-time", "wgpu", diff --git a/crates/effects/src/lib.rs b/crates/effects/src/lib.rs index c865d4f96..9130b061c 100644 --- a/crates/effects/src/lib.rs +++ b/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. diff --git a/crates/effects/src/stylize3.rs b/crates/effects/src/stylize3.rs index e28f9a4fa..e8a141cf4 100644 --- a/crates/effects/src/stylize3.rs +++ b/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); } }); diff --git a/crates/effects/src/transition.rs b/crates/effects/src/transition.rs index 23e54bb6b..e465ffd1b 100644 --- a/crates/effects/src/transition.rs +++ b/crates/effects/src/transition.rs @@ -314,53 +314,86 @@ pub(crate) fn card_wipe_plan(ctx: &EffectCtx, b: &Buf) -> Option { 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 diff --git a/crates/engine/src/background/offload.rs b/crates/engine/src/background/offload.rs index 23cff4da6..f9038accb 100644 --- a/crates/engine/src/background/offload.rs +++ 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 |_| {}, ) diff --git a/crates/engine/src/background/remote.rs b/crates/engine/src/background/remote.rs index 6539ac767..27f9396e1 100644 --- a/crates/engine/src/background/remote.rs +++ b/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)), diff --git a/crates/engine/src/background/remote/tests.rs b/crates/engine/src/background/remote/tests.rs index 85f7a1807..e2db10b8e 100644 --- a/crates/engine/src/background/remote/tests.rs +++ b/crates/engine/src/background/remote/tests.rs @@ -192,6 +192,8 @@ struct Deferred { asked: std::sync::Mutex>, gate: Arc, 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)); +} diff --git a/crates/engine/src/cmd/content_fill.rs b/crates/engine/src/cmd/content_fill.rs index d96be073d..967651a37 100644 --- a/crates/engine/src/cmd/content_fill.rs +++ b/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 &'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, 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 { 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 { &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 { #[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] diff --git a/crates/engine/src/settings/prefs/document.rs b/crates/engine/src/settings/prefs/document.rs index 1f68c4a06..5db41e554 100644 --- a/crates/engine/src/settings/prefs/document.rs +++ b/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 = 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::(Value::Object(trial)).is_ok() { diff --git a/crates/engine/src/tests/settings.rs b/crates/engine/src/tests/settings.rs index 5e38c90a3..7d65061ea 100644 --- a/crates/engine/src/tests/settings.rs +++ b/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#"{ diff --git a/crates/export/src/lib.rs b/crates/export/src/lib.rs index e04c465f9..1d1b3de1f 100644 --- a/crates/export/src/lib.rs +++ b/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 { 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 { let comp = cx.comp().ok_or(ExportError::NoComp)?; let total = cx.settings.frame_count(comp); diff --git a/crates/export/src/pipeline.rs b/crates/export/src/pipeline.rs index 8e8c90561..b0d1d96c9 100644 --- a/crates/export/src/pipeline.rs +++ b/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())); } diff --git a/crates/export/tests/it/h264_gpu_yuv.rs b/crates/export/tests/it/h264_gpu_yuv.rs index 0f99a7b54..02963ef46 100644 --- a/crates/export/tests/it/h264_gpu_yuv.rs +++ b/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); } diff --git a/crates/gpu/Cargo.toml b/crates/gpu/Cargo.toml index 49fe51c16..01154dd9f 100644 --- a/crates/gpu/Cargo.toml +++ b/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] diff --git a/crates/gpu/src/adv3d.rs b/crates/gpu/src/adv3d.rs index c96e57d2d..b88322316 100644 --- a/crates/gpu/src/adv3d.rs +++ b/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> { 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 GpuImage { + e.uses(canvas); let out = e.image(canvas.width, canvas.height); e.copy_into(canvas, &out, 0, 0); if pixels.is_empty() { diff --git a/crates/gpu/src/context.rs b/crates/gpu/src/context.rs index 44e777b5a..cc861c85b 100644 --- a/crates/gpu/src/context.rs +++ b/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>, /// 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>, + /// 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, 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, + /// 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, /// 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) -> 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 { 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 { - 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>, readbacks: Arc, + /// 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>, + 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 { 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 { + 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, 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, 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 { + 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> { + 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> { 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> { + 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}; diff --git a/crates/gpu/src/effects.rs b/crates/gpu/src/effects.rs index 5a2b80f1c..b33c15e99 100644 --- a/crates/gpu/src/effects.rs +++ b/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; diff --git a/crates/gpu/src/fx_pixel2.rs b/crates/gpu/src/fx_pixel2.rs index acba936d0..a62cbaff0 100644 --- a/crates/gpu/src/fx_pixel2.rs +++ b/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 { - 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 { diff --git a/crates/gpu/src/fx_tone.rs b/crates/gpu/src/fx_tone.rs index cb7144114..f6e7062e2 100644 --- a/crates/gpu/src/fx_tone.rs +++ b/crates/gpu/src/fx_tone.rs @@ -174,6 +174,62 @@ fn parse_map(s: &str) -> Vec { 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> { + 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 { 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 { } "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]; diff --git a/crates/gpu/src/fx_transition.rs b/crates/gpu/src/fx_transition.rs index 60378491e..b5cb07012 100644 --- a/crates/gpu/src/fx_transition.rs +++ b/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 { 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 { diff --git a/crates/gpu/src/inpaint.rs b/crates/gpu/src/inpaint.rs new file mode 100644 index 000000000..25cbc31b9 --- /dev/null +++ b/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 { + 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::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 = Vec::with_capacity(knowns.iter().map(|k| k.len()).sum()); + let mut offs: Vec = 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> = 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 = 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 { + 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 = 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", ¶ms([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 = 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", + ¶ms([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", + ¶ms([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 { + 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 = 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 = 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", ¶ms([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 = 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", ¶ms([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", ¶ms([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 { + 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 = (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", ¶ms([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 { + 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 = 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 +} diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs index c58ccd59e..0844a4335 100644 --- a/crates/gpu/src/lib.rs +++ b/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) -> Option { 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 { @@ -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(); diff --git a/crates/gpu/src/memcap.rs b/crates/gpu/src/memcap.rs new file mode 100644 index 000000000..1c2ce6456 --- /dev/null +++ b/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::().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"); + } +} diff --git a/crates/gpu/src/shaders/common.wgsl b/crates/gpu/src/shaders/common.wgsl index 2c65c9d4a..1d8a701b6 100644 --- a/crates/gpu/src/shaders/common.wgsl +++ b/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 { + if (!(abs(x) <= 8192.0)) { + return vec2(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(cs, -sn); } + case 2: { return vec2(-sn, -cs); } + case 3: { return vec2(-cs, sn); } + default: { return vec2(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(x); + let ab = bitcast(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(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(q); } + if (q == 0x1000000u) { q >>= 1u; exponent += 1; } + return bitcast((u32(exponent + 127) << 23u) | (q & 0x7fffffu)); +} + // ---------------------------------------------------------------- sampling (raster::Image) fn tex_get(t: texture_2d, x: i32, y: i32) -> vec4 { @@ -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). diff --git a/crates/gpu/src/shaders/fx_pixel2.wgsl b/crates/gpu/src/shaders/fx_pixel2.wgsl index 2c6f5e50c..583225299 100644 --- a/crates/gpu/src/shaders/fx_pixel2.wgsl +++ b/crates/gpu/src/shaders/fx_pixel2.wgsl @@ -633,52 +633,9 @@ fn fp2_hextile(@builtin(global_invocation_id) gid: vec3) { 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) diff --git a/crates/gpu/src/shaders/fx_sim.wgsl b/crates/gpu/src/shaders/fx_sim.wgsl index 0c6023f32..7deeac048 100644 --- a/crates/gpu/src/shaders/fx_sim.wgsl +++ b/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; diff --git a/crates/gpu/src/shaders/fx_stylize.wgsl b/crates/gpu/src/shaders/fx_stylize.wgsl index 7fe42df7a..1fa8134fc 100644 --- a/crates/gpu/src/shaders/fx_stylize.wgsl +++ b/crates/gpu/src/shaders/fx_stylize.wgsl @@ -42,7 +42,11 @@ fn fxs_qdiv(x: vec3, a: f32) -> vec3 { let r = fma(-q, vec3(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(1.0), q); + let corrected = fma(r / a, vec3(1.0), q); + return vec3( + 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(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) { switch mode { case 1u: { let lv = P.f[0].x; - let c = clamp(fxs_qdiv(px.xyz, a), vec3(0.0), vec3(1.0)) * lv; - let o = vec3(round_away(c.x), round_away(c.y), round_away(c.z)) / lv; + let c = vec3(unit_div(px.x, a), unit_div(px.y, a), unit_div(px.z, a)); + let o = vec3(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(o * a, a)); } case 2u: { diff --git a/crates/gpu/src/shaders/fx_tone.wgsl b/crates/gpu/src/shaders/fx_tone.wgsl index c99d199c5..bb206adbd 100644 --- a/crates/gpu/src/shaders/fx_tone.wgsl +++ b/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) { case 12u: { let n = u0.z; let aa = max(a, 0.0); - let r = vec3(fxt_lookup(c.x, n), fxt_lookup(c.y, n), fxt_lookup(c.z, n)); + var straight = vec3(0.0); + if (a > 1e-6) { + straight = vec3(unit_div(px.x, a), unit_div(px.y, a), unit_div(px.z, a)); + } + let r = vec3(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); diff --git a/crates/gpu/src/shaders/fx_transition.wgsl b/crates/gpu/src/shaders/fx_transition.wgsl index 5037ec9a6..0e9553d58 100644 --- a/crates/gpu/src/shaders/fx_transition.wgsl +++ b/crates/gpu/src/shaders/fx_transition.wgsl @@ -224,28 +224,8 @@ fn ftr_grid(@builtin(global_invocation_id) gid: vec3) { 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(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); } diff --git a/crates/gpu/src/shaders/fx_vr.wgsl b/crates/gpu/src/shaders/fx_vr.wgsl index a2bb364c8..9e2d4390e 100644 --- a/crates/gpu/src/shaders/fx_vr.wgsl +++ b/crates/gpu/src/shaders/fx_vr.wgsl @@ -16,7 +16,7 @@ fn fxv_unpremul(p: vec4) -> vec4 { fn fxv_equi_dir(x: f32, y: f32, w: f32, h: f32) -> vec3 { let lon = (x / w - 0.5) * 2.0 * FXV_PI; let lat = (0.5 - y / h) * FXV_PI; - return vec3(cos(lat) * sin(lon), sin(lat), cos(lat) * cos(lon)); + return vec3(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 { cu = u / r; cv = v / r; } - return vec4(sin(th) * cu, cos(th), sin(th) * cv, 1.0); + return vec4(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 { cu = u / r; cv = v / r; } - return vec4(sin(th) * cu, sin(th) * cv, cos(th), 1.0); + return vec4(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) { 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); } diff --git a/crates/gpu/src/shaders/inpaint.wgsl b/crates/gpu/src/shaders/inpaint.wgsl new file mode 100644 index 000000000..56c468f85 --- /dev/null +++ b/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(tx2, ty2), 0); + let b = textureLoad(aux, vec2(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) { + 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(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(fx, fy), 0); + n = n + 1.0; + } + } + } + var v = vec4(0.0); + if (n > 0.0) { + v = s / n; + } + textureStore(out, vec2(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) { + 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(x, y), textureLoad(src, vec2(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) { + 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(x, y), 0); + if (data[i32(P.u[3].x) + y * w + x] <= 0.5) { + var s = vec4(0.0); + var n = 0.0; + if (x > 0) { + s = s + textureLoad(src, vec2(x - 1, y), 0); + n = n + 1.0; + } + if (x + 1 < w) { + s = s + textureLoad(src, vec2(x + 1, y), 0); + n = n + 1.0; + } + if (y > 0) { + s = s + textureLoad(src, vec2(x, y - 1), 0); + n = n + 1.0; + } + if (y + 1 < h) { + s = s + textureLoad(src, vec2(x, y + 1), 0); + n = n + 1.0; + } + if (n > 0.0) { + v = s / n; + } + } + textureStore(out, vec2(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) { + 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(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(fx, fy), 0); + n = n + 1.0; + } + } + } + var v = vec4(0.0); + if (n > 0.0) { + v = s / n; + } + textureStore(out, vec2(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) { + 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(x0, y0), 0); + let b = textureLoad(src, vec2(x1, y0), 0); + let c = textureLoad(src, vec2(x0, y1), 0); + let d = textureLoad(src, vec2(x1, y1), 0); + let top = a + (b - a) * tx; + let bot = c + (d - c) * tx; + textureStore(out, vec2(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) { + 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(cx, cy), 0); + var o = vec4(0.0); + if (c.a > 0.5) { + o = vec4(c.r * 2.0 + f32(x % 2), c.g * 2.0 + f32(y % 2), 0.0, 1.0); + } + textureStore(out, vec2(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 { + 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(x, y); + } + } + return vec2(-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) { + 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(0.0, 0.0, IMP_INF, 0.0); + if (imp_near(i)) { + let prev = textureLoad(src, vec2(x, y), 0); + var s = vec2(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(s.x), f32(s.y), d, 1.0); + } + } + textureStore(out, vec2(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) { + 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(x, y), 0); + let i = y * w + x; + if (!imp_near(i) || ((x + y) & 1) != i32(P.u[0].w)) { + textureStore(out, vec2(x, y), cur); + return; + } + var best = vec2(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(nx, ny), 0); + if (n.a <= 0.5) { + continue; + } + let c = vec2(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(x, y), vec4(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) { + 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(x, y), 0); + let i = y * w + x; + if (!imp_near(i) || cur.a <= 0.5) { + textureStore(out, vec2(x, y), cur); + return; + } + var best = vec2(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(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(x, y), vec4(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) { + 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(x, y), 0); + if (imp_near(i)) { + let f = textureLoad(src, vec2(x, y), 0); + o = vec4(f.r, f.g, 0.0, 1.0); + } + textureStore(out, vec2(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) { + 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(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(px, py), 0); + if (f.a <= 0.5 || f.b >= IMP_INF) { + continue; + } + let s = vec2(i32(f.r) - dx, i32(f.g) - dy); + let wt = exp(-f.b / (2.0 * sigma2)); + acc = acc + textureLoad(aux, vec2(s.x, s.y), 0) * wt; + wsum = wsum + wt; + } + } + if (wsum > 0.0) { + textureStore(out, vec2(x, y), acc / wsum); + } else { + textureStore(out, vec2(x, y), textureLoad(aux, vec2(x, y), 0)); + } +} + +// The workgroup's 1024 distance bins (`imp_hist`); the name is unique across the concatenated +// shader sources because `var` is module scope. +var imp_hist_bins: array, 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 imp_hist_max: atomic; + +// 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) { + 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(x, y), 0); + if (f.a <= 0.5 || f.b >= IMP_INF) { + continue; + } + atomicMax(&imp_hist_max, bitcast(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(bin), 0), vec4(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(0, 0), vec4(f32(atomicLoad(&imp_hist_bins[0])), bitcast(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) { + 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(min(2 * cx + dx, w - 1), min(2 * cy + dy, h - 1)), 0); + sx = sx + v.x; + sy = sy + v.y; + } + } + textureStore(out, vec2(cx, cy), vec4(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 { + let v = textureLoad(src, clamp(vec2(x, y), vec2(0, 0), vec2(w - 1, h - 1)), 0); + return vec2(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 { + 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) { + 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, 6>; + var ns = 0; + seeds[ns] = vec2(0, 0); + ns = ns + 1; + if (P.u[2].x != 0u) { + var nbr = array, 5>(vec2(0, 0), vec2(-1, 0), vec2(1, 0), vec2(0, -1), vec2(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(nc, nr), 0); + seeds[ns] = vec2(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; + var la8: array; + 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(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(d0, d1); + } + } + } + } + textureStore(out, vec2(c, r), vec4(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) { + 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(c, r), 0); + let d = vec2(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; + 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(c, r), vec4(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) { + 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; + var ys: array; + 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(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(c, r), vec4(xs[n / 2], ys[n / 2], 0.0, 0.0)); +} diff --git a/crates/gpu/src/shaders/particles.wgsl b/crates/gpu/src/shaders/particles.wgsl index 694d6fb67..ac12b8fc7 100644 --- a/crates/gpu/src/shaders/particles.wgsl +++ b/crates/gpu/src/shaders/particles.wgsl @@ -53,12 +53,12 @@ fn ps_norm3(v: vec3) -> vec3 { fn ps_rand_dir(i: u32, k: u32, flat: bool) -> vec3 { if (flat) { let a = ps_h(i, k) * PS_TAU; - return vec3(cos(a), sin(a), 0.0); + return vec3(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(r * cos(a), r * sin(a), z); + return vec3(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(sin(a), -cos(a)); + let d = vec2(sin_p(a), -cos_p(a)); let speed = max(vel + ps_hs(id, 3u) * vel_spread, 0.0); var off: vec2; 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(cos(aa) * rr, sin(aa) * rr); + off = vec2(cos_p(aa) * rr, sin_p(aa) * rr); } let v = d * speed; let dt = S.f[2].z; diff --git a/crates/gpu/src/tests.rs b/crates/gpu/src/tests.rs index 44e8922a6..35ad402ba 100644 --- a/crates/gpu/src/tests.rs +++ b/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) { + let p = vec2(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(q, fxs_posterize_bin(q, P.u[0].y), 0.0, 1.0)); + return; + } + textureStore(out, p, vec4(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", ¶ms, &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()?; + 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) = (¶ms, &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"); + }); +} diff --git a/crates/gpu/src/tests/fx_stylize.rs b/crates/gpu/src/tests/fx_stylize.rs index df99e1ced..d9ca95712 100644 --- a/crates/gpu/src/tests/fx_stylize.rs +++ b/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: ¶ms, 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::>().join(","))), + ("phase".into(), n(phase)), + ("applyPhaseMapToAlpha".into(), on()), + ] + .into(), + }; + let ctx = + || effectcraft_effects::EffectCtx { params: ¶ms, 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: ¶ms, 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). diff --git a/crates/gpu/src/tests/inpaint.rs b/crates/gpu/src/tests/inpaint.rs new file mode 100644 index 000000000..e00d3cc18 --- /dev/null +++ b/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 { + (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(>, &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(>, &hole); + let mem = membrane_fill(&src, &hole); + let pm = g.patch(&src, &hole, None, 7, 1); + let (pm_db, mem_db) = (psnr(&pm, >, Some(&hole)), psnr(&mem, >, 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(>, &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, >, 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(>, &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, >, Some(&hole)) > 25.0, "{:.1} dB", psnr(&gpu_pm, >, 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(>, &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, >, 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(>, &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], >s[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 = 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(>, &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(>, &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, >, 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(>, &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"); + } + } +} diff --git a/crates/gpu/src/walk.rs b/crates/gpu/src/walk.rs index fdcb58dd4..c7212dc84 100644 --- a/crates/gpu/src/walk.rs +++ b/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(()); diff --git a/crates/gpu/src/yuv.rs b/crates/gpu/src/yuv.rs index aba3ec700..27df90110 100644 --- a/crates/gpu/src/yuv.rs +++ b/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); diff --git a/crates/host/tests/it/plugins.rs b/crates/host/tests/it/plugins.rs index 83891148d..593e1ffcd 100644 --- a/crates/host/tests/it/plugins.rs +++ b/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 { 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; diff --git a/crates/raster/src/flow.rs b/crates/raster/src/flow.rs index feb033d65..180b235bf 100644 --- a/crates/raster/src/flow.rs +++ b/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, diff --git a/crates/raster/src/inpaint.rs b/crates/raster/src/inpaint.rs index c64de3375..592ecd221 100644 --- a/crates/raster/src/inpaint.rs +++ b/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> { + 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> { 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 = frames.par_iter().zip(&holes).map(|(f, hl)| membrane_fill(f, hl)).collect(); + let pre: Vec = 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, Vec) = 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; diff --git a/crates/render/src/accel.rs b/crates/render/src/accel.rs index e6bbd2113..fe5a8f3fe 100644 --- a/crates/render/src/accel.rs +++ b/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 } diff --git a/crates/render/src/passes.rs b/crates/render/src/passes.rs index 5f0373b84..37c94fff4 100644 --- a/crates/render/src/passes.rs +++ b/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 { diff --git a/crates/render/src/tests.rs b/crates/render/src/tests.rs index 865a4f59f..c77fc2d94 100644 --- a/crates/render/src/tests.rs +++ b/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 { + 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) -> Option { + 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(); diff --git a/crates/ui-egui/src/frames.rs b/crates/ui-egui/src/frames.rs index 87cd9a17b..6a98abfe6 100644 --- a/crates/ui-egui/src/frames.rs +++ b/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, + gpu_adjusted: Option, /// The playing preview, if any: evictions spare the frames it shows next. playhead: Option, /// 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 { 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 { - 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) diff --git a/crates/ui-egui/src/frames/tests.rs b/crates/ui-egui/src/frames/tests.rs index 692df16f7..ace51b0bb 100644 --- a/crates/ui-egui/src/frames/tests.rs +++ b/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] diff --git a/crates/ui-egui/src/theme.rs b/crates/ui-egui/src/theme.rs index aa1b30350..e67bd2d7d 100644 --- a/crates/ui-egui/src/theme.rs +++ b/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 diff --git a/docs/hardware-parity.md b/docs/hardware-parity.md index c25da4928..998fe4453 100644 --- a/docs/hardware-parity.md +++ b/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 |