diff --git a/Cargo.lock b/Cargo.lock index 6f31e68d..3def0e94 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1157,6 +1157,7 @@ dependencies = [ "cpu-time", "effectcraft-automation", "effectcraft-engine", + "effectcraft-gpu", "effectcraft-host", "effectcraft-render", "effectcraft-time", @@ -1256,6 +1257,25 @@ dependencies = [ "serde", ] +[[package]] +name = "effectcraft-gpu" +version = "0.1.0" +dependencies = [ + "bytemuck", + "effectcraft-color", + "effectcraft-effects", + "effectcraft-geom", + "effectcraft-keyframe", + "effectcraft-project", + "effectcraft-raster", + "effectcraft-render", + "effectcraft-text", + "effectcraft-time", + "log", + "pollster 0.4.0", + "wgpu", +] + [[package]] name = "effectcraft-host" version = "0.1.0" @@ -1406,6 +1426,7 @@ dependencies = [ "effectcraft-color", "effectcraft-engine", "effectcraft-geom", + "effectcraft-gpu", "effectcraft-host", "effectcraft-keyframe", "effectcraft-path", diff --git a/Cargo.toml b/Cargo.toml index 8212d74a..a1941cf4 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -23,6 +23,7 @@ effectcraft-text = { path = "crates/text" } effectcraft-effects = { path = "crates/effects" } effectcraft-track = { path = "crates/track" } effectcraft-render = { path = "crates/render" } +effectcraft-gpu = { path = "crates/gpu" } effectcraft-media = { path = "crates/media" } effectcraft-export = { path = "crates/export" } effectcraft-expr = { path = "crates/expr" } @@ -66,6 +67,11 @@ filmcraft-prores = { git = "https://github.com/storytold/filmcraft", rev = "89cf filmcraft-aac = { git = "https://github.com/storytold/filmcraft", rev = "89cfbf24f698b0f0e538eb7795929f19bbfb1dbb" } filmcraft-isobmff = { git = "https://github.com/storytold/filmcraft", rev = "89cfbf24f698b0f0e538eb7795929f19bbfb1dbb" } +# gpu (the same wgpu eframe / egui-wgpu use) +wgpu = { version = "30.0.1", features = ["fragile-send-sync-non-atomic-wasm"] } +bytemuck = "1" +pollster = "0.4" + # ui egui = "0.36" eframe = { version = "0.36", default-features = false } @@ -108,6 +114,9 @@ opt-level = 3 [profile.dev.package.effectcraft-render] opt-level = 3 +[profile.dev.package.effectcraft-gpu] +opt-level = 3 + [profile.dev.package.effectcraft-effects] opt-level = 3 diff --git a/apps/effectcraft-cli/Cargo.toml b/apps/effectcraft-cli/Cargo.toml index 2adf0873..ff8842d9 100644 --- a/apps/effectcraft-cli/Cargo.toml +++ b/apps/effectcraft-cli/Cargo.toml @@ -17,6 +17,7 @@ effectcraft-engine = { workspace = true } effectcraft-host = { workspace = true } effectcraft-automation = { workspace = true } effectcraft-render = { workspace = true } +effectcraft-gpu = { workspace = true } effectcraft-time = { workspace = true } serde_json = { workspace = true } image = { workspace = true } diff --git a/apps/effectcraft-cli/src/main.rs b/apps/effectcraft-cli/src/main.rs index 43e75837..92e5c1fc 100644 --- a/apps/effectcraft-cli/src/main.rs +++ b/apps/effectcraft-cli/src/main.rs @@ -13,11 +13,14 @@ //! [--work-area] [--fps N] [--resolution full|half|third|quarter|K] [--quality best|draft] [--channels rgb|rgba] //! [--jpeg-quality N] [--bitrate KBPS] [--prores proxy|lt|standard|hq|4444|4444xq] [--audio auto|on|off] //! effectcraft-cli render F.ecproj --queue render the project's Render Queue -//! effectcraft-cli bench [--comp C] [--time S] [--scale K] [--n N] [--play N] render timings +//! effectcraft-cli bench [--comp C] [--time S] [--scale K] [--n N] [--play N] [--gpu] render timings +//! (--gpu: CPU vs GPU ms/frame for every comp at Full and Half) //! effectcraft-cli mcp [--bridge PORT] MCP server on stdio //! //! Project: --project F.ecproj (or a positional *.ecproj) | --demo | --empty (default: demo; mcp: empty) //! Saving: --save (back to --project) | --save-as F.ecproj +//! GPU: --gpu renders on the GPU compositor (Mercury GPU Acceleration) when an adapter exists; +//! the default is the CPU (Mercury Software Only) //! Bridge: --bridge PORT drives a running `effectcraft --control PORT` instead of a headless session //! Output: --json for one compact JSON document on stdout (errors: {"error": …}, exit 1) //! ``` @@ -41,10 +44,11 @@ const USAGE: &str = "usage: effectcraft-cli : id or name, '-' = active comp; : id, '#n' or name; : JSON or bare string"; /// Options that take a value. @@ -196,7 +200,7 @@ fn backend(args: &Args, default_demo: bool) -> Result { } return Ok(Backend::bridge(addr)?); } - let mut b = Backend::headless(effectcraft_host::session()); + let mut b = Backend::headless(session(args)?); if let Some(p) = &args.project { b.exec("file.open", json!({"path": p}))?; } else if args.flag("--demo") || (default_demo && !args.flag("--empty")) { @@ -205,6 +209,17 @@ fn backend(args: &Args, default_demo: bool) -> Result { Ok(b) } +/// A wired session; `--gpu` attaches the GPU compositor (renders then follow the project's +/// renderer setting, Mercury GPU Acceleration by default). +fn session(args: &Args) -> Result { + let mut s = effectcraft_host::session(); + if args.flag("--gpu") { + let g = effectcraft_gpu::Gpu::headless().ok_or_else(|| Failure::Error("--gpu: no usable GPU adapter".into()))?; + s.accel = Some(std::sync::Arc::new(g)); + } + Ok(s) +} + /// `--save` / `--save-as` after a mutating command. fn maybe_save(b: &mut Backend, args: &Args) -> Result, Failure> { let r = if let Some(p) = args.opt("--save-as") { @@ -436,7 +451,7 @@ fn render(args: &Args, json_out: bool) -> Result<(), Failure> { if args.opt("--bridge").is_some() { return usage_err("render runs headless; use `exec renderQueue.add` / `renderQueue.render` with --bridge"); } - let mut s = effectcraft_host::session(); + let mut s = session(args)?; match &args.project { Some(p) => s.execute("file.open", json!({"path": p})).map_err(|e| Failure::Error(e.to_string()))?, None => s.execute("file.openDemoProject", json!({})).map_err(|e| Failure::Error(e.to_string()))?, @@ -549,6 +564,9 @@ fn bench_cmd(args: &Args) -> Result<(), Failure> { if let Some(c) = args.opt("--comp") { s.execute("comp.open", json!({"comp": c})).map_err(|e| Failure::Error(e.to_string()))?; } + if args.flag("--gpu") { + return bench_gpu(&s, args); + } let cid = s.active_comp_id().ok_or_else(|| Failure::Error("no composition".into()))?; let t = Tick::from_seconds_f64(args.num("--time")?.unwrap_or(3.0)); let opts = RenderOpts { scale: args.num("--scale")?.unwrap_or(1.0), ..Default::default() }; @@ -664,3 +682,70 @@ fn bench_play(s: &Session, cid: ItemId, t: Tick, opts: RenderOpts, n: usize) { let st = cache.stats(); eprintln!(" layer cache: {} hits, {} misses, {} entries, {:.1} MB", st.hits, st.misses, st.entries, st.bytes as f64 / 1e6); } + +/// Median wall time (ms) of `n` runs of `f` after one warm-up run. +fn time_ms(n: usize, mut f: impl FnMut()) -> f64 { + f(); + let mut v: Vec = (0..n) + .map(|_| { + let t0 = std::time::Instant::now(); + f(); + t0.elapsed().as_secs_f64() * 1e3 + }) + .collect(); + median(&mut v) +} + +/// `bench --gpu`: CPU vs GPU ms/frame for every comp (or `--comp`) at Full and Half. "cold" +/// renders everything (no layer cache); "warm" reuses the layer cache (playback / scrubbing of +/// unchanged layers: compositing cost); "viewer" is the GPU display path without readback. +fn bench_gpu(s: &Session, args: &Args) -> Result<(), Failure> { + use effectcraft_render::Backend; + let gpu = effectcraft_gpu::Gpu::headless().ok_or_else(|| Failure::Error("--gpu: no usable GPU adapter".into()))?; + let n = args.num("--n")?.unwrap_or(10.0).max(1.0) as usize; + let comps: Vec = match args.opt("--comp") { + Some(_) => s.active_comp_id().into_iter().collect(), + None => s.project.comps().map(|(id, _)| *id).collect(), + }; + eprintln!("GPU: {} — median of {n} runs (ms/frame)", effectcraft_render::Accelerator::name(&gpu)); + eprintln!("{:<28} {:>5} {:>10} {:>9} {:>9} {:>9} {:>9} {:>9}", "comp", "res", "size", "cpu cold", "gpu cold", "cpu warm", "gpu warm", "gpu view"); + for cid in comps { + let Some(comp) = s.project.comp(cid) else { continue }; + let name = s.project.item(cid).map(|i| i.name.clone()).unwrap_or_default(); + let t = Tick::from_seconds_f64(args.num("--time")?.unwrap_or(3.0).min(comp.duration.seconds() * 0.5)); + for (label, scale) in [("Full", 1.0), ("Half", 0.5)] { + let mk = |backend: Backend, cache: Option<&LayerCache>| { + let mut r = Renderer::new(&s.project, s.footage.as_ref(), RenderOpts { scale, backend, ..Default::default() }); + r.expr = s.expr.as_deref(); + r.cache = cache; + r.accel = Some(&gpu); + r.comp_frame(cid, t) + }; + let cpu_cold = time_ms(n, || { + std::hint::black_box(mk(Backend::Cpu, None)); + }); + let gpu_cold = time_ms(n, || { + std::hint::black_box(mk(Backend::Gpu, None)); + }); + let cache = LayerCache::default(); + let cpu_warm = time_ms(n, || { + std::hint::black_box(mk(Backend::Cpu, Some(&cache))); + }); + let gcache = LayerCache::default(); + let gpu_warm = time_ms(n, || { + std::hint::black_box(mk(Backend::Gpu, Some(&gcache))); + }); + let view = time_ms(n, || { + let mut r = Renderer::new(&s.project, s.footage.as_ref(), RenderOpts { scale, backend: Backend::Gpu, ..Default::default() }); + r.expr = s.expr.as_deref(); + r.cache = Some(&gcache); + r.accel = Some(&gpu); + std::hint::black_box(gpu.render_display(&r, cid, t)); + gpu.wait(); + }); + let size = format!("{}x{}", (comp.width as f64 * scale).round(), (comp.height as f64 * scale).round()); + eprintln!("{name:<28} {label:>5} {size:>10} {cpu_cold:>9.2} {gpu_cold:>9.2} {cpu_warm:>9.2} {gpu_warm:>9.2} {view:>9.2}"); + } + } + Ok(()) +} diff --git a/crates/effects/src/color2.rs b/crates/effects/src/color2.rs index 8d707c68..8ed52d2b 100644 --- a/crates/effects/src/color2.rs +++ b/crates/effects/src/color2.rs @@ -119,6 +119,12 @@ impl Curve { h00 * self.ys[k] + h10 * h * self.ms[k] + h01 * self.ys[k + 1] + h11 * h * self.ms[k + 1] } + /// Control points (x, y), Hermite tangents and the 1025-entry table over 0..1 that + /// [`Curve::eval`] interpolates (the GPU path uploads them). + pub fn tables(&self) -> (&[f32], &[f32], &[f32], &[f32]) { + (&self.xs, &self.ys, &self.ms, &self.lut) + } + pub fn eval(&self, x: f32) -> f32 { if !(0.0..=1.0).contains(&x) { return self.eval_exact(x); diff --git a/crates/effects/src/lib.rs b/crates/effects/src/lib.rs index 7c694622..2e29061a 100644 --- a/crates/effects/src/lib.rs +++ b/crates/effects/src/lib.rs @@ -353,6 +353,11 @@ pub fn registry() -> &'static [EffectSpec] { v.extend(paint::specs()); v.extend(puppet::specs()); v.sort_by(|a, b| a.category.cmp(b.category).then(a.name.cmp(b.name))); + for s in v.iter_mut() { + if GPU_EFFECTS.contains(&s.id) { + s.gpu = true; + } + } v }) } @@ -411,6 +416,28 @@ pub fn instantiate(spec: &EffectSpec, ids: &mut Ids, instance_name: &str, layer_ g } +/// Effects the GPU compositor (`effectcraft-gpu`) implements with the CPU effect's semantics +/// (Mercury GPU Acceleration; the Effects & Presets GPU badge). Expression controls are +/// pass-throughs and count as GPU effects too. +pub const GPU_EFFECTS: &[&str] = &[ + "ec.blur.gaussian", + "ec.blur.fastbox", + "ec.blur.directional", + "ec.stylize.glow", + "ec.color.levels", + "ec.color.curves", + "ec.color.huesaturation", + "ec.color.tint", + "ec.generate.fill", + "ec.generate.gradientramp", + "ec.noise.fractal", + "ec.perspective.dropshadow", + "ec.color.brightnesscontrast", + "ec.color.exposure", + "ec.channel.invert", + "ec.distort.transform", +]; + /// Effects whose output depends on [`EffectCtx::time`] directly (not only through animated /// parameters). The renderer's layer cache folds the layer time into the key for these. /// The `time_dependence_is_declared` test checks the list against every registered effect. diff --git a/crates/engine/src/commands/comp_more.rs b/crates/engine/src/commands/comp_more.rs index 6a066b5a..e3057f67 100644 --- a/crates/engine/src/commands/comp_more.rs +++ b/crates/engine/src/commands/comp_more.rs @@ -89,7 +89,7 @@ fn save_frame(s: &mut Session, p: &Value) -> Result { let t = f_p(p, "time").map(Tick::from_seconds_f64).unwrap_or(s.time()); let scale = f_p(p, "scale").unwrap_or(1.0).clamp(0.01, 4.0); let bg = comp.background; - let img = s.render(cid, t, effectcraft_render::RenderOpts { scale, ..Default::default() }); + let img = s.render(cid, t, effectcraft_render::RenderOpts { scale, backend: effectcraft_render::Backend::Auto, ..Default::default() }); let rgba = img.to_rgba8_over(bg); let png = encode_png(&rgba, img.width, img.height).map_err(|e| EngineError::Other(format!("PNG encoding failed: {e}")))?; s.services.write_file(&path, &png).map_err(|e| EngineError::Other(format!("cannot write {path}: {e}")))?; diff --git a/crates/engine/src/commands/file.rs b/crates/engine/src/commands/file.rs index ee61f900..a85a391f 100644 --- a/crates/engine/src/commands/file.rs +++ b/crates/engine/src/commands/file.rs @@ -152,6 +152,17 @@ fn project_settings(s: &mut Session, p: &Value) -> Result { ), }; } + if let Some(r) = p.get("renderer").or_else(|| p.get("gpuAcceleration")) { + proj.settings.gpu_acceleration = match r { + Value::Bool(b) => *b, + Value::String(s) => match effectcraft_render::Backend::parse(s) { + Some(effectcraft_render::Backend::Cpu) => false, + Some(_) => true, + None => return Err(bad("file.projectSettings", format!("renderer: gpu|software, not `{s}`"))), + }, + _ => return Err(bad("file.projectSettings", "renderer: gpu|software")), + }; + } if let Some(t) = str_p(p, "timeDisplay") { proj.settings.time_display = if t.eq_ignore_ascii_case("frames") { effectcraft_project::TimeDisplayStyle::Frames } else { effectcraft_project::TimeDisplayStyle::Timecode }; @@ -161,6 +172,41 @@ fn project_settings(s: &mut Session, p: &Value) -> Result { Ok(serde_json::to_value(&s.project.settings).unwrap_or_default()) } +/// Video Rendering and Effects: report or set the renderer (Mercury GPU Acceleration when a GPU +/// adapter exists, else Mercury Software Only). +fn render_backend(s: &mut Session, p: &Value) -> Result { + if let Some(b) = p.get("backend").or_else(|| p.get("renderer")) { + let gpu = match b { + Value::Bool(b) => *b, + Value::String(v) => match effectcraft_render::Backend::parse(v) { + Some(effectcraft_render::Backend::Cpu) => false, + Some(_) => true, + None => return Err(bad("render.backend", format!("backend: gpu|cpu, not `{v}`"))), + }, + _ => return Err(bad("render.backend", "backend: gpu|cpu")), + }; + if gpu != s.project.settings.gpu_acceleration { + s.edit("Project Settings", None, |proj, _| { + proj.settings.gpu_acceleration = gpu; + Ok(()) + })?; + } + } + Ok(backend_status(s)) +} + +/// The renderer setting and what renders with it. +pub fn backend_status(s: &Session) -> Value { + let adapter = s.accel.as_ref().map(|a| a.name()); + let gpu = s.project.settings.gpu_acceleration; + json!({ + "renderer": if gpu { "Mercury GPU Acceleration" } else { "Mercury Software Only" }, + "gpuAcceleration": gpu, + "adapter": adapter, + "active": if gpu && adapter.is_some() { "gpu" } else { "cpu" }, + }) +} + fn cycle_depth(s: &mut Session, _: &Value) -> Result { s.edit("Project Bit Depth", None, |proj, _| { proj.settings.bit_depth = proj.settings.bit_depth.next(); @@ -184,10 +230,11 @@ pub fn specs() -> Vec { "Project Settings...", ["File"], Some("Cmd+Alt+Shift+K"), - "{bitDepth?: 8|16|32, workingSpace?: none|srgb|rec709|rec2020|p3, linearize?, blendLinear?, timeDisplay?: timecode|frames}", + "{bitDepth?: 8|16|32, workingSpace?: none|srgb|rec709|rec2020|p3, linearize?, blendLinear?, renderer?: gpu|software, timeDisplay?: timecode|frames}", always, project_settings ), cmd!("file.cycleBitDepth", "Cycle Project Bit Depth", [], None, "{}", always, cycle_depth), + cmd!("render.backend", "Video Rendering and Effects", [], None, "{backend?: gpu|cpu}", always, render_backend), ] } diff --git a/crates/engine/src/commands/query.rs b/crates/engine/src/commands/query.rs index f3823150..69a9dc3b 100644 --- a/crates/engine/src/commands/query.rs +++ b/crates/engine/src/commands/query.rs @@ -47,6 +47,7 @@ fn summary(s: &mut Session, _: &Value) -> Result { "path": s.path, "dirty": s.is_dirty(), "bitDepth": s.project.settings.bit_depth.label(), + "renderer": super::file::backend_status(s), "items": items, "activeComp": s.state.active_comp.map(|c| c.0), "time": s.time().seconds(), diff --git a/crates/engine/src/commands/render_queue.rs b/crates/engine/src/commands/render_queue.rs index d22384dd..c905ce6c 100644 --- a/crates/engine/src/commands/render_queue.rs +++ b/crates/engine/src/commands/render_queue.rs @@ -495,7 +495,8 @@ fn save_current_preview(s: &mut Session, p: &Value) -> Result { let mut item = RenderQueueItem::new(0, cid); item.output = OutputModule::for_format(format); item.output.output = path.clone(); - let job = crate::ExportJob { project: &s.project, footage: s.footage.as_ref(), expr: s.expr.as_deref(), item: &item, path: &path }; + let job = + crate::ExportJob { project: &s.project, footage: s.footage.as_ref(), expr: s.expr.as_deref(), accel: s.accel.as_deref(), item: &item, path: &path }; let r = exporter.export(&job, &mut |_, _| true).map_err(EngineError::Other)?; s.toast(format!("Saved preview to {}", r.path)); Ok(serde_json::to_value(&r).unwrap_or_default()) diff --git a/crates/engine/src/lib.rs b/crates/engine/src/lib.rs index 564dbc1b..adb510a4 100644 --- a/crates/engine/src/lib.rs +++ b/crates/engine/src/lib.rs @@ -225,6 +225,10 @@ pub struct Session { pub services: Arc, pub footage: Arc, pub expr: Option>, + /// GPU compositor (Mercury GPU Acceleration), when the frontend has one: used by renders + /// that ask for [`effectcraft_render::Backend::Gpu`]/`Auto`, and by Render Queue exports + /// when the project's renderer is the GPU. + pub accel: Option>, /// Expression syntax checker (set by the host that links the expression engine). pub expr_check: Option std::result::Result<(), String>>, pub importer: Option>, @@ -268,6 +272,7 @@ impl Default for Session { services: Arc::new(FsServices), footage: Arc::new(NoFootage), expr: None, + accel: None, expr_check: None, importer: None, exporter: None, @@ -464,6 +469,7 @@ impl Session { let mut r = Renderer::new(&self.project, self.footage.as_ref(), opts); r.expr = self.expr.as_deref(); r.cache = Some(&self.layer_cache); + r.accel = self.accel.as_deref(); r.comp_frame(comp, t) } @@ -482,7 +488,7 @@ impl Session { let c = self.project.comp(comp).ok_or(EngineError::NoComp)?; let long = c.width.max(c.height).max(1) as f64; let scale = if max_side == 0 { 1.0 } else { (max_side as f64 / long).min(1.0) }; - let img = self.render(comp, t, RenderOpts { scale, ..Default::default() }); + let img = self.render(comp, t, RenderOpts { scale, backend: effectcraft_render::Backend::Auto, ..Default::default() }); Ok((img.width, img.height, img.to_rgba8_over(c.background))) } diff --git a/crates/engine/src/render_queue.rs b/crates/engine/src/render_queue.rs index 95c90824..bc490f47 100644 --- a/crates/engine/src/render_queue.rs +++ b/crates/engine/src/render_queue.rs @@ -23,6 +23,8 @@ pub struct ExportJob<'a> { pub footage: &'a dyn FootageSource, pub expr: Option<&'a dyn ExprHost>, pub item: &'a RenderQueueItem, + /// GPU compositor to render with when the project's renderer is Mercury GPU Acceleration. + pub accel: Option<&'a dyn effectcraft_render::Accelerator>, /// Resolved output path (templates expanded except the `#` frame-number run). pub path: &'a str, } @@ -120,6 +122,7 @@ struct Work { project: Arc, footage: Arc, expr: Option>, + accel: Option>, exporter: Arc, /// Per queue item: one (item with that output module, resolved path) per output module. items: Vec>, @@ -155,7 +158,7 @@ fn run_work(w: Work, shared: &JobShared) { // first module's file is the item's output. let mut r: Result = Err("no output module".into()); for (mi, (mitem, path)) in modules.iter().enumerate() { - let job = ExportJob { project: &w.project, footage: w.footage.as_ref(), expr: w.expr.as_deref(), item: mitem, path }; + let job = ExportJob { project: &w.project, footage: w.footage.as_ref(), expr: w.expr.as_deref(), accel: w.accel.as_deref(), item: mitem, path }; let rr = w.exporter.export(&job, &mut |done, total| { let mut s = lock(&shared.state); s.done = done; @@ -288,7 +291,7 @@ impl Session { return Err("nothing is queued: add a composition (Composition ▸ Add to Render Queue) and tick Render".into()); } let ids: Vec = items.iter().filter_map(|m| m.first().map(|(i, _)| i.id)).collect(); - let work = Work { project: self.project.clone(), footage: self.footage.clone(), expr: self.expr.clone(), exporter, items }; + let work = Work { project: self.project.clone(), footage: self.footage.clone(), expr: self.expr.clone(), accel: self.accel.clone(), exporter, items }; let shared = Arc::new(JobShared::default()); if wait { run_work(work, &shared); diff --git a/crates/export/src/lib.rs b/crates/export/src/lib.rs index 1616926d..87c9f80b 100644 --- a/crates/export/src/lib.rs +++ b/crates/export/src/lib.rs @@ -56,6 +56,8 @@ pub struct Job<'a> { pub project: &'a Project, pub footage: &'a dyn FootageSource, pub expr: Option<&'a dyn ExprHost>, + /// GPU compositor: used when the project's renderer is Mercury GPU Acceleration. + pub accel: Option<&'a dyn effectcraft_render::Accelerator>, pub comp: ItemId, pub settings: &'a RenderSettings, pub output: &'a OutputModule, @@ -172,9 +174,11 @@ pub(crate) fn render_frame(job: &Job, comp: &Comp, i: u64) -> Image { draft: job.settings.quality == RenderQuality::Draft, // Output renders always look through the comp's active camera. view: None, + backend: effectcraft_render::Backend::Auto, }; let mut r = Renderer::new(job.project, job.footage, opts); r.expr = job.expr; + r.accel = job.accel; r.comp_frame(job.comp, t) } diff --git a/crates/export/tests/roundtrip.rs b/crates/export/tests/roundtrip.rs index 761b2c67..186aa1a9 100644 --- a/crates/export/tests/roundtrip.rs +++ b/crates/export/tests/roundtrip.rs @@ -55,7 +55,7 @@ fn settings() -> RenderSettings { fn run(p: &Project, cid: ItemId, footage: &dyn FootageSource, om: &OutputModule, path: &Path) -> effectcraft_export::Report { let s = settings(); let path = path.to_string_lossy().to_string(); - let job = Job { project: p, footage, expr: None, comp: cid, settings: &s, output: om, path: &path, sink: None }; + let job = Job { project: p, footage, expr: None, accel: None, comp: cid, settings: &s, output: om, path: &path, sink: None }; let mut last = Progress::default(); let mut calls = 0; let r = export(&job, &mut |pr| { @@ -167,7 +167,7 @@ fn h264_odd_size_is_cropped_even() { let comp = p.comp(cid).expect("comp"); assert_eq!(s.output_size(comp).0 % 2, 1); let path = d.join("odd.mp4").to_string_lossy().to_string(); - let job = Job { project: &p, footage: &NoFootage, expr: None, comp: cid, settings: &s, output: &om, path: &path, sink: None }; + let job = Job { project: &p, footage: &NoFootage, expr: None, accel: None, comp: cid, settings: &s, output: &om, path: &path, sink: None }; let r = export(&job, &mut |_| true).expect("export"); assert_eq!((r.width % 2, r.height % 2), (0, 0)); } @@ -287,7 +287,7 @@ fn cancel_stops_early() { let om = OutputModule::for_format(OutputFormat::PngSequence); let s = settings(); let path = d.join("c_[#####].png").to_string_lossy().to_string(); - let job = Job { project: &p, footage: &NoFootage, expr: None, comp: cid, settings: &s, output: &om, path: &path, sink: None }; + let job = Job { project: &p, footage: &NoFootage, expr: None, accel: None, comp: cid, settings: &s, output: &om, path: &path, sink: None }; let r = export(&job, &mut |_| false); assert!(matches!(r, Err(effectcraft_export::ExportError::Cancelled))); let written = std::fs::read_dir(&d).map(|r| r.count()).unwrap_or(0); diff --git a/crates/gpu/Cargo.toml b/crates/gpu/Cargo.toml new file mode 100644 index 00000000..36c751df --- /dev/null +++ b/crates/gpu/Cargo.toml @@ -0,0 +1,29 @@ +[package] +name = "effectcraft-gpu" +description = "The GPU compositor (wgpu compute): layer transforms, masks/mattes, the 38 blend modes, motion blur, bit-depth quantisation and GPU effects, with the CPU renderer as the reference" +version.workspace = true +edition.workspace = true +license.workspace = true +rust-version.workspace = true + +[lints] +workspace = true + +[dependencies] +effectcraft-time = { workspace = true } +effectcraft-geom = { workspace = true } +effectcraft-color = { workspace = true } +effectcraft-project = { workspace = true } +effectcraft-raster = { workspace = true } +effectcraft-effects = { workspace = true } +effectcraft-render = { workspace = true } +effectcraft-keyframe = { workspace = true } +wgpu = { workspace = true } +bytemuck = { workspace = true } +log = { workspace = true } + +[target.'cfg(not(target_arch = "wasm32"))'.dependencies] +pollster = { workspace = true } + +[dev-dependencies] +effectcraft-text = { workspace = true } diff --git a/crates/gpu/src/context.rs b/crates/gpu/src/context.rs new file mode 100644 index 00000000..9f814377 --- /dev/null +++ b/crates/gpu/src/context.rs @@ -0,0 +1,528 @@ +//! Device, pipelines, textures, uploads and readback. + +use std::collections::HashMap; +use std::sync::{Arc, Mutex, Weak}; + +use effectcraft_effects::Buf; +use effectcraft_raster::Image; +use wgpu::util::DeviceExt; + +/// Compute entry points in `kernels.wgsl` (one pipeline each). +const ENTRIES: &[&str] = &[ + "warp_blend", + "blend_full", + "matte", + "preserve", + "knockout", + "channel_mix", + "quantize", + "convert", + "half", + "box_h", + "box_v", + "directional", + "glow_bright", + "glow_combine", + "shadow_make", + "shadow_combine", + "pointwise", +]; + +/// Pixel format of every working texture: premultiplied RGBA, 32-bit float (32 bpc headroom; +/// 8/16 bpc are emulated by clamping and quantising, as on the CPU). +pub const FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Rgba32Float; + +/// Uploaded layer buffers kept on the GPU (bytes). +const UPLOAD_BUDGET: usize = 1536 << 20; + +/// A GPU image: premultiplied RGBA f32 in a texture (cheap to clone; immutable once written). +#[derive(Clone, Debug)] +pub struct GpuImage { + pub texture: wgpu::Texture, + pub width: u32, + pub height: u32, +} + +/// Uniform parameter block shared by every kernel (see `shaders/common.wgsl`). +#[derive(Clone, Copy, Debug, Default)] +pub(crate) struct Params { + pub u: [[u32; 4]; 4], + pub f: [[f32; 4]; 12], +} + +impl Params { + fn bytes(&self) -> Vec { + let mut v = Vec::with_capacity(256); + for r in &self.u { + for x in r { + v.extend_from_slice(&x.to_le_bytes()); + } + } + for r in &self.f { + for x in r { + v.extend_from_slice(&x.to_le_bytes()); + } + } + v + } +} + +struct Upload { + buf: Weak, + img: GpuImage, + bytes: usize, + last_use: u64, +} + +#[derive(Default)] +struct Uploads { + map: HashMap, + bytes: usize, + clock: u64, +} + +/// The device and everything compiled for it. +pub struct GpuContext { + pub device: wgpu::Device, + pub queue: wgpu::Queue, + pub(crate) name: String, + pub(crate) max_dim: u32, + bgl: wgpu::BindGroupLayout, + pipelines: HashMap<&'static str, wgpu::ComputePipeline>, + display_bgl: wgpu::BindGroupLayout, + display: wgpu::ComputePipeline, + dummy_tex: wgpu::TextureView, + dummy_buf: wgpu::Buffer, + uploads: Mutex, +} + +fn tex_entry(binding: u32) -> wgpu::BindGroupLayoutEntry { + wgpu::BindGroupLayoutEntry { + binding, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: false }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, + count: None, + } +} + +fn storage_tex_entry(binding: u32, format: wgpu::TextureFormat) -> wgpu::BindGroupLayoutEntry { + wgpu::BindGroupLayoutEntry { + binding, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::StorageTexture { access: wgpu::StorageTextureAccess::WriteOnly, format, view_dimension: wgpu::TextureViewDimension::D2 }, + count: None, + } +} + +impl GpuContext { + /// Build on an existing device (the desktop app shares egui-wgpu's). `Err` when the + /// adapter cannot run the compositor (no compute shaders, e.g. WebGL2; no float storage + /// textures). + pub fn new(adapter: &wgpu::Adapter, device: wgpu::Device, queue: wgpu::Queue) -> Result { + let info = adapter.get_info(); + if !adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) { + return Err(format!("{} ({:?}): no compute shaders", info.name, info.backend)); + } + for f in [FORMAT, wgpu::TextureFormat::Rgba8Unorm] { + if !adapter.get_texture_format_features(f).allowed_usages.contains(wgpu::TextureUsages::STORAGE_BINDING) { + return Err(format!("{} ({:?}): {f:?} storage textures unsupported", info.name, info.backend)); + } + } + let name = format!("{} ({:?})", info.name, info.backend); + let src = [include_str!("shaders/common.wgsl"), include_str!("shaders/kernels.wgsl")].concat(); + let module = + device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("effectcraft kernels"), source: wgpu::ShaderSource::Wgsl(src.into()) }); + let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("effectcraft kernels"), + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None }, + count: None, + }, + tex_entry(1), + tex_entry(2), + storage_tex_entry(3, FORMAT), + wgpu::BindGroupLayoutEntry { + binding: 4, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: true }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + ], + }); + let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("effectcraft kernels"), + bind_group_layouts: &[Some(&bgl)], + immediate_size: 0, + }); + let pipelines = ENTRIES + .iter() + .map(|e| { + let p = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some(e), + layout: Some(&layout), + module: &module, + entry_point: Some(e), + compilation_options: Default::default(), + cache: None, + }); + (*e, p) + }) + .collect(); + let dmodule = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("effectcraft display"), + source: wgpu::ShaderSource::Wgsl(include_str!("shaders/display.wgsl").into()), + }); + let display_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("effectcraft display"), + entries: &[tex_entry(1), storage_tex_entry(3, wgpu::TextureFormat::Rgba8Unorm)], + }); + let dlayout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("effectcraft display"), + bind_group_layouts: &[Some(&display_bgl)], + immediate_size: 0, + }); + let display = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("display"), + layout: Some(&dlayout), + module: &dmodule, + entry_point: Some("display"), + compilation_options: Default::default(), + cache: None, + }); + let dummy = device.create_texture(&wgpu::TextureDescriptor { + label: Some("dummy"), + size: wgpu::Extent3d { width: 1, height: 1, depth_or_array_layers: 1 }, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu::TextureDimension::D2, + format: FORMAT, + usage: wgpu::TextureUsages::TEXTURE_BINDING, + view_formats: &[], + }); + let dummy_buf = + device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("dummy"), contents: &[0u8; 16], usage: wgpu::BufferUsages::STORAGE }); + let max_dim = device.limits().max_texture_dimension_2d; + Ok(GpuContext { + device, + queue, + name, + max_dim, + bgl, + pipelines, + display_bgl, + display, + dummy_tex: dummy.create_view(&Default::default()), + dummy_buf, + uploads: Mutex::new(Uploads::default()), + }) + } + + /// A device of its own on the best adapter (CLI, tests, benchmarks). `None` when no + /// adapter qualifies. + #[cfg(not(target_arch = "wasm32"))] + pub fn headless() -> Option { + let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle()); + let adapter = pollster::block_on( + instance.request_adapter(&wgpu::RequestAdapterOptions { power_preference: wgpu::PowerPreference::HighPerformance, ..Default::default() }), + ) + .ok()?; + let limits = adapter.limits(); + let required_limits = wgpu::Limits { + max_texture_dimension_2d: limits.max_texture_dimension_2d.min(16384), + max_buffer_size: limits.max_buffer_size, + ..wgpu::Limits::default() + }; + let (device, queue) = + pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { label: Some("effectcraft gpu"), required_limits, ..Default::default() })) + .ok()?; + device.on_uncaptured_error(Arc::new(|e| log::error!("wgpu: {e}"))); + GpuContext::new(&adapter, device, queue).map_err(|e| log::info!("gpu: {e}")).ok() + } + + /// Readback (GPU → CPU) waits for the device; impossible on the browser's main thread. + pub fn can_readback(&self) -> bool { + cfg!(not(target_arch = "wasm32")) + } + + pub(crate) fn fits(&self, w: u32, h: u32) -> bool { + w > 0 && h > 0 && w <= self.max_dim && h <= self.max_dim + } + + /// A new (zeroed) working texture. + pub(crate) fn image(&self, w: u32, h: u32) -> GpuImage { + let texture = self.device.create_texture(&wgpu::TextureDescriptor { + label: None, + size: wgpu::Extent3d { width: w.max(1), height: h.max(1), depth_or_array_layers: 1 }, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu::TextureDimension::D2, + format: FORMAT, + usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::COPY_SRC | wgpu::TextureUsages::COPY_DST, + view_formats: &[], + }); + GpuImage { texture, width: w, height: h } + } + + /// Upload a CPU image (`None` when it does not fit the device). + pub fn upload_image(&self, img: &Image) -> Option { + if !self.fits(img.width, img.height) { + return None; + } + let g = self.image(img.width, img.height); + self.queue.write_texture( + wgpu::TexelCopyTextureInfo { texture: &g.texture, mip_level: 0, origin: wgpu::Origin3d::ZERO, aspect: wgpu::TextureAspect::All }, + bytemuck::cast_slice(&img.data), + wgpu::TexelCopyBufferLayout { offset: 0, bytes_per_row: Some(img.width * 16), rows_per_image: Some(img.height) }, + wgpu::Extent3d { width: img.width, height: img.height, depth_or_array_layers: 1 }, + ); + Some(g) + } + + /// Upload a (layer cache) buffer, reusing the texture while the same `Arc` is alive: static + /// layers upload once, not every frame. + pub(crate) fn upload_buf(&self, buf: &Arc) -> Option { + let key = Arc::as_ptr(buf) as usize; + if let Ok(mut u) = self.uploads.lock() { + u.clock += 1; + let now = u.clock; + if let Some(e) = u.map.get_mut(&key) + && e.buf.upgrade().is_some_and(|b| Arc::ptr_eq(&b, buf)) + { + e.last_use = now; + return Some(e.img.clone()); + } + } + let img = self.upload_image(&buf.img)?; + if let Ok(mut u) = self.uploads.lock() { + let bytes = buf.img.data.len() * 16; + let now = u.clock; + if let Some(old) = u.map.insert(key, Upload { buf: Arc::downgrade(buf), img: img.clone(), bytes, last_use: now }) { + u.bytes -= old.bytes; + } + u.bytes += bytes; + // Drop textures of freed buffers, then the least recently used over budget. + if u.bytes > UPLOAD_BUDGET { + let dead: Vec = u.map.iter().filter(|(_, e)| e.buf.strong_count() == 0).map(|(k, _)| *k).collect(); + for k in dead { + if let Some(e) = u.map.remove(&k) { + u.bytes -= e.bytes; + } + } + } + while u.bytes > UPLOAD_BUDGET { + let Some(k) = u.map.iter().min_by_key(|(_, e)| e.last_use).map(|(k, _)| *k) else { break }; + if let Some(e) = u.map.remove(&k) { + u.bytes -= e.bytes; + } + } + } + Some(img) + } + + /// Forget uploaded layer textures. + pub fn clear_uploads(&self) { + if let Ok(mut u) = self.uploads.lock() { + u.map.clear(); + u.bytes = 0; + } + } +} + +/// Records GPU work for one render (a command encoder submitted at readback or the end). +pub(crate) struct Enc<'g> { + pub g: &'g GpuContext, + enc: Option, + pending: usize, +} + +impl<'g> Enc<'g> { + pub fn new(g: &'g GpuContext) -> Enc<'g> { + Enc { g, enc: None, pending: 0 } + } + + fn encoder(&mut self) -> &mut wgpu::CommandEncoder { + let g = self.g; + self.enc.get_or_insert_with(|| g.device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("effectcraft") })) + } + + /// A new zeroed image. + pub fn image(&self, w: u32, h: u32) -> GpuImage { + self.g.image(w, h) + } + + /// Run kernel `entry` writing `out` (dispatched over `groups` workgroups). + pub fn dispatch( + &mut self, + entry: &str, + p: &Params, + src: &GpuImage, + aux: Option<&GpuImage>, + out: &GpuImage, + data: Option<&wgpu::Buffer>, + groups: (u32, u32), + ) { + let g = self.g; + let Some(pipe) = g.pipelines.get(entry) else { + log::error!("gpu: no kernel {entry}"); + return; + }; + let ub = g.device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: None, contents: &p.bytes(), usage: wgpu::BufferUsages::UNIFORM }); + let sv = src.texture.create_view(&Default::default()); + let av = aux.map(|a| a.texture.create_view(&Default::default())); + let ov = out.texture.create_view(&Default::default()); + let bg = g.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout: &g.bgl, + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: ub.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(&sv) }, + wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::TextureView(av.as_ref().unwrap_or(&g.dummy_tex)) }, + wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&ov) }, + wgpu::BindGroupEntry { binding: 4, resource: data.unwrap_or(&g.dummy_buf).as_entire_binding() }, + ], + }); + let enc = self.encoder(); + { + let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some(entry), timestamp_writes: None }); + pass.set_pipeline(pipe); + pass.set_bind_group(0, &bg, &[]); + pass.dispatch_workgroups(groups.0.max(1), groups.1.max(1), 1); + } + self.pending += 1; + // Keep command buffers (and the transient textures they hold) bounded. + if self.pending >= 256 { + self.submit(); + } + } + + /// Per-pixel kernel over `out`'s size (16×16 workgroups). + pub fn pixels(&mut self, entry: &str, p: &Params, src: &GpuImage, aux: Option<&GpuImage>, out: &GpuImage, data: Option<&wgpu::Buffer>) { + let groups = (out.width.div_ceil(16), out.height.div_ceil(16)); + self.dispatch(entry, p, src, aux, out, data, groups); + } + + /// 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(); + 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 }, + wgpu::Extent3d { width: src.width, height: src.height, depth_or_array_layers: 1 }, + ); + } + + /// Submit recorded work. + pub fn submit(&mut self) { + if let Some(enc) = self.enc.take() { + self.g.queue.submit([enc.finish()]); + } + self.pending = 0; + } + + /// Read an image back to the CPU (submits and waits). `None` where readback is impossible. + pub fn download(&mut self, img: &GpuImage) -> Option { + let bytes = self.read_texture(&img.texture, img.width, img.height, 16)?; + let mut out = Image::new(img.width, img.height); + for (p, c) in out.data.iter_mut().zip(bytes.chunks_exact(16)) { + for (k, v) in p.iter_mut().enumerate() { + *v = f32::from_le_bytes([c[4 * k], c[4 * k + 1], c[4 * k + 2], c[4 * k + 3]]); + } + } + 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> { + if !self.g.can_readback() { + return None; + } + let row = w * bpp; + let padded = row.div_ceil(wgpu::COPY_BYTES_PER_ROW_ALIGNMENT) * wgpu::COPY_BYTES_PER_ROW_ALIGNMENT; + let size = padded as u64 * h as u64; + let buf = self.g.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("readback"), + size, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + self.encoder().copy_texture_to_buffer( + wgpu::TexelCopyTextureInfo { texture, mip_level: 0, origin: wgpu::Origin3d::ZERO, 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 }, + ); + self.submit(); + let (tx, rx) = std::sync::mpsc::channel(); + buf.map_async(wgpu::MapMode::Read, .., move |r| { + let _ = tx.send(r); + }); + if let Err(e) = self.g.device.poll(wgpu::PollType::wait_indefinitely()) { + log::error!("gpu: poll: {e}"); + return None; + } + rx.recv().ok()?.ok()?; + let view = buf.get_mapped_range(..).ok()?; + let mut out = Vec::with_capacity((row * h) as usize); + for y in 0..h as usize { + out.extend_from_slice(&view[y * padded as usize..y * padded as usize + row as usize]); + } + drop(view); + buf.unmap(); + Some(out) + } + + /// A data buffer for kernels that read tables (curves). + pub fn data(&self, v: &[f32]) -> wgpu::Buffer { + let mut bytes = Vec::with_capacity(v.len().max(4) * 4); + for x in v { + bytes.extend_from_slice(&x.to_le_bytes()); + } + while bytes.len() < 16 { + bytes.push(0); + } + self.g.device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("data"), contents: &bytes, usage: wgpu::BufferUsages::STORAGE }) + } + + /// Convert to the viewer's premultiplied RGBA8 texture (egui-wgpu's native texture format). + pub fn display(&mut self, img: &GpuImage) -> wgpu::Texture { + let g = self.g; + let tex = g.device.create_texture(&wgpu::TextureDescriptor { + label: Some("viewer frame"), + size: wgpu::Extent3d { width: img.width, height: img.height, depth_or_array_layers: 1 }, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu::TextureDimension::D2, + format: wgpu::TextureFormat::Rgba8Unorm, + usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::COPY_SRC, + view_formats: &[], + }); + let sv = img.texture.create_view(&Default::default()); + let ov = tex.create_view(&Default::default()); + let bg = g.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout: &g.display_bgl, + entries: &[ + wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(&sv) }, + wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&ov) }, + ], + }); + let enc = self.encoder(); + { + let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("display"), timestamp_writes: None }); + pass.set_pipeline(&g.display); + pass.set_bind_group(0, &bg, &[]); + pass.dispatch_workgroups(img.width.div_ceil(16), img.height.div_ceil(16), 1); + } + tex + } +} diff --git a/crates/gpu/src/effects.rs b/crates/gpu/src/effects.rs new file mode 100644 index 00000000..4addbaa8 --- /dev/null +++ b/crates/gpu/src/effects.rs @@ -0,0 +1,347 @@ +//! GPU effects: the same steps as the CPU effects in `effectcraft-effects` (padding, blur +//! radii, parameter conversions), with the pixel loops as compute kernels. + +use effectcraft_effects::{Buf, EffectCtx}; +use effectcraft_geom::{Mat3, vec2}; +use effectcraft_raster::Sampling; +use effectcraft_render::FxStep; + +use crate::context::{Enc, GpuImage, Params}; +use crate::ops; + +/// Effects with a GPU implementation (must equal `effectcraft_effects::GPU_EFFECTS`), plus the +/// expression controls (pass-throughs). +pub fn supports(id: &str) -> bool { + effectcraft_effects::GPU_EFFECTS.contains(&id) || id.starts_with("ec.control.") +} + +/// A layer buffer on the GPU (see [`Buf`]). +struct GBuf { + img: GpuImage, + offset: [f64; 2], + scale: f64, +} + +impl GBuf { + fn to_px(&self, p: [f64; 2]) -> (f64, f64) { + (p[0] * self.scale + self.offset[0], p[1] * self.scale + self.offset[1]) + } + + /// Buf::pad: grow by `pad` transparent pixels on every side. + fn pad(&mut self, e: &mut Enc, pad: u32) -> Option<()> { + if pad == 0 { + return Some(()); + } + let (w, h) = (self.img.width + 2 * pad, self.img.height + 2 * pad); + if !e.g.fits(w, h) { + return None; + } + let out = e.image(w, h); + e.copy_into(&self.img, &out, pad, pad); + self.img = out; + self.offset[0] += pad as f64; + self.offset[1] += pad as f64; + Some(()) + } +} + +/// Run a chain of GPU effects (one upload, one readback), quantising after each when `levels` +/// is set. `None` when anything cannot run here. +pub(crate) fn run_chain(e: &mut Enc, chain: &[FxStep], buf: &Buf, levels: Option) -> Option { + if !e.g.can_readback() || buf.img.is_empty() { + return None; + } + let mut b = GBuf { img: e.g.upload_image(&buf.img)?, offset: buf.offset, scale: buf.scale }; + for step in chain { + b = apply(e, step.spec.id, &step.ctx, b)?; + if let Some(l) = levels { + b.img = ops::quantize(e, &b.img, l); + } + } + let img = e.download(&b.img)?; + Some(Buf { img, offset: b.offset, scale: b.scale }) +} + +fn apply(e: &mut Enc, id: &str, ctx: &EffectCtx, b: GBuf) -> Option { + match id { + "ec.blur.gaussian" => gaussian(e, ctx, b), + "ec.blur.fastbox" => box_blur(e, ctx, b), + "ec.blur.directional" => directional(e, ctx, b), + "ec.stylize.glow" => glow(e, ctx, b), + "ec.perspective.dropshadow" => drop_shadow(e, ctx, b), + "ec.distort.transform" => transform(e, ctx, b), + "ec.color.curves" => curves(e, ctx, b), + _ if id.starts_with("ec.control.") => Some(b), + _ => pointwise(e, id, ctx, b), + } +} + +// ---------------------------------------------------------------- blurs + +/// raster::blur::box_radii. +fn box_radii(sigma: f64, n: usize) -> Vec { + if sigma <= 0.0 { + return vec![0; n]; + } + let w_ideal = (12.0 * sigma * sigma / n as f64 + 1.0).sqrt(); + let mut wl = w_ideal.floor() as i64; + if wl % 2 == 0 { + wl -= 1; + } + let wu = wl + 2; + let m_ideal = (12.0 * sigma * sigma - (n as i64 * wl * wl) as f64 - 4.0 * n as f64 * wl as f64 - 3.0 * n as f64) / (-4.0 * wl as f64 - 4.0); + let m = m_ideal.round() as i64; + (0..n as i64).map(|i| (((if i < m { wl } else { wu }) - 1) / 2).max(0) as usize).collect() +} + +/// Horizontal passes with radii `rx`, then vertical passes with radii `ry` (raster::blur). +fn box_passes(e: &mut Enc, img: &GpuImage, rx: &[usize], ry: &[usize], repeat: bool) -> GpuImage { + let mut cur = img.clone(); + let passes = rx.iter().map(|&r| (false, r)).chain(ry.iter().map(|&r| (true, r))).filter(|p| p.1 > 0); + for (vertical, r) in passes { + let block = if vertical { (4 * r).clamp(32, 256) } else { 256 } as u32; + let (lines, n) = if vertical { (cur.width, cur.height) } else { (cur.height, cur.width) }; + let mut p = Params::default(); + p.u[0] = [r as u32, repeat as u32, block, 0]; + let out = e.image(cur.width, cur.height); + e.dispatch(if vertical { "box_v" } else { "box_h" }, &p, &cur, None, &out, None, (lines.div_ceil(64), n.div_ceil(block))); + cur = out; + } + cur +} + +/// raster::gaussian_blur (3 box passes per axis). +fn gaussian_blur(e: &mut Enc, img: &GpuImage, sx: f64, sy: f64, repeat: bool) -> GpuImage { + let rx = if sx > 0.05 { box_radii(sx, 3) } else { vec![] }; + let ry = if sy > 0.05 { box_radii(sy, 3) } else { vec![] }; + box_passes(e, img, &rx, &ry, repeat) +} + +fn dims_xy(dim: u32) -> (f64, f64) { + match dim { + 1 => (1.0, 0.0), + 2 => (0.0, 1.0), + _ => (1.0, 1.0), + } +} + +fn gaussian(e: &mut Enc, ctx: &EffectCtx, mut b: GBuf) -> Option { + let s = ctx.params.f("blurriness").max(0.0) * 0.5 * b.scale; + if s <= 0.0 { + return Some(b); + } + let (kx, ky) = dims_xy(ctx.params.e("dimensions")); + let repeat = ctx.params.b("repeatEdge"); + if !repeat && !ctx.adjustment { + b.pad(e, (s * 3.0).ceil() as u32)?; + } + b.img = gaussian_blur(e, &b.img, s * kx, s * ky, repeat || ctx.adjustment); + Some(b) +} + +fn box_blur(e: &mut Enc, ctx: &EffectCtx, mut b: GBuf) -> Option { + let r = (ctx.params.f("radius") * b.scale).round().max(0.0) as usize; + if r == 0 { + return Some(b); + } + let it = ctx.params.f("iterations").round().clamp(1.0, 50.0) as usize; + let (kx, ky) = dims_xy(ctx.params.e("dimensions")); + let repeat = ctx.params.b("repeatEdge"); + if !repeat && !ctx.adjustment { + b.pad(e, (r * it) as u32 + 1)?; + } + b.img = box_passes(e, &b.img, &vec![r * kx as usize; it], &vec![r * ky as usize; it], repeat || ctx.adjustment); + Some(b) +} + +fn directional(e: &mut Enc, ctx: &EffectCtx, mut b: GBuf) -> Option { + let len = ctx.params.f("length") * b.scale; + if len < 0.5 { + return Some(b); + } + if !ctx.adjustment { + b.pad(e, len.ceil() as u32 + 1)?; + } + let a = ctx.params.f("direction").to_radians(); + let n = (len.ceil() as usize * 2 + 1).clamp(3, 257); + let mut p = Params::default(); + p.f[0] = [a.sin() as f32, -a.cos() as f32, len as f32, n as f32]; + let out = e.image(b.img.width, b.img.height); + e.pixels("directional", &p, &b.img, None, &out, None); + b.img = out; + Some(b) +} + +fn glow(e: &mut Enc, ctx: &EffectCtx, mut b: GBuf) -> Option { + let thr = ctx.params.f("threshold") as f32 / 100.0; + let radius = ctx.params.f("radius") * b.scale; + let intensity = ctx.params.f("intensity") as f32; + let use_colors = ctx.params.e("colors") == 1; + let ca = ctx.params.color("colorA"); + let cb = ctx.params.color("colorB"); + if !ctx.adjustment { + b.pad(e, (radius * 1.5).ceil() as u32 + 2)?; + } + let mut p = Params::default(); + p.u[0][0] = use_colors as u32; + p.f[0][0] = thr; + p.f[1] = ca; + p.f[2] = cb; + let bright = e.image(b.img.width, b.img.height); + e.pixels("glow_bright", &p, &b.img, None, &bright, None); + let s = (radius / 2.0).max(0.5); + let blurred = gaussian_blur(e, &bright, s, s, false); + let mut p = Params::default(); + p.u[0][0] = (ctx.params.e("operation") == 1) as u32; + p.f[0][0] = intensity; + let out = e.image(b.img.width, b.img.height); + e.pixels("glow_combine", &p, &b.img, Some(&blurred), &out, None); + b.img = out; + Some(b) +} + +fn drop_shadow(e: &mut Enc, ctx: &EffectCtx, mut b: GBuf) -> Option { + let color = ctx.params.color("color"); + let opacity = ctx.params.f("opacity") as f32 / 100.0; + let dir = ctx.params.f("direction").to_radians(); + let dist = ctx.params.f("distance") * b.scale; + let soft = ctx.params.f("softness") * b.scale; + let only = ctx.params.b("shadowOnly"); + let pad = (dist + soft * 1.5).ceil() as u32 + 2; + if !ctx.adjustment { + b.pad(e, pad)?; + } + let (dx, dy) = (dir.sin() * dist, -dir.cos() * dist); + let mut p = Params::default(); + p.f[0] = color; + p.f[1] = [opacity, dx as f32, dy as f32, 0.0]; + let mut sh = e.image(b.img.width, b.img.height); + e.pixels("shadow_make", &p, &b.img, None, &sh, None); + if soft > 0.0 { + sh = gaussian_blur(e, &sh, soft / 2.0, soft / 2.0, false); + } + if !only { + let out = e.image(b.img.width, b.img.height); + e.pixels("shadow_combine", &Params::default(), &sh, Some(&b.img), &out, None); + sh = out; + } + b.img = sh; + Some(b) +} + +fn transform(e: &mut Enc, ctx: &EffectCtx, mut b: GBuf) -> Option { + let anchor = b.to_px(ctx.params.v2("anchor")); + let pos = b.to_px(ctx.params.v2("position")); + let sh = ctx.params.f("scaleHeight"); + let sw = if ctx.params.b("uniform") { sh } else { ctx.params.f("scaleWidth") }; + let opacity = ctx.params.f("opacity") as f32 / 100.0; + let m = Mat3::translate(vec2(pos.0, pos.1)) + * Mat3::rotate_deg(ctx.params.f("rotation")) + * Mat3::skew_deg(ctx.params.f("skew"), ctx.params.f("skewAxis")) + * Mat3::scale(vec2(sw / 100.0, sh / 100.0)) + * Mat3::translate(vec2(-anchor.0, -anchor.1)); + let empty = e.image(b.img.width, b.img.height); + b.img = ops::warp(e, &empty, &b.img, &m, Sampling::Bilinear, effectcraft_color::BlendMode::Normal, opacity, 0, None); + Some(b) +} + +// ---------------------------------------------------------------- per-pixel + +fn curves(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option { + let parsed = ["rgb", "red", "green", "blue", "alpha"].map(|id| effectcraft_effects::Curve::parse(ctx.params.s(id))); + if parsed.iter().all(Option::is_none) { + return Some(b); + } + let mut data = vec![]; + let mut offs = [-1.0f32; 5]; + for (i, c) in parsed.iter().enumerate() { + let Some(c) = c else { continue }; + let (xs, ys, ms, lut) = c.tables(); + offs[i] = data.len() as f32; + data.push(xs.len() as f32); + data.extend_from_slice(xs); + data.extend_from_slice(ys); + data.extend_from_slice(ms); + data.extend_from_slice(lut); + } + let buf = e.data(&data); + let mut p = Params::default(); + p.u[0][0] = 8; + p.f[0] = [offs[0], offs[1], offs[2], offs[3]]; + p.f[1][0] = offs[4]; + let out = e.image(b.img.width, b.img.height); + e.pixels("pointwise", &p, &b.img, None, &out, Some(&buf)); + Some(GBuf { img: out, ..b }) +} + +fn rgb(c: [f32; 4]) -> [f32; 4] { + [c[0], c[1], c[2], 0.0] +} + +fn pointwise(e: &mut Enc, id: &str, ctx: &EffectCtx, b: GBuf) -> Option { + let f = |k: &str| ctx.params.f(k); + let mut p = Params::default(); + match id { + "ec.color.tint" => { + p.u[0][0] = 1; + p.f[0] = rgb(ctx.params.color("black")); + p.f[1] = rgb(ctx.params.color("white")); + p.f[2][0] = f("amount") as f32 / 100.0; + } + "ec.color.brightnesscontrast" => { + let br = f("brightness") as f32 / 100.0; + let ct = f("contrast") as f32 / 100.0; + let k = if ct >= 0.0 { 1.0 / (1.0 - ct * 0.99) } else { 1.0 + ct }; + p.u[0][0] = 2; + p.f[0] = [br, k, 0.0, 0.0]; + } + "ec.color.huesaturation" => { + p.u[0] = [3, ctx.params.b("colorize") as u32, 0, 0]; + p.f[0] = [f("hue") as f32 / 360.0, f("saturation") as f32 / 100.0, f("lightness") as f32 / 100.0, 0.0]; + p.f[1] = [f("colorizeHue") as f32 / 360.0, f("colorizeSaturation") as f32 / 100.0, f("colorizeLightness") as f32 / 100.0, 0.0]; + } + "ec.color.levels" => { + p.u[0] = [4, !ctx.params.b("noClip") as u32, 0, 0]; + p.f[0] = [f("inBlack") as f32, f("inWhite") as f32, f("gamma").max(0.01) as f32, f("outBlack") as f32]; + p.f[1][0] = f("outWhite") as f32; + } + "ec.color.exposure" => { + p.u[0][0] = 5; + p.f[0] = [2f32.powf(f("exposure") as f32), f("offset") as f32, f("gamma").max(0.01) as f32, 0.0]; + } + "ec.channel.invert" => { + p.u[0] = [6, ctx.params.e("channel"), 0, 0]; + p.f[0][0] = 1.0 - f("blend") as f32 / 100.0; + } + "ec.generate.fill" => { + p.u[0] = [7, ctx.params.b("invert") as u32, 0, 0]; + p.f[0] = ctx.params.color("color"); + p.f[1][0] = f("opacity") as f32 / 100.0; + } + "ec.generate.gradientramp" => { + let s = b.to_px(ctx.params.v2("start")); + let en = b.to_px(ctx.params.v2("end")); + let (dx, dy) = (en.0 - s.0, en.1 - s.1); + let len2 = (dx * dx + dy * dy).max(1e-9); + p.u[0] = [9, (ctx.params.e("shape") == 1) as u32, ctx.seed, 0]; + p.f[0] = [s.0 as f32, s.1 as f32, en.0 as f32, en.1 as f32]; + p.f[1] = rgb(ctx.params.color("startColor")); + p.f[2] = rgb(ctx.params.color("endColor")); + p.f[3] = [len2 as f32, f("scatter") as f32 / 512.0, f("blend") as f32 / 100.0, 0.0]; + } + "ec.noise.fractal" => { + let octaves = f("complexity").clamp(1.0, 20.0); + let (sr, cr) = (f("rotation") as f32).to_radians().sin_cos(); + let (ox, oy) = b.to_px(ctx.params.v2("offset")); + p.u[0] = [10, ctx.params.e("fractalType"), octaves.ceil() as u32, f("seed") as u32 ^ 0x51ed]; + p.u[1][0] = ctx.params.b("invert") as u32; + p.f[0] = [f("contrast") as f32 / 100.0, f("brightness") as f32 / 100.0, (f("scale") * b.scale).max(1.0) as f32, (octaves - octaves.floor()) as f32]; + p.f[1] = [ox as f32, oy as f32, sr, cr]; + p.f[2] = [f("evolution") as f32 / 360.0, f("blend") as f32 / 100.0, f("opacity") as f32 / 100.0, 0.0]; + } + _ => return None, + } + let out = e.image(b.img.width, b.img.height); + e.pixels("pointwise", &p, &b.img, None, &out, None); + Some(GBuf { img: out, ..b }) +} diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs new file mode 100644 index 00000000..e7b54ab2 --- /dev/null +++ b/crates/gpu/src/lib.rs @@ -0,0 +1,127 @@ +//! The EffectCraft GPU compositor (Mercury GPU Acceleration's counterpart), on wgpu compute +//! shaders: Metal, Vulkan, Direct3D 12 and WebGPU. +//! +//! The CPU [`Renderer`] stays the reference and keeps rendering layer *content* (sources, masks, +//! CPU effects, layer styles) into its layer cache. [`Gpu`] composites: it uploads the cached +//! layer buffers (once per buffer), applies the layer transforms with the CPU's sampling +//! (nearest / bilinear / bicubic, minification pre-filter), motion-blur sub-samples, track +//! mattes, Preserve Transparency, layer styles' passes, all 38 blend modes and the 8/16 bpc +//! clamping and quantisation, and converts colour spaces. 3D runs, adjustment layers and +//! wireframes run on the CPU between GPU steps (read back, draw, upload). +//! +//! GPU effects ([`effectcraft_effects::GPU_EFFECTS`]) run as compute kernels with the CPU +//! effect's exact steps (padding, box-blur radii, parameter conversions); chains of them are +//! uploaded and read back once. +//! +//! Plug a [`Gpu`] into [`Renderer::accel`] (it implements [`Accelerator`]); renders then use it +//! when [`RenderOpts::backend`](effectcraft_render::RenderOpts) asks for it. The viewer can +//! skip readback entirely with [`Gpu::render_display`], which leaves an RGBA8 texture for +//! egui-wgpu to draw. + +mod context; +mod effects; +mod ops; +mod walk; + +use std::sync::Arc; + +pub use context::{GpuContext, GpuImage}; +use effectcraft_effects::Buf; +use effectcraft_project::ItemId; +use effectcraft_raster::Image; +use effectcraft_render::{Accelerator, FxStep, Renderer}; +use effectcraft_time::Tick; +pub use wgpu; + +use crate::context::Enc; + +/// The GPU compositor (cheap to clone; one device shared by all clones). +#[derive(Clone)] +pub struct Gpu { + ctx: Arc, +} + +/// A viewer frame left on the GPU: premultiplied RGBA8 (`wgpu::TextureFormat::Rgba8Unorm`), +/// ready for `egui_wgpu::Renderer::register_native_texture`. +#[derive(Clone, Debug)] +pub struct DisplayFrame { + pub texture: wgpu::Texture, + pub width: u32, + pub height: u32, +} + +impl Gpu { + /// On a device of its own (CLI, tests, benchmarks); `None` without a usable adapter. + #[cfg(not(target_arch = "wasm32"))] + pub fn headless() -> Option { + GpuContext::headless().map(Gpu::from_context) + } + + /// On an existing device (the desktop app shares egui-wgpu's). + pub fn new(adapter: &wgpu::Adapter, device: wgpu::Device, queue: wgpu::Queue) -> Result { + GpuContext::new(adapter, device, queue).map(Gpu::from_context) + } + + pub fn from_context(ctx: GpuContext) -> Gpu { + Gpu { ctx: Arc::new(ctx) } + } + + pub fn context(&self) -> &GpuContext { + &self.ctx + } + + pub fn device(&self) -> &wgpu::Device { + &self.ctx.device + } + + /// Render a top-level comp frame and leave it on the GPU as a display texture (no + /// readback unless a CPU fallback step needs one). `None` = render on the CPU. + pub fn render_display(&self, r: &Renderer, comp: ItemId, t: Tick) -> Option { + let mut e = Enc::new(&self.ctx); + let img = walk::render(&mut e, r, comp, t)?; + let texture = e.display(&img); + e.submit(); + Some(DisplayFrame { texture, width: img.width, height: img.height }) + } + + /// Read a display frame back as premultiplied RGBA8 (Info panel sampling, eyedroppers). + pub fn read_display(&self, f: &DisplayFrame) -> Option> { + Enc::new(&self.ctx).read_texture(&f.texture, f.width, f.height, 4) + } + + /// Render a frame on the GPU and read it back (`None` = not handled). + pub fn render(&self, r: &Renderer, comp: ItemId, t: Tick) -> Option { + if !self.ctx.can_readback() { + return None; + } + let mut e = Enc::new(&self.ctx); + let img = walk::render(&mut e, r, comp, t)?; + e.download(&img) + } + + /// Wait for all submitted GPU work (benchmarks). + pub fn wait(&self) { + let _ = self.ctx.device.poll(wgpu::PollType::wait_indefinitely()); + } +} + +impl Accelerator for Gpu { + fn name(&self) -> String { + self.ctx.name.clone() + } + + fn comp_frame(&self, r: &Renderer, comp: ItemId, t: Tick) -> Option { + self.render(r, comp, t) + } + + fn supports_effect(&self, id: &str) -> bool { + self.ctx.can_readback() && effects::supports(id) + } + + fn effects(&self, chain: &[FxStep], buf: &Buf, levels: Option) -> Option { + effects::run_chain(&mut Enc::new(&self.ctx), chain, buf, levels) + } +} + +#[cfg(test)] +mod tests; diff --git a/crates/gpu/src/ops.rs b/crates/gpu/src/ops.rs new file mode 100644 index 00000000..775e86c1 --- /dev/null +++ b/crates/gpu/src/ops.rs @@ -0,0 +1,167 @@ +//! Compositing primitives: each mirrors a CPU operation in `effectcraft-raster` / +//! `effectcraft-render` and returns a new image (GPU images are immutable). + +use effectcraft_color::{BlendMode, Conversion}; +use effectcraft_geom::{Mat3, Rect, vec2}; +use effectcraft_project::MatteKind; +use effectcraft_raster::Sampling; + +use crate::context::{Enc, GpuImage, Params}; + +/// Index of a blend mode in `BlendMode::ALL` (the kernels' mode numbers). +pub fn mode_id(m: BlendMode) -> u32 { + BlendMode::ALL.iter().position(|x| *x == m).unwrap_or(0) as u32 +} + +fn sampling_id(s: Sampling) -> u32 { + match s { + Sampling::Nearest => 0, + Sampling::Bilinear => 1, + Sampling::Bicubic => 2, + } +} + +fn rows(m: &Mat3) -> [[f32; 4]; 3] { + m.0.map(|r| [r[0] as f32, r[1] as f32, r[2] as f32, 0.0]) +} + +/// Box-filter halving (minification pre-filter). +pub fn half(e: &mut Enc, src: &GpuImage) -> GpuImage { + let out = e.image(src.width.div_ceil(2).max(1), src.height.div_ceil(2).max(1)); + e.pixels("half", &Params::default(), src, None, &out, None); + out +} + +/// `raster::composite_warp` (and, with `accumulate = Some(weight)`, `accumulate_warp`): draw +/// `src` transformed by `m` (src pixels → dst pixels) onto `dst`. +pub fn warp( + e: &mut Enc, + dst: &GpuImage, + src: &GpuImage, + m: &Mat3, + sampling: Sampling, + mode: BlendMode, + opacity: f32, + seed: u32, + accumulate: Option, +) -> GpuImage { + if src.width == 0 || src.height == 0 || opacity <= 0.0 { + return dst.clone(); + } + let Some(inv0) = m.inverse() else { return dst.clone() }; + // Minification: pick a pre-filtered level so sampling doesn't alias. + let scale = m.mean_scale(); + let mut img = src.clone(); + let mut inv = inv0; + if scale < 0.5 && m.is_affine() { + let mut s = scale * 2.0; + let mut f = 0.5; + img = half(e, src); + while s < 0.5 && img.width > 1 && img.height > 1 { + img = half(e, &img); + f *= 0.5; + s *= 2.0; + } + inv = Mat3::scale(vec2(f, f)) * inv; + } + let bounds = m.map_rect(&Rect::from_size(src.width as f64, src.height as f64)).inflate(1.0).round_out(); + let area = bounds.intersect(&Rect::from_size(dst.width as f64, dst.height as f64)); + if area.is_empty() { + return dst.clone(); + } + let (x0, x1) = (area.x0.max(0.0) as u32, area.x1.min(dst.width as f64) as u32); + let (y0, y1) = (area.y0.max(0.0) as u32, area.y1.min(dst.height as f64) as u32); + let mut p = Params::default(); + let clear = matches!(mode, BlendMode::StencilAlpha | BlendMode::StencilLuma); + let flags = u32::from(accumulate.is_some()) | if clear && accumulate.is_none() { 2 } else { 0 }; + p.u[0] = [mode_id(mode), sampling_id(sampling), seed, flags]; + let r = rows(&inv); + p.f[0] = r[0]; + p.f[1] = r[1]; + p.f[2] = r[2]; + p.f[3] = [opacity, accumulate.unwrap_or(1.0), 0.0, 0.0]; + p.f[4] = [x0 as f32, y0 as f32, x1 as f32, y1 as f32]; + let out = e.image(dst.width, dst.height); + e.pixels("warp_blend", &p, &img, Some(dst), &out, None); + out +} + +/// `Image::blend_from`: blend `src` (same size) onto `dst`. +pub fn blend_full(e: &mut Enc, dst: &GpuImage, src: &GpuImage, mode: BlendMode, opacity: f32, seed: u32) -> GpuImage { + let mut p = Params::default(); + p.u[0] = [mode_id(mode), 0, seed, 0]; + p.f[0] = [opacity, 0.0, 0.0, 0.0]; + let out = e.image(dst.width, dst.height); + e.pixels("blend_full", &p, src, Some(dst), &out, None); + out +} + +/// Multiply an isolated layer by its track matte. +pub fn matte(e: &mut Enc, iso: &GpuImage, matte: &GpuImage, kind: MatteKind) -> GpuImage { + let mut p = Params::default(); + p.u[0][0] = match kind { + MatteKind::Alpha => 0, + MatteKind::AlphaInverted => 1, + MatteKind::Luma => 2, + MatteKind::LumaInverted => 3, + }; + let out = e.image(iso.width, iso.height); + e.pixels("matte", &p, iso, Some(matte), &out, None); + out +} + +/// Preserve Transparency: multiply by the alpha below. +pub fn preserve(e: &mut Enc, iso: &GpuImage, canvas: &GpuImage) -> GpuImage { + let out = e.image(iso.width, iso.height); + e.pixels("preserve", &Params::default(), iso, Some(canvas), &out, None); + out +} + +/// Knockout: clear what lies below the layer's content. +pub fn knockout(e: &mut Enc, canvas: &GpuImage, shape: &GpuImage) -> GpuImage { + let out = e.image(canvas.width, canvas.height); + e.pixels("knockout", &Params::default(), canvas, Some(shape), &out, None); + out +} + +/// Styled-layer finish: move `canvas` toward `tmp` by `opacity` on the enabled channels. +pub fn channel_mix(e: &mut Enc, canvas: &GpuImage, tmp: &GpuImage, channels: [bool; 3], opacity: f32) -> GpuImage { + let mut p = Params::default(); + p.u[0] = [channels[0] as u32, channels[1] as u32, channels[2] as u32, 0]; + p.f[0][0] = opacity; + let out = e.image(canvas.width, canvas.height); + e.pixels("channel_mix", &p, canvas, Some(tmp), &out, None); + out +} + +/// Clamp and quantise to `levels` steps (8/16 bpc). +pub fn quantize(e: &mut Enc, img: &GpuImage, levels: f32) -> GpuImage { + let mut p = Params::default(); + p.f[0][0] = levels; + let out = e.image(img.width, img.height); + e.pixels("quantize", &p, img, None, &out, None); + out +} + +fn curve_id(s: Option) -> u32 { + use effectcraft_color::ColorSpace::*; + match s { + None => 0, + Some(Srgb | DisplayP3) => 1, + Some(Rec709 | Rec2020) => 2, + } +} + +/// Colour conversion on straight colour (`render::color::convert`). +pub fn convert(e: &mut Enc, img: &GpuImage, c: &Conversion) -> GpuImage { + let mut p = Params::default(); + p.u[0] = [curve_id(c.decode), c.matrix.is_some() as u32, curve_id(c.encode), 0]; + if let Some(m) = c.matrix { + for i in 0..3 { + p.f[i] = [m[i][0], m[i][1], m[i][2], 0.0]; + } + } + let out = e.image(img.width, img.height); + e.pixels("convert", &p, img, None, &out, None); + out +} diff --git a/crates/gpu/src/shaders/common.wgsl b/crates/gpu/src/shaders/common.wgsl new file mode 100644 index 00000000..dee86b16 --- /dev/null +++ b/crates/gpu/src/shaders/common.wgsl @@ -0,0 +1,452 @@ +// EffectCraft GPU compositor: shared bindings and helpers. +// +// Every kernel uses the same bind group: a uniform block of parameters, two input textures +// (premultiplied RGBA f32, read with textureLoad), one output storage texture and a read-only +// storage buffer of extra data (curve tables). Each helper mirrors a CPU function in +// effectcraft-raster / effectcraft-color operation for operation, so results match the CPU +// reference to float rounding. + +struct Params { + u: array, 4>, + f: array, 12>, +} + +@group(0) @binding(0) var P: Params; +@group(0) @binding(1) var src: texture_2d; +@group(0) @binding(2) var aux: texture_2d; +@group(0) @binding(4) var data: array; + +// ---------------------------------------------------------------- sampling (raster::Image) + +fn tex_get(t: texture_2d, x: i32, y: i32) -> vec4 { + let d = vec2(textureDimensions(t)); + if (x < 0 || y < 0 || x >= d.x || y >= d.y) { + return vec4(0.0); + } + return textureLoad(t, vec2(x, y), 0); +} + +fn tex_get_clamped(t: texture_2d, x: i32, y: i32) -> vec4 { + let d = vec2(textureDimensions(t)); + return textureLoad(t, vec2(clamp(x, 0, d.x - 1), clamp(y, 0, d.y - 1)), 0); +} + +// Image::sample_bilinear (pixel centres at +0.5; transparent outside). +fn sample_bilinear(t: texture_2d, x: f32, y: f32) -> vec4 { + let fx = x - 0.5; + let fy = y - 0.5; + let x0 = floor(fx); + let y0 = floor(fy); + let tx = fx - x0; + let ty = fy - y0; + let xi = i32(x0); + let yi = i32(y0); + let a = tex_get(t, xi, yi); + let b = tex_get(t, xi + 1, yi); + let c = tex_get(t, xi, yi + 1); + let d = tex_get(t, xi + 1, yi + 1); + let top = a + (b - a) * tx; + let bot = c + (d - c) * tx; + return top + (bot - top) * ty; +} + +// Image::sample_bilinear_clamped (edge pixels repeat). +fn sample_bilinear_clamped(t: texture_2d, x: f32, y: f32) -> vec4 { + let fx = x - 0.5; + let fy = y - 0.5; + let x0 = floor(fx); + let y0 = floor(fy); + let tx = fx - x0; + let ty = fy - y0; + let xi = i32(x0); + let yi = i32(y0); + let a = tex_get_clamped(t, xi, yi); + let b = tex_get_clamped(t, xi + 1, yi); + let c = tex_get_clamped(t, xi, yi + 1); + let d = tex_get_clamped(t, xi + 1, yi + 1); + let top = a + (b - a) * tx; + let bot = c + (d - c) * tx; + return top + (bot - top) * ty; +} + +fn cubic_w(t: f32) -> vec4 { + let t2 = t * t; + let t3 = t2 * t; + return vec4(-0.5 * t3 + t2 - 0.5 * t, 1.5 * t3 - 2.5 * t2 + 1.0, -1.5 * t3 + 2.0 * t2 + 0.5 * t, 0.5 * t3 - 0.5 * t2); +} + +// Image::sample_bicubic (Catmull-Rom; negatives clamped after ringing). +fn sample_bicubic(t: texture_2d, x: f32, y: f32) -> vec4 { + let fx = x - 0.5; + let fy = y - 0.5; + let x0 = floor(fx); + let y0 = floor(fy); + let wx = cubic_w(fx - x0); + let wy = cubic_w(fy - y0); + let xi = i32(x0); + let yi = i32(y0); + var o = vec4(0.0); + for (var j = 0; j < 4; j++) { + for (var i = 0; i < 4; i++) { + let p = tex_get(t, xi - 1 + i, yi - 1 + j); + o += p * (wx[i] * wy[j]); + } + } + return max(o, vec4(0.0)); +} + +// 0 nearest, 1 bilinear, 2 bicubic (raster::Sampling). +fn sample(t: texture_2d, sampling: u32, x: f32, y: f32) -> vec4 { + if (sampling == 0u) { + return tex_get(t, i32(floor(x)), i32(floor(y))); + } + if (sampling == 2u) { + return sample_bicubic(t, x, y); + } + return sample_bilinear(t, x, y); +} + +// raster::hash_noise. +fn hash_noise(x: u32, y: u32, seed: u32) -> f32 { + var h = (x * 0x8da6b343u) ^ (y * 0xd8163841u) ^ (seed * 0xcb1ab31fu); + h = h ^ (h >> 13u); + h = h * 0x5bd1e995u; + h = h ^ (h >> 15u); + return f32(h & 0x00ffffffu) / 16777216.0; +} + +// powf for x >= 0 (0^y = 0 for y > 0; WGSL leaves pow(0, y) to the implementation). +fn powz(x: f32, y: f32) -> f32 { + if (x <= 0.0) { + return select(0.0, 1.0, y == 0.0); + } + return pow(x, y); +} + +// Rust's f32::round (half away from zero; WGSL's round() is half to even). +fn round_away(x: f32) -> f32 { + let t = trunc(x); + if (abs(x - t) >= 0.5) { + return t + sign(x); + } + return t; +} + +// f32::rem_euclid(1.0). +fn fract_euclid(x: f32) -> f32 { + return x - floor(x); +} + +fn luminance(c: vec3) -> f32 { + return 0.2126 * c.x + 0.7152 * c.y + 0.0722 * c.z; +} + +// ---------------------------------------------------------------- blend modes (color::blend) +// Mode numbers are indices into BlendMode::ALL. + +fn burn(cb: f32, cs: f32) -> f32 { + if (cb >= 1.0) { + return 1.0; + } + if (cs <= 0.0) { + return 0.0; + } + return 1.0 - min((1.0 - cb) / cs, 1.0); +} + +fn dodge(cb: f32, cs: f32) -> f32 { + if (cb <= 0.0) { + return 0.0; + } + if (cs >= 1.0) { + return 1.0; + } + return min(cb / (1.0 - cs), 1.0); +} + +fn hard_light(cb: f32, cs: f32) -> f32 { + if (cs <= 0.5) { + return cb * 2.0 * cs; + } + let s = 2.0 * cs - 1.0; + return cb + s - cb * s; +} + +fn soft_light(cb: f32, cs: f32) -> f32 { + if (cs <= 0.5) { + return cb - (1.0 - 2.0 * cs) * cb * (1.0 - cb); + } + var d: f32; + if (cb <= 0.25) { + d = ((16.0 * cb - 12.0) * cb + 4.0) * cb; + } else { + d = sqrt(max(cb, 0.0)); + } + return cb + (2.0 * cs - 1.0) * (d - cb); +} + +fn vivid(cb: f32, cs: f32) -> f32 { + if (cs <= 0.5) { + return burn(cb, 2.0 * cs); + } + return dodge(cb, 2.0 * (cs - 0.5)); +} + +fn separable(mode: u32, cb: f32, cs: f32) -> f32 { + switch mode { + case 3u: { return min(cb, cs); } + case 4u: { return cb * cs; } + case 5u: { return burn(cb, cs); } + case 6u: { + if (cs <= 0.0) { + return 0.0; + } + return max(1.0 - (1.0 - cb) / cs, 0.0); + } + case 7u: { return max(cb + cs - 1.0, 0.0); } + case 9u, 14u: { return cb + cs; } + case 10u: { return max(cb, cs); } + case 11u: { return cb + cs - cb * cs; } + case 12u: { return dodge(cb, cs); } + case 13u: { + if (cs >= 1.0) { + return 1.0; + } + return min(cb / (1.0 - cs), 1.0); + } + case 16u: { return hard_light(cs, cb); } + case 17u: { return soft_light(cb, cs); } + case 18u: { return hard_light(cb, cs); } + case 19u: { return clamp(cb + 2.0 * cs - 1.0, 0.0, 1.0); } + case 20u: { return vivid(cb, cs); } + case 21u: { + if (cs <= 0.5) { + return min(cb, 2.0 * cs); + } + return max(cb, 2.0 * cs - 1.0); + } + case 22u: { + if (cb + cs >= 1.0) { + return 1.0; + } + return 0.0; + } + case 23u, 24u: { return abs(cb - cs); } + case 25u: { return cb + cs - 2.0 * cb * cs; } + case 26u: { return max(cb - cs, 0.0); } + case 27u: { + if (cs <= 0.0) { + if (cb > 0.0) { + return 1.0; + } + return 0.0; + } + return cb / cs; + } + default: { return cs; } + } +} + +fn blum(c: vec3) -> f32 { + return 0.3 * c.x + 0.59 * c.y + 0.11 * c.z; +} + +fn clip_color(c: vec3) -> vec3 { + let l = blum(c); + let n = min(min(c.x, c.y), c.z); + let x = max(max(c.x, c.y), c.z); + var o = c; + if (n < 0.0) { + o = l + (o - l) * l / max(l - n, 1e-9); + } + if (x > 1.0) { + o = l + (o - l) * (1.0 - l) / max(x - l, 1e-9); + } + return o; +} + +fn set_lum(c: vec3, l: f32) -> vec3 { + let d = l - blum(c); + return clip_color(c + d); +} + +fn bsat(c: vec3) -> f32 { + return max(max(c.x, c.y), c.z) - min(min(c.x, c.y), c.z); +} + +fn set_sat(c: vec3, s: f32) -> vec3 { + let mx = max(max(c.x, c.y), c.z); + let mn = min(min(c.x, c.y), c.z); + if (mx - mn <= 1e-9) { + return vec3(0.0); + } + return (c - mn) * s / (mx - mn); +} + +fn non_separable(mode: u32, cb: vec3, cs: vec3) -> vec3 { + switch mode { + case 28u: { return set_lum(set_sat(cs, bsat(cb)), blum(cb)); } + case 29u: { return set_lum(set_sat(cb, bsat(cs)), blum(cb)); } + case 30u: { return set_lum(cs, blum(cb)); } + case 31u: { return set_lum(cb, blum(cs)); } + case 8u: { + if (luminance(cs) < luminance(cb)) { + return cs; + } + return cb; + } + case 15u: { + if (luminance(cs) > luminance(cb)) { + return cs; + } + return cb; + } + default: { + return vec3(separable(mode, cb.x, cs.x), separable(mode, cb.y, cs.y), separable(mode, cb.z, cs.z)); + } + } +} + +fn supports_hdr(mode: u32) -> bool { + switch mode { + case 0u, 1u, 2u, 3u, 4u, 7u, 8u, 9u, 10u, 11u, 14u, 15u, 23u, 24u, 26u, 27u, 32u, 33u, 34u, 35u, 36u, 37u: { return true; } + default: { return false; } + } +} + +fn is_stencil(mode: u32) -> bool { + return mode >= 32u && mode <= 35u; +} + +// color::blend_pixel. +fn blend_pixel(mode: u32, dst: vec4, src: vec4, noise: f32) -> vec4 { + let sa = src.w; + let da = dst.w; + switch mode { + case 0u: { + let k = 1.0 - sa; + return vec4(src.xyz + dst.xyz * k, sa + da * k); + } + case 1u, 2u: { + if (sa <= 0.0 || noise >= sa) { + return dst; + } + let inv = 1.0 / sa; + return vec4(src.xyz * inv, 1.0); + } + case 32u: { return dst * sa; } + case 34u: { return dst * (1.0 - sa); } + case 33u: { return dst * luminance(src.xyz); } + case 35u: { return dst * (1.0 - luminance(src.xyz)); } + case 36u: { return vec4(src.xyz + dst.xyz * (1.0 - sa), min(sa + da, 1.0)); } + case 37u: { return vec4(src.xyz + dst.xyz * (1.0 - sa), sa + da - sa * da); } + default: {} + } + if (sa <= 0.0) { + return dst; + } + if (da <= 0.0) { + return src; + } + var cs = src.xyz / sa; + var cb = dst.xyz / da; + if (!supports_hdr(mode)) { + cs = clamp(cs, vec3(0.0), vec3(1.0)); + cb = clamp(cb, vec3(0.0), vec3(1.0)); + } + let b = non_separable(mode, cb, cs); + let ao = sa + da - sa * da; + return vec4((1.0 - da) * src.xyz + (1.0 - sa) * dst.xyz + sa * da * b, ao); +} + +// ---------------------------------------------------------------- colour (color crate) + +fn srgb_to_linear(v: f32) -> f32 { + if (v <= 0.04045) { + return v / 12.92; + } + return pow((v + 0.055) / 1.055, 2.4); +} + +fn linear_to_srgb(v: f32) -> f32 { + if (v <= 0.0031308) { + return v * 12.92; + } + return 1.055 * pow(max(v, 0.0), 1.0 / 2.4) - 0.055; +} + +fn rgb_to_hsl(c: vec3) -> vec3 { + let mx = max(max(c.x, c.y), c.z); + let mn = min(min(c.x, c.y), c.z); + let l = (mx + mn) * 0.5; + if (abs(mx - mn) < 1e-7) { + return vec3(0.0, 0.0, l); + } + let d = mx - mn; + var s: f32; + if (l > 0.5) { + s = d / (2.0 - mx - mn); + } else { + s = d / (mx + mn); + } + var h: f32; + if (mx == c.x) { + h = ((c.y - c.z) / d + select(0.0, 6.0, c.y < c.z)) / 6.0; + } else if (mx == c.y) { + h = ((c.z - c.x) / d + 2.0) / 6.0; + } else { + h = ((c.x - c.y) / d + 4.0) / 6.0; + } + return vec3(h, s, l); +} + +fn hsl_channel(t0: f32, p: f32, q: f32) -> f32 { + let t = fract_euclid(t0); + if (t < 1.0 / 6.0) { + return p + (q - p) * 6.0 * t; + } + if (t < 0.5) { + return q; + } + if (t < 2.0 / 3.0) { + return p + (q - p) * (2.0 / 3.0 - t) * 6.0; + } + return p; +} + +fn hsl_to_rgb(h: f32, s: f32, l: f32) -> vec3 { + if (s <= 0.0) { + return vec3(l); + } + var q: f32; + if (l < 0.5) { + q = l * (1.0 + s); + } else { + q = l + s - l * s; + } + let p = 2.0 * l - q; + return vec3(hsl_channel(h + 1.0 / 3.0, p, q), hsl_channel(h, p, q), hsl_channel(h - 1.0 / 3.0, p, q)); +} + +// ColorSpace transfer curves: 1 = sRGB (sRGB, Display P3), 2 = 2.4 power (Rec. 709 / 2020). +fn curve_decode(curve: u32, v: f32) -> f32 { + let a = abs(v); + var l = a; + if (curve == 1u) { + l = srgb_to_linear(a); + } else if (curve == 2u) { + l = powz(a, 2.4); + } + return select(l, -l, v < 0.0); +} + +fn curve_encode(curve: u32, v: f32) -> f32 { + let a = abs(v); + var e = a; + if (curve == 1u) { + e = linear_to_srgb(a); + } else if (curve == 2u) { + e = powz(a, 1.0 / 2.4); + } + return select(e, -e, v < 0.0); +} diff --git a/crates/gpu/src/shaders/display.wgsl b/crates/gpu/src/shaders/display.wgsl new file mode 100644 index 00000000..68261875 --- /dev/null +++ b/crates/gpu/src/shaders/display.wgsl @@ -0,0 +1,19 @@ +// Viewer display: premultiplied f32 → premultiplied RGBA8 (ui-egui's frames::to_color_image), +// written straight into a texture egui samples. + +@group(0) @binding(1) var dsrc: texture_2d; +@group(0) @binding(3) var dout: texture_storage_2d; + +@compute @workgroup_size(16, 16) +fn display(@builtin(global_invocation_id) gid: vec3) { + let dims = vec2(textureDimensions(dout)); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + let s = textureLoad(dsrc, p, 0); + let a = clamp(s.w, 0.0, 1.0); + // (v · 255 + 0.5) truncated, as the CPU conversion. + let q = floor(vec4(clamp(s.xyz, vec3(0.0), vec3(a)), a) * 255.0 + 0.5) / 255.0; + textureStore(dout, p, q); +} diff --git a/crates/gpu/src/shaders/kernels.wgsl b/crates/gpu/src/shaders/kernels.wgsl new file mode 100644 index 00000000..98ead549 --- /dev/null +++ b/crates/gpu/src/shaders/kernels.wgsl @@ -0,0 +1,750 @@ +// Compositing and effect kernels (rgba32float output). Parameter layouts are documented next to +// each entry point and mirrored by the Rust dispatch code (src/ops.rs, src/effects.rs). + +@group(0) @binding(3) var out: texture_storage_2d; + +fn out_dims() -> vec2 { + return vec2(textureDimensions(out)); +} + +// ---------------------------------------------------------------- compositing + +// Warp `src` by the inverse matrix (output pixel → source pixel) and blend over `aux` +// (raster::composite_warp / accumulate_warp). +// u[0] = (mode, sampling, seed, flags: 1 accumulate, 2 clear outside bounds) +// f[0..3] = inverse rows; f[3] = (opacity, weight); f[4] = bounds (x0, y0, x1, y1) +@compute @workgroup_size(16, 16) +fn warp_blend(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let x = i32(gid.x); + let y = i32(gid.y); + if (x >= dims.x || y >= dims.y) { + return; + } + let mode = P.u[0].x; + let sampling = P.u[0].y; + let seed = P.u[0].z; + let flags = P.u[0].w; + let b = P.f[4]; + var d = textureLoad(aux, vec2(x, y), 0); + let fx = f32(x) + 0.5; + let fy = f32(y) + 0.5; + if (fx < b.x || fy < b.y || fx > b.z || fy > b.w) { + if ((flags & 2u) != 0u) { + d = vec4(0.0); + } + textureStore(out, vec2(x, y), d); + return; + } + let r0 = P.f[0]; + let r1 = P.f[1]; + let r2 = P.f[2]; + var px = r0.x * fx + r0.y * fy + r0.z; + var py = r1.x * fx + r1.y * fy + r1.z; + let w = r2.x * fx + r2.y * fy + r2.z; + if (w != 1.0 && w != 0.0) { + px = px / w; + py = py / w; + } + let sz = vec2(textureDimensions(src)); + let stencil = is_stencil(mode); + if (px < -1.0 || py < -1.0 || px > sz.x + 1.0 || py > sz.y + 1.0) { + if (stencil && (flags & 1u) == 0u) { + d = blend_pixel(mode, d, vec4(0.0), 0.5); + } + textureStore(out, vec2(x, y), d); + return; + } + var s = sample(src, sampling, px, py); + if ((flags & 1u) != 0u) { + if (s.w > 0.0) { + d = d + s * P.f[3].y; + } + textureStore(out, vec2(x, y), d); + return; + } + if (s.w <= 0.0 && !stencil) { + textureStore(out, vec2(x, y), d); + return; + } + let op = P.f[3].x; + if (op < 1.0) { + s = s * op; + } + var n = 0.5; + if (mode == 1u || mode == 2u) { + n = hash_noise(u32(x), u32(y), seed); + } + textureStore(out, vec2(x, y), blend_pixel(mode, d, s, n)); +} + +// Blend `src` (same size) onto `aux` (Image::blend_from). +// u[0] = (mode, _, seed); f[0].x = opacity +@compute @workgroup_size(16, 16) +fn blend_full(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + let mode = P.u[0].x; + let d = textureLoad(aux, p, 0); + var s = textureLoad(src, p, 0); + if (s.w <= 0.0 && !is_stencil(mode)) { + textureStore(out, p, d); + return; + } + s = s * P.f[0].x; + var n = 0.5; + if (mode == 1u || mode == 2u) { + n = hash_noise(gid.x, gid.y, P.u[0].z); + } + textureStore(out, p, blend_pixel(mode, d, s, n)); +} + +// Track matte: `src` × matte factor of `aux`. u[0].x = 0 alpha, 1 alpha inverted, 2 luma, +// 3 luma inverted. +@compute @workgroup_size(16, 16) +fn matte(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + let q = textureLoad(aux, p, 0); + var k: f32; + switch P.u[0].x { + case 0u: { k = q.w; } + case 1u: { k = 1.0 - q.w; } + case 2u: { k = luminance(q.xyz); } + default: { k = 1.0 - luminance(q.xyz); } + } + k = clamp(k, 0.0, 1.0); + textureStore(out, p, textureLoad(src, p, 0) * k); +} + +// Preserve Transparency: `src` × alpha of `aux`. +@compute @workgroup_size(16, 16) +fn preserve(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + textureStore(out, p, textureLoad(src, p, 0) * textureLoad(aux, p, 0).w); +} + +// Knockout: `src` × (1 − alpha of `aux`). +@compute @workgroup_size(16, 16) +fn knockout(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + let keep = 1.0 - clamp(textureLoad(aux, p, 0).w, 0.0, 1.0); + textureStore(out, p, textureLoad(src, p, 0) * keep); +} + +// Styled layer finish: canvas `src` moves toward `aux` by the opacity, per enabled channel. +// u[0].xyz = R/G/B switches; f[0].x = opacity. +@compute @workgroup_size(16, 16) +fn channel_mix(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + let c = textureLoad(src, p, 0); + let t = textureLoad(aux, p, 0); + let op = P.f[0].x; + var v = t; + if (P.u[0].x == 0u) { + v.x = c.x; + } + if (P.u[0].y == 0u) { + v.y = c.y; + } + if (P.u[0].z == 0u) { + v.z = c.z; + } + textureStore(out, p, c + (v - c) * op); +} + +// 8/16 bpc: clamp premultiplied pixels to 0..1 (colour ≤ alpha) and round. f[0].x = levels. +@compute @workgroup_size(16, 16) +fn quantize(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + let levels = P.f[0].x; + let inv = 1.0 / levels; + let s = textureLoad(src, p, 0); + let a = round_away(clamp(s.w, 0.0, 1.0) * levels) * inv; + var o = vec4(0.0, 0.0, 0.0, a); + o.x = round_away(clamp(s.x, 0.0, a) * levels) * inv; + o.y = round_away(clamp(s.y, 0.0, a) * levels) * inv; + o.z = round_away(clamp(s.z, 0.0, a) * levels) * inv; + textureStore(out, p, o); +} + +// Colour conversion on straight colour (render::color::convert). +// u[0] = (decode curve, has matrix, encode curve); f[0..3] = matrix rows. +@compute @workgroup_size(16, 16) +fn convert(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + let s = textureLoad(src, p, 0); + let a = s.w; + if (a <= 0.0) { + textureStore(out, p, s); + return; + } + var c = s.xyz / a; + let dc = P.u[0].x; + if (dc != 0u) { + c = vec3(curve_decode(dc, c.x), curve_decode(dc, c.y), curve_decode(dc, c.z)); + } + if (P.u[0].y != 0u) { + let m0 = P.f[0].xyz; + let m1 = P.f[1].xyz; + let m2 = P.f[2].xyz; + c = vec3(m0.x * c.x + m0.y * c.y + m0.z * c.z, m1.x * c.x + m1.y * c.y + m1.z * c.z, m2.x * c.x + m2.y * c.y + m2.z * c.z); + } + let ec = P.u[0].z; + if (ec != 0u) { + c = vec3(curve_encode(ec, c.x), curve_encode(ec, c.y), curve_encode(ec, c.z)); + } + textureStore(out, p, vec4(c * a, a)); +} + +// Box-filter halving for minification (raster::warp::half). +@compute @workgroup_size(16, 16) +fn half(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + let sx = 2 * p.x; + let sy = 2 * p.y; + var acc = vec4(0.0); + acc += tex_get_clamped(src, sx, sy); + acc += tex_get_clamped(src, sx + 1, sy); + acc += tex_get_clamped(src, sx, sy + 1); + acc += tex_get_clamped(src, sx + 1, sy + 1); + textureStore(out, p, acc * 0.25); +} + +// ---------------------------------------------------------------- blurs (raster::blur) + +fn line_get(horizontal: bool, line: i32, i: i32, n: i32, repeat: bool) -> vec4 { + var j = i; + if (i < 0 || i >= n) { + if (!repeat) { + return vec4(0.0); + } + j = clamp(i, 0, n - 1); + } + if (horizontal) { + return textureLoad(src, vec2(j, line), 0); + } + return textureLoad(src, vec2(line, j), 0); +} + +// One box pass of radius r along rows (horizontal) or columns, as running sums over blocks of +// `block` pixels (each block restarts its sum, as the CPU's vertical pass does). +fn box_pass(gid: vec3, horizontal: bool) { + let dims = out_dims(); + let r = i32(P.u[0].x); + let repeat = P.u[0].y != 0u; + let block = i32(P.u[0].z); + var lines = dims.x; + var n = dims.y; + if (horizontal) { + lines = dims.y; + n = dims.x; + } + let line = i32(gid.x); + let b0 = i32(gid.y) * block; + if (line >= lines || b0 >= n) { + return; + } + let norm = 1.0 / f32(2 * r + 1); + var acc = vec4(0.0); + for (var i = b0 - r; i <= b0 + r; i++) { + acc += line_get(horizontal, line, i, n, repeat); + } + let b1 = min(b0 + block, n); + for (var i = b0; i < b1; i++) { + if (horizontal) { + textureStore(out, vec2(i, line), acc * norm); + } else { + textureStore(out, vec2(line, i), acc * norm); + } + let add = line_get(horizontal, line, i + r + 1, n, repeat); + let sub = line_get(horizontal, line, i - r, n, repeat); + acc += add - sub; + } +} + +// u[0] = (radius, repeat edges, block); dispatched over (line, block). +@compute @workgroup_size(64, 1) +fn box_h(@builtin(global_invocation_id) gid: vec3) { + box_pass(gid, true); +} + +@compute @workgroup_size(64, 1) +fn box_v(@builtin(global_invocation_id) gid: vec3) { + box_pass(gid, false); +} + +// Directional blur (raster::directional_blur). f[0] = (dx, dy, length, samples). +@compute @workgroup_size(16, 16) +fn directional(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + let dx = P.f[0].x; + let dy = P.f[0].y; + let len = P.f[0].z; + let n = i32(P.f[0].w); + var acc = vec4(0.0); + for (var i = 0; i < n; i++) { + let t = (f32(i) / f32(n - 1) - 0.5) * len; + acc += sample_bilinear(src, f32(p.x) + 0.5 + dx * t, f32(p.y) + 0.5 + dy * t); + } + textureStore(out, p, acc / f32(n)); +} + +// ---------------------------------------------------------------- Glow / Drop Shadow + +// Glow bright pass. u[0].x = A & B colours; f[0].x = threshold; f[1] = colour A; f[2] = colour B. +@compute @workgroup_size(16, 16) +fn glow_bright(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + let px = textureLoad(src, p, 0); + let a = px.w; + if (a <= 0.0) { + textureStore(out, p, vec4(0.0)); + return; + } + let thr = P.f[0].x; + let l = luminance(px.xyz / a); + let k = clamp((l - thr) / max(1.0 - thr, 1e-3), 0.0, 1.0); + if (P.u[0].x != 0u) { + let t = clamp(l, 0.0, 1.0); + let ca = P.f[1].xyz; + let cb = P.f[2].xyz; + let c = ca + (cb - ca) * t; + textureStore(out, p, vec4(c * k * a, k * a)); + } else { + textureStore(out, p, px * k); + } +} + +// Glow composite: original `src` with the blurred glow `aux`. u[0].x = behind; f[0].x = intensity. +@compute @workgroup_size(16, 16) +fn glow_combine(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + var o = textureLoad(src, p, 0); + let g = textureLoad(aux, p, 0) * P.f[0].x; + if (P.u[0].x != 0u) { + let k = 1.0 - o.w; + o = o + g * k; + } else { + o = vec4(o.xyz + g.xyz, min(o.w + g.w * (1.0 - o.w), 1.0)); + } + o.w = clamp(o.w, 0.0, 1.0); + textureStore(out, p, o); +} + +// Drop shadow silhouette. f[0] = shadow colour; f[1] = (opacity, dx, dy). +@compute @workgroup_size(16, 16) +fn shadow_make(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + let c = P.f[0]; + let a = sample_bilinear(src, f32(p.x) + 0.5 - P.f[1].y, f32(p.y) + 0.5 - P.f[1].z).w * P.f[1].x * c.w; + textureStore(out, p, vec4(c.xyz * a, a)); +} + +// The layer over its shadow: `aux` (original) over `src` (shadow). +@compute @workgroup_size(16, 16) +fn shadow_combine(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + let s = textureLoad(src, p, 0); + let o = textureLoad(aux, p, 0); + textureStore(out, p, o + s * (1.0 - o.w)); +} + +// ---------------------------------------------------------------- per-pixel effects + +fn mix3(a: vec3, b: vec3, t: f32) -> vec3 { + return a + (b - a) * t; +} + +// Curves: data = [n, xs.., ys.., ms.., lut (1025)] per channel at offsets f[c]. +fn curve_exact(base: u32, x: f32) -> f32 { + let n = u32(data[base]); + let xs = base + 1u; + let ys = xs + n; + let ms = ys + n; + if (x <= data[xs]) { + return data[ys] + (x - data[xs]) * data[ms]; + } + if (x >= data[xs + n - 1u]) { + return data[ys + n - 1u] + (x - data[xs + n - 1u]) * data[ms + n - 1u]; + } + // partition_point(|v| v <= x) - 1, clamped to n - 2. + var k = 0u; + for (var i = 0u; i < n; i++) { + if (data[xs + i] <= x) { + k = i + 1u; + } + } + k = min(max(k, 1u) - 1u, n - 2u); + let h = data[xs + k + 1u] - data[xs + k]; + let t = (x - data[xs + k]) / h; + let t2 = t * t; + let t3 = t * t * t; + let h00 = 2.0 * t3 - 3.0 * t2 + 1.0; + let h10 = t3 - 2.0 * t2 + t; + let h01 = -2.0 * t3 + 3.0 * t2; + let h11 = t3 - t2; + return h00 * data[ys + k] + h10 * h * data[ms + k] + h01 * data[ys + k + 1u] + h11 * h * data[ms + k + 1u]; +} + +fn curve_eval(base_f: f32, x: f32) -> f32 { + if (base_f < 0.0) { + return x; + } + let base = u32(base_f); + if (!(x >= 0.0 && x <= 1.0)) { + return curve_exact(base, x); + } + let n = u32(data[base]); + let lut = base + 1u + 3u * n; + let f = x * 1024.0; + let i = min(u32(f), 1023u); + let t = f - f32(i); + return data[lut + i] + (data[lut + i + 1u] - data[lut + i]) * t; +} + +// Fractal Noise lattice (effects::generate::lattice / value_noise). +fn lattice(ix: i32, iy: i32, iz: i32, seed: u32) -> f32 { + var h = (bitcast(ix) * 0x27d4eb2du) ^ (bitcast(iy) * 0x165667b1u) ^ (bitcast(iz) * 0x9e3779b9u) ^ (seed * 0x85ebca6bu); + h = h ^ (h >> 15u); + h = h * 0x2c1b3c6du; + h = h ^ (h >> 12u); + h = h * 0x297a2d39u; + h = h ^ (h >> 15u); + return f32(h & 0xffffu) / 65535.0; +} + +fn fade(t: f32) -> f32 { + return t * t * t * (t * (t * 6.0 - 15.0) + 10.0); +} + +fn value_noise(x: f32, y: f32, z: f32, seed: u32) -> f32 { + let x0 = floor(x); + let y0 = floor(y); + let z0 = floor(z); + let fx = fade(x - x0); + let fy = fade(y - y0); + let fz = fade(z - z0); + let ix = i32(x0); + let iy = i32(y0); + let iz = i32(z0); + let c000 = lattice(ix, iy, iz, seed); + let c100 = lattice(ix + 1, iy, iz, seed); + let c010 = lattice(ix, iy + 1, iz, seed); + let c110 = lattice(ix + 1, iy + 1, iz, seed); + let c001 = lattice(ix, iy, iz + 1, seed); + let c101 = lattice(ix + 1, iy, iz + 1, seed); + let c011 = lattice(ix, iy + 1, iz + 1, seed); + let c111 = lattice(ix + 1, iy + 1, iz + 1, seed); + let l00 = c000 + (c100 - c000) * fx; + let l01 = c010 + (c110 - c010) * fx; + let a = l00 + (l01 - l00) * fy; + let l10 = c001 + (c101 - c001) * fx; + let l11 = c011 + (c111 - c011) * fx; + let b = l10 + (l11 - l10) * fy; + return a + (b - a) * fz; +} + +// generate::put: blend a generated straight colour (alpha ga) over the pixel. +fn put(px: vec4, g: vec3, ga: f32, blend_orig: f32) -> vec4 { + let g2 = vec4(g * ga, ga); + let k = 1.0 - blend_orig; + return px * blend_orig + g2 * k; +} + +// Straight-colour map helper: skips transparent pixels (Image::map_straight). +fn straight(px: vec4) -> vec3 { + return px.xyz / px.w; +} + +// u[0].x = op: +// 1 Tint f[0] = black, f[1] = white, f[2].x = amount +// 2 Bright/Contrast f[0] = (brightness, k) +// 3 Hue/Saturation u[0].y = colorize; f[0] = (hue, sat, light); f[1] = (ch, cs, cl) +// 4 Levels f[0] = (ib, iw, g, ob); f[1] = (ow); u[0].y = clip +// 5 Exposure f[0] = (e, offset, gamma) +// 6 Invert u[0].y = channel; f[0].x = blend +// 7 Fill f[0] = colour; f[1] = (opacity); u[0].y = invert +// 8 Curves f[0] = data offsets (rgb, red, green, blue), f[1].x = alpha (−1 = none) +// 9 Gradient Ramp f[0] = (sx, sy, ex, ey); f[1] = start colour; f[2] = end colour; +// f[3] = (len2, scatter, blend); u[0] = (_, radial, seed) +// 10 Fractal Noise f[0] = (contrast, brightness, scale, frac); f[1] = (ox, oy, sr, cr); +// f[2] = (evo, blend, opacity); u[0] = (_, kind, n_oct, seed); u[1].x = invert +@compute @workgroup_size(16, 16) +fn pointwise(@builtin(global_invocation_id) gid: vec3) { + let dims = out_dims(); + let p = vec2(gid.xy); + if (p.x >= dims.x || p.y >= dims.y) { + return; + } + var px = textureLoad(src, p, 0); + let a = px.w; + let op = P.u[0].x; + switch op { + case 1u: { + if (a > 0.0) { + let c = straight(px); + let l = clamp(luminance(c), 0.0, 1.0); + let blk = P.f[0].xyz; + let wht = P.f[1].xyz; + let t = blk + (wht - blk) * l; + px = vec4(mix3(c, t, P.f[2].x) * a, a); + } + } + case 2u: { + if (a > 0.0) { + let c = straight(px); + let br = P.f[0].x; + let k = P.f[0].y; + px = vec4(max((c + br - 0.5) * k + 0.5, vec3(0.0)) * a, a); + } + } + case 3u: { + if (a > 0.0) { + let c = straight(px); + let hsl = rgb_to_hsl(c); + var h = hsl.x; + var s = hsl.y; + var l = hsl.z; + if (P.u[0].y != 0u) { + let cl = P.f[1].z; + var k = l; + if (cl > 0.0) { + k = 1.0 - l; + } + h = P.f[1].x; + s = P.f[1].y; + l = clamp(l + cl * k, 0.0, 1.0); + } else { + let hue = P.f[0].x; + let sat = P.f[0].y; + let light = P.f[0].z; + var s2: f32; + if (sat >= 0.0) { + s2 = s + (1.0 - s) * sat * min(s, 1.0); + } else { + s2 = s * (1.0 + sat); + } + var l2: f32; + if (light >= 0.0) { + l2 = l + (1.0 - l) * light; + } else { + l2 = l * (1.0 + light); + } + h = fract_euclid(h + hue); + s = clamp(s2, 0.0, 1.0); + l = l2; + } + px = vec4(hsl_to_rgb(h, s, l) * a, a); + } + } + case 4u: { + if (a > 0.0) { + let c = straight(px); + let ib = P.f[0].x; + let iw = P.f[0].y; + let g = P.f[0].z; + let ob = P.f[0].w; + let ow = P.f[1].x; + var o = vec3(0.0); + for (var i = 0; i < 3; i++) { + var t = (c[i] - ib) / max(iw - ib, 1e-6); + if (P.u[0].y != 0u) { + t = clamp(t, 0.0, 1.0); + } + t = powz(max(t, 0.0), 1.0 / g); + o[i] = ob + (ow - ob) * t; + } + px = vec4(o * a, a); + } + } + case 5u: { + if (a > 0.0) { + let c = straight(px); + var o = vec3(0.0); + for (var i = 0; i < 3; i++) { + let lin = srgb_to_linear(max(c[i], 0.0)); + let v = powz(max(lin * P.f[0].x + P.f[0].y, 0.0), 1.0 / P.f[0].z); + o[i] = linear_to_srgb(v); + } + px = vec4(o * a, a); + } + } + case 6u: { + let ch = P.u[0].y; + let blend = P.f[0].x; + if (ch == 4u) { + let na = 1.0 - a; + var k = 0.0; + if (a > 0.0) { + k = na / a; + } + let inv = vec4(px.xyz * k, na); + px = px + (inv - px) * blend; + } else if (a > 0.0) { + let s = straight(px); + var o = s; + switch ch { + case 1u: { o.x = 1.0 - s.x; } + case 2u: { o.y = 1.0 - s.y; } + case 3u: { o.z = 1.0 - s.z; } + default: { o = 1.0 - s; } + } + px = vec4((s + (o - s) * blend) * a, a); + } + } + case 7u: { + let c = P.f[0]; + let invert = P.u[0].y != 0u; + let opacity = P.f[1].x; + var fa = px.w; + if (invert) { + fa = 1.0 - px.w; + } + let f = c.xyz * fa; + let rgb = px.xyz + (f - px.xyz) * opacity; + var alpha = px.w; + if (invert) { + alpha = px.w + (fa - px.w) * opacity; + } + px = vec4(rgb, alpha); + } + case 8u: { + // effects::color2::map_ca: transparent pixels only change through the alpha curve. + if (a <= 0.0 && px.x == 0.0 && px.y == 0.0 && px.z == 0.0) { + let na = curve_eval(P.f[1].x, 0.0); + if (na > 0.0) { + px = vec4(0.0, 0.0, 0.0, clamp(na, 0.0, 1.0)); + } + } else { + var c = vec3(0.0); + var aa = max(a, 0.0); + if (a > 1e-6) { + c = px.xyz / a; + aa = a; + } + var o = vec3(0.0); + for (var i = 0; i < 3; i++) { + o[i] = curve_eval(P.f[0][i + 1], curve_eval(P.f[0].x, c[i])); + } + let na = clamp(curve_eval(P.f[1].x, aa), 0.0, 1.0); + px = vec4(o * na, na); + } + } + case 9u: { + let s = P.f[0].xy; + let e = P.f[0].zw; + let d = e - s; + let v = vec2(f32(p.x) + 0.5, f32(p.y) + 0.5) - s; + let len2 = P.f[3].x; + var t: f32; + if (P.u[0].y != 0u) { + t = sqrt((v.x * v.x + v.y * v.y) / len2); + } else { + t = (v.x * d.x + v.y * d.y) / len2; + } + let scatter = P.f[3].y; + if (scatter > 0.0) { + t += (hash_noise(gid.x, gid.y, P.u[0].z) - 0.5) * scatter; + } + t = clamp(t, 0.0, 1.0); + let c0 = P.f[1].xyz; + let c1 = P.f[2].xyz; + px = put(px, c0 + (c1 - c0) * t, 1.0, P.f[3].z); + } + case 10u: { + let contrast = P.f[0].x; + let brightness = P.f[0].y; + let scale = P.f[0].z; + let frac = P.f[0].w; + let ox = P.f[1].x; + let oy = P.f[1].y; + let sr = P.f[1].z; + let cr = P.f[1].w; + let evo = P.f[2].x; + let kind = P.u[0].y; + let n_oct = P.u[0].z; + let seed = P.u[0].w; + let dx = f32(p.x) + 0.5 - ox; + let dy = f32(p.y) + 0.5 - oy; + let u = (dx * cr + dy * sr) / scale; + let v = (-dx * sr + dy * cr) / scale; + var sum = 0.0; + var amp = 1.0; + var norm = 0.0; + var f = 1.0; + for (var o = 0u; o < n_oct; o++) { + var w = 1.0; + if (o + 1u == n_oct && frac > 0.0) { + w = frac; + } + var n = value_noise(u * f, v * f, evo + f32(o) * 7.31, seed + o); + if (kind == 1u) { + n = 1.0 - abs(n * 2.0 - 1.0); + } + sum += n * amp * w; + norm += amp * w; + amp *= 0.5; + f *= 2.0; + } + var val = sum / max(norm, 1e-6); + val = clamp((val - 0.5) * contrast + 0.5 + brightness, 0.0, 1.0); + if (P.u[1].x != 0u) { + val = 1.0 - val; + } + px = put(px, vec3(val), P.f[2].z, P.f[2].y); + } + default: {} + } + textureStore(out, p, px); +} diff --git a/crates/gpu/src/tests.rs b/crates/gpu/src/tests.rs new file mode 100644 index 00000000..ca2fb1b0 --- /dev/null +++ b/crates/gpu/src/tests.rs @@ -0,0 +1,517 @@ +//! GPU vs CPU: the CPU renderer is the oracle. Every scene renders on both and must agree +//! within ≤ 1/255 at 8 bpc and ≤ 1e-3 at 16/32 bpc (per channel, premultiplied; relative for +//! over-range 32 bpc values above 1). +//! +//! Two documented exceptions, both from float rounding (the GPU computes in f32 with fast math, +//! the CPU's transforms in f64): +//! * 8/16 bpc quantise after every layer and effect, so a value within ~1e-6 of a rounding +//! boundary can land one step apart and later steps (Overlay's ×2 slope, colour-space curves) +//! can widen it: up to 0.3 % of pixels may differ by at most 4/255. +//! * Discontinuous modes (Hard Mix, Dissolve, Darker/Lighter Color) may flip pixels sitting +//! exactly on a threshold: up to 0.5 % of the frame. +//! +//! The tests skip (pass with a note) when no GPU adapter is available, as on headless CI. + +use std::sync::{Arc, OnceLock}; + +use effectcraft_color::{BlendMode, ColorSpace, Label}; +use effectcraft_keyframe::{Keyframe, ShapePath, Value}; +use effectcraft_project::build::{self, Ids}; +use effectcraft_project::{ + AlphaMode, BitDepth, Comp, Footage, FootageKind, ItemId, ItemKind, Layer, LayerSource, MaskMode, MatteKind, Project, Quality, Sampling, Solid, TrackMatte, +}; +use effectcraft_raster::Image; +use effectcraft_render::{Backend, FootageSource, RenderOpts, Renderer}; +use effectcraft_time::{FrameRate, Tick}; + +use crate::Gpu; + +fn gpu() -> Option<&'static Gpu> { + static G: OnceLock> = OnceLock::new(); + G.get_or_init(|| { + let g = Gpu::headless(); + if g.is_none() { + eprintln!("effectcraft-gpu tests: no GPU adapter, skipping GPU comparisons"); + } + g + }) + .as_ref() +} + +/// Procedural footage: smooth colour gradients, fine noise and an alpha with soft, opaque and +/// fully transparent regions (different per item). +struct Pattern; + +fn pattern(seed: u32, w: u32, h: u32) -> Image { + let mut img = Image::new(w, h); + let k = seed as f32 * 0.37; + for y in 0..h { + for x in 0..w { + let (u, v) = (x as f32 / w as f32, y as f32 / h as f32); + let n = effectcraft_raster::hash_noise(x, y, seed) * 0.25; + let r = (0.5 + 0.5 * (u * 7.0 + k).sin()) * 0.75 + n; + let g = (0.5 + 0.5 * (v * 5.0 - k).cos()) * 0.75 + n * 0.5; + let b = ((u + v) * 0.5 + k).fract() * 0.8 + 0.1; + let d = ((u - 0.5).powi(2) + (v - 0.5).powi(2)).sqrt(); + let a = ((0.48 - d) * 6.0).clamp(0.0, 1.0) * if (x / 9 + y / 7 + seed).is_multiple_of(5) { 0.6 } else { 1.0 }; + img.set(x, y, [r.min(1.0) * a, g.min(1.0) * a, b * a, a]); + } + } + img +} + +impl FootageSource for Pattern { + fn frame(&self, item: ItemId, f: &Footage, _t: Tick) -> Option> { + Some(Arc::new(pattern(item.0 as u32, f.width, f.height))) + } +} + +struct Scene { + p: Project, + cid: ItemId, + comp: Comp, +} + +impl Scene { + fn new(depth: BitDepth) -> Scene { + let mut p = Project::default(); + p.settings.bit_depth = depth; + let comp = Comp::new(97, 61, FrameRate::FPS_30, Tick::from_seconds_f64(2.0)); + let cid = p.add_item("Comp", Label::Sandstone, None, ItemKind::Comp(comp.clone().into())); + Scene { p, cid, comp } + } + + fn footage(&mut self, w: u32, h: u32) -> Layer { + let f = Footage { + path: "pattern.png".into(), + kind: FootageKind::Still, + width: w, + height: h, + pixel_aspect: 1.0, + frame_rate: FrameRate::FPS_30, + native_rate: None, + duration: Tick::ZERO, + has_video: true, + has_audio: false, + alpha: AlphaMode::Premultiplied, + premul_color: [0.0; 3], + loop_count: 1, + codec: String::new(), + missing: false, + sequence: vec![], + color_profile: None, + }; + let fid = self.p.add_item("Pattern", Label::Lavender, None, ItemKind::Footage(f)); + build::layer(&mut self.p, &self.comp, "Pattern", LayerSource::Footage { item: fid }, (w, h), None) + } + + fn solid(&mut self, color: [f32; 3], w: u32, h: u32) -> Layer { + let sid = self.p.add_item("Solid", Label::Red, None, ItemKind::Solid(Solid { color, width: w, height: h, pixel_aspect: 1.0 })); + build::layer(&mut self.p, &self.comp, "Solid", LayerSource::Solid { item: sid }, (w, h), None) + } + + /// Add on top of the stack. + fn push(&mut self, l: Layer) -> effectcraft_project::LayerId { + let id = l.id; + self.p.comp_mut(self.cid).unwrap().layers.insert(0, l); + id + } + + fn effect(&mut self, l: &mut Layer, id: &str, vals: &[(&str, Value)]) { + let spec = effectcraft_effects::find(id).unwrap(); + let mut next = self.p.next_id; + let size = effectcraft_render::source_size(&self.p, l); + let mut g = effectcraft_effects::instantiate(spec, &mut Ids(&mut next), spec.name, [size.0 as f64, size.1 as f64]); + self.p.next_id = next; + for (k, v) in vals { + g.prop_mut(k).unwrap_or_else(|| panic!("{id}: no {k}")).value = v.clone(); + } + l.props.sub_mut("effects").unwrap().children.push(g.into()); + } +} + +fn set(l: &mut Layer, path: &str, v: Value) { + l.props.prop_mut(path).unwrap_or_else(|| panic!("no {path}")).value = v; +} + +fn v3(x: f64, y: f64) -> Value { + Value::Vec3([x, y, 0.0]) +} + +/// Differences of a GPU render against the CPU render. +#[derive(Debug)] +struct Diff { + max: f32, + /// Pixels over the tolerance. + over: usize, + total: usize, + worst: (u32, u32, [f32; 4], [f32; 4]), + /// 8/16 bpc render. + quantized: bool, +} + +fn diff(cpu: &Image, gpu: &Image, tol: f32) -> Diff { + assert_eq!((cpu.width, cpu.height), (gpu.width, gpu.height), "size"); + let mut d = Diff { max: 0.0, over: 0, total: cpu.data.len(), worst: (0, 0, [0.0; 4], [0.0; 4]), quantized: false }; + for (i, (a, b)) in cpu.data.iter().zip(&gpu.data).enumerate() { + // Absolute below 1, relative above (32 bpc over-range values, e.g. Divide by ~0). + let m = (0..4).map(|c| (a[c] - b[c]).abs() / a[c].abs().max(1.0)).fold(0.0f32, f32::max); + let m = if m.is_nan() { f32::INFINITY } else { m }; + if m > tol { + d.over += 1; + } + if m > d.max { + d.max = m; + d.worst = (i as u32 % cpu.width, i as u32 / cpu.width, *a, *b); + } + } + d +} + +fn tolerance(depth: BitDepth) -> f32 { + match depth { + BitDepth::Bpc8 => 1.0 / 255.0 + 1e-6, + _ => 1e-3, + } +} + +/// Render on the CPU and the GPU; `Some(diff)` (None without an adapter). +fn compare_at(s: &Scene, opts: RenderOpts, t: Tick) -> Option { + let g = gpu()?; + let cpu = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Cpu, ..opts }).comp_frame_cpu(s.cid, t); + let mut r = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Gpu, ..opts }); + r.accel = Some(g); + let img = g.render(&r, s.cid, t).expect("the GPU renders the frame"); + let mut d = diff(&cpu, &img, tolerance(s.p.settings.bit_depth)); + d.quantized = s.p.settings.bit_depth != BitDepth::Bpc32; + Some(d) +} + +fn opts() -> RenderOpts { + RenderOpts { motion_blur: true, ..Default::default() } +} + +/// Assert agreement; `allow` = fraction of pixels allowed over the tolerance. +fn check(label: &str, d: Option, allow: f64) { + let Some(d) = d else { return }; + let frac = d.over as f64 / d.total as f64; + // Quantised depths: rounding-boundary flips (see the module docs). + let (allow, cap) = if d.quantized && allow == 0.0 { (0.003, 4.0 / 255.0 + 1e-6) } else { (allow, f32::INFINITY) }; + assert!(frac <= allow && d.max <= cap, "{label}: {} of {} pixels over tolerance (max {:.6}); worst {:?}", d.over, d.total, d.max, d.worst); + if d.over > 0 { + eprintln!("{label}: {} of {} pixels over tolerance (max {:.6})", d.over, d.total, d.max); + } +} + +fn discontinuous(m: BlendMode) -> bool { + matches!(m, BlendMode::HardMix | BlendMode::Dissolve | BlendMode::DancingDissolve | BlendMode::DarkerColor | BlendMode::LighterColor) +} + +/// Background, an Add layer (over-range values in 32 bpc) and a rotated, scaled top layer in +/// `mode` at 70 % opacity. +fn blend_scene(depth: BitDepth, mode: BlendMode) -> Scene { + let mut s = Scene::new(depth); + let bg = s.footage(97, 61); + s.push(bg); + let mut add = s.footage(80, 50); + add.blend_mode = BlendMode::Add; + set(&mut add, "transform/opacity", Value::Scalar(50.0)); + s.push(add); + let mut top = s.footage(64, 48); + top.blend_mode = mode; + set(&mut top, "transform/opacity", Value::Scalar(70.0)); + set(&mut top, "transform/rotation", Value::Scalar(17.0)); + set(&mut top, "transform/scale", Value::Vec3([115.0, 90.0, 100.0])); + set(&mut top, "transform/position", v3(52.3, 29.6)); + s.push(top); + s +} + +#[test] +fn all_blend_modes_at_every_bit_depth() { + for depth in [BitDepth::Bpc8, BitDepth::Bpc16, BitDepth::Bpc32] { + for mode in BlendMode::ALL { + let s = blend_scene(depth, mode); + let allow = if discontinuous(mode) { 0.005 } else { 0.0 }; + check(&format!("{mode:?} {depth:?}"), compare_at(&s, opts(), Tick::ZERO), allow); + } + } +} + +#[test] +fn transforms_sampling_and_resolution() { + for depth in [BitDepth::Bpc8, BitDepth::Bpc32] { + for (i, (sampling, quality, scale, rot, sc)) in [ + (Sampling::Bilinear, Quality::Best, 1.0, 0.0, 100.0), + (Sampling::Bilinear, Quality::Best, 1.0, 33.0, 140.0), + (Sampling::Bicubic, Quality::Best, 1.0, -21.0, 160.0), + (Sampling::Bilinear, Quality::Best, 1.0, 12.0, 30.0), + (Sampling::Bicubic, Quality::Best, 0.5, 45.0, 80.0), + (Sampling::Bilinear, Quality::Draft, 1.0, 0.0, 200.0), + (Sampling::Bilinear, Quality::Best, 0.25, 5.0, 100.0), + ] + .into_iter() + .enumerate() + { + let mut s = Scene::new(depth); + let bg = s.solid([0.1, 0.2, 0.3], 97, 61); + s.push(bg); + let mut l = s.footage(90, 70); + l.switches.sampling = sampling; + l.switches.quality = quality; + set(&mut l, "transform/rotation", Value::Scalar(rot)); + set(&mut l, "transform/scale", Value::Vec3([sc, sc * 0.9, 100.0])); + set(&mut l, "transform/position", v3(47.25, 31.5)); + s.push(l); + check(&format!("transform #{i} {depth:?}"), compare_at(&s, RenderOpts { scale, ..opts() }, Tick::ZERO), 0.0); + } + } +} + +#[test] +fn masks_mattes_and_preserve_transparency() { + for depth in [BitDepth::Bpc8, BitDepth::Bpc32] { + for kind in [MatteKind::Alpha, MatteKind::AlphaInverted, MatteKind::Luma, MatteKind::LumaInverted] { + let mut s = Scene::new(depth); + let bg = s.footage(97, 61); + s.push(bg); + let mut matte = s.footage(60, 40); + matte.switches.video = false; + set(&mut matte, "transform/rotation", Value::Scalar(-12.0)); + let mut next = s.p.next_id; + let m = build::mask(&mut Ids(&mut next), "Mask 1", ShapePath::ellipse([30.0, 20.0], 50.0, 34.0), MaskMode::Add, [255, 255, 0]); + s.p.next_id = next; + matte.props.sub_mut("masks").unwrap().children.push(m.into()); + let mid = s.push(matte); + let mut fill = s.footage(97, 61); + fill.track_matte = Some(TrackMatte { layer: mid, kind }); + fill.blend_mode = BlendMode::Screen; + s.push(fill); + let mut pt = s.solid([0.9, 0.4, 0.1], 50, 30); + pt.preserve_transparency = true; + set(&mut pt, "transform/opacity", Value::Scalar(60.0)); + s.push(pt); + check(&format!("matte {kind:?} {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.0); + } + } +} + +#[test] +fn motion_blur_accumulates_like_the_cpu() { + for depth in [BitDepth::Bpc8, BitDepth::Bpc32] { + let mut s = Scene::new(depth); + s.p.comp_mut(s.cid).unwrap().enable_motion_blur = true; + s.comp.enable_motion_blur = true; + let bg = s.solid([0.05, 0.05, 0.1], 97, 61); + s.push(bg); + let mut l = s.footage(30, 30); + l.switches.motion_blur = true; + l.blend_mode = BlendMode::Add; + let pos = l.props.prop_mut("transform/position").unwrap(); + pos.keys = vec![Keyframe::new(Tick::ZERO, v3(10.0, 20.0)), Keyframe::new(Tick::from_seconds_f64(0.5), v3(90.0, 45.0))]; + let rot = l.props.prop_mut("transform/rotation").unwrap(); + rot.keys = vec![Keyframe::new(Tick::ZERO, Value::Scalar(0.0)), Keyframe::new(Tick::from_seconds_f64(0.5), Value::Scalar(180.0))]; + s.push(l); + check(&format!("motion blur {depth:?}"), compare_at(&s, opts(), Tick::from_seconds_f64(0.2)), 0.0); + } +} + +#[test] +fn colour_management_and_linear_blending() { + for (ws, linearize, blend_linear) in + [(None, false, true), (Some(ColorSpace::Rec2020), false, false), (Some(ColorSpace::Srgb), true, false), (Some(ColorSpace::Rec709), false, true)] + { + for depth in [BitDepth::Bpc8, BitDepth::Bpc32] { + let mut s = blend_scene(depth, BlendMode::Overlay); + s.p.settings.working_space = ws; + s.p.settings.linearize = linearize; + s.p.settings.blend_linear = blend_linear; + check(&format!("colour {ws:?} {linearize} {blend_linear} {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.0); + } + } +} + +#[test] +fn layer_styles_collapsed_precomps_and_cpu_fallbacks() { + for depth in [BitDepth::Bpc8, BitDepth::Bpc32] { + // Layer styles (drop shadow pass + stroke body). + let mut s = Scene::new(depth); + let bg = s.footage(97, 61); + s.push(bg); + let mut l = s.solid([0.8, 0.3, 0.2], 40, 30); + let mut next = s.p.next_id; + effectcraft_project::styles::add_style(&mut l, &mut Ids(&mut next), "dropShadow", &s.comp.global_light, true).unwrap(); + effectcraft_project::styles::add_style(&mut l, &mut Ids(&mut next), "stroke", &s.comp.global_light, true).unwrap(); + s.p.next_id = next; + l.blend_mode = BlendMode::Multiply; + set(&mut l, "transform/rotation", Value::Scalar(10.0)); + s.push(l); + check(&format!("styles {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.0); + + // A collapsed precomp (nested blend modes act on the parent) and an adjustment layer + // (CPU fallback between GPU steps). + let mut s = Scene::new(depth); + let inner = Comp::new(60, 40, FrameRate::FPS_30, Tick::from_seconds_f64(2.0)); + let inner_id = s.p.add_item("Inner", Label::Sandstone, None, ItemKind::Comp(inner.clone().into())); + let (a, b) = { + let mut a = s.footage(60, 40); + a.blend_mode = BlendMode::Normal; + let mut b = s.solid([0.2, 0.7, 0.4], 30, 20); + b.blend_mode = BlendMode::Difference; + (a, b) + }; + s.p.comp_mut(inner_id).unwrap().layers = vec![b, a]; + let bg = s.footage(97, 61); + s.push(bg); + let mut pre = build::layer(&mut s.p, &s.comp, "Inner", LayerSource::Comp { item: inner_id }, (60, 40), None); + pre.switches.collapse = true; + set(&mut pre, "transform/rotation", Value::Scalar(25.0)); + set(&mut pre, "transform/opacity", Value::Scalar(80.0)); + s.push(pre); + let mut adj = s.solid([1.0, 1.0, 1.0], 50, 61); + adj.switches.adjustment = true; + s.effect(&mut adj, "ec.channel.invert", &[]); + s.push(adj); + check(&format!("collapse + adjustment {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.0); + } +} + +/// The effect alone on a buffer: GPU kernels vs the CPU effect (32-bit float), including the +/// buffer geometry (padding, offset). +fn effect_direct(id: &str, vals: &[(&str, Value)], adjustment: bool) { + let Some(g) = gpu() else { return }; + let spec = effectcraft_effects::find(id).unwrap(); + let size = [70.0, 44.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 (k, v) in vals { + params.values.insert(k.to_string(), v.clone()); + } + for scale in [1.0, 0.5] { + let img = pattern(7, (size[0] * scale) as u32, (size[1] * scale) as u32); + let buf = effectcraft_effects::Buf { img, offset: [0.0, 0.0], scale }; + let ctx = || effectcraft_effects::EffectCtx { params: ¶ms, time: 0.25, layer_size: size, seed: 11, adjustment, env: Default::default() }; + let cpu = (spec.render)(&ctx(), buf.clone()); + let out = effectcraft_render::Accelerator::effects(g, &[effectcraft_render::FxStep { spec, ctx: ctx() }], &buf, None).expect("the GPU runs the effect"); + assert_eq!((out.offset, out.scale), (cpu.offset, cpu.scale), "{id}: geometry"); + let d = diff(&cpu.img, &out.img, 1e-3); + assert!(d.over == 0, "{id} {vals:?} adj {adjustment} scale {scale}: {} of {} pixels over 1e-3 (max {}); worst {:?}", d.over, d.total, d.max, d.worst); + } +} + +/// One footage layer with `fx` applied, composited over a background, at 8 and 32 bpc. +fn effect_case(id: &str, vals: &[(&str, Value)]) { + effect_direct(id, vals, false); + effect_direct(id, vals, true); + for depth in [BitDepth::Bpc8, BitDepth::Bpc32] { + let mut s = Scene::new(depth); + let bg = s.solid([0.15, 0.1, 0.2], 97, 61); + s.push(bg); + let mut l = s.footage(70, 44); + s.effect(&mut l, id, vals); + set(&mut l, "transform/rotation", Value::Scalar(8.0)); + s.push(l); + check(&format!("{id} {vals:?} {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.0); + check(&format!("{id} {vals:?} {depth:?} half"), compare_at(&s, RenderOpts { scale: 0.5, ..opts() }, Tick::ZERO), 0.0); + } +} + +fn n(v: f64) -> Value { + Value::Scalar(v) +} + +fn c(r: f64, g: f64, b: f64) -> Value { + Value::Color([r, g, b, 1.0]) +} + +#[test] +fn gpu_effects_match_cpu() { + effect_case("ec.blur.gaussian", &[("blurriness", n(9.0))]); + effect_case("ec.blur.gaussian", &[("blurriness", n(4.0)), ("dimensions", Value::Enum(1)), ("repeatEdge", Value::Bool(true))]); + effect_case("ec.blur.fastbox", &[("radius", n(5.0)), ("iterations", n(2.0))]); + effect_case("ec.blur.fastbox", &[("radius", n(3.0)), ("dimensions", Value::Enum(2)), ("repeatEdge", Value::Bool(true))]); + effect_case("ec.blur.directional", &[("direction", n(30.0)), ("length", n(12.0))]); + effect_case("ec.stylize.glow", &[("threshold", n(40.0)), ("radius", n(14.0)), ("intensity", n(1.5))]); + effect_case("ec.stylize.glow", &[("colors", Value::Enum(1)), ("operation", Value::Enum(1)), ("colorA", c(1.0, 0.8, 0.1))]); + effect_case("ec.color.levels", &[("inBlack", n(0.1)), ("inWhite", n(0.8)), ("gamma", n(1.6)), ("outBlack", n(0.05))]); + effect_case("ec.color.levels", &[("inWhite", n(0.5)), ("noClip", Value::Bool(true))]); + effect_case( + "ec.color.curves", + &[("rgb", Value::Str("0,0 0.3,0.45 1,1".into())), ("red", Value::Str("0,0.1 1,0.9".into())), ("alpha", Value::Str("0,0 0.5,0.7 1,1".into()))], + ); + effect_case("ec.color.huesaturation", &[("hue", n(70.0)), ("saturation", n(40.0)), ("lightness", n(-20.0))]); + effect_case("ec.color.huesaturation", &[("colorize", Value::Bool(true)), ("colorizeHue", n(200.0)), ("colorizeLightness", n(15.0))]); + effect_case("ec.color.tint", &[("black", c(0.1, 0.0, 0.3)), ("white", c(1.0, 0.9, 0.5)), ("amount", n(80.0))]); + effect_case("ec.generate.fill", &[("color", c(0.2, 0.6, 1.0)), ("opacity", n(70.0))]); + effect_case("ec.generate.fill", &[("invert", Value::Bool(true))]); + effect_case("ec.generate.gradientramp", &[("start", Value::Vec2([5.0, 3.0])), ("end", Value::Vec2([60.0, 40.0])), ("blend", n(25.0))]); + effect_case("ec.generate.gradientramp", &[("shape", Value::Enum(1)), ("scatter", n(40.0))]); + effect_case("ec.noise.fractal", &[]); + effect_case( + "ec.noise.fractal", + &[("fractalType", Value::Enum(1)), ("complexity", n(4.5)), ("rotation", n(30.0)), ("evolution", n(90.0)), ("scale", n(40.0)), ("blend", n(30.0))], + ); + effect_case("ec.perspective.dropshadow", &[("distance", n(6.0)), ("softness", n(8.0)), ("opacity", n(80.0))]); + effect_case("ec.perspective.dropshadow", &[("shadowOnly", Value::Bool(true)), ("color", c(0.0, 0.2, 0.6))]); + effect_case("ec.color.brightnesscontrast", &[("brightness", n(20.0)), ("contrast", n(35.0))]); + effect_case("ec.color.brightnesscontrast", &[("brightness", n(-10.0)), ("contrast", n(-40.0))]); + effect_case("ec.color.exposure", &[("exposure", n(1.3)), ("offset", n(0.02)), ("gamma", n(1.2))]); + effect_case("ec.channel.invert", &[]); + effect_case("ec.channel.invert", &[("channel", Value::Enum(4)), ("blend", n(30.0))]); + effect_case("ec.channel.invert", &[("channel", Value::Enum(2))]); + effect_case("ec.distort.transform", &[("rotation", n(20.0)), ("scaleHeight", n(80.0)), ("skew", n(10.0)), ("opacity", n(70.0))]); +} + +/// A chain of GPU effects (one upload, quantised between steps) and a non-GPU effect in the +/// middle (split into two GPU chains around the CPU effect). +#[test] +fn effect_chains_and_mixed_stacks() { + for depth in [BitDepth::Bpc8, BitDepth::Bpc32] { + let mut s = Scene::new(depth); + let mut l = s.footage(80, 50); + s.effect(&mut l, "ec.color.exposure", &[("exposure", n(0.7))]); + s.effect(&mut l, "ec.blur.gaussian", &[("blurriness", n(5.0))]); + s.effect(&mut l, "ec.stylize.posterize", &[]); + s.effect(&mut l, "ec.stylize.glow", &[]); + s.effect(&mut l, "ec.channel.invert", &[("blend", n(50.0))]); + s.push(l); + check(&format!("chain {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.0); + } +} + +#[test] +fn registry_badges_match_the_gpu_implementation() { + for s in effectcraft_effects::registry() { + let implemented = crate::effects::supports(s.id); + assert_eq!(s.gpu, implemented, "{}: registry gpu = {}, GPU implementation = {implemented}", s.id, s.gpu); + } +} + +#[test] +fn backend_selection_and_display_frames() { + let Some(g) = gpu() else { return }; + let s = blend_scene(BitDepth::Bpc8, BlendMode::Screen); + // Auto follows the project's Mercury GPU / Software Only setting; Cpu never uses the GPU. + let mut r = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Auto, ..opts() }); + r.accel = Some(g); + assert!(r.active_accel().is_some()); + let mut p2 = s.p.clone(); + p2.settings.gpu_acceleration = false; + let mut r2 = Renderer::new(&p2, &Pattern, RenderOpts { backend: Backend::Auto, ..opts() }); + r2.accel = Some(g); + assert!(r2.active_accel().is_none()); + let mut r3 = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Cpu, ..opts() }); + r3.accel = Some(g); + assert!(r3.active_accel().is_none()); + // The display texture holds what the CPU viewer path shows (8-bit premultiplied). + let f = g.render_display(&r, s.cid, Tick::ZERO).expect("display frame"); + let bytes = g.read_display(&f).expect("readback"); + let cpu = Renderer::new(&s.p, &Pattern, opts()).comp_frame_cpu(s.cid, Tick::ZERO); + for (p, q) in cpu.data.iter().zip(bytes.chunks_exact(4)) { + let a = p[3].clamp(0.0, 1.0); + let want = [p[0].clamp(0.0, a), p[1].clamp(0.0, a), p[2].clamp(0.0, a), a].map(|v| (v * 255.0 + 0.5) as u8); + for k in 0..4 { + assert!((want[k] as i32 - q[k] as i32).abs() <= 1, "{want:?} vs {q:?}"); + } + } +} diff --git a/crates/gpu/src/walk.rs b/crates/gpu/src/walk.rs new file mode 100644 index 00000000..d43c7b1d --- /dev/null +++ b/crates/gpu/src/walk.rs @@ -0,0 +1,188 @@ +//! The per-frame compositing walk on the GPU: the same steps as `Renderer::draw_comp` (layer +//! order, 3D runs, collapsed precomps, track mattes, Preserve Transparency, layer styles, motion +//! blur, bit-depth quantisation, colour conversions), with layer pixels taken from the CPU +//! renderer and its layer cache. Steps without a GPU implementation (3D runs, adjustment and +//! wireframe layers) read the canvas back, run on the CPU and upload the result. + +use std::sync::Arc; + +use effectcraft_color::BlendMode; +use effectcraft_effects::Buf; +use effectcraft_project::{ItemId, Layer, Quality}; +use effectcraft_raster::Image; +use effectcraft_render::{EvalCtx, Renderer, styles}; +use effectcraft_time::Tick; + +use crate::context::{Enc, GpuImage}; +use crate::ops; + +/// Render a comp frame into a GPU image (the canvas in the output space, as +/// `Renderer::comp_frame_cpu` returns it). `None` when something cannot run here (the caller +/// renders on the CPU instead). +pub(crate) fn render<'g>(e: &mut Enc<'g>, r: &Renderer, comp_id: ItemId, t: Tick) -> Option { + let ctx = r.eval_ctx(comp_id, t)?; + let s = r.opts.scale; + let w = ((ctx.comp.width as f64 * s).round() as u32).max(1); + let h = ((ctx.comp.height as f64 * s).round() as u32).max(1); + if r.depth() > 16 || !e.g.fits(w, h) { + return None; + } + let mut canvas = e.image(w, h); + draw_comp(e, r, &ctx, &mut canvas)?; + let pipe = r.pipe(); + if let Some(c) = pipe.from_blend() { + canvas = ops::convert(e, &canvas, &c); + } + if let Some(l) = pipe.levels { + canvas = ops::quantize(e, &canvas, l); + } + if r.depth() == 0 + && let Some(c) = pipe.output() + { + canvas = ops::convert(e, &canvas, &c); + if let Some(l) = pipe.levels { + canvas = ops::quantize(e, &canvas, l); + } + } + Some(canvas) +} + +/// Run a CPU drawing step on the canvas (read back, draw, upload). +fn on_cpu(e: &mut Enc, canvas: &mut GpuImage, f: impl FnOnce(&mut Image)) -> Option<()> { + let mut img = e.download(canvas)?; + f(&mut img); + *canvas = e.g.upload_image(&img)?; + Some(()) +} + +fn draw_comp<'a>(e: &mut Enc, r: &Renderer<'a>, ctx: &EvalCtx<'a>, canvas: &mut GpuImage) -> Option<()> { + let visible = r.visible_layers(ctx); + let levels = r.pipe().levels; + let mut i = 0; + while i < visible.len() { + if visible[i].is_3d() { + let mut j = i; + while j < visible.len() && visible[j].is_3d() { + j += 1; + } + on_cpu(e, canvas, |img| r.draw_3d_run(ctx, &visible[i..j], img))?; + i = j; + } else { + match r.collapsed(ctx, visible[i]) { + Some(item) => draw_collapsed(e, r, ctx, visible[i], item, canvas)?, + None => draw_layer(e, r, ctx, visible[i], canvas)?, + } + i += 1; + } + // 8/16 bpc: the comp is an integer buffer after every layer. + if let Some(l) = levels { + *canvas = ops::quantize(e, canvas, l); + } + } + Some(()) +} + +fn draw_collapsed<'a>(e: &mut Enc, r: &Renderer<'a>, ctx: &EvalCtx<'a>, layer: &Layer, item: ItemId, canvas: &mut GpuImage) -> Option<()> { + let opacity = r.layer_opacity(ctx, layer); + if opacity <= 0.0 { + return Some(()); + } + let matte = r.track_matte(ctx, layer).is_some(); + if layer.blend_mode == BlendMode::Normal && !matte && !layer.preserve_transparency { + if let Some((sub, nctx)) = r.collapse_into(ctx, layer, item, opacity) { + draw_comp(e, &sub, &nctx, canvas)?; + } + return Some(()); + } + let Some((sub, nctx)) = r.collapse_into(ctx, layer, item, 1.0) else { return Some(()) }; + let mut iso = e.image(canvas.width, canvas.height); + draw_comp(e, &sub, &nctx, &mut iso)?; + composite_iso(e, r, ctx, layer, iso, canvas, opacity) +} + +fn draw_layer(e: &mut Enc, r: &Renderer, ctx: &EvalCtx, layer: &Layer, canvas: &mut GpuImage) -> Option<()> { + if layer.switches.adjustment || layer.switches.quality == Quality::Wireframe { + // Adjustment layers run their effects on the comp below (CPU, with GPU effects where + // available); wireframes draw outlines. + return on_cpu(e, canvas, |img| r.draw_layer_cpu(ctx, layer, img)); + } + let opacity = r.layer_opacity(ctx, layer); + if opacity <= 0.0 { + return Some(()); + } + let Some(st) = r.styled_layer(ctx, layer) else { return Some(()) }; + if st.plain { + return composite_layer(e, r, ctx, layer, &st.body, canvas, opacity); + } + // Layer styles: exterior passes with their own modes, then the body; the layer's opacity + // fades the whole stack; knockout and channel switches as on the CPU. + let bl = styles::blending(ctx, layer); + let (w, h) = (canvas.width, canvas.height); + if r.track_matte(ctx, layer).is_some() || layer.preserve_transparency { + let mut iso = e.image(w, h); + for (b, m) in &st.passes { + iso = place(e, r, ctx, layer, b, &iso, *m, 1.0)?; + } + iso = place(e, r, ctx, layer, &st.body, &iso, BlendMode::Normal, 1.0)?; + return composite_iso(e, r, ctx, layer, iso, canvas, opacity); + } + let mut tmp = canvas.clone(); + if bl.knockout > 0 { + let k = place(e, r, ctx, layer, &st.content, &e.image(w, h), BlendMode::Normal, 1.0)?; + tmp = ops::knockout(e, &tmp, &k); + } + for (b, m) in &st.passes { + tmp = place(e, r, ctx, layer, b, &tmp, *m, 1.0)?; + } + tmp = place(e, r, ctx, layer, &st.body, &tmp, layer.blend_mode, 1.0)?; + *canvas = ops::channel_mix(e, canvas, &tmp, bl.channels, opacity); + Some(()) +} + +fn composite_layer(e: &mut Enc, r: &Renderer, ctx: &EvalCtx, layer: &Layer, buf: &Arc, canvas: &mut GpuImage, opacity: f32) -> Option<()> { + let matte = r.track_matte(ctx, layer).is_some(); + if !matte && !layer.preserve_transparency { + *canvas = place(e, r, ctx, layer, buf, canvas, layer.blend_mode, opacity)?; + return Some(()); + } + // Render in isolation, then matte / preserve transparency, then blend. + let iso = place(e, r, ctx, layer, buf, &e.image(canvas.width, canvas.height), BlendMode::Normal, 1.0)?; + composite_iso(e, r, ctx, layer, iso, canvas, opacity) +} + +/// Track matte / Preserve Transparency on an isolated layer render, then blend. +fn composite_iso(e: &mut Enc, r: &Renderer, ctx: &EvalCtx, layer: &Layer, mut iso: GpuImage, canvas: &mut GpuImage, opacity: f32) -> Option<()> { + if let Some((m, kind)) = r.track_matte(ctx, layer) { + let mut mimg = e.image(canvas.width, canvas.height); + if m.is_active_at(ctx.time) + && let Some(mb) = r.layer_buf(ctx, m) + { + let mo = ctx.opacity(m) as f32; + mimg = place(e, r, ctx, m, &mb, &mimg, BlendMode::Normal, mo)?; + } + iso = ops::matte(e, &iso, &mimg, kind); + } + if layer.preserve_transparency { + iso = ops::preserve(e, &iso, canvas); + } + *canvas = ops::blend_full(e, canvas, &iso, layer.blend_mode, opacity, layer.id.0 as u32); + Some(()) +} + +/// Draw a processed layer buffer onto `target` (transform, motion blur sub-samples). +fn place(e: &mut Enc, r: &Renderer, ctx: &EvalCtx, layer: &Layer, buf: &Arc, target: &GpuImage, mode: BlendMode, opacity: f32) -> Option { + if buf.img.is_empty() || opacity <= 0.0 { + return Some(target.clone()); + } + let src = e.g.upload_buf(buf)?; + let pl = r.placement(ctx, layer, buf); + if let [m] = pl.matrices.as_slice() { + return Some(ops::warp(e, target, &src, m, pl.sampling, mode, opacity, pl.seed, None)); + } + let mut acc = e.image(target.width, target.height); + let k = 1.0 / pl.matrices.len() as f32; + for m in &pl.matrices { + acc = ops::warp(e, &acc, &src, m, pl.sampling, BlendMode::Normal, 1.0, 0, Some(k)); + } + Some(ops::blend_full(e, target, &acc, mode, opacity, pl.seed)) +} diff --git a/crates/host/src/lib.rs b/crates/host/src/lib.rs index df563fc7..16199dfb 100644 --- a/crates/host/src/lib.rs +++ b/crates/host/src/lib.rs @@ -32,6 +32,7 @@ impl Exporter for FileExporter { project: job.project, footage: job.footage, expr: job.expr, + accel: job.accel, comp: job.item.comp, settings: &job.item.settings, output: &job.item.output, diff --git a/crates/project/src/lib.rs b/crates/project/src/lib.rs index 88fc6d02..d6bbf5e4 100644 --- a/crates/project/src/lib.rs +++ b/crates/project/src/lib.rs @@ -113,6 +113,10 @@ pub struct ProjectSettings { /// Frame numbering starts at 0 (or 1). pub frame_start: i64, pub audio_sample_rate: u32, + /// Video Rendering and Effects ▸ Use: Mercury GPU Acceleration (`true`, the default; used + /// when a GPU adapter exists) or Mercury Software Only (`false`, the CPU compositor). + #[serde(default = "yes")] + pub gpu_acceleration: bool, } impl Default for ProjectSettings { @@ -125,6 +129,7 @@ impl Default for ProjectSettings { time_display: TimeDisplayStyle::Timecode, frame_start: 0, audio_sample_rate: 48_000, + gpu_acceleration: true, } } } diff --git a/crates/render/src/lib.rs b/crates/render/src/lib.rs index 05c05efa..31df5966 100644 --- a/crates/render/src/lib.rs +++ b/crates/render/src/lib.rs @@ -128,6 +128,53 @@ impl FootageSource for NoFootage { } } +/// Which compositor renders a frame (Project Settings ▸ Video Rendering and Effects). +#[derive(Clone, Copy, Debug, Default, PartialEq, Eq, Hash, serde::Serialize, serde::Deserialize)] +#[serde(rename_all = "lowercase")] +pub enum Backend { + /// The CPU reference compositor (Mercury Software Only). + #[default] + Cpu, + /// The GPU compositor whenever an [`Accelerator`] is attached, whatever the project says. + Gpu, + /// The GPU compositor when an [`Accelerator`] is attached and the project's renderer is + /// Mercury GPU Acceleration (the viewer's choice); the CPU otherwise. + Auto, +} + +impl Backend { + pub fn parse(s: &str) -> Option { + match s.to_ascii_lowercase().as_str() { + "cpu" | "software" | "mercury software only" => Some(Backend::Cpu), + "gpu" | "mercury gpu acceleration" => Some(Backend::Gpu), + "auto" => Some(Backend::Auto), + _ => None, + } + } +} + +/// One effect of a GPU effect chain (see [`Accelerator::effects`]). +pub struct FxStep<'x> { + pub spec: &'static effectcraft_effects::EffectSpec, + pub ctx: EffectCtx<'x>, +} + +/// A GPU (or other hardware) backend for the compositor. The CPU [`Renderer`] stays the +/// reference: an accelerator renders what it supports and returns `None` for the rest, which +/// then runs on the CPU. Implemented by `effectcraft-gpu`. +pub trait Accelerator: Send + Sync { + /// Adapter / backend name for status readouts ("Apple M2 (Metal)"). + fn name(&self) -> String; + /// Render a top-level comp frame (what [`Renderer::comp_frame`] returns). `None` = not + /// handled (the CPU renders it). + fn comp_frame(&self, r: &Renderer, comp: ItemId, t: Tick) -> Option; + /// The accelerator implements effect `id` with the CPU effect's semantics. + fn supports_effect(&self, id: &str) -> bool; + /// Run a chain of supported effects on `buf`, clamping and quantising to `levels` after + /// every effect when set (8/16 bpc). `None` = not handled. + fn effects(&self, chain: &[FxStep], buf: &Buf, levels: Option) -> Option; +} + #[derive(Clone, Copy, Debug)] pub struct RenderOpts { /// Output scale relative to comp pixels (1 = Full, 0.5 = Half…). @@ -139,11 +186,13 @@ pub struct RenderOpts { pub draft: bool, /// 3D view camera override (viewer Front/Top/Custom views); None = the comp's active camera. pub view: Option, + /// CPU or GPU compositor (the GPU needs [`Renderer::accel`]). + pub backend: Backend, } impl Default for RenderOpts { fn default() -> Self { - RenderOpts { scale: 1.0, motion_blur: true, guides: false, draft: false, view: None } + RenderOpts { scale: 1.0, motion_blur: true, guides: false, draft: false, view: None, backend: Backend::Cpu } } } @@ -156,6 +205,8 @@ pub struct Renderer<'a> { pub cache: Option<&'a LayerCache>, /// Optional per-layer timing sink (`frame --bench` / perf readouts). pub profile: Option<&'a std::sync::Mutex>>, + /// GPU backend used when [`RenderOpts::backend`] asks for it. + pub accel: Option<&'a dyn Accelerator>, /// Current nesting depth (precomp recursion guard). depth: usize, /// The project's colour pipeline (bit depth, working space, linear blending). @@ -204,6 +255,7 @@ impl<'a> Renderer<'a> { opts, cache: None, profile: None, + accel: None, depth: 0, pipe: color::Pipe::of(&project.settings), outer: None, @@ -220,7 +272,7 @@ impl<'a> Renderer<'a> { /// The nested comp of a precomp layer whose transformations collapse into this comp: the /// Collapse Transformations switch is on and the layer has no masks, effects or layer /// styles (those force a flattened render of the precomp, as in After Effects). - pub(crate) fn collapsed(&self, ctx: &EvalCtx, layer: &Layer) -> Option { + pub fn collapsed(&self, ctx: &EvalCtx, layer: &Layer) -> Option { if !layer.switches.collapse || layer.switches.adjustment || self.depth > 12 { return None; } @@ -244,7 +296,7 @@ impl<'a> Renderer<'a> { /// their transforms are concatenated with the precomp layer's (one resample, no /// rasterisation at the precomp's bounds), their opacity multiplied by `opacity`, and their /// 3D layers use this comp's camera and lights. - pub(crate) fn collapse_into(&self, ctx: &EvalCtx<'a>, layer: &Layer, item: ItemId, opacity: f32) -> Option<(Renderer<'a>, EvalCtx<'a>)> { + pub fn collapse_into(&self, ctx: &EvalCtx<'a>, layer: &Layer, item: ItemId, opacity: f32) -> Option<(Renderer<'a>, EvalCtx<'a>)> { let nc = self.project.comp(item)?; let nctx = EvalCtx { comp_id: item, comp: nc, time: ctx.source_time(layer), ..*ctx }; let (l2c, _) = ctx.layer_to_comp(layer); @@ -336,8 +388,30 @@ impl<'a> Renderer<'a> { EvalCtx { project: self.project, comp_id, comp, time: t, expr: self.expr } } + /// The accelerator to use for this render, if any (see [`Backend`]). + pub fn active_accel(&self) -> Option<&'a dyn Accelerator> { + let a = self.accel?; + match self.opts.backend { + Backend::Cpu => None, + Backend::Gpu => Some(a), + Backend::Auto => self.project.settings.gpu_acceleration.then_some(a), + } + } + /// Render a composition at comp time `t` (transparent background, comp size × scale). + /// Top-level frames go to the GPU when [`Self::active_accel`] is set and handles them. pub fn comp_frame(&self, comp_id: ItemId, t: Tick) -> Image { + if self.depth == 0 + && let Some(a) = self.active_accel() + && let Some(img) = a.comp_frame(self, comp_id, t) + { + return img; + } + self.comp_frame_cpu(comp_id, t) + } + + /// [`Self::comp_frame`] on the CPU compositor only. + pub fn comp_frame_cpu(&self, comp_id: ItemId, t: Tick) -> Image { let Some(comp) = self.project.comp(comp_id) else { return Image::new(1, 1) }; let s = self.opts.scale; let w = ((comp.width as f64 * s).round() as u32).max(1); @@ -364,15 +438,8 @@ impl<'a> Renderer<'a> { /// Draw a comp's layers (bottom to top) into `canvas` (blending space). fn draw_comp(&self, ctx: &EvalCtx<'a>, canvas: &mut Image) { - let (comp, t) = (ctx.comp, ctx.time); - let any_solo = comp.layers.iter().any(|l| l.switches.solo && l.source.is_av() && l.is_active_at(t)); // Bottom-to-top, with runs of consecutive 3D layers depth-sorted (farthest first). - let visible: Vec<&Layer> = comp - .layers - .iter() - .rev() - .filter(|l| l.is_active_at(t) && l.has_video() && (!any_solo || l.switches.solo) && (self.opts.guides || !l.switches.guide)) - .collect(); + let visible = self.visible_layers(ctx); let mut i = 0; while i < visible.len() { if visible[i].is_3d() { @@ -425,18 +492,49 @@ impl<'a> Renderer<'a> { let host = FxHost { r: self, ctx, layer, index: Default::default() }; let env = EffectEnv { masks: &mask_shapes, host: Some(&host), comp_time: ctx.time.seconds(), frame_rate: ctx.comp.frame_rate.as_f64(), effect_index: 0 }; - for (i, g) in fx.groups().enumerate() { - if i >= limit { - break; - } - if !g.enabled { - continue; - } - let GroupKind::Effect { effect } = &g.kind else { continue }; - let Some(spec) = effectcraft_effects::find(effect) else { continue }; - if effectcraft_effects::audio_fx::is_audio_effect(spec.id) { - continue; + // Video effects in stack order (index, group, spec); disabled and audio effects skipped. + let stack: Vec<(usize, &effectcraft_project::PropGroup, &'static effectcraft_effects::EffectSpec)> = fx + .groups() + .enumerate() + .take_while(|(i, _)| *i < limit) + .filter(|(_, g)| g.enabled) + .filter_map(|(i, g)| match &g.kind { + GroupKind::Effect { effect } => effectcraft_effects::find(effect).map(|s| (i, g, s)), + _ => None, + }) + .filter(|(_, _, s)| !effectcraft_effects::audio_fx::is_audio_effect(s.id)) + .collect(); + let accel = self.active_accel(); + let mut k = 0; + while k < stack.len() { + let (i, g, spec) = stack[k]; + // A run of GPU effects goes to the accelerator in one upload / readback. + if let Some(a) = accel + && spec.gpu + && a.supports_effect(spec.id) + { + let run: Vec<_> = stack[k..].iter().take_while(|(_, _, s)| s.gpu && a.supports_effect(s.id)).collect(); + let params: Vec = run.iter().map(|(_, g, _)| self.effect_params(ctx, layer, g)).collect(); + let steps: Vec = run + .iter() + .zip(¶ms) + .map(|((i, g, s), params)| FxStep { + spec: s, + ctx: EffectCtx { params, time: lt.seconds(), layer_size, seed: g.uid as u32, adjustment, env: EffectEnv { effect_index: *i, ..env } }, + }) + .collect(); + let t0 = web_time::Instant::now(); + if let Some(out) = a.effects(&steps, &buf, self.pipe.levels) { + buf = out; + if let Some(v) = timing.as_deref_mut() { + let ms = t0.elapsed().as_secs_f64() * 1e3 / run.len() as f64; + v.extend(run.iter().map(|(_, _, s)| (format!("{} (gpu)", s.id), ms))); + } + k += run.len(); + continue; + } } + k += 1; let params = self.effect_params(ctx, layer, g); host.index.store(i, std::sync::atomic::Ordering::Relaxed); let env = EffectEnv { effect_index: i, ..env }; @@ -574,7 +672,7 @@ impl<'a> Renderer<'a> { } /// [`Self::layer_buf`] in the blending space (linear with Blend Colors Using 1.0 Gamma). - pub(crate) fn blend_layer_buf(&self, ctx: &EvalCtx, layer: &Layer) -> Option> { + pub fn blend_layer_buf(&self, ctx: &EvalCtx, layer: &Layer) -> Option> { let (b, k) = self.layer_buf_keyed(ctx, layer)?; Some(self.to_blend(b, k)) } @@ -733,32 +831,45 @@ impl<'a> Renderer<'a> { Mat3::scale(vec2(s, s)) * l2c * Mat3::scale(vec2(1.0 / buf.scale, 1.0 / buf.scale)) * Mat3::translate(vec2(-buf.offset[0], -buf.offset[1])) } - /// Draw a processed layer into `target` (transform, motion blur) with `mode`/`opacity`. - fn place(&self, ctx: &EvalCtx, layer: &Layer, buf: &Buf, target: &mut Image, mode: BlendMode, opacity: f32) { + /// How a processed layer buffer lands in the output: one buffer → output matrix, or one per + /// motion-blur sub-sample (averaged), with the layer's sampling filter and dissolve seed. + pub fn placement(&self, ctx: &EvalCtx, layer: &Layer, buf: &Buf) -> Placement { let samples = if self.opts.motion_blur && ctx.comp.enable_motion_blur && layer.switches.motion_blur { if self.opts.draft { 4 } else { ctx.comp.motion_blur_samples.clamp(2, 64) as usize } } else { 1 }; - let opts = WarpOpts { sampling: self.sampling(layer), opacity, mode, seed: layer.id.0 as u32, clip: None }; + let (sampling, seed) = (self.sampling(layer), layer.id.0 as u32); if samples <= 1 { - let m = self.buf_matrix(ctx, layer, buf); - composite_warp(target, &buf.img, &m, &opts); - return; + return Placement { matrices: vec![self.buf_matrix(ctx, layer, buf)], sampling, seed }; } let fd = ctx.comp.frame_duration().seconds(); let angle = ctx.comp.shutter_angle / 360.0; let phase = ctx.comp.shutter_phase / 360.0; + let matrices = (0..samples) + .map(|i| { + let f = phase + angle * i as f64 / (samples - 1) as f64; + self.buf_matrix(&ctx.at(ctx.time + Tick::from_seconds_f64(f * fd)), layer, buf) + }) + .collect(); + Placement { matrices, sampling, seed } + } + + /// Draw a processed layer into `target` (transform, motion blur) with `mode`/`opacity`. + fn place(&self, ctx: &EvalCtx, layer: &Layer, buf: &Buf, target: &mut Image, mode: BlendMode, opacity: f32) { + let pl = self.placement(ctx, layer, buf); + let opts = WarpOpts { sampling: pl.sampling, opacity, mode, seed: pl.seed, clip: None }; + if let [m] = pl.matrices.as_slice() { + composite_warp(target, &buf.img, m, &opts); + return; + } // Sub-samples sum straight into one buffer (each warp is row-parallel). let mut acc = Image::new(target.width, target.height); - let k = 1.0 / samples as f32; - for i in 0..samples { - let f = phase + angle * i as f64 / (samples - 1) as f64; - let sub = ctx.at(ctx.time + Tick::from_seconds_f64(f * fd)); - let m = self.buf_matrix(&sub, layer, buf); - effectcraft_raster::accumulate_warp(&mut acc, &buf.img, &m, opts.sampling, k); + let k = 1.0 / pl.matrices.len() as f32; + for m in &pl.matrices { + effectcraft_raster::accumulate_warp(&mut acc, &buf.img, m, opts.sampling, k); } - target.blend_from(&acc, mode, opacity, layer.id.0 as u32); + target.blend_from(&acc, mode, opacity, pl.seed); } fn draw_layer(&self, ctx: &EvalCtx, layer: &Layer, canvas: &mut Image) { @@ -906,18 +1017,99 @@ impl<'a> Renderer<'a> { } } -/// A layer ready to composite (see [`Renderer::styled`]). -struct StyledLayer { - passes: Vec<(Arc, BlendMode)>, - body: Arc, +/// A layer ready to composite (see [`Renderer::styled_layer`]). +pub struct StyledLayer { + /// Exterior style passes (Drop Shadow, Outer Glow) with their blend modes. + pub passes: Vec<(Arc, BlendMode)>, + pub body: Arc, /// Pre-style pixels (knockout shape). - content: Arc, + pub content: Arc, /// Styled cache key (the content key when plain). key: Option, /// Content cache key. ckey: Option, /// No layer styles: composite `body` the usual way. - plain: bool, + pub plain: bool, +} + +/// Where a layer buffer lands (see [`Renderer::placement`]). +#[derive(Clone, Debug)] +pub struct Placement { + /// Buffer pixel → output pixel; several = motion-blur sub-samples (equal weights). + pub matrices: Vec, + pub sampling: effectcraft_raster::Sampling, + /// Dissolve noise seed. + pub seed: u32, +} + +/// Hooks for accelerated compositors (`effectcraft-gpu`): the pieces of the CPU walk they reuse +/// or fall back to. The CPU compositor itself is [`Renderer::comp_frame_cpu`]. +impl<'a> Renderer<'a> { + /// The project's colour pipeline. + pub fn pipe(&self) -> color::Pipe { + self.pipe + } + + /// Nesting depth (0 = the top-level comp). + pub fn depth(&self) -> usize { + self.depth + } + + /// Evaluation context of comp `comp_id` at comp time `t`. + pub fn eval_ctx(&self, comp_id: ItemId, t: Tick) -> Option> { + Some(self.ctx(comp_id, self.project.comp(comp_id)?, t)) + } + + /// The layers drawn at the context time, bottom to top (solo, guides and visibility + /// applied). + pub fn visible_layers(&self, ctx: &EvalCtx<'a>) -> Vec<&'a Layer> { + let (comp, t) = (ctx.comp, ctx.time); + let any_solo = comp.layers.iter().any(|l| l.switches.solo && l.source.is_av() && l.is_active_at(t)); + comp.layers + .iter() + .rev() + .filter(|l| l.is_active_at(t) && l.has_video() && (!any_solo || l.switches.solo) && (self.opts.guides || !l.switches.guide)) + .collect() + } + + /// Draw a run of consecutive 3D layers into `canvas` (CPU). + pub fn draw_3d_run(&self, ctx: &EvalCtx<'a>, run: &[&Layer], canvas: &mut Image) { + three_d::compose::draw_run(self, ctx, run, canvas); + } + + /// Draw one 2D layer into `canvas` on the CPU (any kind: adjustment, wireframe, styled…). + pub fn draw_layer_cpu(&self, ctx: &EvalCtx, layer: &Layer, canvas: &mut Image) { + self.draw_layer(ctx, layer, canvas); + } + + /// Draw a collapsed precomp layer into `canvas` on the CPU. + pub fn draw_collapsed_cpu(&self, ctx: &EvalCtx<'a>, layer: &Layer, item: ItemId, canvas: &mut Image) { + self.draw_collapsed(ctx, layer, item, canvas); + } + + /// The layer's finished buffers (styles included) in the blending space. + pub fn styled_layer(&self, ctx: &EvalCtx, layer: &Layer) -> Option { + let mut timing = self.profile.map(|_| LayerTiming { depth: self.depth, layer: layer.name.clone(), ..Default::default() }); + let t0 = web_time::Instant::now(); + let st = self.styled(ctx, layer, timing.as_mut()).map(|st| self.styled_to_blend(st)); + if let (Some(prof), Some(mut timing)) = (self.profile, timing) { + timing.process_ms = t0.elapsed().as_secs_f64() * 1e3; + if let Ok(mut v) = prof.lock() { + v.push(timing); + } + } + st + } + + /// Opacity the layer composites with (its Opacity × collapsed precomps' opacity). + pub fn layer_opacity(&self, ctx: &EvalCtx, layer: &Layer) -> f32 { + ctx.opacity(layer) as f32 * self.opacity_mul + } + + /// The layer's track matte layer and kind, if it has a valid one. + pub fn track_matte<'c>(&self, ctx: &EvalCtx<'c>, layer: &Layer) -> Option<(&'c Layer, MatteKind)> { + layer.track_matte.and_then(|tm| ctx.comp.layer(tm.layer).filter(|m| m.id != layer.id).map(|m| (m, tm.kind))) + } } /// The layer's frame blending mode, if the comp's Enable Frame Blending switch is on. diff --git a/crates/ui-egui/src/lib.rs b/crates/ui-egui/src/lib.rs index 7627c096..9a032442 100644 --- a/crates/ui-egui/src/lib.rs +++ b/crates/ui-egui/src/lib.rs @@ -319,7 +319,14 @@ impl EffectcraftApp { let Some(c) = self.session.project.comp(comp) else { return }; let key = self.frame_key(comp, frame, scale); let t = c.frame_rate.tick_of(frame); - let opts = RenderOpts { scale, motion_blur: true, guides: true, draft: self.ui.viewer.fast_preview, view: self.session.view_camera(comp) }; + let opts = RenderOpts { + scale, + motion_blur: true, + guides: true, + draft: self.ui.viewer.fast_preview, + view: self.session.view_camera(comp), + backend: effectcraft_engine::render::Backend::Auto, + }; if urgent { self.frames.request_urgent(&self.render_source(), key, comp, t, opts); } else { diff --git a/docs/architecture.md b/docs/architecture.md index dc95f35b..c3d478e3 100644 --- a/docs/architecture.md +++ b/docs/architecture.md @@ -27,6 +27,7 @@ cpal or muda. Everything in L0 to L4, the egui UI and the web app also build for | L3 | `media` | Footage decoding (FilmCraft's pure-Rust codecs), image sequences, frame cache | | L3 | `expr` | The expression engine (JavaScript via boa) with the After Effects object model | | L3 | `export` | Render queue encoding: H.264, ProRes, PNG/JPEG/TIFF/EXR sequences, GIF, audio | +| L3 | `gpu` | The GPU compositor and GPU effects on wgpu compute shaders (Metal, Vulkan, Direct3D 12, WebGPU), checked against the CPU renderer | | L3 | `lottie` | Lottie JSON / dotLottie import and export (layers, precomps, eased and spatial keyframes, shapes, masks, mattes) with a warnings list for what Lottie cannot express | | L4 | `engine` | `Session`: project, undo history, editor state, the command registry and menus | | L4 | `host` | A fully wired `Session` (media, expressions, exporter) for the frontends | @@ -35,7 +36,7 @@ cpal or muda. Everything in L0 to L4, the egui UI and the web app also build for | L6 | apps `effectcraft`, `effectcraft-cli`, `effectcraft-web` | Desktop app; command-line tool (render, exec, get/set, MCP); the browser app (wasm32, [web.md](web.md)) | Allowed same-layer edges: `path → keyframe, raster`, `text → path`, `effects → project, text, path`, -`media / expr / export → render`, `export → media`, `lottie → format`, `host → engine`. +`media / expr / export / gpu → render`, `export → media`, `lottie → format`, `host → engine`. ## 2. Time @@ -103,6 +104,27 @@ groups into `Params` keys (`effects::flatten_params`). Commands: `paint.*`, `pup Half, Third and Quarter resolution render proportionally fewer pixels end to end. +**GPU compositor** (`crates/gpu`; Project Settings ▸ Video Rendering and Effects ▸ Mercury GPU +Acceleration, the default, or Mercury Software Only; `render.backend`). The CPU renderer is the +reference and keeps rendering layer content (sources, masks, CPU effects, layer styles) into the +layer cache. The render crate defines an `Accelerator` trait; `effectcraft_gpu::Gpu` implements it +on wgpu compute shaders. `RenderOpts::backend` picks `Cpu`, `Gpu` or `Auto` (GPU when an +accelerator is attached and the project's renderer is the GPU). A GPU frame walks the comp like +`draw_comp`: cached layer buffers are uploaded once per buffer, then transformed with the CPU's +sampling (nearest, bilinear, Catmull-Rom bicubic, the same minification pre-filter), motion-blur +sub-samples are accumulated, and track mattes, Preserve Transparency, layer style passes, +knockout, all 38 blend modes (a WGSL port of `color::blend`), the 8/16 bpc clamp-and-quantise steps +and colour-space conversions run on the GPU. 3D runs, adjustment layers and wireframes run on the +CPU between GPU steps (read back, draw, upload). GPU effects (`effects::GPU_EFFECTS`: Gaussian, +Fast Box and Directional Blur, Glow, Levels, Curves, Hue/Saturation, Tint, Fill, Gradient Ramp, +Fractal Noise, Drop Shadow, Brightness & Contrast, Exposure, Invert, Transform) repeat the CPU +effect's steps (padding, box radii, parameters) as kernels; consecutive GPU effects run as one +chain with one upload and one readback. Tests render scenes on both paths and compare them +(≤ 1/255 at 8 bpc, ≤ 1e-3 at 32 bpc); they skip without an adapter. The desktop viewer builds the +`Gpu` on egui-wgpu's device and shows frames from GPU textures without reading them back +(`ui-egui::frames`); headless renders, the CLI (unless `--gpu`) and CI use the CPU. On the web +(WebGPU) the GPU composites viewer frames; steps that need a readback fall back to the CPU. + **Colour and bit depth** (`crates/render/src/color.rs`, `crates/color/src/space.rs`). Pixels are `f32`, but 8 and 16 bpc projects clamp and quantise each layer after its source and masks and after every effect, and the comp after every layer (16 bpc uses 0..32768), so over-range values diff --git a/docs/parity.md b/docs/parity.md index df6b963a..31e73c87 100644 --- a/docs/parity.md +++ b/docs/parity.md @@ -69,11 +69,11 @@ rendering, on-canvas text editing) are where the conservative figure is most lik | Text | 63% | 5.8 | on-canvas editing with a cursor, per-character styles, paragraph spacing and indents | | Web | 60% | 2.5 | threads, browser storage, audio | | 3D | 58% | 20.5 | Advanced 3D (models, PBR, image-based light), 3D camera tracker, multi-view layouts | -| Effects | 56% | 14.6 | Effect Controls widgets (angle dial, point crosshair, eyedropper, curves), GPU effects, 41 missing effects | +| Effects | 58% | 13.6 | Effect Controls widgets (angle dial, point crosshair, eyedropper, curves), GPU versions of more effects, 41 missing effects | | Interface | 55% | 10.4 | drag-to-dock and floating panels, viewer rulers/snapping/channels/snapshots, preferences, native macOS menus | | Project | 53% | 8.1 | 8/16/32-bit pipeline and colour management, auto-save, proxies, folder moves | | Masks & roto | 44% | 13.0 | mask tracking, Roto Brush, free transform and RotoBezier | -| Preview | 43% | 8.1 | GPU compositor, disk cache, region of interest, snapshots and exposure | +| Preview | 55% | 6.1 | GPU 3D and adjustment layers, disk cache, region of interest, snapshots and exposure | | Tracking | 0% | 13.0 | point tracker, Warp Stabilizer, mask and face tracking | | Paint | 0% | 4.0 | Brush, Clone Stamp, Eraser | @@ -90,7 +90,8 @@ Mocha and Cineware are third-party and not counted. 1. On-canvas text editing and per-character styles. 2. Effect Controls widgets: angle dial, point crosshair, eyedropper, curves and levels editors. 3. Viewer basics: snapping, rulers, channel view, snapshots, exposure, a drawable region of interest. -4. A GPU (wgpu) compositor, then GPU effects. +4. ~~A GPU (wgpu) compositor, then GPU effects~~ (done: 2D compositing and 16 GPU effects; 3D runs + and adjustment layers still composite on the CPU). 5. Pen tool for shape paths, shape vertex editing, the Layer viewer, free transform and RotoBezier. 6. Motion-path handles in the viewer; graph editor transform box and snapping. 7. Point tracking and stabilization (in progress).