LR-EDIT-DETAIL-DENOISE: avoid shared vector write races (#479)

crates/gpu/src/nn.rs

@@ -509,7 +509,6 @@ fn conv_source(bm: u32, bn: u32, bk: u32) -> String {
         .replace("{{BM}}", &bm.to_string())
         .replace("{{BN}}", &bn.to_string())
         .replace("{{BK}}", &bk.to_string())
-        .replace("{{A_VECS}}", &(bk * bm / 4).to_string())
         .replace("{{B_VECS}}", &(bk * bn / 4).to_string())
 }
 
@@ -900,6 +899,7 @@ mod tests {
 
     fn worst(a: &[f32], b: &[f32]) -> f32 {
         assert_eq!(a.len(), b.len());
+        assert!(a.iter().chain(b).all(|v| v.is_finite()), "non-finite denoise output");
         let scale = b.iter().fold(1e-3f32, |m, v| m.max(v.abs()));
         a.iter().zip(b).fold(0f32, |m, (x, y)| m.max((x - y).abs())) / scale
     }
@@ -923,6 +923,21 @@ mod tests {
         }
     }
 
+    // Issue #479: several pixels loaded by different invocations shared a vec4,
+    // so lane writes raced on Metal. Exercise a partial block repeatedly with new
+    // inputs; stale shared-memory values must never leak into the next dispatch.
+    #[test]
+    fn partial_tiles_remain_correct_when_reusing_a_runner() {
+        let n = net(3, 4, 0, 8);
+        let Some(g) = gpu(&n, 3) else { return };
+        for seed in 100..112 {
+            let x = picture(4 * 3 * 3, seed);
+            let want = reference::run(&n, 3, &x).unwrap();
+            let e = worst(&g.run(&x).unwrap(), &want);
+            assert!(e < 2e-4, "seed {seed}: off by {e:e}");
+        }
+    }
+
     #[test]
     fn a_runner_serves_many_tiles_from_many_threads() {
         let n = net(32, 8, 2, 7);
@@ -1020,7 +1035,6 @@ mod tests {
             .replace("{{BM}}", &bm.to_string())
             .replace("{{BN}}", &bn.to_string())
             .replace("{{BK}}", &bk.to_string())
-            .replace("{{A_VECS}}", &(bk * bm / 4).to_string())
             .replace("{{B_VECS}}", &(bk * bn / 4).to_string());
         let d = dev().unwrap();
         let layout = conv_bgl(d);

crates/gpu/src/wgsl/nn_conv.wgsl

@@ -44,8 +44,11 @@ const TX: u32 = BN / 4u;
 const A_ITEMS: u32 = BM * (BK / 4u);
 const B_ITEMS: u32 = BK * (BN / 4u);
 
-// A[k][m] four pixels to a vec4, B[k][n] four channels to a vec4
-var<workgroup> As: array<vec4<f32>, {{A_VECS}}>;
+// A[k][m] has one scalar per pixel; B[k][n] packs four channels per vec4.
+// Different invocations must not write lanes of the same shared vector: WGSL permits
+// a component write to access the entire vector. A later barrier cannot undo that race.
+// Each scalar below has one writer; gather four pixels only after the barrier.
+var<workgroup> As: array<f32, BK * BM>;
 var<workgroup> Bs: array<vec4<f32>, {{B_VECS}}>;
 
 @compute @workgroup_size(256)
@@ -88,12 +91,11 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) t
                         }
                     }
                 }
-                let base = (q * 4u) * (BM / 4u) + ml / 4u;
-                let lane = ml % 4u;
-                As[base][lane] = v.x;
-                As[base + BM / 4u][lane] = v.y;
-                As[base + 2u * (BM / 4u)][lane] = v.z;
-                As[base + 3u * (BM / 4u)][lane] = v.w;
+                let base = q * 4u * BM + ml;
+                As[base] = v.x;
+                As[base + BM] = v.y;
+                As[base + 2u * BM] = v.z;
+                As[base + 3u * BM] = v.w;
             }
             for (var item = tid; item < B_ITEMS; item += 256u) {
                 let r = item / (BN / 4u);
@@ -103,7 +105,8 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) t
             }
             workgroupBarrier();
             for (var kk = 0u; kk < BK; kk++) {
-                let a = As[kk * (BM / 4u) + ty];
+                let ai = kk * BM + ty * 4u;
+                let a = vec4<f32>(As[ai], As[ai + 1u], As[ai + 2u], As[ai + 3u]);
                 let b = Bs[kk * TX + tx];
                 acc0 += a.x * b;
                 acc1 += a.y * b;

docs/denoise.md

@@ -88,6 +88,13 @@ from a check tile; this estimate can choose imperfectly on unmeasured hardware.
 Settings offers Automatic, Graphics card and Processor. GPU tests cover supported
 synthetic networks against the scalar executor; a real-model test compares CPU/GPU.
 
+The convolution kernel stores its shared input tile as scalar elements, each with
+one writer, and gathers vectors after a workgroup barrier. Writing separate lanes
+of the same shared vector from different invocations is not safe under WGSL
+([component references](https://www.w3.org/TR/WGSL/#vector-component-reference));
+this caused incorrect Metal results (issue #479). Synthetic equivalence tests include
+repeated partial tiles and concurrent calls, and reject non-finite outputs.
+
 ### Historical performance, measured 2026-10-07
 
 These are the contributor's measurements **before the current upstream integration**,

docs/parity.md

@@ -267,7 +267,7 @@ under the UI command that wraps it (a no-op error); the UI command now wins.
 | LR-EDIT-DETAIL-SHARPEN | Sharpening | P0 | ✅ | `ctl:detail.sharpenAmount`, `ctl:detail.sharpenRadius`, `ctl:detail.sharpenDetail`, `ctl:detail.sharpenMasking` | no Alt-drag mask preview |
 | LR-EDIT-DETAIL-NR | Luminance noise reduction | P0 | ✅ | `ctl:detail.nrLuminance`, `ctl:detail.nrDetail`, `ctl:detail.nrContrast` | |
 | LR-EDIT-DETAIL-CNR | Colour noise reduction | P0 | ✅ | `ctl:detail.nrColor`, `ctl:detail.nrColorDetail`, `ctl:detail.nrColorSmoothness` | |
-| LR-EDIT-DETAIL-DENOISE | AI denoise | P2 | 🟡 | cmd:denoise.toggle, cmd:denoise.models.install, cmd:denoise.queue, ctl:enhance.denoise, `crates/denoise/`, `docs/denoise.md` | non-destructive Bayer RAW, opt-in weights, pure-Rust CPU/GPU; X-Trans/linear RGB, model policy and Lightroom fidelity remain |
+| LR-EDIT-DETAIL-DENOISE | AI denoise | P2 | 🟡 | cmd:denoise.toggle, cmd:denoise.models.install, cmd:denoise.queue, ctl:enhance.denoise, `crates/denoise/`, `docs/denoise.md` | non-destructive Bayer RAW, opt-in weights, pure-Rust CPU/GPU; GPU convolution uses independently written shared scalars (Metal equivalence regression, issue #479); X-Trans/linear RGB, model policy and Lightroom fidelity remain |
 | LR-EDIT-DETAIL-RAWDETAILS | Improved demosaic toggle | P2 | ⬜ | | |
 | LR-EDIT-DETAIL-SUPERRES | Super resolution | P2 | ⬜ | | |
 | LR-EDIT-DETAIL-AISHARPEN | AI sharpen | OOS | 🚫 | | |