From bec18415611df47cbb60109f3f64efa02c494717 Mon Sep 17 00:00:00 2001 From: Elliot Graham Date: Fri, 9 Oct 2026 13:12:47 +0100 Subject: [PATCH] LR-EDIT-DETAIL-DENOISE: avoid shared vector write races (#479) --- crates/gpu/src/nn.rs | 18 ++++++++++++++++-- crates/gpu/src/wgsl/nn_conv.wgsl | 21 ++++++++++++--------- docs/denoise.md | 7 +++++++ docs/parity.md | 2 +- 4 files changed, 36 insertions(+), 12 deletions(-) diff --git a/crates/gpu/src/nn.rs b/crates/gpu/src/nn.rs index 4231dae8c..2aba039ba 100644 --- a/crates/gpu/src/nn.rs +++ b/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); diff --git a/crates/gpu/src/wgsl/nn_conv.wgsl b/crates/gpu/src/wgsl/nn_conv.wgsl index d991e9e76..20a0ce996 100644 --- a/crates/gpu/src/wgsl/nn_conv.wgsl +++ b/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 As: array, {{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 As: array; var Bs: array, {{B_VECS}}>; @compute @workgroup_size(256) @@ -88,12 +91,11 @@ fn main(@builtin(workgroup_id) wg: vec3, @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, @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(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; diff --git a/docs/denoise.md b/docs/denoise.md index f1196e9b5..be208a9b8 100644 --- a/docs/denoise.md +++ 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**, diff --git a/docs/parity.md b/docs/parity.md index c06dad1b0..be103874a 100644 --- a/docs/parity.md +++ b/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 | 🚫 | | |