Merge PR #1126

crates/gpu/src/render.rs

@@ -229,10 +229,19 @@ pub(crate) fn rgb_words(img: &Rgb32f) -> &[f32] {
 // ---------------------------------------------------------------------------------------------
 // Filters (twins of `lightcraft_raster::blur` / `resample` and `lightcraft_pipeline::local`)
 
-/// Pixels each blur thread slides its running sum over (at least; longer for large radii, so the
+/// Rows each `box_v` thread slides its running sum over (at least; longer for large radii, so the
 /// priming sum stays a small share of the work).
 const CHUNK: usize = 16;
 
+/// One box pass of radius `r` (clamped edges) over a `w × h` image of `nc` interleaved channels,
+/// `src → dst`: horizontal or vertical.
+fn box_pass(cx: &mut Cx<'_>, horizontal: bool, src: &Buf, dst: &Buf, (w, h): (usize, usize), nc: usize, r: usize) {
+    let chunk = (2 * r + 1).clamp(CHUNK, 128);
+    let p = [w as u32, h as u32, nc as u32, r as u32, chunk as u32];
+    let (k, groups) = if horizontal { ("box_h", groups2(w, h, [64, 4])) } else { ("box_v", groups2(w, h.div_ceil(chunk), [64, 4])) };
+    cx.run(k, &p, &[Some(src), Some(dst)], groups);
+}
+
 /// Gaussian blur of a `w × h` image of `nc` interleaved channels: three box passes each way
 /// (`lightcraft_raster::blur::gaussian`).
 pub(crate) fn gaussian(cx: &mut Cx<'_>, src: &Buf, w: usize, h: usize, nc: usize, sigma: f32) -> Buf {
@@ -240,18 +249,15 @@ pub(crate) fn gaussian(cx: &mut Cx<'_>, src: &Buf, w: usize, h: usize, nc: usize
         return cx.copy(src);
     }
     let radii = lightcraft_raster::blur::box_radii(sigma);
-    let passes: Vec<(&str, usize)> = radii.iter().map(|r| ("box_h", *r)).chain(radii.iter().map(|r| ("box_v", *r))).filter(|p| p.1 > 0).collect();
+    let passes: Vec<(bool, usize)> = radii.iter().map(|r| (true, *r)).chain(radii.iter().map(|r| (false, *r))).filter(|p| p.1 > 0).collect();
     if passes.is_empty() {
         return cx.copy(src);
     }
     // ping-pong between two buffers
     let mut bufs = [cx.gpu.buffer(w * h * nc), cx.gpu.buffer(w * h * nc)];
-    for (i, (k, r)) in passes.iter().enumerate() {
-        let chunk = (2 * r + 1).clamp(CHUNK, 128);
-        let groups = if *k == "box_h" { groups2(w.div_ceil(chunk), h, [64, 4]) } else { groups2(w, h.div_ceil(chunk), [64, 4]) };
-        let p = [w as u32, h as u32, nc as u32, *r as u32, chunk as u32];
+    for (i, (horizontal, r)) in passes.iter().enumerate() {
         let (out, prev) = (&bufs[i % 2], &bufs[(i + 1) % 2]);
-        cx.run(k, &p, &[Some(if i == 0 { src } else { prev }), Some(out)], groups);
+        box_pass(cx, *horizontal, if i == 0 { src } else { prev }, out, (w, h), nc, *r);
     }
     let last = (passes.len() - 1) % 2;
     let [a, b] = std::mem::replace(&mut bufs, [cx.gpu.buffer(0), cx.gpu.buffer(0)]);
@@ -1153,6 +1159,33 @@ mod tests {
         assert!(covered > 0 && covered < w * h, "{covered} px covered");
     }
 
+    /// The blur against `lightcraft_raster::blur::gaussian`: widths that do and don't fill the last
+    /// workgroup, windows wider than the image, and 1, 2 (as two planes) and 3 channels.
+    #[test]
+    fn gaussian_matches_cpu() {
+        use lightcraft_raster::blur;
+        let Some(gpu) = crate::device() else { return };
+        let _scope = crate::ctx::RenderScope::new(gpu);
+        let v = |x: usize, y: usize, c: usize| ((x * 7 + y * 13 + c * 5) % 17) as f32 / 16.0 + if x % 61 < 3 { 2.0 } else { 0.0 };
+        for (w, h) in [(1000, 5), (257, 3), (5, 4), (1, 9)] {
+            for sigma in [0.8f32, 2.0, 4.5, 30.0, 250.0, 1000.0] {
+                let planes: Vec<Plane> = (0..3).map(|c| Plane::from_fn(w, h, |x, y| v(x, y, c))).collect();
+                let want: Vec<Plane> = planes.iter().map(|p| blur::gaussian(p, sigma)).collect();
+                for nc in [1, 2, 3] {
+                    let words: Vec<f32> = (0..w * h).flat_map(|i| (0..nc).map(move |c| (i, c))).map(|(i, c)| planes[c].data[i]).collect();
+                    let mut cx = Cx::new(gpu);
+                    let src = gpu.upload(&words);
+                    let out = gaussian(&mut cx, &src, w, h, nc, sigma);
+                    let got: Vec<f32> = cx.read(&out, w * h * nc);
+                    for (i, g) in got.iter().enumerate() {
+                        let want = want[i % nc].data[i / nc];
+                        assert!((g - want).abs() < 1e-5 * (1.0 + sigma / 10.0), "{w}×{h} nc={nc} sigma={sigma}: word {i} {g} vs {want}");
+                    }
+                }
+            }
+        }
+    }
+
     /// Kernel timings on a 24 MP plane: `cargo test --release -p lightcraft-gpu -- --ignored --nocapture`.
     #[test]
     #[ignore]
@@ -1160,6 +1193,7 @@ mod tests {
         let Some(gpu) = crate::device() else { return };
         let _scope = crate::ctx::RenderScope::new(gpu);
         let (w, h) = (6000, 4000);
+        eprintln!("{} ({:?})", gpu.info.name, gpu.info.backend);
         let src = gpu.upload(&vec![0.5f32; w * h * 3]);
         let time = |name: &str, f: &mut dyn FnMut(&mut Cx<'_>)| {
             let mut best = f64::MAX;
@@ -1173,9 +1207,19 @@ mod tests {
             }
             eprintln!("{name:<40} {best:.1} ms");
         };
+        for nc in [1, 2, 3] {
+            let out = gpu.buffer(w * h * nc);
+            // noise reduction's radii, Texture's and Dehaze's at 6000 px
+            for r in [2, 11, 120] {
+                for (k, horizontal) in [("box_h", true), ("box_v", false)] {
+                    time(&format!("{k} nc={nc} r={r}"), &mut |cx| box_pass(cx, horizontal, &src, &out, (w, h), nc, r));
+                }
+            }
+        }
         for nc in [1, 2, 3] {
             time(&format!("gaussian nc={nc} sigma=2"), &mut |cx| drop(gaussian(cx, &src, w, h, nc, 2.0)));
         }
+        time("gaussian nc=3 sigma=4.5 (colour NR)", &mut |cx| drop(gaussian(cx, &src, w, h, 3, 4.5)));
         time("map log_lum", &mut |cx| {
             let o = cx.gpu.buffer(w * h);
             map(cx, "log_lum_k", w * h, &[], [Some(&src), None, None], &o);

crates/gpu/src/wgsl/blur.wgsl

@@ -1,6 +1,6 @@
 // Box passes (clamped edges) over an image of NC interleaved channels; three of each approximate a
-// Gaussian (`lightcraft_raster::blur::gaussian`). Each thread slides a running sum over CH pixels.
-// P: w, h, nc, r, ch. Bindings: src, dst.
+// Gaussian (`lightcraft_raster::blur::gaussian`).
+// P: w, h, nc, r, ch (box_v only). Bindings: src, dst.
 
 // Pixel `i` (all nc ≤ 3 channels at once: one contiguous read per pixel).
 fn ld(i: u32, nc: u32) -> vec3<f32> {
@@ -25,31 +25,27 @@ fn st(i: u32, nc: u32, v: vec3<f32>) {
     }
 }
 
-// Horizontal: one thread per CH consecutive pixels of a row.
+// Horizontal: one thread per pixel, summing its 2r + 1 window; adjacent threads read adjacent
+// pixels, which the cache serves. (Threads that each slid a running sum over 16 pixels read pixels
+// 16 apart across a workgroup: on a Radeon 780M that pass took 9–15× as long as `box_v`.)
 @compute @workgroup_size(64, 4)
 fn box_h(@builtin(global_invocation_id) g: vec3<u32>) {
     let w = pu(0u);
     let h = pu(1u);
     let nc = pu(2u);
     let r = i32(pu(3u));
-    let ch = pu(4u);
+    let x = g.x;
     let y = g.y;
-    let x0 = g.x * ch;
-    if (y >= h || x0 >= w) {
+    if (x >= w || y >= h) {
         return;
     }
     let last = i32(w) - 1;
-    let inv = 1.0 / f32(2 * r + 1);
     let row = y * w;
-    let x1 = min(x0 + ch, w);
     var acc = vec3<f32>(0.0);
     for (var k = -r; k <= r; k++) {
-        acc += ld(row + u32(clamp(i32(x0) + k, 0, last)), nc);
-    }
-    for (var x = x0; x < x1; x++) {
-        st(row + x, nc, acc * inv);
-        acc = acc + ld(row + u32(min(i32(x) + r + 1, last)), nc) - ld(row + u32(max(i32(x) - r, 0)), nc);
+        acc += ld(row + u32(clamp(i32(x) + k, 0, last)), nc);
     }
+    st(row + x, nc, acc * (1.0 / f32(2 * r + 1)));
 }
 
 // Vertical: one thread per CH consecutive rows of a column (the CPU re-primes every BAND rows the

docs/gpu-pipeline.md

@@ -139,7 +139,8 @@ walks the drive.
 - Kernels (`crates/gpu/src/wgsl/`): `geom` (orientation pixel map, bilinear through the
   crop/straighten/flip affine or the full lens + perspective warp), `resize` (separable resample
   with the CPU's taps; Mitchell prefilter, box/bilinear for the fast guided filter), `blur` (box
-  passes with running sums over pixel chunks; three each way = the CPU's Gaussian), `map`
+  passes, three each way = the CPU's Gaussian: horizontal ones sum each pixel's window, so adjacent
+  threads read adjacent pixels; vertical ones slide running sums down chunks of a column), `map`
   (log luminance, dark channel, guided-filter steps, white balance, luminance / colour NR, airlight
   sampling), `mask` (linear / radial / luminance / colour range / brush shapes, combine, finalize),
   `finish` (the whole per-pixel stage) — each mirrors a named CPU function.
@@ -239,7 +240,7 @@ mean 0.0003 LSB. That is the same picture as on Metal, inside the bounds
 (`cargo test -p lightcraft-gpu --release --test equivalence -- --nocapture`).
 
 Remaining differences come from f32 vs f64 coordinate math, fast-math transcendental functions on
-Metal and running-sum order in the box filters — all far below one 8-bit step.
+Metal and summation order in the box filters — all far below one 8-bit step.
 
 ## Measurements
 `render_bench` on `corpus/raw/arw-sony-a7m3-compressed.arw` (24 MP; loupe = 2560 px preview →
@@ -257,6 +258,14 @@ so CPU numbers are pessimistic; the GPU numbers are less affected):
 | per-pixel stage 1920×1280, + vibrance/saturation | 53–65 ms | 5 ms |
 | export render 6000×4000 | 1037–1414 ms | 291–330 ms |
 
+On Windows (Radeon 780M iGPU, DX12), a full-size JPEG export of the same raw with heavy edits (NR 30
+/ 25, clarity, texture, dehaze, sharpening, highlights / shadows) took 3.6–4.0 s on the GPU against
+1.9–2.0 s on the CPU: the horizontal box pass had each thread slide a running sum over 16 pixels, so
+the threads of a workgroup read pixels 16 apart, and it took 9–15× as long as the vertical pass.
+Summing each pixel's window instead (adjacent threads, adjacent pixels): 1.9–2.0 s, with the GPU's
+NR stage 1.5 → 0.4 s and spatial planes 0.89 → 0.29 s (`LIGHTCRAFT_PROFILE`). The ignored test
+`bench_kernels` (`crates/gpu/src/render.rs`) times the passes on a 24 MP plane.
+
 GPU timings include the readback of the 8-bit result and the histogram. Device creation + kernel
 compilation: ~0.4 s once per process — the desktop app starts it on a background thread at launch
 (`lightcraft_gpu::warm_up`), so the first loupe render doesn't wait for it; other processes create

docs/whats-new.md

@@ -141,6 +141,11 @@
   the other editor (Edit in Other Editor, now ⌥⌘E), and where PhotoCraft is if you use a portable copy.
 - Export can write PSD too (flat, 8- or 16-bit).
 
+### Faster GPU exports
+- Exports with noise reduction, Texture, Clarity, Dehaze or Highlights and Shadows can be much faster on the GPU,
+  depending on the graphics card. On a Windows laptop with an AMD Radeon 780M, a full-size export of a 24 MP raw
+  with those edits went from about 3.7 s to 1.9 s. Photos look the same as before.
+
 ### Reliability
 - Upgrading from a Windows installer older than 0.6 (such as 0.5.0) keeps LightCraft in the folder it was installed
   in, instead of moving it back to Program Files (issue #399). If that folder has been deleted since, the upgrade