From a61a52ed8b37880244d1e62a3b3d2fd4c3b71144 Mon Sep 17 00:00:00 2001 From: THEtheChad <880190+THEtheChad@users.noreply.github.com> Date: Sun, 11 Oct 2026 08:08:14 +0000 Subject: [PATCH] perf(gpu): LR-BEHAV-GPU: horizontal box blur reads adjacent pixels The horizontal box pass gave each thread 16+ consecutive pixels to slide a running sum over, so a workgroup's threads read pixels 16 apart. On a Radeon 780M (DX12) that pass took 9-15x as long as the vertical one, and it sits under every GPU Gaussian and guided filter (NR, texture, dehaze, clarity/base). Each thread now sums one pixel's 2r+1 window: adjacent threads read adjacent pixels. Same 64x4 workgroups, parameters and bindings; box_v unchanged. Measured on the Radeon 780M (6000x4000): guided filter at sigma 1.75 828 -> 146 ms, colour-NR Gaussian 495 -> 104 ms. Full-size heavy-edit export of the CC0 Sony ARW: GPU 3.6-4.0 s -> 1.9-2.0 s (CPU 1.9-2.1 s); GPU-vs-CPU output distance unchanged. Adds gaussian_matches_cpu (GPU blur vs lightcraft_raster::blur::gaussian, 1-3 channels, partial workgroups, windows wider than the image) and per-pass rows to bench_kernels. Co-Authored-By: Claude Opus 5.5 --- crates/gpu/src/render.rs | 58 ++++++++++++++++++++++++++++++----- crates/gpu/src/wgsl/blur.wgsl | 22 ++++++------- docs/gpu-pipeline.md | 13 ++++++-- docs/whats-new.md | 5 +++ 4 files changed, 76 insertions(+), 22 deletions(-) diff --git a/crates/gpu/src/render.rs b/crates/gpu/src/render.rs index 44458350f..acafbdf5d 100644 --- a/crates/gpu/src/render.rs +++ b/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)]); @@ -1132,6 +1138,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 = (0..3).map(|c| Plane::from_fn(w, h, |x, y| v(x, y, c))).collect(); + let want: Vec = planes.iter().map(|p| blur::gaussian(p, sigma)).collect(); + for nc in [1, 2, 3] { + let words: Vec = (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 = 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] @@ -1139,6 +1172,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; @@ -1152,9 +1186,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); diff --git a/crates/gpu/src/wgsl/blur.wgsl b/crates/gpu/src/wgsl/blur.wgsl index b0e28370d..375fb6d55 100644 --- a/crates/gpu/src/wgsl/blur.wgsl +++ b/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 { @@ -25,31 +25,27 @@ fn st(i: u32, nc: u32, v: vec3) { } } -// 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) { 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(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 diff --git a/docs/gpu-pipeline.md b/docs/gpu-pipeline.md index 883c3a0bf..edfdf55d6 100644 --- a/docs/gpu-pipeline.md +++ b/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. @@ -233,7 +234,7 @@ noise-reduction size come from the whole frame, and a window that spots or Auto cut back after the readback. The tests skip (pass with a note) when no adapter exists. 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 → @@ -251,6 +252,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 diff --git a/docs/whats-new.md b/docs/whats-new.md index fb8b285a3..6f93109dd 100644 --- a/docs/whats-new.md +++ b/docs/whats-new.md @@ -41,6 +41,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