M5.2e: renders and exports use the GPU; render_bench GPU column

`media::develop` renders on the GPU when available (views, render_now, exports;
thumbnails stay on the CPU) and falls back to the CPU otherwise. `app.gpu
{enabled?}` toggles it at runtime, `LIGHTCRAFT_GPU=0` disables it for the process,
`ui.inspect` reports the adapter under perf.gpu.

Performance: dropped device buffers are recycled once their thread has submitted
(fresh 100-300 MB allocations per pass cost as much as the passes), Gaussian
passes ping-pong between two buffers and read all channels of a pixel at once,
and long box radii slide over longer chunks. LIGHTCRAFT_PROFILE prints real GPU
stage timings. render_bench prints CPU and GPU columns and the CPU/GPU difference.

24 MP ARW, loaded machine, min of 7: cold loupe 415 -> 32 ms, exposure drag
22-42 -> 3.9 ms, NR drag (1920) 301+ -> 18 ms, export render 1037+ -> 291 ms;
max 1 LSB difference. docs/gpu-pipeline.md describes the implementation.

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

Cargo.lock

@@ -2085,6 +2085,7 @@ dependencies = [
  "lightcraft-color",
  "lightcraft-develop",
  "lightcraft-geom",
+ "lightcraft-gpu",
  "lightcraft-meta",
  "lightcraft-pipeline",
  "lightcraft-preview",

crates/engine/Cargo.toml

@@ -15,6 +15,7 @@ lightcraft-raster = { workspace = true }
 lightcraft-develop = { workspace = true }
 lightcraft-scenes = { workspace = true }
 lightcraft-pipeline = { workspace = true }
+lightcraft-gpu = { workspace = true }
 lightcraft-catalog = { workspace = true }
 lightcraft-preview = { workspace = true }
 lightcraft-codecs = { workspace = true }

crates/engine/examples/render_bench.rs

@@ -8,11 +8,15 @@
 //! and a full-size render + JPEG encode (export); `ONLY=batch` with several files times a full-size
 //! export of each (decode + render + encode). Prints minimum wall-clock and minimum process CPU
 //! time: on a shared machine the CPU time shows the work done, wall-clock also the wait for cores.
+//! Each scenario runs on the CPU pipeline and, when a GPU adapter exists, on `lightcraft-gpu`
+//! (second column; `LIGHTCRAFT_GPU=0` for the CPU only).
+use std::sync::Arc;
 use std::time::Instant;
 
 use lightcraft_develop::DevelopSettings;
 use lightcraft_engine::export::{ExportOptions, encode_image};
-use lightcraft_pipeline::{Quality, RenderRequest, SourceInfo, StageCache, render, render_cached};
+use lightcraft_engine::media::develop;
+use lightcraft_pipeline::{Quality, RenderRequest, SourceInfo, StageCache};
 use lightcraft_raster::Rgb32f;
 use lightcraft_raster::resample::{Filter, fit};
 
@@ -66,6 +70,19 @@ fn typical() -> DevelopSettings {
     s
 }
 
+/// Max and mean |Δ| (8-bit, RGB) between two renders.
+fn diff(a: &lightcraft_raster::Rgba8, b: &lightcraft_raster::Rgba8) -> (u8, f64) {
+    let (mut max, mut sum) = (0u8, 0u64);
+    for (p, q) in a.data.iter().zip(&b.data) {
+        for c in 0..3 {
+            let d = p[c].abs_diff(q[c]);
+            max = max.max(d);
+            sum += d as u64;
+        }
+    }
+    (max, sum as f64 / (a.data.len() * 3).max(1) as f64)
+}
+
 fn main() {
     let args: Vec<String> = std::env::args().skip(1).collect();
     let n: usize = std::env::var("N").ok().and_then(|v| v.parse().ok()).unwrap_or(5);
@@ -87,50 +104,76 @@ fn main() {
             SourceInfo { raw: true, ..Default::default() },
         ),
     };
+    let full = Arc::new(full);
     println!("source {}×{} ({:.1} MP), {n} runs each", full.width, full.height, (full.width * full.height) as f64 / 1e6);
-    let preview = std::sync::Arc::new(fit(&full, 2560, 2560, Filter::Box));
+    // GPU column: `LIGHTCRAFT_GPU=0` (or no adapter) prints only the CPU column.
+    let t = Instant::now();
+    let gpu = lightcraft_engine::gpu::available();
+    match lightcraft_engine::gpu::adapter_name().filter(|_| gpu) {
+        Some(a) => println!("gpu: {a} (device + kernels: {:.0} ms)", t.elapsed().as_secs_f64() * 1e3),
+        None => println!("gpu: none"),
+    }
+    let backends: &[bool] = if gpu { &[false, true] } else { &[false] };
+    let row = |name: &str, f: &mut dyn FnMut(bool)| {
+        let cols: Vec<String> = backends.iter().map(|&g| format!("{}", best(n, || f(g)))).collect();
+        println!("{name:<38} cpu: {}{}", cols[0], cols.get(1).map(|g| format!("  |  gpu: {g}")).unwrap_or_default());
+    };
+    let preview = Arc::new(fit(&full, 2560, 2560, Filter::Box));
     let s = typical();
     let view = RenderRequest::fit(1920, 1280);
     let draft = RenderRequest { max_w: 1152, max_h: 768, quality: Quality::Draft, apply_crop: true };
+    if gpu {
+        let c = develop(&preview, &info, &s, &view, None, false).image;
+        let g = develop(&preview, &info, &s, &view, None, true).image;
+        let (max, mean) = diff(&c, &g);
+        println!("gpu vs cpu, loupe 1920×1280 typical: max {max} LSB, mean {mean:.4} LSB");
+    }
 
     if run("cold") {
-        println!("loupe 1920×1280 cold:                 {}", best(n, || drop(render(&preview, &info, &s, &view))));
-        println!("loupe draft 1152×768 cold:            {}", best(n, || drop(render(&preview, &info, &s, &draft))));
+        // a new photo: no cached stages (the GPU also uploads the preview)
+        row("loupe 1920×1280 cold:", &mut |g| drop(develop(&preview, &info, &s, &view, None, g)));
+        row("loupe draft 1152×768 cold:", &mut |g| drop(develop(&preview, &info, &s, &draft, None, g)));
     }
-    let cache = StageCache::default();
+    let caches = [StageCache::default(), StageCache::default()];
     let mut k = 0.0;
     if run("drag") {
-        drop(render_cached(&preview, &info, &s, &view, &cache));
-        let ms = best(n, || {
+        for g in backends {
+            drop(develop(&preview, &info, &s, &view, Some(&caches[*g as usize]), *g));
+        }
+        row("loupe 1920×1280 exposure drag (warm):", &mut |g| {
             let mut t = s.clone();
             k += 1.0;
-            t.light.contrast = k;
+            t.light.contrast = k % 50.0;
             t.light.exposure = 0.3 + k * 0.01;
-            drop(render_cached(&preview, &info, &t, &view, &cache));
+            drop(develop(&preview, &info, &t, &view, Some(&caches[g as usize]), g));
         });
-        println!("loupe 1920×1280 exposure drag (warm): {ms}");
-        drop(render_cached(&preview, &info, &s, &draft, &cache));
-        let ms = best(n, || {
+        for g in backends {
+            drop(develop(&preview, &info, &s, &draft, Some(&caches[*g as usize]), *g));
+        }
+        row("loupe draft highlights drag (warm):", &mut |g| {
             let mut t = s.clone();
             k += 1.0;
-            t.light.highlights = -40.0 + k;
-            drop(render_cached(&preview, &info, &t, &draft, &cache));
+            t.light.highlights = -40.0 + k % 50.0;
+            drop(develop(&preview, &info, &t, &draft, Some(&caches[g as usize]), g));
         });
-        println!("loupe draft highlights drag (warm):   {ms}");
-        let ms = best(n, || {
+        row("loupe draft clarity drag (warm):", &mut |g| {
             let mut t = s.clone();
             k += 1.0;
-            t.effects.clarity = 15.0 + k;
-            drop(render_cached(&preview, &info, &t, &draft, &cache));
+            t.effects.clarity = 15.0 + k % 50.0;
+            drop(develop(&preview, &info, &t, &draft, Some(&caches[g as usize]), g));
         });
-        println!("loupe draft clarity drag (warm):      {ms}");
-        let ms = best(n, || {
+        row("loupe draft NR drag (warm):", &mut |g| {
             let mut t = s.clone();
             k += 1.0;
-            t.detail.nr_luminance = 30.0 + k;
-            drop(render_cached(&preview, &info, &t, &draft, &cache));
+            t.detail.nr_luminance = 30.0 + k % 50.0;
+            drop(develop(&preview, &info, &t, &draft, Some(&caches[g as usize]), g));
+        });
+        row("loupe 1920×1280 NR drag (warm):", &mut |g| {
+            let mut t = s.clone();
+            k += 1.0;
+            t.detail.nr_luminance = 30.0 + k % 50.0;
+            drop(develop(&preview, &info, &t, &view, Some(&caches[g as usize]), g));
         });
-        println!("loupe draft NR drag (warm):           {ms}");
     }
     if run("finish") {
         // per-pixel stage cost by feature (warm cache: exposure changes only)
@@ -154,13 +197,14 @@ fn main() {
         for (name, f) in variants {
             let mut t = s.clone();
             f(&mut t);
-            drop(render_cached(&preview, &info, &t, &view, &cache));
-            let ms = best(n, || {
+            for g in backends {
+                drop(develop(&preview, &info, &t, &view, Some(&caches[*g as usize]), *g));
+            }
+            row(&format!("finish 1920×1280 {name}"), &mut |g| {
                 k += 1.0;
                 t.light.exposure = 0.3 + k * 0.001;
-                drop(render_cached(&preview, &info, &t, &view, &cache));
+                drop(develop(&preview, &info, &t, &view, Some(&caches[g as usize]), g));
             });
-            println!("finish 1920×1280 {name:<22} {ms}");
         }
     }
     if run("batch") && args.len() > 1 {
@@ -171,22 +215,39 @@ fn main() {
         let ids: Vec<_> = r["imported"].as_array().expect("ids").iter().filter_map(|v| v.as_u64()).map(lightcraft_engine::catalog::PhotoId).collect();
         let items: Vec<_> = ids.iter().enumerate().map(|(i, id)| (*id, i + 1)).collect();
         let o = ExportOptions::default();
-        let ms = best(n.min(3), || {
+        row(&format!("batch export of {} files (decode+render+JPEG):", items.len()), &mut |g| {
+            lightcraft_engine::gpu::set_enabled(g);
             for &(id, seq) in &items {
                 drop(export_photo(&mut session, id, &o, seq).expect("export"));
                 session.media.forget(id); // as in a batch: every original is decoded once
             }
         });
-        println!("batch export of {} files (decode+render+JPEG): {ms}", items.len());
+        lightcraft_engine::gpu::set_enabled(true);
         return;
     }
     if run("export") {
         let big = RenderRequest::fit(full.width, full.height);
         let ne = n.min(3);
         let mut img = None;
-        let ms = best(ne, || img = Some(render(&full, &info, &s, &big).image));
-        println!("export render {}×{}:            {ms}", full.width, full.height);
+        let cols: Vec<String> = backends
+            .iter()
+            .map(|&g| {
+                let ms = best(ne, || img = Some(develop(&full, &info, &s, &big, None, g).image));
+                format!("{ms}")
+            })
+            .collect();
+        println!(
+            "{:<38} cpu: {}{}",
+            format!("export render {}×{}:", full.width, full.height),
+            cols[0],
+            cols.get(1).map(|g| format!("  |  gpu: {g}")).unwrap_or_default()
+        );
+        if gpu {
+            let c = develop(&full, &info, &s, &big, None, false).image;
+            let (max, mean) = diff(&c, img.as_ref().expect("rendered"));
+            println!("gpu vs cpu, export {}×{} typical: max {max} LSB, mean {mean:.4} LSB", full.width, full.height);
+        }
         let img = img.expect("rendered");
-        println!("export JPEG encode:                   {}", best(ne, || drop(encode_image(&img, &ExportOptions::default()))));
+        println!("export JPEG encode:                    {}", best(ne, || drop(encode_image(&img, &ExportOptions::default()))));
     }
 }

crates/engine/src/cmd/query.rs

@@ -129,6 +129,12 @@ pub fn specs() -> Vec<CommandSpec> {
                 "versions": ph.versions.iter().map(|v| json!({"name": v.name, "created": v.created})).collect::<Vec<_>>(),
             }))
         }),
+        cmd!(query "app.gpu", "GPU Rendering", [], None, "{enabled?: bool} — allow/forbid GPU rendering (CPU fallback; LIGHTCRAFT_GPU=0 forbids it for the process)", always, |_, p| {
+            if let Some(on) = p.get("enabled").and_then(Value::as_bool) {
+                lightcraft_gpu::set_enabled(on);
+            }
+            Ok(json!({"enabled": lightcraft_gpu::enabled(), "available": lightcraft_gpu::available(), "adapter": lightcraft_gpu::adapter_name()}))
+        }),
         cmd!(query "journal.list", "Command Journal", [], None, "{limit?}", always, |s, p| {
             let lim = p.get("limit").and_then(Value::as_u64).unwrap_or(100) as usize;
             let start = s.journal.len().saturating_sub(lim);

crates/engine/src/lib.rs

@@ -30,7 +30,7 @@ use lightcraft_develop::DevelopSettings;
 pub use media::{RenderJob, SourceLevel};
 use serde_json::Value;
 pub use view::{LibrarySource, Selection};
-pub use {lightcraft_catalog as catalog, lightcraft_develop as develop, lightcraft_pipeline as pipeline};
+pub use {lightcraft_catalog as catalog, lightcraft_develop as develop, lightcraft_gpu as gpu, lightcraft_pipeline as pipeline};
 
 #[derive(Debug, thiserror::Error)]
 pub enum EngineError {

crates/engine/src/media.rs

@@ -236,10 +236,10 @@ impl RenderJob {
         let was_loaded = matches!(self.source, SourceRef::Loaded(_));
         match self.source.load() {
             Ok(src) => {
-                let rendered = match &self.stages {
-                    Some(st) => lightcraft_pipeline::render_cached(&src, &self.info, &self.settings, &self.request, st),
-                    None => lightcraft_pipeline::render(&src, &self.info, &self.settings, &self.request),
-                };
+                // Thumbnails (many small jobs side by side) stay on the CPU; views and exports use
+                // the GPU when there is one.
+                let gpu = self.cache.is_none();
+                let rendered = develop(&src, &self.info, &self.settings, &self.request, self.stages.as_deref(), gpu);
                 if let Some((cache, key)) = &self.cache {
                     cache.put(*key, Arc::new(rendered.image.clone()));
                 }
@@ -250,6 +250,19 @@ impl RenderJob {
     }
 }
 
+/// Render `src`: on the GPU when `gpu` is set and a GPU is available (`lightcraft_gpu`), else on
+/// the CPU, reusing `stages` either way. Both produce the same image within 1–3 LSB (see
+/// `docs/gpu-pipeline.md`); `LIGHTCRAFT_GPU=0` or [`lightcraft_gpu::set_enabled`] forces the CPU.
+pub fn develop(src: &Arc<Rgb32f>, info: &SourceInfo, s: &DevelopSettings, req: &RenderRequest, stages: Option<&StageCache>, gpu: bool) -> Rendered {
+    if gpu && let Some(r) = lightcraft_gpu::render(src, info, s, req, stages) {
+        return r;
+    }
+    match stages {
+        Some(st) => lightcraft_pipeline::render_cached(src, info, s, req, st),
+        None => lightcraft_pipeline::render(src, info, s, req),
+    }
+}
+
 /// What identifies a photo's pixels for caching: its content hash, else its source.
 pub fn content_key(p: &Photo) -> String {
     match (&p.content_hash, &p.source) {

crates/gpu/src/ctx.rs

@@ -58,17 +58,46 @@ struct Kernel {
 }
 
 /// A device buffer of `len` 32-bit values.
+///
+/// Dropped buffers are recycled: they first wait in a per-thread list (commands recorded on this
+/// thread may still use them) and become reusable once this thread submits ([`Gpu::submitted`]),
+/// as the queue then runs every earlier use before any later one. Reuse matters: allocating (and
+/// zero-filling) a fresh 100–300 MB buffer per pass costs as much as the pass itself.
 pub struct Buf {
-    pub buf: wgpu::Buffer,
+    buf: Option<wgpu::Buffer>,
     pub len: usize,
 }
 
+impl Buf {
+    pub fn raw(&self) -> &wgpu::Buffer {
+        self.buf.as_ref().expect("live buffer")
+    }
+}
+
+impl Drop for Buf {
+    fn drop(&mut self) {
+        if let Some(b) = self.buf.take() {
+            RETIRED.with(|r| r.borrow_mut().push(b));
+        }
+    }
+}
+
+thread_local! {
+    /// Buffers dropped on this thread since its last submit.
+    static RETIRED: std::cell::RefCell<Vec<wgpu::Buffer>> = const { std::cell::RefCell::new(Vec::new()) };
+}
+
+/// Most bytes kept in the free pool.
+const POOL_BYTES: u64 = 2 << 30;
+
 pub struct Gpu {
     pub device: wgpu::Device,
     pub queue: wgpu::Queue,
     pub info: wgpu::AdapterInfo,
     kernels: HashMap<&'static str, Kernel>,
     dummy: wgpu::Buffer,
+    /// Recycled buffers (see [`Buf`]).
+    free: std::sync::Mutex<Vec<wgpu::Buffer>>,
     /// Largest storage buffer the device accepts (bytes).
     pub max_buffer: u64,
 }
@@ -180,19 +209,37 @@ impl Gpu {
         });
         let max_buffer = limits.max_storage_buffer_binding_size.min(limits.max_buffer_size);
         log::info!("gpu: {} ({:?})", info.name, info.backend);
-        Some(Gpu { device, queue, info, kernels, dummy, max_buffer })
+        Some(Gpu { device, queue, info, kernels, dummy, free: Default::default(), max_buffer })
     }
 
-    /// A zero-initialized buffer of `len` 32-bit values.
+    /// A buffer of `len` 32-bit values with undefined contents (a recycled one when possible).
     pub fn buffer(&self, len: usize) -> Buf {
         let size = (len.max(1) * 4) as u64;
-        let buf = self.device.create_buffer(&wgpu::BufferDescriptor {
-            label: None,
-            size,
-            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::COPY_DST,
-            mapped_at_creation: false,
+        let recycled = {
+            let mut f = self.free.lock().unwrap_or_else(|e| e.into_inner());
+            f.iter().position(|b| b.size() == size).map(|i| f.swap_remove(i))
+        };
+        let buf = recycled.unwrap_or_else(|| {
+            self.device.create_buffer(&wgpu::BufferDescriptor {
+                label: None,
+                size,
+                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::COPY_DST,
+                mapped_at_creation: false,
+            })
         });
-        Buf { buf, len }
+        Buf { buf: Some(buf), len }
+    }
+
+    /// Submit command buffers, then make the buffers this thread dropped before reusable.
+    pub fn submit(&self, cmds: impl IntoIterator<Item = wgpu::CommandBuffer>) {
+        self.queue.submit(cmds);
+        let retired = RETIRED.with(|r| std::mem::take(&mut *r.borrow_mut()));
+        let mut f = self.free.lock().unwrap_or_else(|e| e.into_inner());
+        f.extend(retired);
+        let mut total: u64 = f.iter().map(|b| b.size()).sum();
+        while total > POOL_BYTES && !f.is_empty() {
+            total -= f.remove(0).size();
+        }
     }
 
     /// Upload 32-bit values.
@@ -207,7 +254,7 @@ impl Gpu {
             contents: bytes,
             usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::COPY_DST,
         });
-        Buf { buf, len }
+        Buf { buf: Some(buf), len }
     }
 
     /// Whether a buffer of `len` 32-bit values fits the device limits.
@@ -230,7 +277,7 @@ impl Gpu {
         });
         let mut entries = vec![wgpu::BindGroupEntry { binding: 0, resource: params.as_entire_binding() }];
         for (i, b) in bufs.iter().enumerate() {
-            let buf = b.map(|b| &b.buf).unwrap_or(&self.dummy);
+            let buf = b.map(|b| b.raw()).unwrap_or(&self.dummy);
             entries.push(wgpu::BindGroupEntry { binding: i as u32 + 1, resource: buf.as_entire_binding() });
         }
         let bg = self.device.create_bind_group(&wgpu::BindGroupDescriptor { label: None, layout: &k.layout, entries: &entries });
@@ -250,8 +297,8 @@ impl Gpu {
             usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
             mapped_at_creation: false,
         });
-        enc.copy_buffer_to_buffer(&src.buf, 0, &staging, 0, size);
-        self.queue.submit([enc.finish()]);
+        enc.copy_buffer_to_buffer(src.raw(), 0, &staging, 0, size);
+        self.submit([enc.finish()]);
         let slice = staging.slice(..size);
         let (tx, rx) = std::sync::mpsc::channel();
         slice.map_async(wgpu::MapMode::Read, move |r| {

crates/gpu/src/lib.rs

@@ -56,7 +56,7 @@ fn env_disabled() -> bool {
 }
 
 #[cfg(not(target_arch = "wasm32"))]
-fn device() -> Option<&'static ctx::Gpu> {
+pub(crate) fn device() -> Option<&'static ctx::Gpu> {
     static GPU: std::sync::OnceLock<Option<ctx::Gpu>> = std::sync::OnceLock::new();
     if env_disabled() {
         return None;

crates/gpu/src/render.rs

@@ -111,11 +111,27 @@ impl<'a> Cx<'a> {
         self.gpu.finish_and_read(enc, b, len)
     }
 
+    /// Submit what is recorded and wait for it (profiling only: stage timings).
+    pub fn sync(&mut self) {
+        if let Some(enc) = self.enc.take() {
+            self.gpu.submit([enc.finish()]);
+        }
+        let _ = self.gpu.device.poll(wgpu::PollType::wait_indefinitely());
+    }
+
+    /// A zero-filled buffer (the clear is recorded).
+    pub fn zeroed(&mut self, len: usize) -> Buf {
+        let out = self.gpu.buffer(len);
+        let enc = self.enc.get_or_insert_with(|| self.gpu.encoder());
+        enc.clear_buffer(out.raw(), 0, None);
+        out
+    }
+
     /// A copy of `b` (recorded).
     pub fn copy(&mut self, b: &Buf) -> Buf {
         let out = self.gpu.buffer(b.len);
         let enc = self.enc.get_or_insert_with(|| self.gpu.encoder());
-        enc.copy_buffer_to_buffer(&b.buf, 0, &out.buf, 0, (b.len * 4) as u64);
+        enc.copy_buffer_to_buffer(b.raw(), 0, out.raw(), 0, (b.len * 4) as u64);
         out
     }
 
@@ -135,7 +151,8 @@ 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.
+/// Pixels each blur 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;
 
 /// Gaussian blur of a `w × h` image of `nc` interleaved channels: three box passes each way
@@ -145,19 +162,22 @@ 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 mut cur: Option<Buf> = None;
-    let passes = radii.iter().map(|r| ("box_h", *r)).chain(radii.iter().map(|r| ("box_v", *r)));
-    for (k, r) in passes {
-        if r == 0 {
-            continue;
-        }
-        let out = cx.gpu.buffer(w * h * nc);
-        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];
-        cx.run(k, &p, &[Some(cur.as_ref().unwrap_or(src)), Some(&out)], groups);
-        cur = Some(out);
+    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();
+    if passes.is_empty() {
+        return cx.copy(src);
     }
-    cur.unwrap_or_else(|| 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];
+        let (out, prev) = (&bufs[i % 2], &bufs[(i + 1) % 2]);
+        cx.run(k, &p, &[Some(if i == 0 { src } else { prev }), Some(out)], groups);
+    }
+    let last = (passes.len() - 1) % 2;
+    let [a, b] = std::mem::replace(&mut bufs, [cx.gpu.buffer(0), cx.gpu.buffer(0)]);
+    if last == 0 { a } else { b }
 }
 
 /// Resample taps of `lightcraft_raster::resample::resize`, packed for the `resize_*` kernels.
@@ -377,9 +397,11 @@ pub fn render(
     stages: Option<&GpuStages>,
 ) -> Option<Rendered> {
     let mut t = profiling().then(std::time::Instant::now);
-    let lap = |what: &str, t: &mut Option<std::time::Instant>| {
+    // Under `LIGHTCRAFT_PROFILE` each stage is submitted and waited for, so the timings are real.
+    let lap = |what: &str, t: &mut Option<std::time::Instant>, cx: &mut Cx<'_>| {
         if let Some(t) = t {
-            eprintln!("  gpu {what}: {:.1} ms (recorded)", t.elapsed().as_secs_f64() * 1e3);
+            cx.sync();
+            eprintln!("  gpu {what}: {:.1} ms", t.elapsed().as_secs_f64() * 1e3);
             *t = std::time::Instant::now();
         }
     };
@@ -393,6 +415,7 @@ pub fn render(
     let cached = stages.and_then(|c| c.get(src, plan.geo));
     let mut host = Host::default();
     let mut cx = Cx::new(gpu);
+    lap("plan", &mut t, &mut cx);
 
     // 1. geometry (on the CPU when the source exceeds the device's buffer limit)
     let sampled = match &cached {
@@ -412,14 +435,14 @@ pub fn render(
             b
         }
     };
-    lap("sample", &mut t);
+    lap("sample", &mut t, &mut cx);
 
     // 2. white balance, defringe, spots, noise reduction
     let lin = match cached.as_ref().and_then(|e| e.lin.clone()).filter(|(k, _)| *k == plan.lin_key) {
         Some((_, b)) => b,
         None => Arc::new(linear(&mut cx, &sampled, info, &plan, &mut host)),
     };
-    lap("wb/nr", &mut t);
+    lap("wb/nr", &mut t, &mut cx);
 
     // 3. spatial planes
     let mut planes = match cached.map(|e| e.planes) {
@@ -427,14 +450,14 @@ pub fn render(
         _ => Planes { key: plan.lin_key, ..Default::default() },
     };
     let prep = prepare(&mut cx, &lin, &plan, req, &mut planes);
-    lap("planes", &mut t);
+    lap("planes", &mut t, &mut cx);
     if let Some(c) = stages {
         c.put(Entry { src: src.clone(), geo: plan.geo, sampled, lin: Some((plan.lin_key, lin.clone())), planes });
     }
 
     // 4. masks
     let (masks, terms) = masks(&mut cx, &lin, &prep, &plan, &mut host);
-    lap("masks", &mut t);
+    lap("masks", &mut t, &mut cx);
 
     // 5. per-pixel stage
     let fp = FinishParams::new(s, &plan.frame, info, w, h, plan.px_per_long, prep.air);
@@ -460,8 +483,9 @@ pub fn render(
     );
     let data: Vec<[u8; 4]> = cx.read(&out, n);
     let image = Rgba8 { width: w, height: h, data };
-    lap("finish + readback", &mut t);
+    lap("finish + readback", &mut t, &mut cx);
     let histogram = Histogram::of_srgb8(&image);
+    lap("histogram", &mut t, &mut cx);
     Some(Rendered { image, histogram })
 }
 
@@ -598,7 +622,7 @@ fn masks(cx: &mut Cx<'_>, lin: &Buf, prep: &Prep, plan: &Plan<'_>, host: &mut Ho
     let [a, b, c, d, e, f] = frame.out_to_norm(w, h).0;
     let (kx, ky) = (frame.ow / long, frame.oh / long);
     let aff = [a * kx, b * ky, c * kx, d * ky, e * kx, f * ky].map(|v| (v as f32).to_bits());
-    let alpha = cx.gpu.buffer(n * list.len());
+    let alpha = cx.zeroed(n * list.len());
     let comp_plane = cx.gpu.buffer(n);
     for (mi, m) in list.iter().enumerate() {
         let mut first = true;
@@ -701,6 +725,35 @@ fn masks(cx: &mut Cx<'_>, lin: &Buf, prep: &Prep, plan: &Plan<'_>, host: &mut Ho
 mod tests {
     use super::*;
 
+    /// Kernel timings on a 24 MP plane: `cargo test --release -p lightcraft-gpu -- --ignored --nocapture`.
+    #[test]
+    #[ignore]
+    fn bench_kernels() {
+        let Some(gpu) = crate::device() else { return };
+        let (w, h) = (6000, 4000);
+        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;
+            for _ in 0..5 {
+                let mut cx = Cx::new(gpu);
+                cx.sync();
+                let t = std::time::Instant::now();
+                f(&mut cx);
+                cx.sync();
+                best = best.min(t.elapsed().as_secs_f64() * 1e3);
+            }
+            eprintln!("{name:<40} {best:.1} ms");
+        };
+        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("map log_lum", &mut |cx| {
+            let o = cx.gpu.buffer(w * h);
+            map(cx, "log_lum_k", w * h, &[], [Some(&src), None, None], &o);
+        });
+        time("guided sigma=1.75", &mut |cx| drop(guided(cx, &src, w, h, 1.75, 0.01)));
+    }
+
     #[test]
     fn orient_map_matches_image_oriented() {
         use Orientation::*;

crates/gpu/src/wgsl/blur.wgsl

@@ -2,6 +2,29 @@
 // Gaussian (`lightcraft_raster::blur::gaussian`). Each thread slides a running sum over CH pixels.
 // P: w, h, nc, r, ch. Bindings: src, dst.
 
+// Pixel `i` (all nc ≤ 3 channels at once: one contiguous read per pixel).
+fn ld(i: u32, nc: u32) -> vec3<f32> {
+    let j = i * nc;
+    if (nc == 1u) {
+        return vec3<f32>(src[j], 0.0, 0.0);
+    }
+    if (nc == 2u) {
+        return vec3<f32>(src[j], src[j + 1u], 0.0);
+    }
+    return vec3<f32>(src[j], src[j + 1u], src[j + 2u]);
+}
+
+fn st(i: u32, nc: u32, v: vec3<f32>) {
+    let j = i * nc;
+    dst[j] = v.x;
+    if (nc > 1u) {
+        dst[j + 1u] = v.y;
+    }
+    if (nc > 2u) {
+        dst[j + 2u] = v.z;
+    }
+}
+
 // Horizontal: one thread per CH consecutive pixels of a row.
 @compute @workgroup_size(64, 4)
 fn box_h(@builtin(global_invocation_id) g: vec3<u32>) {
@@ -19,15 +42,13 @@ fn box_h(@builtin(global_invocation_id) g: vec3<u32>) {
     let inv = 1.0 / f32(2 * r + 1);
     let row = y * w;
     let x1 = min(x0 + ch, w);
-    for (var c = 0u; c < nc; c++) {
-        var acc = 0.0;
-        for (var k = -r; k <= r; k++) {
-            acc += src[(row + u32(clamp(i32(x0) + k, 0, last))) * nc + c];
-        }
-        for (var x = x0; x < x1; x++) {
-            dst[(row + x) * nc + c] = acc * inv;
-            acc = acc + src[(row + u32(min(i32(x) + r + 1, last))) * nc + c] - src[(row + u32(max(i32(x) - r, 0))) * nc + c];
-        }
+    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);
     }
 }
 
@@ -48,14 +69,12 @@ fn box_v(@builtin(global_invocation_id) g: vec3<u32>) {
     let last = i32(h) - 1;
     let inv = 1.0 / f32(2 * r + 1);
     let y1 = min(y0 + ch, h);
-    for (var c = 0u; c < nc; c++) {
-        var acc = 0.0;
-        for (var k = -r; k <= r; k++) {
-            acc += src[(u32(clamp(i32(y0) + k, 0, last)) * w + x) * nc + c];
-        }
-        for (var y = y0; y < y1; y++) {
-            dst[(y * w + x) * nc + c] = acc * inv;
-            acc = acc + src[(u32(min(i32(y) + r + 1, last)) * w + x) * nc + c] - src[(u32(max(i32(y) - r, 0)) * w + x) * nc + c];
-        }
+    var acc = vec3<f32>(0.0);
+    for (var k = -r; k <= r; k++) {
+        acc += ld(u32(clamp(i32(y0) + k, 0, last)) * w + x, nc);
+    }
+    for (var y = y0; y < y1; y++) {
+        st(y * w + x, nc, acc * inv);
+        acc = acc + ld(u32(min(i32(y) + r + 1, last)) * w + x, nc) - ld(u32(max(i32(y) - r, 0)) * w + x, nc);
     }
 }

crates/ui-egui/src/control.rs

@@ -77,7 +77,7 @@ pub fn inspect(app: &LightcraftApp, ctx: &egui::Context) -> Value {
         "selection": app.session.selection.ids.iter().map(|p| p.0).collect::<Vec<_>>(),
         "activeMask": app.session.active_mask,
         "widgetCount": app.widgets.len(),
-        "perf": {"frameMs": app.perf.frame_ms, "fps": app.perf.fps, "lastRenderMs": app.renderer.last_main_ms, "renderQueue": app.renderer.queued(), "rendersInFlight": app.renderer.in_flight(), "rendersDone": app.renderer.completed, "thumbTextures": app.renderer.thumb_textures()},
+        "perf": {"frameMs": app.perf.frame_ms, "fps": app.perf.fps, "lastRenderMs": app.renderer.last_main_ms, "renderQueue": app.renderer.queued(), "rendersInFlight": app.renderer.in_flight(), "rendersDone": app.renderer.completed, "thumbTextures": app.renderer.thumb_textures(), "gpu": lightcraft_engine::gpu::available().then(lightcraft_engine::gpu::adapter_name).flatten()},
         "status": app.ui.status,
     })
 }

docs/gpu-pipeline.md

@@ -1,67 +1,88 @@
-# GPU develop pipeline — design note (M5.2, not implemented yet)
+# GPU develop pipeline (M5.2)
 
-The CPU pipeline (`crates/pipeline`) is the reference ("oracle"). A `lightcraft-gpu` crate (L3) will
-evaluate the same stages with wgpu compute shaders (WGSL), on native (Metal/Vulkan/DX12) and on the
-web (WebGPU). This note records the plan so the CPU work in M5 stays compatible with it.
+The CPU pipeline (`crates/pipeline`) is the reference ("oracle"). `lightcraft-gpu` (`crates/gpu`, L3)
+evaluates the same stages with wgpu compute shaders (WGSL) on Metal / Vulkan / DX12. Everything is
+pure Rust (wgpu, naga); the drivers are the system's. No GL backend is compiled in.
 
-## Why, and what the CPU path already does
-- CPU numbers (24 MP source, 2560 px preview, ~2.5 MP loupe; see `crates/engine/examples/render_bench.rs`):
-  the per-pixel stage (`finish`) costs ~50 ns/px of CPU time with typical edits (+~65 ns/px with
-  vibrance/saturation: the OkLCh round trip), noise reduction + spatial planes another ~80 ns/px
-  when they have to be recomputed. With all cores and the stage cache this meets the CPU budgets
-  for tone/colour drags (~7 ms draft, ~18 ms full at 2.5 MP); a GPU makes cold full-quality
-  renders, NR drags, 4–8K displays and 1:1 zoom of 60 MP files interactive, and frees the CPU for
-  previews/import.
-- The CPU path is already structured the way the GPU path wants it:
-  - **Stage cache** (`StageCache`): resampled source → white-balanced/denoised image → spatial
-    planes (log luminance, guided base, clarity/texture bands, dark channel) → per-pixel stage. Each
-    stage is keyed by the settings that feed it. On the GPU the same keys decide which textures are
-    still valid; a tone/colour/exposure drag re-runs one full-screen compute pass.
-  - **Exposure after the spatial filters** (shift-equivariant on log luminance), so exposure is a
-    per-pixel uniform, not a cache invalidation.
-  - **Draft quality** during drags (smaller output, capped base radius), full on release.
+## Where it is used
+- `lightcraft_engine::media::develop` (called by every `RenderJob`): loupe / before / compare views,
+  `render_now` (CLI, MCP, control channel renders) and exports render on the GPU when one is
+  available; grid/filmstrip thumbnails (many small jobs in parallel) stay on the CPU.
+- Anything the GPU path cannot do returns `None` and the CPU renders instead: no adapter (CI
+  machines, software-only adapters), `LIGHTCRAFT_GPU=0` (whole process), the `app.gpu {enabled}`
+  command (runtime preference; `ui.inspect` → `perf.gpu` shows the adapter), an output larger than
+  the device's storage-buffer limit, or a device error / panic during a render (the process then
+  stays on the CPU).
+- wasm32: the crate compiles to the CPU fallback (WebGPU needs asynchronous device creation and a
+  device shared with the canvas — open item below).
 
-## Architecture
-- `lightcraft-gpu` owns a `wgpu::Device` (shared with eframe's when running in the desktop app, so
-  the loupe texture never leaves the GPU; headless device for CLI/MCP/export).
-- Resident per-photo state: the source pyramid as `rgba16float` textures (levels /2, /4, … down to
-  the thumbnail), uploaded once per photo; the stage textures mirror `StageCache` entries
-  (`r32float` planes, `rgba16float` images).
-- Kernels (one WGSL module each, workgroups of 16×16, storage textures):
-  1. `sample`: orientation + crop/rotate + flips from the nearest pyramid level (bilinear on the
-     level ≥ output scale — matches the CPU's area-prefiltered resample within tolerance).
-  2. `wb`: 3×3 matrix (+ spots as instanced quads: copy/blend from the source offset).
-  3. `nr`: guided filter on log luminance (box sums via two-pass separable running sums in shared
-     memory, or a summed-area table for large radii) + chroma Gaussian.
-  4. `planes`: log luminance; guided base / clarity (fast guided filter: coefficients on a /s
-     subsampled grid, bilinear upsample — exactly the CPU `guided_fast` structure); texture
-     Gaussian (3 box passes); dark channel Gaussian + airlight (parallel reduction: histogram of the
-     dark channel → 99.5th percentile).
-  5. `masks`: one pass per mask shape (analytic shapes evaluate per pixel; brushes rasterized as
-     instanced soft discs; range masks read the image/planes).
-  6. `finish`: the per-pixel stage, a straight port of `finish.rs` (tone LUT and curve LUTs as 1D
-     textures, colour ops in OkLCh, gamut map, sRGB encode, grain hash). Output `rgba8unorm` for
-     display, `rgba16float`/`rgba32float` for export.
-  7. `histogram`: atomics into a storage buffer (256 bins × 4).
-- Readback only for export and for `ui.screenshot`/MCP renders.
+## Structure
+- Shared parameters, so the two implementations cannot drift apart: `pipeline::plan` (effective
+  settings, frame, output size, stage-cache keys), `Frame::sample_plan`, `local::{wb_matrix_for,
+  nr_params, plane_sigmas, guided_fast_step, airlight_of}`, `finish::{FinishParams, mask_terms}`,
+  `masks::brush_dabs`; exact tables (tone LUT, sRGB LUT, curve LUTs, resample taps) and the OkLab
+  matrices are uploaded / generated into the WGSL prelude from the CPU values.
+- Buffers, not textures: images are `array<f32>` with the CPU's interleaved layout (RGB, 1–3
+  channels), so upload / readback are plain copies of `Rgb32f` / `Plane` / `Rgba8` data.
+- 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`
+  (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.
+- Per-stage hybrid: defringe and spot removal (rare, CPU-only for now) download the resampled image,
+  run on the CPU and upload; heuristic mask shapes (Sky, Subject, Background, …) and brushes of more
+  than 4096 dabs are evaluated on the CPU and uploaded; sources over the buffer limit are resampled
+  on the CPU.
+- Stage cache: `GpuStages` mirrors `StageCache` (same keys: source identity + `geo`, `lin_key`,
+  per-plane radii) and lives in it as an extension, so the UI needs no change. The uploaded source is
+  kept per view. A tone / colour / exposure drag re-runs one dispatch of `finish` plus the readback.
+- One command encoder per render; the only mid-render readback is the airlight samples (dehaze, when
+  the dark channel is recomputed). Dropped buffers are recycled (exact size, ≤ 2 GB pool) once the
+  thread that dropped them has submitted — allocating and zero-filling fresh 100–300 MB buffers
+  per pass cost as much as the passes.
+- `LIGHTCRAFT_PROFILE=1` prints GPU stage timings (each stage is then submitted and waited for).
 
 ## Correctness: CPU oracle and equivalence tests
-- Every kernel has a CPU twin; tests render the synthetic scenes (`lightcraft-scenes`) on both and
-  require max |Δ| ≤ 1/1023 in display-encoded values (architecture §1.6), plus the existing pipeline
-  tests run against the GPU renderer behind a feature flag. Tests skip (not fail) when no adapter is
-  available (CI machines), with a software adapter (wgpu's GL/lavapipe) where possible.
-- Float determinism: avoid fast-math-sensitive constructs (fused multiply-add differences are within
-  tolerance); reductions (airlight, histogram) use integer atomics.
+`crates/gpu/tests/equivalence.rs` renders the same settings on both and compares 8-bit sRGB:
+bounds **mean |Δ| < 0.5 LSB and max |Δ| ≤ 3 LSB** per channel. Measured (Apple M4 Pro, Metal):
+**max 1 LSB, mean ≤ 0.0003 LSB** for every case — 26 settings cases (all tone/colour/effects tools,
+3 vignette styles, grain, curves, grading, B&W, WB, NR, crop/straighten/flip, lens + perspective,
+three mask sets incl. CPU-evaluated shapes, spots + defringe), all 8 orientations ± crop, embedded
+DNG lens data, display-referred sources, draft quality, full-size renders; the 24 MP export in
+`render_bench` also differs by max 1 LSB. Cached (slider-drag) GPU renders are bit-identical to
+fresh ones. The tests skip (pass with a note) when no adapter exists.
 
-## Scheduling
-- One render queue per view; a new revision cancels queued work for the old one (CPU already drops
-  superseded jobs per slot). Drafts during drags at display size (GPU can afford full size), full
-  quality (larger NR/guided radii at source resolution for 1:1) on release.
-- Fallback order: GPU → CPU (`render_cached`) when no adapter, on device loss, or for wasm builds
-  without WebGPU.
+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.
 
-## Open questions
-- f16 vs f32 storage for planes (bandwidth vs. guided-filter precision: start with f32).
-- Sharing the device with eframe on the web (single WebGPU device per canvas).
-- Large exports (60 MP+) exceed texture limits on some adapters (8192/16384): tile with overlap
-  equal to the largest filter radius (the planes' radii are known per settings).
+## Measurements
+`render_bench` on `corpus/raw/arw-sony-a7m3-compressed.arw` (24 MP; loupe = 2560 px preview →
+1920×1280), min of 7 runs, wall clock, on a heavily shared machine (load average 60–110 on 14 cores,
+so CPU numbers are pessimistic; the GPU numbers are less affected):
+
+| scenario | CPU | GPU |
+|---|---|---|
+| loupe 1920×1280 cold (new photo, incl. upload) | 415 ms | 32 ms |
+| loupe draft 1152×768 cold | 169 ms | 21 ms |
+| exposure drag 1920×1280 (warm) | 22–42 ms | 3.9 ms |
+| highlights drag, draft (warm) | 14 ms | 3.2 ms |
+| clarity drag, draft (warm) | 16–23 ms | 3.0 ms |
+| NR drag, draft / 1920×1280 (warm) | 186 / 301–471 ms | 10 / 18 ms |
+| per-pixel stage 1920×1280, + vibrance/saturation | 53–65 ms | 5 ms |
+| export render 6000×4000 | 1037–1414 ms | 291–330 ms |
+
+GPU timings include the readback of the 8-bit result and the histogram. Device creation + kernel
+compilation: ~0.4 s once per process (lazily, on the first GPU render).
+
+## Open items
+- Keep the loupe texture on the GPU (render into an `egui-wgpu` texture on eframe's device instead of
+  reading back); share the device with eframe.
+- WebGPU in the browser build (async init, single device per canvas).
+- Defringe and spot removal kernels; Sky/Subject heuristics (replaced by the segmenter in M12).
+- Tiling for sources / outputs beyond the buffer limit (60 MP+ on smaller adapters): tile with
+  overlap equal to the largest filter radius.
+- Large-radius blurs at full size (dehaze dark channel at 24 MP) dominate the export: a summed-area
+  table or a downsampled dark channel would cut them further.
+- GPU histogram (atomics) to skip the CPU pass over the readback.