M5.2c: mask evaluation on the GPU

Linear and radial gradients, luminance and colour ranges and brush strokes are
evaluated by a WGSL kernel per component, combined (add/subtract/intersect,
component and mask invert, amount) into one device buffer of alpha planes that
the per-pixel stage reads directly. Brushes use the CPU's dabs
(`masks::brush_dabs`, shared with the rasterizer) and the closed form of the
flow accumulation (1 − Π(1 − a·flow), capped at density). Sky / Subject /
Background and other heuristic shapes, and brushes of more than 4096 dabs, are
evaluated on the CPU inside the same render and uploaded.

New equivalence case with CPU-evaluated shapes, subtract/intersect and amount;
all cases stay at max 1 LSB.

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

crates/gpu/src/ctx.rs

@@ -39,6 +39,11 @@ const MODULES: &[Module] = &[
         bindings: &[("a", false, "f32"), ("b", false, "f32"), ("c", false, "f32"), ("dst", true, "f32")],
         entries: &["log_lum_k", "dark_k", "guided_pre", "guided_ab", "guided_apply", "wb_k", "nr_lum", "chroma_k", "nr_col", "subsample"],
     },
+    Module {
+        src: include_str!("wgsl/mask.wgsl"),
+        bindings: &[("img", false, "f32"), ("log_l", false, "f32"), ("aux", false, "f32"), ("c", true, "f32"), ("alpha", true, "f32")],
+        entries: &["shape", "combine", "finalize"],
+    },
 ];
 
 struct Kernel {

crates/gpu/src/render.rs

@@ -435,21 +435,121 @@ fn prepare(cx: &mut Cx<'_>, lin: &Buf, plan: &Plan<'_>, req: &RenderRequest, pla
     Prep { log_l, base, clarity, texture, dark, air }
 }
 
-/// Evaluate the masks: their alpha planes (concatenated) and their adjustment terms.
+/// Most brush dabs the mask kernel evaluates per pixel (more: the CPU rasterizes the brush).
+const MAX_DABS: usize = 4096;
+
+/// Evaluate the masks (`lightcraft_pipeline::masks::evaluate`): their alpha planes (concatenated)
+/// and their adjustment terms. Shapes without a kernel (Sky, Subject, …) run on the CPU.
 fn masks(cx: &mut Cx<'_>, lin: &Buf, prep: &Prep, plan: &Plan<'_>, host: &mut Host) -> (Option<Buf>, Vec<[f32; MASK_TERMS]>) {
+    use lightcraft_develop::{MaskOp, MaskShape};
     let s = &*plan.settings;
-    if !s.masks.iter().any(|m| m.visible && !m.components.is_empty()) {
+    let list: Vec<_> = s.masks.iter().filter(|m| m.visible && !m.components.is_empty()).collect();
+    if list.is_empty() {
         return (None, Vec::new());
     }
     let (w, h) = (plan.w, plan.h);
-    let img = host.lin.get_or_insert_with(|| Arc::new(cx.read_rgb(lin, w, h))).clone();
-    let l = host.log_l.get_or_insert_with(|| Arc::new(cx.read_plane(&prep.log_l, w, h))).clone();
+    let n = w * h;
+    let frame = &plan.frame;
     let ev = s.light.exposure as f32;
-    let ev = lightcraft_pipeline::masks::evaluate(&s.masks, &plan.frame, w, h, &img, &l, ev);
-    let terms = ev.iter().map(|m| mask_terms(&m.adjust)).collect();
-    let mut all = Vec::with_capacity(ev.len() * w * h);
-    for m in &ev {
-        all.extend_from_slice(&m.alpha.data);
+    let long = frame.ow.max(frame.oh);
+    let [a, b, c, d, e, f] = frame.out_to_norm(w, h).0;
+    let (kx, ky) = (frame.ow / long, frame.oh / long);
+    let aff = [a * kx, b * ky, c * kx, d * ky, e * kx, f * ky].map(|v| (v as f32).to_bits());
+    let alpha = cx.gpu.buffer(n * list.len());
+    let comp_plane = cx.gpu.buffer(n);
+    for (mi, m) in list.iter().enumerate() {
+        let mut first = true;
+        for comp in &m.components {
+            let op = match comp.op {
+                _ if first && comp.op != MaskOp::Intersect => 0u32,
+                MaskOp::Add => 1,
+                MaskOp::Subtract => 2,
+                MaskOp::Intersect => 3,
+            };
+            first = false;
+            let mut p = vec![w as u32, h as u32, 0, comp.invert as u32];
+            p.extend_from_slice(&aff);
+            let f32s = |v: &[f64]| v.iter().map(|x| (*x as f32).to_bits()).collect::<Vec<u32>>();
+            let mut aux: Vec<f32> = Vec::new();
+            let kind = match &comp.shape {
+                MaskShape::Linear { start, end } => {
+                    let (a, b) = (frame.norm_to_long(*start), frame.norm_to_long(*end));
+                    let d = b - a;
+                    p.extend(f32s(&[a.x, a.y, d.x, d.y, d.dot(d).max(1e-12)]));
+                    Some(0)
+                }
+                MaskShape::Radial { center, rx, ry, angle, feather, invert } => {
+                    let c = frame.norm_to_long(*center);
+                    let (sn, co) = (-angle.to_radians()).sin_cos();
+                    p.extend(f32s(&[c.x, c.y, co, sn, rx.max(1e-6), ry.max(1e-6), (*feather / 100.0).clamp(0.0, 1.0)]));
+                    p.push(*invert as u32);
+                    Some(1)
+                }
+                MaskShape::LuminanceRange { lo, hi, lo_feather, hi_feather } => {
+                    p.extend(f32s(&[*lo, *hi, lo_feather.max(1e-3), hi_feather.max(1e-3)]));
+                    p.push(ev.to_bits());
+                    Some(2)
+                }
+                MaskShape::ColorRange { samples, refine } => {
+                    let tol = 0.04 + 0.16 * (*refine as f32 / 100.0);
+                    p.extend([tol.to_bits(), ev.exp2().to_bits(), samples.len() as u32]);
+                    aux.extend(samples.iter().flat_map(|s| s.map(|v| v as f32)));
+                    Some(3)
+                }
+                MaskShape::Brush { strokes } => {
+                    let dabs: Vec<_> = strokes.iter().map(|st| lightcraft_pipeline::masks::brush_dabs(st, frame, w, h)).collect();
+                    if dabs.iter().map(|d| d.dabs.len()).sum::<usize>() <= MAX_DABS {
+                        p.push(dabs.len() as u32);
+                        let mut off = 11 * dabs.len();
+                        for d in &dabs {
+                            let (mut x0, mut y0, mut x1, mut y1) = (f64::MAX, f64::MAX, f64::MIN, f64::MIN);
+                            for q in &d.dabs {
+                                (x0, y0, x1, y1) = (x0.min(q.x - d.r), y0.min(q.y - d.r), x1.max(q.x + d.r), y1.max(q.y + d.r));
+                            }
+                            aux.extend([off as f32, d.dabs.len() as f32, d.r as f32, d.hard as f32, d.flow, d.density, d.erase as u8 as f32]);
+                            aux.extend([x0, y0, x1, y1].map(|v| v as f32));
+                            off += 2 * d.dabs.len();
+                        }
+                        for d in &dabs {
+                            aux.extend(d.dabs.iter().flat_map(|q| [q.x as f32, q.y as f32]));
+                        }
+                        Some(4)
+                    } else {
+                        None
+                    }
+                }
+                _ => None,
+            };
+            let plane = match kind {
+                Some(k) => {
+                    p[2] = k;
+                    let aux = (!aux.is_empty()).then(|| cx.gpu.upload(&aux));
+                    cx.run("shape", &p, &[Some(lin), Some(&prep.log_l), aux.as_ref(), Some(&comp_plane), Some(&alpha)], groups2(w, h, [16, 16]));
+                    None
+                }
+                None => {
+                    // no kernel: evaluate on the CPU
+                    let img = host.lin.get_or_insert_with(|| Arc::new(cx.read_rgb(lin, w, h))).clone();
+                    let l = host.log_l.get_or_insert_with(|| Arc::new(cx.read_plane(&prep.log_l, w, h))).clone();
+                    let mut v = lightcraft_pipeline::masks::shape_alpha(&comp.shape, frame, w, h, &img, &l, ev);
+                    if comp.invert {
+                        v.data.iter_mut().for_each(|x| *x = 1.0 - *x);
+                    }
+                    Some(cx.gpu.upload(&v.data))
+                }
+            };
+            let cbuf = plane.as_ref().unwrap_or(&comp_plane);
+            let k = [n as u32, 0, mi as u32, op];
+            cx.run("combine", &k, &[None, None, None, Some(cbuf), Some(&alpha)], groups1(n));
+        }
+        let amt = (m.adjust.amount / 100.0) as f32;
+        cx.run(
+            "finalize",
+            &[n as u32, 0, mi as u32, m.invert as u32, amt.to_bits()],
+            &[None, None, None, Some(&comp_plane), Some(&alpha)],
+            groups1(n),
+        );
     }
-    (Some(cx.gpu.upload(&all)), terms)
+    let terms = list.iter().map(|m| mask_terms(&m.adjust)).collect();
+    (Some(alpha), terms)
 }

crates/gpu/src/wgsl/mask.wgsl

@@ -0,0 +1,149 @@
+// Mask evaluation (`lightcraft_pipeline::masks`): one component shape → `c`, then combined into
+// the mask's alpha plane. Bindings: img (rgb, pre-exposure), log_l, aux (shape data), c, alpha.
+
+// Position of output pixel (x, y) in long-edge units: P[4..10] = affine (a, b, c, d, e, f).
+fn pos(x: u32, y: u32) -> vec2<f32> {
+    let px = f32(x) + 0.5;
+    let py = f32(y) + 0.5;
+    return vec2<f32>(pf(4u) * px + pf(6u) * py + pf(8u), pf(5u) * px + pf(7u) * py + pf(9u));
+}
+
+// P: w, h, kind, invert, affine[6], then per kind:
+//   0 linear: a.xy, d.xy, |d|²         1 radial: centre.xy, cos, sin, rx, ry, feather, invert
+//   2 luminance range: lo, hi, lo feather, hi feather, ev
+//   3 colour range: tol, gain, sample count (OkLab triples in aux)
+//   4 brush: stroke count (stroke records + dabs in aux)
+@compute @workgroup_size(16, 16)
+fn shape(@builtin(global_invocation_id) g: vec3<u32>) {
+    let w = pu(0u);
+    let h = pu(1u);
+    if (g.x >= w || g.y >= h) {
+        return;
+    }
+    let i = g.y * w + g.x;
+    let kind = pu(2u);
+    var v = 0.0;
+    if (kind == 0u) {
+        let p = pos(g.x, g.y);
+        let a = vec2<f32>(pf(10u), pf(11u));
+        let d = vec2<f32>(pf(12u), pf(13u));
+        let t = dot(p - a, d) / pf(14u);
+        v = 1.0 - sstep(0.0, 1.0, t);
+    } else if (kind == 1u) {
+        let p = pos(g.x, g.y);
+        let d = p - vec2<f32>(pf(10u), pf(11u));
+        let co = pf(12u);
+        let s = pf(13u);
+        let x = d.x * co - d.y * s;
+        let y = d.x * s + d.y * co;
+        let qx = x / pf(14u);
+        let qy = y / pf(15u);
+        let r = sqrt(qx * qx + qy * qy);
+        let f = pf(16u);
+        v = 1.0 - sstep(1.0 - max(f, 0.01), 1.0, r);
+        if (pu(17u) != 0u) {
+            v = 1.0 - v;
+        }
+    } else if (kind == 2u) {
+        let lo = pf(10u);
+        let hi = pf(11u);
+        let y = clamp((log_l[i] + pf(14u) + 8.0) / 12.0, 0.0, 1.0);
+        v = sstep(lo - pf(12u), lo, y) * (1.0 - sstep(hi, hi + pf(13u), y));
+    } else if (kind == 3u) {
+        let tol = pf(10u);
+        let gain = pf(11u);
+        let c = vec3<f32>(img[3u * i], img[3u * i + 1u], img[3u * i + 2u]) * gain;
+        let lab = oklab(c / (1.0 + c));
+        var best = 3.402823e38;
+        for (var k = 0u; k < pu(12u); k++) {
+            let s = vec3<f32>(aux[3u * k], aux[3u * k + 1u], aux[3u * k + 2u]);
+            let da = lab.y - s.y;
+            let db = lab.z - s.z;
+            let dl = lab.x - s.x;
+            best = min(best, sqrt(da * da + db * db + 0.25 * (dl * dl)));
+        }
+        v = 1.0 - sstep(tol * 0.5, tol, best);
+    } else if (kind == 4u) {
+        // stroke record: dab offset, dab count, r, hard, flow, density, erase, bbox x0 y0 x1 y1
+        let px = f32(g.x) + 0.5;
+        let py = f32(g.y) + 0.5;
+        for (var k = 0u; k < pu(10u); k++) {
+            let o = 11u * k;
+            if (px < aux[o + 7u] || py < aux[o + 8u] || px > aux[o + 9u] || py > aux[o + 10u]) {
+                continue;
+            }
+            let off = u32(aux[o]);
+            let cnt = u32(aux[o + 1u]);
+            let r = aux[o + 2u];
+            let hard = aux[o + 3u];
+            let flow = aux[o + 4u];
+            var keep = 1.0;
+            for (var j = 0u; j < cnt; j++) {
+                let dx = px - aux[off + 2u * j];
+                let dy = py - aux[off + 2u * j + 1u];
+                let dd = sqrt(dx * dx + dy * dy);
+                if (dd > r) {
+                    continue;
+                }
+                var a = 1.0;
+                if (dd > hard) {
+                    a = 1.0 - sstep(hard, r, dd);
+                }
+                keep *= 1.0 - a * flow;
+            }
+            let sa = min(1.0 - keep, aux[o + 5u]);
+            if (aux[o + 6u] != 0.0) {
+                v = v * (1.0 - sa);
+            } else {
+                v = max(v, sa);
+            }
+        }
+    }
+    if (pu(3u) != 0u) {
+        v = 1.0 - v;
+    }
+    c[i] = v;
+}
+
+// Combine `c` into mask plane P[2]: P[3] = 0 copy, 1 add (max), 2 subtract, 3 intersect.
+@compute @workgroup_size(256)
+fn combine(@builtin(global_invocation_id) g: vec3<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
+    let i = lin_index(g, nw);
+    let n = pu(0u);
+    if (i >= n) {
+        return;
+    }
+    let j = pu(2u) * n + i;
+    let op = pu(3u);
+    let a = alpha[j];
+    let v = c[i];
+    if (op == 0u) {
+        alpha[j] = v;
+    } else if (op == 1u) {
+        alpha[j] = max(a, v);
+    } else if (op == 2u) {
+        alpha[j] = a * (1.0 - v);
+    } else {
+        alpha[j] = a * v;
+    }
+}
+
+// Finish mask plane P[2]: P[3] = invert, P[4] = amount (applied when ≠ 1).
+@compute @workgroup_size(256)
+fn finalize(@builtin(global_invocation_id) g: vec3<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
+    let i = lin_index(g, nw);
+    let n = pu(0u);
+    if (i >= n) {
+        return;
+    }
+    let j = pu(2u) * n + i;
+    var a = alpha[j];
+    if (pu(3u) != 0u) {
+        a = 1.0 - a;
+    }
+    let amt = pf(4u);
+    if (abs(amt - 1.0) > 1e-6) {
+        a = a * amt;
+    }
+    alpha[j] = a;
+}

crates/gpu/tests/equivalence.rs

@@ -226,6 +226,33 @@ fn cases() -> Vec<(&'static str, Edit)> {
                 },
             ];
         }),
