diff --git a/crates/gpu/src/context.rs b/crates/gpu/src/context.rs index d549149d..66bee99d 100644 --- a/crates/gpu/src/context.rs +++ b/crates/gpu/src/context.rs @@ -26,6 +26,7 @@ const ENTRIES: &[&str] = &[ "shadow_make", "shadow_combine", "pointwise", + "adjust_mix", ]; /// Entry points that also bind group 1 (four read-only storage buffers; see @@ -582,6 +583,27 @@ impl<'g> Enc<'g> { self.bytes(bytes) } + /// Copy an image into a storage buffer (RGBA f32 rows) for kernels that read a third image + /// through `data`: (buffer, row length in pixels). + pub fn to_buffer(&mut self, img: &GpuImage) -> (wgpu::Buffer, u32) { + let row = (img.width * 16).div_ceil(wgpu::COPY_BYTES_PER_ROW_ALIGNMENT) * wgpu::COPY_BYTES_PER_ROW_ALIGNMENT; + let buf = self.g.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("image rows"), + size: row as u64 * img.height as u64, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + self.encoder().copy_texture_to_buffer( + wgpu::TexelCopyTextureInfo { texture: &img.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(row), rows_per_image: Some(img.height) }, + }, + wgpu::Extent3d { width: img.width, height: img.height, depth_or_array_layers: 1 }, + ); + (buf, row / 16) + } + /// A read-only storage buffer holding `bytes` (padded to 16 bytes). pub fn bytes(&self, mut bytes: Vec) -> wgpu::Buffer { while bytes.len() < 16 || bytes.len() % 4 != 0 { diff --git a/crates/gpu/src/effects.rs b/crates/gpu/src/effects.rs index fdc317b3..36b692f1 100644 --- a/crates/gpu/src/effects.rs +++ b/crates/gpu/src/effects.rs @@ -62,6 +62,48 @@ pub(crate) fn run_chain(e: &mut Enc, chain: &[FxStep], buf: &Buf, levels: Option Some(Buf { img, offset: b.offset, scale: b.scale }) } +/// An effect stack running on a GPU-resident image (adjustment layers): GPU chains stay on +/// the device; CPU effects read the image back, run, and upload the result. +pub(crate) struct GpuFx<'e, 'g> { + e: &'e mut Enc<'g>, + b: Option, + levels: Option, +} + +impl<'e, 'g> GpuFx<'e, 'g> { + pub fn new(e: &'e mut Enc<'g>, img: GpuImage, offset: [f64; 2], scale: f64, levels: Option) -> Self { + GpuFx { e, b: Some(GBuf { img, offset, scale }), levels } + } + + /// The result: (image, offset, scale); `None` when a step failed. + pub fn finish(self) -> Option<(GpuImage, [f64; 2], f64)> { + self.b.map(|b| (b.img, b.offset, b.scale)) + } +} + +impl effectcraft_render::FxTarget for GpuFx<'_, '_> { + fn gpu(&mut self, steps: &[FxStep]) -> bool { + let Some(b) = &self.b else { return false }; + let mut cur = GBuf { img: b.img.clone(), offset: b.offset, scale: b.scale }; + for step in steps { + let Some(next) = apply(self.e, step.spec.id, &step.ctx, cur) else { return false }; + cur = next; + if let Some(l) = self.levels { + cur.img = ops::quantize(self.e, &cur.img, l); + } + } + self.b = Some(cur); + true + } + + fn cpu(&mut self, f: &mut dyn FnMut(Buf) -> Buf) { + let Some(b) = self.b.take() else { return }; + let Some(img) = self.e.download(&b.img) else { return }; + let out = f(Buf { img, offset: b.offset, scale: b.scale }); + self.b = self.e.g.upload_image(&out.img).map(|img| GBuf { img, offset: out.offset, scale: out.scale }); + } +} + fn apply(e: &mut Enc, id: &str, ctx: &EffectCtx, b: GBuf) -> Option { match id { "ec.blur.gaussian" => gaussian(e, ctx, b), diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs index 7a2634b3..a7c9cdc4 100644 --- a/crates/gpu/src/lib.rs +++ b/crates/gpu/src/lib.rs @@ -143,4 +143,6 @@ mod tests; #[cfg(test)] mod tests_3d; #[cfg(test)] +mod tests_adjust; +#[cfg(test)] mod tests_adv3d; diff --git a/crates/gpu/src/ops.rs b/crates/gpu/src/ops.rs index 775e86c1..cd743fbf 100644 --- a/crates/gpu/src/ops.rs +++ b/crates/gpu/src/ops.rs @@ -134,6 +134,17 @@ pub fn channel_mix(e: &mut Enc, canvas: &GpuImage, tmp: &GpuImage, channels: [bo out } +/// Adjustment layer finish: move `canvas` toward `adjusted` by `matte`'s alpha × `opacity`. +pub fn adjust_mix(e: &mut Enc, canvas: &GpuImage, adjusted: &GpuImage, matte: &GpuImage, opacity: f32) -> GpuImage { + let (rows, stride) = e.to_buffer(matte); + let mut p = Params::default(); + p.u[0][0] = stride; + p.f[0][0] = opacity; + let out = e.image(canvas.width, canvas.height); + e.pixels("adjust_mix", &p, canvas, Some(adjusted), &out, Some(&rows)); + 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(); diff --git a/crates/gpu/src/shaders/kernels.wgsl b/crates/gpu/src/shaders/kernels.wgsl index c94076a0..443abcf9 100644 --- a/crates/gpu/src/shaders/kernels.wgsl +++ b/crates/gpu/src/shaders/kernels.wgsl @@ -828,3 +828,19 @@ fn pointwise(@builtin(global_invocation_id) gid: vec3) { } textureStore(out, p, px); } + +// Adjustment layer finish (Renderer::draw_adjustment): canvas `src` moves toward the adjusted +// comp `aux` by the footprint's alpha × opacity. The footprint is `data` (RGBA f32 rows of +// u[0].x pixels); f[0].x = opacity. +@compute @workgroup_size(16, 16) +fn adjust_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 a = textureLoad(aux, p, 0); + let k = data[(gid.y * P.u[0].x + gid.x) * 4u + 3u] * P.f[0].x; + textureStore(out, p, c + (a - c) * k); +} diff --git a/crates/gpu/src/tests_adjust.rs b/crates/gpu/src/tests_adjust.rs new file mode 100644 index 00000000..ae7f88da --- /dev/null +++ b/crates/gpu/src/tests_adjust.rs @@ -0,0 +1,105 @@ +//! Adjustment layers on the GPU vs the CPU: GPU-only stacks (resident on the device), mixed +//! stacks (CPU effects in between read back and upload), masks, opacity, colour management. + +use effectcraft_keyframe::{ShapePath, Value}; +use effectcraft_project::build::{self, Ids}; +use effectcraft_project::{BitDepth, ColorSpace, MaskMode}; +use effectcraft_render::{Backend, FxStep, FxTarget, RenderOpts, Renderer}; +use effectcraft_time::Tick; + +use crate::tests::{Pattern, Scene, c, check, compare_at, gpu, n, opts, set}; + +fn adjustment_scene(depth: BitDepth, stack: &[(&str, Vec<(&str, Value)>)], masked: bool) -> Scene { + let mut s = Scene::new(depth); + let bg = s.footage(97, 61); + s.push(bg); + let mut l = s.footage(60, 40); + set(&mut l, "transform/rotation", Value::Scalar(15.0)); + s.push(l); + let mut adj = s.solid([1.0, 1.0, 1.0], 80, 50); + adj.switches.adjustment = true; + set(&mut adj, "transform/opacity", Value::Scalar(75.0)); + set(&mut adj, "transform/rotation", Value::Scalar(-6.0)); + for (id, vals) in stack { + s.effect(&mut adj, id, vals); + } + if masked { + let mut next = s.p.next_id; + let m = build::mask(&mut Ids(&mut next), "Mask 1", ShapePath::ellipse([40.0, 25.0], 60.0, 40.0), MaskMode::Add, [255, 255, 0]); + s.p.next_id = next; + adj.props.sub_mut("masks").unwrap().children.push(m.into()); + } + s.push(adj); + let top = s.solid([0.2, 0.6, 0.9], 20, 20); + s.push(top); + s +} + +fn gpu_stack() -> Vec<(&'static str, Vec<(&'static str, Value)>)> { + vec![ + ("ec.blur.gaussian", vec![("blurriness", n(6.0))]), + ("ec.color.levels", vec![("inBlack", n(0.1)), ("gamma", n(1.4))]), + ("ec.color.tint", vec![("black", c(0.1, 0.0, 0.2)), ("amount", n(60.0))]), + ] +} + +#[test] +fn gpu_adjustment_chains_match_the_cpu() { + for depth in [BitDepth::Bpc8, BitDepth::Bpc16, BitDepth::Bpc32] { + for masked in [false, true] { + let s = adjustment_scene(depth, &gpu_stack(), masked); + check(&format!("gpu adjustment {depth:?} masked {masked}"), compare_at(&s, opts(), Tick::ZERO), 0.0); + check(&format!("gpu adjustment {depth:?} masked {masked} half"), compare_at(&s, RenderOpts { scale: 0.5, ..opts() }, Tick::ZERO), 0.0); + } + } +} + +#[test] +fn mixed_adjustment_chains_fall_back_per_effect() { + let mut stack = gpu_stack(); + stack.insert(1, ("ec.stylize.posterize", vec![])); + stack.push(("ec.stylize.glow", vec![("threshold", n(30.0))])); + for depth in [BitDepth::Bpc8, BitDepth::Bpc32] { + let s = adjustment_scene(depth, &stack, true); + check(&format!("mixed adjustment {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.0); + } + // Colour-managed: effects run in the working space, compositing in linear light. + for depth in [BitDepth::Bpc8, BitDepth::Bpc32] { + let mut s = adjustment_scene(depth, &stack, false); + s.p.settings.working_space = Some(ColorSpace::Rec709); + s.p.settings.blend_linear = true; + check(&format!("colour-managed adjustment {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.0); + } +} + +/// Counts where each run of the stack goes. +#[derive(Default)] +struct Probe { + gpu: usize, + cpu: usize, +} + +impl FxTarget for Probe { + fn gpu(&mut self, steps: &[FxStep]) -> bool { + self.gpu += steps.len(); + true + } + fn cpu(&mut self, _f: &mut dyn FnMut(effectcraft_effects::Buf) -> effectcraft_effects::Buf) { + self.cpu += 1; + } +} + +#[test] +fn adjustment_stacks_split_into_gpu_runs_and_cpu_steps() { + let Some(g) = gpu() else { return }; + let mut stack = gpu_stack(); + stack.insert(1, ("ec.stylize.posterize", vec![])); + let s = adjustment_scene(BitDepth::Bpc32, &stack, false); + let mut r = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Gpu, ..opts() }); + r.accel = Some(g); + let ctx = r.eval_ctx(s.cid, Tick::ZERO).unwrap(); + let adj = ctx.comp.layers.iter().find(|l| l.switches.adjustment).unwrap(); + let mut p = Probe::default(); + r.run_effects_on(&ctx, adj, true, &mut p); + assert_eq!((p.gpu, p.cpu), (3, 1)); +} diff --git a/crates/gpu/src/walk.rs b/crates/gpu/src/walk.rs index a0b51f1f..cfec880d 100644 --- a/crates/gpu/src/walk.rs +++ b/crates/gpu/src/walk.rs @@ -124,9 +124,14 @@ fn draw_collapsed<'a>(e: &mut Enc, r: &Renderer<'a>, ctx: &EvalCtx<'a>, layer: & } 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. + if layer.switches.adjustment { + if draw_adjustment(e, r, ctx, layer, canvas).is_some() { + return Some(()); + } + return on_cpu(e, canvas, |img| r.draw_layer_cpu(ctx, layer, img)); + } + if layer.switches.quality == Quality::Wireframe { + // Wireframes draw outlines on the CPU. return on_cpu(e, canvas, |img| r.draw_layer_cpu(ctx, layer, img)); } let opacity = r.layer_opacity(ctx, layer); @@ -162,6 +167,33 @@ fn draw_layer(e: &mut Enc, r: &Renderer, ctx: &EvalCtx, layer: &Layer, canvas: & Some(()) } +/// Adjustment layer (`Renderer::draw_adjustment`): the effect stack runs on the comp below, +/// resident on the GPU (non-GPU effects read back, run on the CPU and upload), and the result +/// replaces the canvas inside the layer's footprint × opacity. `None` = draw it on the CPU. +fn draw_adjustment(e: &mut Enc, r: &Renderer, ctx: &EvalCtx, layer: &Layer, canvas: &mut GpuImage) -> Option<()> { + let opacity = r.layer_opacity(ctx, layer); + let Some(foot) = r.adjustment_footprint(ctx, layer) else { return Some(()) }; + let (w, h) = (canvas.width, canvas.height); + let matte = place(e, r, ctx, layer, &Arc::new(foot), &e.image(w, h), BlendMode::Normal, 1.0)?; + let pipe = r.pipe(); + // Effects run in the working space, not the linear blending space. + let below = match pipe.from_blend() { + Some(c) => ops::convert(e, canvas, &c), + None => canvas.clone(), + }; + let mut fx = crate::effects::GpuFx::new(e, below, [0.0; 2], r.opts.scale, pipe.levels); + r.run_effects_on(ctx, layer, true, &mut fx); + let (mut adjusted, _, _) = fx.finish()?; + if adjusted.width != w || adjusted.height != h { + return Some(()); + } + if let Some(c) = pipe.to_blend() { + adjusted = ops::convert(e, &adjusted, &c); + } + *canvas = ops::adjust_mix(e, canvas, &adjusted, &matte, 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 { diff --git a/crates/render/src/lib.rs b/crates/render/src/lib.rs index 5ae4ebc5..ffb55936 100644 --- a/crates/render/src/lib.rs +++ b/crates/render/src/lib.rs @@ -242,6 +242,41 @@ pub trait Accelerator: Send + Sync { } } +/// Where an effect stack runs (see [`Renderer::run_effects_on`]): a CPU buffer, or an image +/// resident on an accelerator. +pub trait FxTarget { + /// Run a chain of accelerator-supported effects (quantising after each to the bit depth). + /// `false` = not handled: the effects then run one by one through [`FxTarget::cpu`]. + fn gpu(&mut self, steps: &[FxStep]) -> bool; + /// Run one CPU effect: `f` maps the buffer. + fn cpu(&mut self, f: &mut dyn FnMut(Buf) -> Buf); +} + +/// The CPU effect target (GPU chains go through [`Accelerator::effects`]: upload, readback). +struct CpuFx<'a> { + accel: Option<&'a dyn Accelerator>, + levels: Option, + buf: Option, +} + +impl FxTarget for CpuFx<'_> { + fn gpu(&mut self, steps: &[FxStep]) -> bool { + let (Some(a), Some(b)) = (self.accel, &self.buf) else { return false }; + match a.effects(steps, b, self.levels) { + Some(out) => { + self.buf = Some(out); + true + } + None => false, + } + } + fn cpu(&mut self, f: &mut dyn FnMut(Buf) -> Buf) { + if let Some(b) = self.buf.take() { + self.buf = Some(f(b)); + } + } +} + #[derive(Clone, Copy, Debug)] pub struct RenderOpts { /// Output scale relative to comp pixels (1 = Full, 0.5 = Half…). @@ -592,19 +627,28 @@ impl<'a> Renderer<'a> { } /// Run the first `limit` effects of the layer's stack (by position; disabled ones count). - fn apply_effects_timed( + fn apply_effects_timed(&self, ctx: &EvalCtx, layer: &Layer, buf: Buf, adjustment: bool, limit: usize, timing: Option<&mut Vec<(String, f64)>>) -> Buf { + let mut t = CpuFx { accel: self.active_accel(), levels: self.pipe.levels, buf: Some(buf) }; + self.run_effect_stack(ctx, layer, adjustment, limit, timing, &mut t); + t.buf.unwrap_or_else(|| Buf { img: Image::new(1, 1), offset: [0.0; 2], scale: self.opts.scale }) + } + + /// Run the layer's video effect stack on `target`: runs of GPU-capable effects go to + /// [`FxTarget::gpu`] (when an accelerator is active and supports them), the rest to + /// [`FxTarget::cpu`] one by one (quantised to the bit depth after each). + fn run_effect_stack( &self, ctx: &EvalCtx, layer: &Layer, - mut buf: Buf, adjustment: bool, limit: usize, mut timing: Option<&mut Vec<(String, f64)>>, - ) -> Buf { + target: &mut dyn FxTarget, + ) { if !layer.switches.effects { - return buf; + return; } - let Some(fx) = layer.effects() else { return buf }; + let Some(fx) = layer.effects() else { return }; let lt = layer.layer_time(ctx.time); let size = source_size(self.project, layer); let layer_size = if size.0 == 0 { [ctx.comp.width as f64, ctx.comp.height as f64] } else { [size.0 as f64, size.1 as f64] }; @@ -635,7 +679,7 @@ impl<'a> Renderer<'a> { 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. + // A run of GPU effects goes to the accelerator in one go. if let Some(a) = accel && spec.gpu && a.supports_effect(spec.id) @@ -651,8 +695,7 @@ impl<'a> Renderer<'a> { }) .collect(); let t0 = web_time::Instant::now(); - if let Some(out) = a.effects(&steps, &buf, self.pipe.levels) { - buf = out; + if target.gpu(&steps) { 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))); @@ -667,14 +710,16 @@ impl<'a> Renderer<'a> { let env = EffectEnv { effect_index: i, ..env }; let ectx = EffectCtx { params: ¶ms, time: lt.seconds(), layer_size, seed: g.uid as u32, adjustment, env }; let t0 = web_time::Instant::now(); - buf = effectcraft_effects::apply(spec, &ectx, buf); - // 8/16 bpc effects write integer pixels. - self.pipe.quantize(&mut buf.img); + target.cpu(&mut |buf| { + let mut buf = effectcraft_effects::apply(spec, &ectx, buf); + // 8/16 bpc effects write integer pixels. + self.pipe.quantize(&mut buf.img); + buf + }); if let Some(v) = timing.as_deref_mut() { v.push((spec.id.to_string(), t0.elapsed().as_secs_f64() * 1e3)); } } - buf } /// Source pixels of a layer in layer space (before masks/effects). @@ -1389,6 +1434,22 @@ impl<'a> Renderer<'a> { three_d::compose::gpu_run(self, ctx, run, out) } + /// Run the layer's effect stack on `target` (an accelerator's resident image): GPU-capable + /// runs through [`FxTarget::gpu`], the others through [`FxTarget::cpu`], with the CPU's + /// parameters and environment. `adjustment`: the stack runs on the comp below (adjustment + /// layers). + pub fn run_effects_on(&self, ctx: &EvalCtx, layer: &Layer, adjustment: bool, target: &mut dyn FxTarget) { + self.run_effect_stack(ctx, layer, adjustment, usize::MAX, None, target); + } + + /// An adjustment layer's footprint (its source with masks applied, layer space): where + /// its effects show. `None` = the layer changes nothing. + pub fn adjustment_footprint(&self, ctx: &EvalCtx, layer: &Layer) -> Option { + let mut foot = self.source(ctx, layer)?; + let _ = masks::apply(ctx, layer, &mut foot); + Some(foot) + } + /// 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, false);