Merge upstream/main into pr/overlay-windows (#1027 landed)

GPU handles(): only output-device proofs stay on the CPU; Show All Masks now renders zoom windows (upstream); proof-boundary tolerance in the GPU equivalence test; device-profile draft test.

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

crates/gpu/src/backend.rs

@@ -213,12 +213,12 @@ fn disarm(marker: &std::path::Path, active: &Mutex<u32>) {
 }
 
 /// Removes the marker on the way out, a panic included (the process survived).
-struct Armed(Option<PathBuf>);
+struct Armed<'a>(Option<PathBuf>, &'a Mutex<u32>);
 
-impl Drop for Armed {
+impl Drop for Armed<'_> {
     fn drop(&mut self) {
         if let Some(m) = &self.0 {
-            disarm(m, &ACTIVE);
+            disarm(m, self.1);
         }
     }
 }
@@ -227,10 +227,16 @@ impl Drop for Armed {
 /// removing the init marker (if one is set).
 pub(crate) fn with_marker<T>(what: &str, f: impl FnOnce() -> T) -> T {
     let marker = MARKER.lock().unwrap_or_else(|e| e.into_inner()).clone();
+    with_marker_on(marker, &ACTIVE, what, f)
+}
+
+/// [`with_marker`] on an explicit marker path and scope counter (the tests use their own, so that
+/// device-creating tests running in parallel cannot join their scopes).
+fn with_marker_on<T>(marker: Option<PathBuf>, active: &Mutex<u32>, what: &str, f: impl FnOnce() -> T) -> T {
     if let Some(m) = &marker {
-        arm(m, &ACTIVE, what);
+        arm(m, active, what);
     }
-    let _armed = Armed(marker);
+    let _armed = Armed(marker, active);
     f()
 }
 
@@ -352,9 +358,6 @@ mod tests {
     }
     use super::*;
 
-    /// The tests that set the process-wide marker take turns.
-    static GLOBAL: Mutex<()> = Mutex::new(());
-
     #[test]
     fn parses_backend_names() {
         assert_eq!(parse_backends("dx12"), Some(BackendChoice::Use(Backends::DX12)));
@@ -421,13 +424,12 @@ mod tests {
 
     #[test]
     fn init_marker_is_written_during_creation_and_removed_after() {
-        let _turn = GLOBAL.lock().unwrap_or_else(|e| e.into_inner());
         let dir = std::env::temp_dir().join(format!("lc-gpu-marker-{}", std::process::id()));
         let m = dir.join("gpu-init.marker");
         let _ = std::fs::remove_dir_all(&dir);
-        set_init_marker(Some(m.clone()));
-        let seen = with_init_marker(Backends::DX12, || m.exists());
-        set_init_marker(None);
+        // a local scope counter: other tests create real devices through the process-wide one
+        let active = Mutex::new(0);
+        let seen = with_marker_on(Some(m.clone()), &active, "device creation", || m.exists());
         assert!(seen, "the marker exists while the device is created");
         assert!(!m.exists(), "and is removed afterwards");
         // a marker left behind by a crashed process is reported once; reading it leaves it in place
@@ -440,6 +442,20 @@ mod tests {
         let _ = std::fs::remove_dir_all(&dir);
     }
 
+    #[test]
+    fn with_init_marker_writes_the_process_wide_marker() {
+        let dir = std::env::temp_dir().join(format!("lc-gpu-global-{}", std::process::id()));
+        let m = dir.join("gpu-init.marker");
+        let _ = std::fs::remove_dir_all(&dir);
+        set_init_marker(Some(m.clone()));
+        let seen = with_init_marker(Backends::DX12, || m.exists());
+        set_init_marker(None);
+        // not checked for removal: device-creating tests running in parallel may still hold the file
+        // (the test above checks removal on its own scope counter)
+        assert!(seen, "the marker set with set_init_marker exists while the device is created");
+        let _ = std::fs::remove_dir_all(&dir);
+    }
+
     /// The window's renderer fallback: forced, remembered, the platform chain, the failed one last.
     #[test]
     fn window_renderer_fallback_order() {
@@ -505,13 +521,11 @@ mod tests {
 
     #[test]
     fn a_panic_inside_the_scope_still_removes_the_marker() {
-        let _turn = GLOBAL.lock().unwrap_or_else(|e| e.into_inner());
         let dir = std::env::temp_dir().join(format!("lc-gpu-panic-{}", std::process::id()));
         let m = dir.join("gpu-init.marker");
         let _ = std::fs::remove_dir_all(&dir);
-        set_init_marker(Some(m.clone()));
-        let r = std::panic::catch_unwind(|| with_marker("kernel compilation", || panic!("compile")));
-        set_init_marker(None);
+        let active = Mutex::new(0);
+        let r = std::panic::catch_unwind(|| with_marker_on(Some(m.clone()), &active, "kernel compilation", || panic!("compile")));
         assert!(r.is_err());
         assert!(!m.exists(), "the process survived, so nothing is left to report");
         let _ = std::fs::remove_dir_all(&dir);

crates/gpu/src/lib.rs

@@ -35,6 +35,8 @@ pub mod backend;
 mod ctx;
 #[cfg(not(target_arch = "wasm32"))]
 mod exit;
+#[cfg(not(target_arch = "wasm32"))]
+pub mod faces;
 #[cfg(all(feature = "denoise", not(target_arch = "wasm32")))]
 pub mod nn;
 #[cfg(not(target_arch = "wasm32"))]
@@ -239,9 +241,9 @@ pub fn wait_idle(timeout: std::time::Duration) -> bool {
     }
 }
 
-/// Why GPU compute that has a device of its own (denoise) is switched off by the user's settings, or `None`: the
+/// Why GPU compute that has a device of its own (denoise, faces) is switched off by the user's settings, or `None`: the
 /// environment switches and the rendering preference, but not a render-device failure (the devices are separate).
-#[cfg(all(feature = "denoise", not(target_arch = "wasm32")))]
+#[cfg(not(target_arch = "wasm32"))]
 pub(crate) fn switched_off() -> Option<String> {
     if env_disabled() {
         return Some(if backend::env_off() { "disabled by LIGHTCRAFT_GPU_BACKEND=off" } else { "disabled by LIGHTCRAFT_GPU=0" }.into());
@@ -437,6 +439,15 @@ pub fn render(src: &Arc<Rgb32f>, info: &SourceInfo, s: &DevelopSettings, req: &R
     try_render(src, info, s, req, stages).ok()
 }
 
+/// Whether the GPU renders requests of this kind at all (a given render can still fall back to
+/// the CPU): the kernel writes 8-bit output, so high-bit-depth exports render on the CPU, and so do
+/// soft proofs against an output-device profile (`device_proof`: its 3D table is CPU-side) and
+/// Visualize HDR, which needs the float HDR render. [`try_render`] follows it, and the loupe sizes
+/// its slider-drag drafts by it (issue #330), so the two can't drift apart.
+pub fn handles(depth: lightcraft_pipeline::OutputDepth, device_proof: bool, overlay: lightcraft_pipeline::Overlay) -> bool {
+    depth == lightcraft_pipeline::OutputDepth::U8 && !device_proof && overlay != lightcraft_pipeline::Overlay::HdrRange
+}
+
 /// [`render`], reporting why this render made no GPU image: [`Fallback::NotTried`] when the GPU
 /// was not used (callers can ask [`unavailable_reason`]), [`Fallback::Failed`] with the reason
 /// when the GPU render failed and the CPU must render it. The reason belongs to this call, unlike
@@ -453,13 +464,8 @@ pub fn try_render(
         // sections switched off with their eye render as if at their defaults (issue #316)
         let effective = s.effective();
         let s: &DevelopSettings = &effective;
-        // the kernel writes 8-bit output: high-bit-depth exports render on the CPU, and so do soft
-        // proofs against an output-device profile (its 3D table); Visualize HDR needs the float HDR render: CPU too
-        if !enabled()
-            || req.depth != lightcraft_pipeline::OutputDepth::U8
-            || req.proof.is_some_and(|p| lightcraft_pipeline::proof_profile(p.profile).is_some())
-            || req.overlay == lightcraft_pipeline::Overlay::HdrRange
-        {
+        let device_proof = req.proof.is_some_and(|p| lightcraft_pipeline::proof_profile(p.profile).is_some());
+        if !enabled() || !handles(req.depth, device_proof, req.overlay) {
             return Err(Fallback::NotTried);
         }
         let s = &*lightcraft_pipeline::settings_for(s, req);
@@ -533,3 +539,19 @@ pub fn try_render(
         Err(Fallback::NotTried)
     }
 }
+
+#[cfg(test)]
+mod tests {
+    use lightcraft_pipeline::{OutputDepth, Overlay};
+
+    // Given a request the 8-bit kernel can't draw (high bit depth, a soft proof, Visualize HDR),
+    // the GPU doesn't take it; the loupe and diagnostic overlays it can draw stay on the GPU
+    #[test]
+    fn the_gpu_handles_8_bit_unproofed_renders_without_the_hdr_range() {
+        assert!(super::handles(OutputDepth::U8, false, Overlay::None));
+        assert!(super::handles(OutputDepth::U8, false, Overlay::Spots(50)));
+        assert!(!super::handles(OutputDepth::U16, false, Overlay::None));
+        assert!(!super::handles(OutputDepth::U8, true, Overlay::None));
+        assert!(!super::handles(OutputDepth::U8, false, Overlay::HdrRange));
+    }
+}

crates/gpu/src/nn.rs

@@ -907,7 +907,7 @@ mod tests {
         match runner(net, tile) {
             Ok(r) => Some(r),
             Err(e) => {
-                eprintln!("skipped: no GPU runner here ({e})");
+                lightcraft_testkit::skip!(Gpu, "no GPU runner here ({e})");
                 None
             }
         }

crates/gpu/src/render.rs

@@ -824,7 +824,7 @@ fn linear(cx: &mut Cx<'_>, sampled: &Buf, info: &SourceInfo, plan: &Plan<'_>, ho
 fn denoise(cx: &mut Cx<'_>, img: Buf, plan: &Plan<'_>) -> Buf {
     let (w, h) = (plan.w, plan.h);
     let n = w * h;
-    let (lum, col) = local::nr_params(&plan.settings, plan.src_long, plan.frame.output_long(w, h));
+    let (lum, col) = local::nr_params(&plan.settings, plan.detail_scale);
     let mut img = img;
     if let Some(nr) = lum {
         let l = cx.gpu.buffer(n);
@@ -1015,7 +1015,8 @@ fn masks(cx: &mut Cx<'_>, lin: &Buf, prep: &Prep, plan: &Plan<'_>, host: &mut Ho
                 }
                 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 {
+                    // pen pressure varies radius and flow per dab, which the kernel doesn't take: the CPU rasterizes it
+                    if dabs.iter().map(|d| d.dabs.len()).sum::<usize>() > MAX_DABS || dabs.iter().any(|d| !d.pressure.is_empty()) {
                         None
                     } else if dabs.iter().any(|d| d.auto) {
                         // Auto Mask: stroke by stroke (each one refined before it's combined)

crates/gpu/src/wgsl/faces_conv.wgsl

@@ -0,0 +1,171 @@
+// Convolution for the face networks (lightcraft-faces), planar (NCHW) float32, batch 1, any stride, padding and grouping.
+//
+// `tiled` is an implicit matrix product per group: rows are the group's output channels (M), columns its output
+// pixels (N), and K runs over input channel x kernel tap. The input is gathered (im2col) straight into workgroup
+// memory, so the lowered matrix is never made. 256 threads make a 64 x 64 block, 4 x 4 values each.
+// `direct` gives one output value per thread, for depthwise convolutions (one output channel per group), where a
+// 64-row tile would be almost all waste.
+//
+//   out[co][p] = bias[co] + sum_k W[co][k] * X[k][p]     X[k][p] = input[g*cg + k/taps][oy*sh + ky - pt][ox*sw + kx - pl]
+
+struct P {
+    m: u32,      // output channels per group
+    n: u32,      // output pixels (oh * ow)
+    k: u32,      // cg * kh * kw
+    cg: u32,     // input channels per group
+    h: u32,
+    w: u32,
+    ow: u32,
+    kh: u32,
+    kw: u32,
+    sh: u32,
+    sw: u32,
+    pt: u32,
+    pl: u32,
+    bias: u32,   // 1 when there is a bias
+    total: u32,  // cout * n
+    row: u32,    // workgroups along x, for `direct` and `reduce` (so a large layer can use y too)
+    splits: u32, // K is cut into this many slices (z = group * splits + slice), summed by `reduce`
+    kchunk: u32, // K values per slice, a multiple of BK
+    _a: u32,
+    _b: u32,
+}
+
+@group(0) @binding(0) var<uniform> p: P;
+@group(0) @binding(1) var<storage, read> src: array<f32>;
+@group(0) @binding(2) var<storage, read> wgt: array<f32>;
+@group(0) @binding(3) var<storage, read> bias: array<f32>;
+@group(0) @binding(4) var<storage, read_write> dst: array<f32>;
+
+const BM: u32 = 64u;
+const BN: u32 = 64u;
+const BK: u32 = 16u;
+
+var<workgroup> As: array<f32, 1024>; // [BK][BM]
+var<workgroup> Bs: array<f32, 1024>; // [BK][BN]
+
+fn gather(g: u32, k: u32, pix: u32) -> f32 {
+    let taps = p.kh * p.kw;
+    let ci = g * p.cg + k / taps;
+    let tap = k % taps;
+    let oy = pix / p.ow;
+    let ox = pix % p.ow;
+    let iy = i32(oy * p.sh + tap / p.kw) - i32(p.pt);
+    let ix = i32(ox * p.sw + tap % p.kw) - i32(p.pl);
+    if (iy < 0 || ix < 0 || iy >= i32(p.h) || ix >= i32(p.w)) {
+        return 0.0;
+    }
+    return src[(ci * p.h + u32(iy)) * p.w + u32(ix)];
+}
+
+@compute @workgroup_size(256)
+fn tiled(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) tid: u32) {
+    let n0 = wg.x * BN;
+    let m0 = wg.y * BM;
+    let g = wg.z / p.splits;
+    let slice = wg.z % p.splits;
+    let k_lo = slice * p.kchunk;
+    let k_hi = min(p.k, k_lo + p.kchunk);
+    let tx = tid % 16u;
+    let ty = tid / 16u;
+    var acc: array<array<f32, 4>, 4>;
+    for (var i = 0u; i < 4u; i++) {
+        for (var j = 0u; j < 4u; j++) {
+            acc[i][j] = 0.0;
+        }
+    }
+    for (var k0 = k_lo; k0 < k_hi; k0 += BK) {
+        // 1024 values of each tile, 4 per thread; every slot has exactly one writer
+        for (var e = tid; e < 1024u; e += 256u) {
+            // weights: neighbouring threads read neighbouring k of one row
+            let mm = e / BK;
+            let ka = k0 + e % BK;
+            let m = m0 + mm;
+            var a = 0.0;
+            if (ka < k_hi && m < p.m) {
+                a = wgt[(g * p.m + m) * p.k + ka];
+            }
+            As[(e % BK) * BM + mm] = a;
+            // input: neighbouring threads read neighbouring pixels
+            let kk = e / BN;
+            let k = k0 + kk;
+            let pix = n0 + e % BN;
+            var b = 0.0;
+            if (k < k_hi && pix < p.n) {
+                b = gather(g, k, pix);
+            }
+            Bs[e] = b;
+        }
+        workgroupBarrier();
+        for (var kk = 0u; kk < BK; kk++) {
+            var a: array<f32, 4>;
+            var b: array<f32, 4>;
+            for (var i = 0u; i < 4u; i++) {
+                a[i] = As[kk * BM + ty * 4u + i];
+                b[i] = Bs[kk * BN + tx + 16u * i];
+            }
+            for (var i = 0u; i < 4u; i++) {
+                for (var j = 0u; j < 4u; j++) {
+                    acc[i][j] += a[i] * b[j];
+                }
+            }
+        }
+        workgroupBarrier();
+    }
+    for (var i = 0u; i < 4u; i++) {
+        let m = m0 + ty * 4u + i;
+        if (m >= p.m) {
+            continue;
+        }
+        let co = g * p.m + m;
+        var b0 = 0.0;
+        if (p.bias != 0u && p.splits == 1u) {
+            b0 = bias[co];
+        }
+        // with slices, each writes a partial sum of its own: [slice][cout][n]
+        let base = slice * p.total + co * p.n;
+        for (var j = 0u; j < 4u; j++) {
+            // columns tx, tx + 16, tx + 32, tx + 48: neighbouring threads write neighbouring pixels
+            let pix = n0 + tx + 16u * j;
+            if (pix < p.n) {
+                dst[base + pix] = acc[i][j] + b0;
+            }
+        }
+    }
+}
+
+@compute @workgroup_size(256)
+fn direct(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) tid: u32) {
+    let at = (wg.y * p.row + wg.x) * 256u + tid;
+    if (at >= p.total) {
+        return;
+    }
+    let co = at / p.n;
+    let pix = at % p.n;
+    let g = co / p.m;
+    var v = 0.0;
+    if (p.bias != 0u) {
+        v = bias[co];
+    }
+    for (var k = 0u; k < p.k; k++) {
+        v += wgt[co * p.k + k] * gather(g, k, pix);
+    }
+    dst[at] = v;
+}
+
+// The sum of the K slices' partial results, plus the bias. `src` holds [splits][cout][n].
+@compute @workgroup_size(256)
+fn reduce(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) tid: u32) {
+    let at = (wg.y * p.row + wg.x) * 256u + tid;
+    if (at >= p.total) {
+        return;
+    }
+    var v = 0.0;
+    if (p.bias != 0u) {
+        v = bias[at / p.n];
+    }
+    for (var s = 0u; s < p.splits; s++) {
+        v += src[s * p.total + at];
+    }
+    dst[at] = v;
+}

crates/gpu/src/wgsl/faces_ops.wgsl

@@ -0,0 +1,188 @@
+// The face networks' other operators (lightcraft-faces `net::program`), so activations stay on the device between layers.
+// Same bindings as faces_conv.wgsl: the parameters, three read-only operands and the output. Element kernels run one
+// thread per output value; `gavg` and `gemv` one workgroup per output value.
+//
+// Parameters are 24 numbers: q(0) = values (or workgroups) to make, q(1) = workgroups along x, then per kernel.
+
+struct P {
+    v: array<vec4<u32>, 6>,
+}
+
+@group(0) @binding(0) var<uniform> p: P;
+@group(0) @binding(1) var<storage, read> a: array<f32>;
+@group(0) @binding(2) var<storage, read> b: array<f32>;
+@group(0) @binding(3) var<storage, read> c: array<f32>;
+@group(0) @binding(4) var<storage, read_write> dst: array<f32>;
+
+fn q(i: u32) -> u32 {
+    return p.v[i / 4u][i % 4u];
+}
+
+fn index(wid: vec3<u32>, lid: u32) -> u32 {
+    return (wid.y * q(1) + wid.x) * 256u + lid;
+}
+
+// q(2): 0 relu, 1 sigmoid
+@compute @workgroup_size(256)
+fn unary(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) lid: u32) {
+    let i = index(wid, lid);
+    if (i >= q(0)) { return; }
+    let x = a[i];
+    if (q(2) == 0u) {
+        dst[i] = max(x, 0.0);
+    } else {
+        dst[i] = 1.0 / (1.0 + exp(-x));
+    }
+}
+
+// q(2) plane, q(3) channels: x * b[c] + c[c]
+@compute @workgroup_size(256)
+fn affine(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) lid: u32) {
+    let i = index(wid, lid);
+    if (i >= q(0)) { return; }
+    let ch = (i / q(2)) % q(3);
+    dst[i] = a[i] * b[ch] + c[ch];
+}
+
+// q(2) plane, q(3) channels: x < 0 ? x * b[c] : x
+@compute @workgroup_size(256)
+fn prelu(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) lid: u32) {
+    let i = index(wid, lid);
+    if (i >= q(0)) { return; }
+    let x = a[i];
+    let ch = (i / q(2)) % q(3);
+    dst[i] = select(x, x * b[ch], x < 0.0);
+}
+
+// q(2): 0 add, 1 sub, 2 mul; q(3..7) output dims; q(7..11) a's strides; q(11..15) b's strides
+@compute @workgroup_size(256)
+fn binary(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) lid: u32) {
+    let i = index(wid, lid);
+    if (i >= q(0)) { return; }
+    var rest = i;
+    var ia = 0u;
+    var ib = 0u;
+    for (var d = 3; d >= 0; d = d - 1) {
+        let du = u32(d);
+        let size = q(3u + du);
+        let at = rest % size;
+        rest = rest / size;
+        ia = ia + at * q(7u + du);
+        ib = ib + at * q(11u + du);
+    }
+    let x = a[ia];
+    let y = b[ib];
+    if (q(2) == 0u) {
+        dst[i] = x + y;
+    } else if (q(2) == 1u) {
+        dst[i] = x - y;
+    } else {
+        dst[i] = x * y;
+    }
+}
+
+// q(3..7) output dims; q(7..11) the input's strides for each output axis
+@compute @workgroup_size(256)
+fn transpose(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) lid: u32) {
+    let i = index(wid, lid);
+    if (i >= q(0)) { return; }
+    var rest = i;
+    var at_in = 0u;
+    for (var d = 3; d >= 0; d = d - 1) {
+        let du = u32(d);
+        let size = q(3u + du);
+        at_in = at_in + (rest % size) * q(7u + du);
+        rest = rest / size;
+    }
+    dst[i] = a[at_in];
+}
+
+// q(2) h, q(3) w, q(4) kh, q(5) kw, q(6) sh, q(7) sw, q(8) oh, q(9) ow
+@compute @workgroup_size(256)
+fn maxpool(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) lid: u32) {
+    let i = index(wid, lid);
+    if (i >= q(0)) { return; }
+    let h = q(2);
+    let w = q(3);
+    let oh = q(8);
+    let ow = q(9);
+    let ox = i % ow;
+    let oy = (i / ow) % oh;
+    let ch = i / (ow * oh);
+    var m = -3.402823e38;
+    for (var ky = 0u; ky < q(4); ky = ky + 1u) {
+        for (var kx = 0u; kx < q(5); kx = kx + 1u) {
+            let y = oy * q(6) + ky;
+            let x = ox * q(7) + kx;
+            if (y < h && x < w) {
+                m = max(m, a[(ch * h + y) * w + x]);
+            }
+        }
+    }
+    dst[i] = m;
+}
+
+// q(2) h, q(3) w, q(4) sh, q(5) sw: nearest upscale by whole numbers
+@compute @workgroup_size(256)
+fn resize(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) lid: u32) {
+    let i = index(wid, lid);
+    if (i >= q(0)) { return; }
+    let h = q(2);
+    let w = q(3);
+    let ow = w * q(5);
+    let oh = h * q(4);
+    let ox = i % ow;
+    let oy = (i / ow) % oh;
+    let ch = i / (ow * oh);
+    dst[i] = a[(ch * h + oy / q(4)) * w + ox / q(5)];
+}
+
+var<workgroup> partial: array<f32, 256>;
+
+fn reduce(lid: u32) {
+    for (var s = 128u; s > 0u; s = s >> 1u) {
+        workgroupBarrier();
+        if (lid < s) {
+            partial[lid] = partial[lid] + partial[lid + s];
+        }
+    }
+    workgroupBarrier();
+}
+
+// one workgroup per channel; q(2) plane
+@compute @workgroup_size(256)
+fn gavg(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) lid: u32) {
+    let ch = wid.y * q(1) + wid.x;
+    let plane = q(2);
+    var sum = 0.0;
+    if (ch < q(0)) {
+        for (var j = lid; j < plane; j = j + 256u) {
+            sum = sum + a[ch * plane + j];
+        }
+    }
+    partial[lid] = sum;
+    reduce(lid);
+    if (lid == 0u && ch < q(0)) {
+        dst[ch] = partial[0] / f32(plane);
+    }
+}
+
+// one workgroup per output; q(2) k, q(3) n, q(4) 1 when B is [n][k], q(5) alpha's bits
+@compute @workgroup_size(256)
+fn gemv(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) lid: u32) {
+    let o = wid.y * q(1) + wid.x;
+    let k = q(2);
+    let n = q(3);
+    var sum = 0.0;
+    if (o < q(0)) {
+        for (var j = lid; j < k; j = j + 256u) {
+            let wi = select(j * n + o, o * k + j, q(4) == 1u);
+            sum = sum + a[j] * b[wi];
+        }
+    }
+    partial[lid] = sum;
+    reduce(lid);
+    if (lid == 0u && o < q(0)) {
+        dst[o] = bitcast<f32>(q(5)) * partial[0] + c[o];
+    }
+}