M12.7: adjustment layers on the GPU (effect stacks resident on the device, CPU fallback per non-GPU effect)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

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<u8>) -> wgpu::Buffer {
         while bytes.len() < 16 || bytes.len() % 4 != 0 {

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<GBuf>,
+    levels: Option<f32>,
+}
+
+impl<'e, 'g> GpuFx<'e, 'g> {
+    pub fn new(e: &'e mut Enc<'g>, img: GpuImage, offset: [f64; 2], scale: f64, levels: Option<f32>) -> 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<GBuf> {
     match id {
         "ec.blur.gaussian" => gaussian(e, ctx, 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;

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();

crates/gpu/src/shaders/kernels.wgsl

@@ -828,3 +828,19 @@ fn pointwise(@builtin(global_invocation_id) gid: vec3<u32>) {
     }
     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<u32>) {
+    let dims = out_dims();
+    let p = vec2<i32>(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);
+}

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));
+}

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<Buf>, canvas: &mut GpuImage, opacity: f32) -> Option<()> {
     let matte = r.track_matte(ctx, layer).is_some();
     if !matte && !layer.preserve_transparency {

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<f32>,
+    buf: Option<Buf>,
+}
+
+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: &params, 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<Buf> {
+        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);