+        ("masks (cpu shapes, ops, amount)", |s| {
+            s.masks = vec![
+                Mask {
+                    components: vec![
+                        MaskComponent { op: MaskOp::Add, invert: false, shape: MaskShape::Sky },
+                        MaskComponent {
+                            op: MaskOp::Subtract,
+                            invert: false,
+                            shape: MaskShape::Radial { center: Point::new(0.7, 0.2), rx: 0.1, ry: 0.1, angle: 0.0, feather: 30.0, invert: false },
+                        },
+                    ],
+                    adjust: lightcraft_develop::LocalAdjustments { exposure: -0.5, dehaze: 40.0, amount: 70.0, ..Default::default() },
+                    ..Default::default()
+                },
+                Mask {
+                    components: vec![MaskComponent { op: MaskOp::Intersect, invert: false, shape: MaskShape::Subject }],
+                    adjust: lightcraft_develop::LocalAdjustments { exposure: 0.5, ..Default::default() },
+                    ..Default::default()
+                },
+                Mask {
+                    components: vec![MaskComponent { op: MaskOp::Add, invert: true, shape: MaskShape::Background }],
+                    adjust: lightcraft_develop::LocalAdjustments { saturation: -60.0, ..Default::default() },
+                    invert: true,
+                    ..Default::default()
+                },
+            ];
+        }),
         ("spots + defringe (cpu stage)", |s| {
             s.spots = vec![Spot { points: vec![Point::new(0.3, 0.3)], size: 0.03, source_offset: Some(Point::new(0.1, 0.0)), ..Default::default() }];
             s.optics.defringe_purple_amount = 5.0;

crates/pipeline/src/masks.rs

@@ -172,29 +172,44 @@ fn tonemap_for_select(c: [f32; 3]) -> [f32; 3] {
     c.map(|v| v / (1.0 + v))
 }
 
-/// Stamp brush strokes into `out` (max-combine; erase strokes subtract).
-fn rasterize_brush(strokes: &[BrushStroke], frame: &Frame, w: usize, h: usize, out: &mut Plane) {
+/// A brush stroke resolved to output pixels: dab centres (spacing r/4 along the path), radius,
+/// hard-core radius, flow and density (0..1), and whether it erases.
+pub struct BrushDabs {
+    pub dabs: Vec<Point>,
+    pub r: f64,
+    pub hard: f64,
+    pub flow: f32,
+    pub density: f32,
+    pub erase: bool,
+}
+
+/// The dabs of stroke `s` in a `w × h` output.
+pub fn brush_dabs(s: &BrushStroke, frame: &Frame, w: usize, h: usize) -> BrushDabs {
     let to_out = frame.norm_to_out(w, h);
     let ppl = frame.px_per_long(w);
-    for s in strokes {
-        let r = (s.size * ppl).max(0.5);
-        let hard = r * (1.0 - (s.feather / 100.0).clamp(0.0, 1.0));
-        let flow = (s.flow / 100.0).clamp(0.0, 1.0) as f32;
-        let dens = (s.density / 100.0).clamp(0.0, 1.0) as f32;
-        // Densify the path so dabs overlap (spacing r/4).
-        let mut dabs: Vec<Point> = Vec::new();
-        for (i, p) in s.points.iter().enumerate() {
-            let q = to_out.apply(*p);
-            if i > 0 {
-                let prev = to_out.apply(s.points[i - 1]);
-                let dist = prev.dist(q);
-                let n = (dist / (r / 4.0).max(0.5)).ceil() as usize;
-                for k in 1..n {
-                    dabs.push(prev.lerp(q, k as f64 / n as f64));
-                }
+    let r = (s.size * ppl).max(0.5);
+    let hard = r * (1.0 - (s.feather / 100.0).clamp(0.0, 1.0));
+    // Densify the path so dabs overlap (spacing r/4).
+    let mut dabs: Vec<Point> = Vec::new();
+    for (i, p) in s.points.iter().enumerate() {
+        let q = to_out.apply(*p);
+        if i > 0 {
+            let prev = to_out.apply(s.points[i - 1]);
+            let dist = prev.dist(q);
+            let n = (dist / (r / 4.0).max(0.5)).ceil() as usize;
+            for k in 1..n {
+                dabs.push(prev.lerp(q, k as f64 / n as f64));
             }
-            dabs.push(q);
         }
+        dabs.push(q);
+    }
+    BrushDabs { dabs, r, hard, flow: (s.flow / 100.0).clamp(0.0, 1.0) as f32, density: (s.density / 100.0).clamp(0.0, 1.0) as f32, erase: s.erase }
+}
+
+/// Stamp brush strokes into `out` (max-combine; erase strokes subtract).
+fn rasterize_brush(strokes: &[BrushStroke], frame: &Frame, w: usize, h: usize, out: &mut Plane) {
+    for s in strokes {
+        let BrushDabs { dabs, r, hard, flow, density: dens, erase } = brush_dabs(s, frame, w, h);
         let mut stroke_alpha = Plane::new(w, h);
         for d in &dabs {
             let (x0, x1) = (((d.x - r).floor().max(0.0)) as usize, ((d.x + r).ceil().max(0.0) as usize).min(w));
@@ -213,7 +228,7 @@ fn rasterize_brush(strokes: &[BrushStroke], frame: &Frame, w: usize, h: usize, o
             }
         }
         for (o, a) in out.data.iter_mut().zip(&stroke_alpha.data) {
-            *o = if s.erase { *o * (1.0 - a) } else { o.max(*a) };
+            *o = if erase { *o * (1.0 - a) } else { o.max(*a) };
         }
     }
 }