diff --git a/Cargo.lock b/Cargo.lock
index 65e80c8d..01e13890 100644
--- a/Cargo.lock
+++ b/Cargo.lock
@@ -2096,6 +2096,7 @@ dependencies = [
"lightcraft-color",
"lightcraft-develop",
"lightcraft-geom",
+ "lightcraft-gpu",
"lightcraft-meta",
"lightcraft-pipeline",
"lightcraft-preview",
@@ -2116,6 +2117,22 @@ dependencies = [
"serde",
]
+[[package]]
+name = "lightcraft-gpu"
+version = "0.1.0"
+dependencies = [
+ "bytemuck",
+ "lightcraft-color",
+ "lightcraft-develop",
+ "lightcraft-geom",
+ "lightcraft-pipeline",
+ "lightcraft-raster",
+ "lightcraft-scenes",
+ "log",
+ "pollster 1.0.1",
+ "wgpu",
+]
+
[[package]]
name = "lightcraft-mcp"
version = "0.1.0"
diff --git a/Cargo.toml b/Cargo.toml
index 6decc70d..76abfa4c 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -21,6 +21,7 @@ lightcraft-meta = { path = "crates/meta" }
lightcraft-develop = { path = "crates/develop" }
lightcraft-scenes = { path = "crates/scenes" }
lightcraft-pipeline = { path = "crates/pipeline" }
+lightcraft-gpu = { path = "crates/gpu" }
lightcraft-catalog = { path = "crates/catalog" }
lightcraft-preview = { path = "crates/preview" }
lightcraft-engine = { path = "crates/engine" }
@@ -57,6 +58,10 @@ moxcms = { version = "0.9", default-features = false, features = ["lut"] }
jxl-oxide = { version = "0.12", default-features = false }
ravif = { version = "0.13", default-features = false, features = ["threading"] } # no "asm": it needs nasm (non-Rust)
+# gpu (crates/gpu; native backends only, no GL: everything stays pure Rust)
+wgpu = { version = "30", default-features = false, features = ["std", "parking_lot", "metal", "vulkan", "dx12", "wgsl"] }
+pollster = "1"
+
# ui
egui = "0.36"
egui_extras = { version = "0.36", default-features = false }
diff --git a/README.md b/README.md
index 4887d5ef..666a29de 100644
--- a/README.md
+++ b/README.md
@@ -202,6 +202,10 @@ lightcraft --control 7980 ~/Pictures/trip
instead of clipping, a filmic shoulder for raw and pixel-exact pass-through for JPEGs you haven't touched.
- **Resolution-independent edits.** Radii and brush sizes are relative to the image, so a 400 px preview, your
5K display and a 60 MP export look the same.
+- **GPU-accelerated, CPU-exact.** The whole develop pipeline runs as wgpu compute kernels (Metal / Vulkan / DX12),
+ checked against the CPU pipeline to within 1/255. On a 24 MP raw (Apple M4 Pro): a slider update re-renders in
+ ~4 ms, a cold 2.5 MP loupe in ~30 ms, and a full-size export in ~0.3 s including a parallel JPEG encode.
+ Without a GPU the same pipeline runs on all CPU cores, redoing only the stages a slider affects.
- **Background rendering.** A worker pool renders the loupe, before/after and every visible thumbnail off the UI
thread; drafts during drags, full quality on release.
- **Local-first.** No account, no cloud, no telemetry, no subscription. Your catalog is an append-only log of
@@ -229,7 +233,7 @@ LightCraft is young and moving fast — see the [roadmap](ROADMAP.md) for estima
| MCP server (headless or live app, persistent libraries), CLI, control channel | ✅ |
| XMP sidecars (read/write, auto-write), reading `crs:` develop settings, preset files (`.lcpreset`, XMP presets) | ✅ |
| Optics (distortion, vignetting, auto + manual CA, defringe, DNG-embedded lens corrections), Geometry (transforms, Constrain Crop), Upright (Auto/Level/Vertical/Full/Guided) | ✅ · camera lens profiles (our own) ⬜ |
-| GPU pipeline | 🚧 |
+| GPU pipeline (wgpu compute, CPU-exact within 1/255) with CPU fallback | ✅ · WebGPU in the browser 🚧 |
| Web build (same UI in the browser via WASM; in-memory imports, export downloads) | ✅ · persistence, workers 🚧 |
diff --git a/crates/codecs/examples/jpeg_bench.rs b/crates/codecs/examples/jpeg_bench.rs
new file mode 100644
index 00000000..487f8619
--- /dev/null
+++ b/crates/codecs/examples/jpeg_bench.rs
@@ -0,0 +1,61 @@
+//! Compare the parallel JPEG encoder with `jpeg-encoder` (size, time, PSNR):
+//! `cargo run --release -p lightcraft-codecs --example jpeg_bench [image]` (default: 6000×4000 synthetic).
+use lightcraft_codecs::ChromaSubsampling;
+use std::time::Instant;
+
+fn psnr(a: &[u8], b: &[u8]) -> f64 {
+ let mse = a.iter().zip(b).map(|(x, y)| (*x as f64 - *y as f64).powi(2)).sum::() / a.len() as f64;
+ 10.0 * (255.0f64.powi(2) / mse.max(1e-9)).log10()
+}
+
+fn best(n: usize, mut f: impl FnMut() -> T) -> (T, f64) {
+ let mut out = None;
+ let mut t = f64::MAX;
+ for _ in 0..n {
+ let s = Instant::now();
+ let r = f();
+ t = t.min(s.elapsed().as_secs_f64() * 1e3);
+ out = Some(r);
+ }
+ (out.unwrap(), t)
+}
+
+fn main() {
+ let (w, h, rgb) = match std::env::args().nth(1) {
+ Some(p) => {
+ let img = image::open(p).expect("open").to_rgb8();
+ (img.width() as usize, img.height() as usize, img.into_raw())
+ }
+ None => {
+ let (w, h) = (6000, 4000);
+ let mut v = Vec::with_capacity(w * h * 3);
+ for y in 0..h {
+ for x in 0..w {
+ let t = (((x as f32 / 37.0).sin() * (y as f32 / 23.0).cos()) * 0.5 + 0.5) * 255.0;
+ v.extend_from_slice(&[(x * 255 / w) as u8, t as u8, ((x ^ y) & 0xff) as u8]);
+ }
+ }
+ (w, h, v)
+ }
+ };
+ for (sub, js) in
+ [(ChromaSubsampling::S444, jpeg_encoder::SamplingFactor::R_4_4_4), (ChromaSubsampling::S420, jpeg_encoder::SamplingFactor::R_4_2_0)]
+ {
+ let (ours, t_ours) = best(5, || lightcraft_codecs::jpeg_par::encode(&rgb, w, h, 3, 90, sub, &[]));
+ let (theirs, t_theirs) = best(3, || {
+ let mut out = Vec::new();
+ let mut e = jpeg_encoder::Encoder::new(&mut out, 90);
+ e.set_sampling_factor(js);
+ e.encode(&rgb, w as u16, h as u16, jpeg_encoder::ColorType::Rgb).unwrap();
+ out
+ });
+ let dec = |b: &[u8]| jpeg_decoder::Decoder::new(b).decode().unwrap();
+ println!(
+ "{w}x{h} {sub:?} q90: parallel {t_ours:.0} ms, {} KB, PSNR {:.2} dB | jpeg-encoder {t_theirs:.0} ms, {} KB, PSNR {:.2} dB",
+ ours.len() / 1024,
+ psnr(&rgb, &dec(&ours)),
+ theirs.len() / 1024,
+ psnr(&rgb, &dec(&theirs))
+ );
+ }
+}
diff --git a/crates/codecs/src/encode.rs b/crates/codecs/src/encode.rs
index 89dadeed..9bfc1f78 100644
--- a/crates/codecs/src/encode.rs
+++ b/crates/codecs/src/encode.rs
@@ -93,6 +93,36 @@ pub fn encode_jpeg(img: &EncodeImage, quality: u8, subsampling: ChromaSubsamplin
3 => (data[..n * 3].into(), jpeg_encoder::ColorType::Rgb),
_ => (data[..n * 4].into(), jpeg_encoder::ColorType::Rgba),
};
+ if subsampling != ChromaSubsampling::S422 {
+ // Parallel encoder (restart-interval bands on all cores).
+ let mut segs: Vec<(u8, Vec)> = Vec::new();
+ if let Some(exif) = meta.exif {
+ if exif.len() > 65_527 {
+ return Err(Error::Encode("EXIF larger than one APP1 segment (64 KiB)".into()));
+ }
+ segs.push((0xE1, [&b"Exif\0\0"[..], exif].concat()));
+ }
+ if let Some(xmp) = meta.xmp {
+ let seg = [XMP_NS, xmp.as_bytes()].concat();
+ if seg.len() > 65_533 {
+ return Err(Error::Encode("XMP larger than one APP1 segment (extended XMP not supported)".into()));
+ }
+ segs.push((0xE1, seg));
+ }
+ if let Some(icc) = meta.icc {
+ // ICC.1 Annex B.4: "ICC_PROFILE\0", 1-based chunk number, chunk count
+ let chunks: Vec<&[u8]> = icc.chunks(65_519).collect();
+ if chunks.len() > 255 {
+ return Err(Error::Encode("ICC profile too large for JPEG".into()));
+ }
+ for (i, c) in chunks.iter().enumerate() {
+ segs.push((0xE2, [&b"ICC_PROFILE\0"[..], &[i as u8 + 1, chunks.len() as u8], c].concat()));
+ }
+ }
+ let ch = if img.channels == 2 { 1 } else { img.channels as usize };
+ // `data` is already gray for 2-channel input (see above)
+ return Ok(crate::jpeg_par::encode(&data, img.width as usize, img.height as usize, ch, quality, subsampling, &segs));
+ }
let mut out = Vec::new();
let mut enc = jpeg_encoder::Encoder::new(&mut out, quality.clamp(1, 100));
enc.set_sampling_factor(match subsampling {
diff --git a/crates/codecs/src/jpeg_par.rs b/crates/codecs/src/jpeg_par.rs
new file mode 100644
index 00000000..7f1e1d20
--- /dev/null
+++ b/crates/codecs/src/jpeg_par.rs
@@ -0,0 +1,495 @@
+//! Parallel baseline JPEG encoder (ITU-T T.81 / JFIF).
+//!
+//! The image is cut into bands of MCU rows separated by restart markers (`DRI` + `RSTn`): DC
+//! prediction restarts at every marker, so bands are entropy-coded independently on all cores and
+//! concatenated. Quantization uses the example tables of T.81 Annex K scaled by the usual 1–100
+//! quality convention; Huffman coding uses the Annex K example tables (interleaved baseline, readable
+//! by every decoder). Supports grayscale and YCbCr 4:4:4 / 4:2:0.
+
+use crate::encode::ChromaSubsampling;
+
+/// T.81 Annex K.1 / K.2 example quantization tables (natural order).
+const Q_LUMA: [u16; 64] = [
+ 16, 11, 10, 16, 24, 40, 51, 61, 12, 12, 14, 19, 26, 58, 60, 55, 14, 13, 16, 24, 40, 57, 69, 56, 14, 17, 22, 29, 51, 87, 80, 62, 18, 22, 37, 56,
+ 68, 109, 103, 77, 24, 35, 55, 64, 81, 104, 113, 92, 49, 64, 78, 87, 103, 121, 120, 101, 72, 92, 95, 98, 112, 100, 103, 99,
+];
+const Q_CHROMA: [u16; 64] = [
+ 17, 18, 24, 47, 99, 99, 99, 99, 18, 21, 26, 66, 99, 99, 99, 99, 24, 26, 56, 99, 99, 99, 99, 99, 47, 66, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99,
+ 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99,
+];
+
+/// Zig-zag position k → natural index.
+const ZIGZAG: [usize; 64] = [
+ 0, 1, 8, 16, 9, 2, 3, 10, 17, 24, 32, 25, 18, 11, 4, 5, 12, 19, 26, 33, 40, 48, 41, 34, 27, 20, 13, 6, 7, 14, 21, 28, 35, 42, 49, 56, 57, 50, 43,
+ 36, 29, 22, 15, 23, 30, 37, 44, 51, 58, 59, 52, 45, 38, 31, 39, 46, 53, 60, 61, 54, 47, 55, 62, 63,
+];
+
+/// T.81 Annex K.3 example Huffman tables: (code counts per length 1..16, symbols).
+const DC_LUMA_BITS: [u8; 16] = [0, 1, 5, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0];
+const DC_CHROMA_BITS: [u8; 16] = [0, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0, 0];
+const DC_VALS: [u8; 12] = [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11];
+const AC_LUMA_BITS: [u8; 16] = [0, 2, 1, 3, 3, 2, 4, 3, 5, 5, 4, 4, 0, 0, 1, 0x7d];
+const AC_LUMA_VALS: [u8; 162] = [
+ 0x01, 0x02, 0x03, 0x00, 0x04, 0x11, 0x05, 0x12, 0x21, 0x31, 0x41, 0x06, 0x13, 0x51, 0x61, 0x07, 0x22, 0x71, 0x14, 0x32, 0x81, 0x91, 0xa1, 0x08,
+ 0x23, 0x42, 0xb1, 0xc1, 0x15, 0x52, 0xd1, 0xf0, 0x24, 0x33, 0x62, 0x72, 0x82, 0x09, 0x0a, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x25, 0x26, 0x27, 0x28,
+ 0x29, 0x2a, 0x34, 0x35, 0x36, 0x37, 0x38, 0x39, 0x3a, 0x43, 0x44, 0x45, 0x46, 0x47, 0x48, 0x49, 0x4a, 0x53, 0x54, 0x55, 0x56, 0x57, 0x58, 0x59,
+ 0x5a, 0x63, 0x64, 0x65, 0x66, 0x67, 0x68, 0x69, 0x6a, 0x73, 0x74, 0x75, 0x76, 0x77, 0x78, 0x79, 0x7a, 0x83, 0x84, 0x85, 0x86, 0x87, 0x88, 0x89,
+ 0x8a, 0x92, 0x93, 0x94, 0x95, 0x96, 0x97, 0x98, 0x99, 0x9a, 0xa2, 0xa3, 0xa4, 0xa5, 0xa6, 0xa7, 0xa8, 0xa9, 0xaa, 0xb2, 0xb3, 0xb4, 0xb5, 0xb6,
+ 0xb7, 0xb8, 0xb9, 0xba, 0xc2, 0xc3, 0xc4, 0xc5, 0xc6, 0xc7, 0xc8, 0xc9, 0xca, 0xd2, 0xd3, 0xd4, 0xd5, 0xd6, 0xd7, 0xd8, 0xd9, 0xda, 0xe1, 0xe2,
+ 0xe3, 0xe4, 0xe5, 0xe6, 0xe7, 0xe8, 0xe9, 0xea, 0xf1, 0xf2, 0xf3, 0xf4, 0xf5, 0xf6, 0xf7, 0xf8, 0xf9, 0xfa,
+];
+const AC_CHROMA_BITS: [u8; 16] = [0, 2, 1, 2, 4, 4, 3, 4, 7, 5, 4, 4, 0, 1, 2, 0x77];
+const AC_CHROMA_VALS: [u8; 162] = [
+ 0x00, 0x01, 0x02, 0x03, 0x11, 0x04, 0x05, 0x21, 0x31, 0x06, 0x12, 0x41, 0x51, 0x07, 0x61, 0x71, 0x13, 0x22, 0x32, 0x81, 0x08, 0x14, 0x42, 0x91,
+ 0xa1, 0xb1, 0xc1, 0x09, 0x23, 0x33, 0x52, 0xf0, 0x15, 0x62, 0x72, 0xd1, 0x0a, 0x16, 0x24, 0x34, 0xe1, 0x25, 0xf1, 0x17, 0x18, 0x19, 0x1a, 0x26,
+ 0x27, 0x28, 0x29, 0x2a, 0x35, 0x36, 0x37, 0x38, 0x39, 0x3a, 0x43, 0x44, 0x45, 0x46, 0x47, 0x48, 0x49, 0x4a, 0x53, 0x54, 0x55, 0x56, 0x57, 0x58,
+ 0x59, 0x5a, 0x63, 0x64, 0x65, 0x66, 0x67, 0x68, 0x69, 0x6a, 0x73, 0x74, 0x75, 0x76, 0x77, 0x78, 0x79, 0x7a, 0x82, 0x83, 0x84, 0x85, 0x86, 0x87,
+ 0x88, 0x89, 0x8a, 0x92, 0x93, 0x94, 0x95, 0x96, 0x97, 0x98, 0x99, 0x9a, 0xa2, 0xa3, 0xa4, 0xa5, 0xa6, 0xa7, 0xa8, 0xa9, 0xaa, 0xb2, 0xb3, 0xb4,
+ 0xb5, 0xb6, 0xb7, 0xb8, 0xb9, 0xba, 0xc2, 0xc3, 0xc4, 0xc5, 0xc6, 0xc7, 0xc8, 0xc9, 0xca, 0xd2, 0xd3, 0xd4, 0xd5, 0xd6, 0xd7, 0xd8, 0xd9, 0xda,
+ 0xe2, 0xe3, 0xe4, 0xe5, 0xe6, 0xe7, 0xe8, 0xe9, 0xea, 0xf2, 0xf3, 0xf4, 0xf5, 0xf6, 0xf7, 0xf8, 0xf9, 0xfa,
+];
+
+/// Huffman code table: symbol → (code, length).
+struct Huff {
+ code: [(u16, u8); 256],
+}
+
+impl Huff {
+ /// Canonical codes from counts (T.81 Annex C).
+ fn new(bits: &[u8; 16], vals: &[u8]) -> Huff {
+ let mut code = [(0u16, 0u8); 256];
+ let (mut c, mut k) = (0u16, 0usize);
+ for (len, &n) in bits.iter().enumerate() {
+ for _ in 0..n {
+ code[vals[k] as usize] = (c, len as u8 + 1);
+ c += 1;
+ k += 1;
+ }
+ c <<= 1;
+ }
+ Huff { code }
+ }
+}
+
+fn scaled_quant(base: &[u16; 64], quality: u8) -> [u16; 64] {
+ let q = quality.clamp(1, 100) as u32;
+ let scale = if q < 50 { 5000 / q } else { 200 - 2 * q };
+ base.map(|b| ((b as u32 * scale + 50) / 100).clamp(1, 255) as u16)
+}
+
+/// Entropy-coded segment writer with 0xFF byte stuffing.
+struct Bits {
+ out: Vec,
+ acc: u32,
+ n: u32,
+}
+
+impl Bits {
+ fn new() -> Bits {
+ Bits { out: Vec::new(), acc: 0, n: 0 }
+ }
+ #[inline]
+ fn put(&mut self, code: u32, len: u32) {
+ if len == 0 {
+ return;
+ }
+ self.acc = (self.acc << len) | (code & ((1 << len) - 1));
+ self.n += len;
+ while self.n >= 8 {
+ let b = (self.acc >> (self.n - 8)) as u8;
+ self.out.push(b);
+ if b == 0xFF {
+ self.out.push(0);
+ }
+ self.n -= 8;
+ }
+ self.acc &= (1 << self.n) - 1;
+ }
+ /// Pad with 1-bits to a byte boundary.
+ fn finish(mut self) -> Vec {
+ if self.n > 0 {
+ let pad = 8 - self.n;
+ self.put((1 << pad) - 1, pad);
+ }
+ self.out
+ }
+}
+
+/// AAN output scale factors: `s[0] = 1`, `s[k] = cos(kπ/16)·√2`.
+fn aan_scale() -> [f32; 8] {
+ std::array::from_fn(|k| if k == 0 { 1.0 } else { ((k as f64 * std::f64::consts::PI / 16.0).cos() * std::f64::consts::SQRT_2) as f32 })
+}
+
+/// Quantization reciprocals with the AAN scaling folded in: coefficient = out / (8·s[u]·s[v]·q).
+fn recip_table(q: &[u16; 64]) -> [f32; 64] {
+ let s = aan_scale();
+ std::array::from_fn(|i| 1.0 / (q[i] as f32 * s[i / 8] * s[i % 8] * 8.0))
+}
+
+/// One 8-point forward DCT (Arai–Agui–Nakajima flowgraph, 5 multiplies) on `d[o + k·stride]`.
+#[inline(always)]
+fn aan_1d(d: &mut [f32; 64], o: usize, stride: usize) {
+ let g = |k: usize| o + k * stride;
+ let (t0, t7) = (d[g(0)] + d[g(7)], d[g(0)] - d[g(7)]);
+ let (t1, t6) = (d[g(1)] + d[g(6)], d[g(1)] - d[g(6)]);
+ let (t2, t5) = (d[g(2)] + d[g(5)], d[g(2)] - d[g(5)]);
+ let (t3, t4) = (d[g(3)] + d[g(4)], d[g(3)] - d[g(4)]);
+ // even part
+ let (t10, t13, t11, t12) = (t0 + t3, t0 - t3, t1 + t2, t1 - t2);
+ d[g(0)] = t10 + t11;
+ d[g(4)] = t10 - t11;
+ let z1 = (t12 + t13) * 0.707_106_77;
+ d[g(2)] = t13 + z1;
+ d[g(6)] = t13 - z1;
+ // odd part
+ let (t10, t11, t12) = (t4 + t5, t5 + t6, t6 + t7);
+ let z5 = (t10 - t12) * 0.382_683_43;
+ let z2 = 0.541_196_1 * t10 + z5;
+ let z4 = 1.306_563 * t12 + z5;
+ let z3 = t11 * 0.707_106_77;
+ let (z11, z13) = (t7 + z3, t7 - z3);
+ d[g(5)] = z13 + z2;
+ d[g(3)] = z13 - z2;
+ d[g(1)] = z11 + z4;
+ d[g(7)] = z11 - z4;
+}
+
+/// Forward DCT of a level-shifted 8×8 block and quantization.
+#[inline]
+fn fdct_quant(block: &[f32; 64], recip: &[f32; 64], out: &mut [i16; 64]) {
+ let mut d = *block;
+ for r in 0..8 {
+ aan_1d(&mut d, r * 8, 1);
+ }
+ for c in 0..8 {
+ aan_1d(&mut d, c, 8);
+ }
+ for i in 0..64 {
+ out[i] = (d[i] * recip[i]).round() as i16;
+ }
+}
+
+#[inline]
+fn magnitude(v: i32) -> (u32, u32) {
+ let a = v.unsigned_abs();
+ let size = 32 - a.leading_zeros();
+ let bits = if v < 0 { (v - 1) as u32 } else { v as u32 };
+ (size, bits & ((1u32 << size) - 1))
+}
+
+fn encode_block(w: &mut Bits, q: &[i16; 64], pred: &mut i32, dc: &Huff, ac: &Huff) {
+ let d = q[0] as i32 - *pred;
+ *pred = q[0] as i32;
+ let (s, b) = magnitude(d);
+ let (c, l) = dc.code[s as usize];
+ w.put(c as u32, l as u32);
+ w.put(b, s);
+ let mut run = 0u32;
+ for &zi in &ZIGZAG[1..] {
+ let v = q[zi] as i32;
+ if v == 0 {
+ run += 1;
+ continue;
+ }
+ while run >= 16 {
+ let (c, l) = ac.code[0xF0];
+ w.put(c as u32, l as u32);
+ run -= 16;
+ }
+ let (s, b) = magnitude(v);
+ let (c, l) = ac.code[((run << 4) | s) as usize];
+ w.put(c as u32, l as u32);
+ w.put(b, s);
+ run = 0;
+ }
+ if run > 0 {
+ let (c, l) = ac.code[0x00];
+ w.put(c as u32, l as u32);
+ }
+}
+
+/// Planar component samples (level-shifted to −128..127) with edge replication on read.
+struct Plane {
+ w: usize,
+ h: usize,
+ data: Vec,
+}
+
+impl Plane {
+ #[inline]
+ fn block(&self, bx: usize, by: usize, out: &mut [f32; 64]) {
+ let (x0, y0) = (bx * 8, by * 8);
+ if x0 + 8 <= self.w && y0 + 8 <= self.h {
+ for y in 0..8 {
+ let i = (y0 + y) * self.w + x0;
+ out[y * 8..y * 8 + 8].copy_from_slice(&self.data[i..i + 8]);
+ }
+ return;
+ }
+ for y in 0..8 {
+ let sy = (by * 8 + y).min(self.h - 1);
+ let row = &self.data[sy * self.w..(sy + 1) * self.w];
+ for x in 0..8 {
+ out[y * 8 + x] = row[(bx * 8 + x).min(self.w - 1)];
+ }
+ }
+ }
+}
+
+fn segment(out: &mut Vec, marker: u8, payload: &[u8]) {
+ out.extend_from_slice(&[0xFF, marker]);
+ out.extend_from_slice(&((payload.len() + 2) as u16).to_be_bytes());
+ out.extend_from_slice(payload);
+}
+
+fn dht(out: &mut Vec, class_id: u8, bits: &[u8; 16], vals: &[u8]) {
+ let mut p = vec![class_id];
+ p.extend_from_slice(bits);
+ p.extend_from_slice(vals);
+ segment(out, 0xC4, &p);
+}
+
+/// Encode interleaved 8-bit samples (`channels` 1, 3 or 4; alpha is dropped).
+pub fn encode(
+ data: &[u8],
+ width: usize,
+ height: usize,
+ channels: usize,
+ quality: u8,
+ sub: ChromaSubsampling,
+ app_segments: &[(u8, Vec)],
+) -> Vec {
+ let gray = channels == 1;
+ let (hs, vs) = match (gray, sub) {
+ (true, _) | (false, ChromaSubsampling::S444) => (1usize, 1usize),
+ (false, _) => (2, 2),
+ };
+ // colour conversion (JFIF YCbCr, full range, level-shifted) in one pass; chroma averaged 2×2 for 4:2:0
+ let n = width * height;
+ #[cfg(feature = "parallel")]
+ use rayon::prelude::*;
+ let ncomp = if gray { 1 } else { 3 };
+ let mut full: Vec> = (0..ncomp).map(|_| vec![0f32; n]).collect();
+ {
+ let rows: Vec<(usize, Vec<&mut [f32]>)> = {
+ let mut its: Vec<_> = full.iter_mut().map(|p| p.chunks_mut(width)).collect();
+ (0..height).map(|y| (y, its.iter_mut().map(|it| it.next().unwrap()).collect())).collect()
+ };
+ let conv = |(y, mut out): (usize, Vec<&mut [f32]>)| {
+ let src = &data[y * width * channels..(y + 1) * width * channels];
+ if gray {
+ for (o, p) in out[0].iter_mut().zip(src.chunks_exact(channels)) {
+ *o = p[0] as f32 - 128.0;
+ }
+ return;
+ }
+ for (x, p) in src.chunks_exact(channels).enumerate() {
+ let (r, g, b) = (p[0] as f32, p[1] as f32, p[2] as f32);
+ out[0][x] = 0.299 * r + 0.587 * g + 0.114 * b - 128.0;
+ out[1][x] = -0.168_736 * r - 0.331_264 * g + 0.5 * b;
+ out[2][x] = 0.5 * r - 0.418_688 * g - 0.081_312 * b;
+ }
+ };
+ #[cfg(feature = "parallel")]
+ rows.into_par_iter().for_each(conv);
+ #[cfg(not(feature = "parallel"))]
+ rows.into_iter().for_each(conv);
+ }
+ let mut it = full.into_iter();
+ let mut planes = vec![Plane { w: width, h: height, data: it.next().unwrap_or_default() }];
+ for c in it {
+ if hs == 1 {
+ planes.push(Plane { w: width, h: height, data: c });
+ continue;
+ }
+ let (cw, ch) = (width.div_ceil(2), height.div_ceil(2));
+ let mut d = vec![0f32; cw * ch];
+ let down = |(cy, out): (usize, &mut [f32])| {
+ let (y0, y1) = (2 * cy, (2 * cy + 1).min(height - 1));
+ for (cx, o) in out.iter_mut().enumerate() {
+ let (x0, x1) = (2 * cx, (2 * cx + 1).min(width - 1));
+ *o = 0.25 * (c[y0 * width + x0] + c[y0 * width + x1] + c[y1 * width + x0] + c[y1 * width + x1]);
+ }
+ };
+ #[cfg(feature = "parallel")]
+ d.par_chunks_mut(cw).enumerate().for_each(down);
+ #[cfg(not(feature = "parallel"))]
+ d.chunks_mut(cw).enumerate().for_each(down);
+ planes.push(Plane { w: cw, h: ch, data: d });
+ }
+
+ let ql = scaled_quant(&Q_LUMA, quality);
+ let qc = scaled_quant(&Q_CHROMA, quality);
+ let (rl, rc) = (recip_table(&ql), recip_table(&qc));
+ let (dcl, acl) = (Huff::new(&DC_LUMA_BITS, &DC_VALS), Huff::new(&AC_LUMA_BITS, &AC_LUMA_VALS));
+ let (dcc, acc) = (Huff::new(&DC_CHROMA_BITS, &DC_VALS), Huff::new(&AC_CHROMA_BITS, &AC_CHROMA_VALS));
+
+ let (mcu_w, mcu_h) = (8 * hs, 8 * vs);
+ let (mx, my) = (width.div_ceil(mcu_w), height.div_ceil(mcu_h));
+ // restart interval: whole MCU rows per band, ≤ 65535 MCUs, enough bands to keep all cores busy
+ let rows_per_band = (my / 64).clamp(1, (65_535 / mx.max(1)).max(1));
+ let bands: Vec = (0..my.div_ceil(rows_per_band)).collect();
+ let encode_band = |band: &usize| -> Vec {
+ let mut w = Bits::new();
+ let mut pred = [0i32; 3];
+ let (mut blk, mut q) = ([0f32; 64], [0i16; 64]);
+ for mcu_y in band * rows_per_band..((band + 1) * rows_per_band).min(my) {
+ for mcu_x in 0..mx {
+ for vy in 0..vs {
+ for hx in 0..hs {
+ planes[0].block(mcu_x * hs + hx, mcu_y * vs + vy, &mut blk);
+ fdct_quant(&blk, &rl, &mut q);
+ encode_block(&mut w, &q, &mut pred[0], &dcl, &acl);
+ }
+ }
+ if !gray {
+ for k in 1..3 {
+ planes[k].block(mcu_x, mcu_y, &mut blk);
+ fdct_quant(&blk, &rc, &mut q);
+ encode_block(&mut w, &q, &mut pred[k], &dcc, &acc);
+ }
+ }
+ }
+ }
+ w.finish()
+ };
+ #[cfg(feature = "parallel")]
+ let coded: Vec> = bands.par_iter().map(encode_band).collect();
+ #[cfg(not(feature = "parallel"))]
+ let coded: Vec> = bands.iter().map(encode_band).collect();
+
+ // headers
+ let mut out = Vec::with_capacity(coded.iter().map(Vec::len).sum::() + 4096);
+ out.extend_from_slice(&[0xFF, 0xD8]);
+ segment(&mut out, 0xE0, &[b'J', b'F', b'I', b'F', 0, 1, 2, 0, 0, 1, 0, 1, 0, 0]);
+ for (marker, payload) in app_segments {
+ segment(&mut out, *marker, payload);
+ }
+ let dqt = |id: u8, q: &[u16; 64]| -> Vec {
+ let mut p = vec![id];
+ p.extend(ZIGZAG.iter().map(|&i| q[i] as u8));
+ p
+ };
+ segment(&mut out, 0xDB, &dqt(0, &ql));
+ if !gray {
+ segment(&mut out, 0xDB, &dqt(1, &qc));
+ }
+ let mut sof = vec![8];
+ sof.extend_from_slice(&(height as u16).to_be_bytes());
+ sof.extend_from_slice(&(width as u16).to_be_bytes());
+ if gray {
+ sof.extend_from_slice(&[1, 1, 0x11, 0]);
+ } else {
+ sof.extend_from_slice(&[3, 1, ((hs as u8) << 4) | vs as u8, 0, 2, 0x11, 1, 3, 0x11, 1]);
+ }
+ segment(&mut out, 0xC0, &sof);
+ dht(&mut out, 0x00, &DC_LUMA_BITS, &DC_VALS);
+ dht(&mut out, 0x10, &AC_LUMA_BITS, &AC_LUMA_VALS);
+ if !gray {
+ dht(&mut out, 0x01, &DC_CHROMA_BITS, &DC_VALS);
+ dht(&mut out, 0x11, &AC_CHROMA_BITS, &AC_CHROMA_VALS);
+ }
+ let interval = (rows_per_band * mx) as u16;
+ if coded.len() > 1 {
+ segment(&mut out, 0xDD, &interval.to_be_bytes());
+ }
+ if gray {
+ segment(&mut out, 0xDA, &[1, 1, 0x00, 0, 63, 0]);
+ } else {
+ segment(&mut out, 0xDA, &[3, 1, 0x00, 2, 0x11, 3, 0x11, 0, 63, 0]);
+ }
+ let last = coded.len().saturating_sub(1);
+ for (i, c) in coded.iter().enumerate() {
+ out.extend_from_slice(c);
+ if i < last {
+ out.extend_from_slice(&[0xFF, 0xD0 + (i % 8) as u8]);
+ }
+ }
+ out.extend_from_slice(&[0xFF, 0xD9]);
+ out
+}
+
+#[cfg(test)]
+mod tests {
+ use super::*;
+
+ fn test_rgb(w: usize, h: usize) -> Vec {
+ let mut v = Vec::with_capacity(w * h * 3);
+ for y in 0..h {
+ for x in 0..w {
+ let t = ((x as f32 / 13.0).sin() * 0.5 + 0.5) * 255.0;
+ v.extend_from_slice(&[(x * 255 / w) as u8, (y * 255 / h) as u8, t as u8]);
+ }
+ }
+ v
+ }
+
+ fn psnr(a: &[u8], b: &[u8]) -> f64 {
+ let mse = a.iter().zip(b).map(|(x, y)| (*x as f64 - *y as f64).powi(2)).sum::() / a.len() as f64;
+ 10.0 * (255.0f64.powi(2) / mse.max(1e-9)).log10()
+ }
+
+ /// The AAN path matches the direct DCT-II definition.
+ #[test]
+ fn aan_matches_reference_dct() {
+ let block: [f32; 64] = std::array::from_fn(|i| ((i * 37 % 255) as f32) - 128.0);
+ let q = [1u16; 64];
+ let mut out = [0i16; 64];
+ fdct_quant(&block, &recip_table(&q), &mut out);
+ for v in 0..8 {
+ for u in 0..8 {
+ let c = |k: usize| if k == 0 { std::f64::consts::FRAC_1_SQRT_2 } else { 1.0 };
+ let mut s = 0.0;
+ for y in 0..8 {
+ for x in 0..8 {
+ s += block[y * 8 + x] as f64
+ * (((2 * x + 1) * u) as f64 * std::f64::consts::PI / 16.0).cos()
+ * (((2 * y + 1) * v) as f64 * std::f64::consts::PI / 16.0).cos();
+ }
+ }
+ let f = 0.25 * c(u) * c(v) * s;
+ assert!((out[v * 8 + u] as f64 - f).abs() <= 0.51, "({u},{v}): {} vs {f}", out[v * 8 + u]);
+ }
+ }
+ }
+
+ #[test]
+ fn huffman_tables_are_consistent() {
+ for (bits, n) in [(&AC_LUMA_BITS, AC_LUMA_VALS.len()), (&AC_CHROMA_BITS, AC_CHROMA_VALS.len()), (&DC_LUMA_BITS, 12), (&DC_CHROMA_BITS, 12)] {
+ assert_eq!(bits.iter().map(|&b| b as usize).sum::(), n);
+ }
+ }
+
+ #[test]
+ fn decodes_with_independent_decoders_at_odd_sizes() {
+ for (w, h, sub) in [
+ (1, 1, ChromaSubsampling::S444),
+ (37, 23, ChromaSubsampling::S420),
+ (640, 1031, ChromaSubsampling::S420),
+ (803, 600, ChromaSubsampling::S444),
+ ] {
+ let src = test_rgb(w, h);
+ let jpg = encode(&src, w, h, 3, 92, sub, &[]);
+ let mut d = jpeg_decoder::Decoder::new(&jpg[..]);
+ let px = d.decode().expect("jpeg-decoder");
+ let info = d.info().unwrap();
+ assert_eq!((info.width as usize, info.height as usize), (w, h));
+ if w * h > 64 {
+ assert!(psnr(&src, &px) > 30.0, "{w}x{h} {sub:?}: psnr {}", psnr(&src, &px));
+ }
+ let z = zune_jpeg::JpegDecoder::new(zune_core::bytestream::ZCursor::new(&jpg[..])).decode().expect("zune-jpeg");
+ assert_eq!(z.len(), px.len());
+ }
+ }
+
+ #[test]
+ fn grayscale_and_quality_monotone() {
+ let (w, h) = (300, 200);
+ let src: Vec = (0..w * h).map(|i| ((i % w) * 255 / w) as u8 ^ ((i / w) as u8 & 0x10)).collect();
+ let lo = encode(&src, w, h, 1, 30, ChromaSubsampling::S444, &[]);
+ let hi = encode(&src, w, h, 1, 95, ChromaSubsampling::S444, &[]);
+ assert!(lo.len() < hi.len());
+ let px = jpeg_decoder::Decoder::new(&hi[..]).decode().unwrap();
+ assert!(psnr(&src, &px) > 35.0);
+ }
+}
diff --git a/crates/codecs/src/lib.rs b/crates/codecs/src/lib.rs
index 5dde3f5c..2c98ea7c 100644
--- a/crates/codecs/src/lib.rs
+++ b/crates/codecs/src/lib.rs
@@ -18,6 +18,7 @@ pub mod encode;
pub mod exif;
pub mod icc;
mod jpeg;
+pub mod jpeg_par;
#[cfg(feature = "jxl")]
mod jxl;
mod other;
diff --git a/crates/color/src/perceptual.rs b/crates/color/src/perceptual.rs
index 40743f80..aa3ffbfa 100644
--- a/crates/color/src/perceptual.rs
+++ b/crates/color/src/perceptual.rs
@@ -44,6 +44,13 @@ fn mul(m: &[[f32; 3]; 3], v: [f32; 3]) -> [f32; 3] {
]
}
+/// The f32 matrices [`oklab_from_2020`] and [`oklab_to_2020`] use (row-major): Rec.2020 → LMS,
+/// LMS → Rec.2020, cube-rooted LMS → Lab, Lab → cube-rooted LMS. GPU kernels embed these.
+pub fn oklab_matrices() -> [[[f32; 3]; 3]; 4] {
+ let m = mats();
+ [m.to_lms, m.from_lms, m.to_lab, m.from_lab]
+}
+
/// Linear Rec.2020 → OkLab `[L, a, b]`.
#[inline]
pub fn oklab_from_2020(rgb: [f32; 3]) -> [f32; 3] {
diff --git a/crates/engine/Cargo.toml b/crates/engine/Cargo.toml
index f075e50b..59922a08 100644
--- a/crates/engine/Cargo.toml
+++ b/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 }
diff --git a/crates/engine/examples/render_bench.rs b/crates/engine/examples/render_bench.rs
index 1fbbb969..8e0027f8 100644
--- a/crates/engine/examples/render_bench.rs
+++ b/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 = 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 = 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 = 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()))));
}
}
diff --git a/crates/engine/src/cmd/query.rs b/crates/engine/src/cmd/query.rs
index f9a018de..5348c351 100644
--- a/crates/engine/src/cmd/query.rs
+++ b/crates/engine/src/cmd/query.rs
@@ -129,6 +129,12 @@ pub fn specs() -> Vec {
"versions": ph.versions.iter().map(|v| json!({"name": v.name, "created": v.created})).collect::>(),
}))
}),
+ 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);
diff --git a/crates/engine/src/lib.rs b/crates/engine/src/lib.rs
index f40c0757..2f971c97 100644
--- a/crates/engine/src/lib.rs
+++ b/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 {
diff --git a/crates/engine/src/media.rs b/crates/engine/src/media.rs
index fa3d72b1..1c0a4070 100644
--- a/crates/engine/src/media.rs
+++ b/crates/engine/src/media.rs
@@ -245,10 +245,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()));
}
@@ -259,6 +259,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, 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) {
diff --git a/crates/gpu/Cargo.toml b/crates/gpu/Cargo.toml
new file mode 100644
index 00000000..765e4c6d
--- /dev/null
+++ b/crates/gpu/Cargo.toml
@@ -0,0 +1,27 @@
+[package]
+name = "lightcraft-gpu"
+version.workspace = true
+edition.workspace = true
+license.workspace = true
+description = "LightCraft GPU develop pipeline: wgpu compute (WGSL) twin of the CPU pipeline, which stays the reference (L3)"
+
+[lints]
+workspace = true
+
+[dependencies]
+lightcraft-geom = { workspace = true }
+lightcraft-color = { workspace = true }
+lightcraft-raster = { workspace = true }
+lightcraft-develop = { workspace = true }
+lightcraft-pipeline = { workspace = true }
+log = { workspace = true }
+
+# Native only: Metal / Vulkan / DX12 through wgpu (pure Rust; the drivers are the system's).
+# The browser build keeps the CPU pipeline (WebGPU needs async device creation: see docs/gpu-pipeline.md).
+[target.'cfg(not(target_arch = "wasm32"))'.dependencies]
+wgpu = { workspace = true }
+pollster = { workspace = true }
+bytemuck = { workspace = true }
+
+[dev-dependencies]
+lightcraft-scenes = { workspace = true }
diff --git a/crates/gpu/src/ctx.rs b/crates/gpu/src/ctx.rs
new file mode 100644
index 00000000..b71471f6
--- /dev/null
+++ b/crates/gpu/src/ctx.rs
@@ -0,0 +1,339 @@
+//! The wgpu device, the compiled kernels, buffers, dispatch and readback.
+
+use std::collections::HashMap;
+
+use wgpu::util::DeviceExt;
+
+/// A kernel module: WGSL source, its storage bindings (name, writable, element type) after the
+/// parameter block `P` (binding 0), and its entry points.
+struct Module {
+ src: &'static str,
+ bindings: &'static [(&'static str, bool, &'static str)],
+ entries: &'static [&'static str],
+}
+
+const MODULES: &[Module] = &[
+ Module {
+ src: include_str!("wgsl/finish.wgsl"),
+ bindings: &[
+ ("img", false, "f32"),
+ ("log_l", false, "f32"),
+ ("base_p", false, "f32"),
+ ("clar", false, "f32"),
+ ("tex", false, "f32"),
+ ("dark", false, "f32"),
+ ("masks", false, "f32"),
+ ("aux", false, "f32"),
+ ("out", true, "u32"),
+ ],
+ entries: &["main"],
+ },
+ Module { src: include_str!("wgsl/blur.wgsl"), bindings: &[("src", false, "f32"), ("dst", true, "f32")], entries: &["box_h", "box_v"] },
+ Module {
+ src: include_str!("wgsl/resize.wgsl"),
+ bindings: &[("src", false, "f32"), ("table", false, "u32"), ("dst", true, "f32")],
+ entries: &["resize_h", "resize_v"],
+ },
+ Module {
+ src: include_str!("wgsl/map.wgsl"),
+ bindings: &[("a", false, "f32"), ("b", false, "f32"), ("c", false, "f32"), ("dst", true, "f32")],
+ entries: &["log_lum_k", "dark_k", "guided_pre", "guided_ab", "guided_apply", "wb_k", "nr_lum", "chroma_k", "nr_col", "subsample"],
+ },
+ Module {
+ src: include_str!("wgsl/mask.wgsl"),
+ bindings: &[("img", false, "f32"), ("log_l", false, "f32"), ("aux", false, "f32"), ("c", true, "f32"), ("alpha", true, "f32")],
+ entries: &["shape", "combine", "finalize"],
+ },
+ Module {
+ src: include_str!("wgsl/geom.wgsl"),
+ bindings: &[("src", false, "f32"), ("dst", true, "f32")],
+ entries: &["orient", "sample_affine", "sample_warp"],
+ },
+];
+
+struct Kernel {
+ pipeline: wgpu::ComputePipeline,
+ layout: wgpu::BindGroupLayout,
+ nbuf: usize,
+}
+
+/// 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 {
+ buf: Option,
+ 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> = 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>,
+ /// Largest storage buffer the device accepts (bytes).
+ pub max_buffer: u64,
+}
+
+/// WGSL constants shared by every module (generated from the CPU pipeline's values).
+fn constants() -> String {
+ use lightcraft_pipeline::finish::{GRAIN_HASH, MASK_TERMS, SRGB_LUT_N};
+ use lightcraft_pipeline::tone::{LUT_MAX_EV, LUT_MIN_EV, LUT_N};
+ let mut s = String::new();
+ let [to_lms, from_lms, to_lab, from_lab] = lightcraft_color::perceptual::oklab_matrices();
+ let to_srgb = lightcraft_color::REC2020.to_space(&lightcraft_color::SRGB).to_f32();
+ for (name, m) in [("OK_TO_LMS", to_lms), ("OK_FROM_LMS", from_lms), ("OK_TO_LAB", to_lab), ("OK_FROM_LAB", from_lab), ("TO_SRGB", to_srgb)] {
+ let rows: Vec = m.iter().map(|r| format!("vec3({:?}, {:?}, {:?})", r[0], r[1], r[2])).collect();
+ s += &format!("const {name} = array, 3>({});\n", rows.join(", "));
+ }
+ s += &format!("const TONE_MIN_EV: f32 = {LUT_MIN_EV:?};\nconst TONE_MAX_EV: f32 = {LUT_MAX_EV:?};\nconst TONE_N: u32 = {LUT_N}u;\n");
+ s += &format!("const TONE_MIN_GAIN: f32 = {:?};\n", 2f32.powf(LUT_MIN_EV));
+ let b = lightcraft_pipeline::geometry::BLANK;
+ s += &format!("const BLANK_R: f32 = {:?};\nconst BLANK_G: f32 = {:?};\nconst BLANK_B: f32 = {:?};\n", b[0], b[1], b[2]);
+ s += &format!("const SRGB_N: u32 = {SRGB_LUT_N}u;\nconst CURVE_N: u32 = {}u;\nconst MASK_TERMS: u32 = {MASK_TERMS}u;\n", crate::params::CURVE_N);
+ for (i, h) in GRAIN_HASH.iter().enumerate() {
+ s += &format!("const GRAIN_H{i}: u32 = {h}u;\n");
+ }
+ for (name, idx) in crate::params::finish_fields() {
+ s += &format!("const F_{name}: u32 = {idx}u;\n");
+ }
+ s
+}
+
+fn module_source(m: &Module, consts: &str) -> String {
+ let mut s = String::from(consts);
+ s += "@group(0) @binding(0) var P: array;\n";
+ for (i, (name, write, ty)) in m.bindings.iter().enumerate() {
+ let access = if *write { "read_write" } else { "read" };
+ s += &format!("@group(0) @binding({}) var {name}: array<{ty}>;\n", i + 1);
+ }
+ s += include_str!("wgsl/common.wgsl");
+ s += m.src;
+ s
+}
+
+impl Gpu {
+ /// Create a device on the best available adapter (no software fallback), or `None`.
+ pub fn new() -> Option {
+ let mut desc = wgpu::InstanceDescriptor::new_without_display_handle();
+ desc.backends = wgpu::Backends::PRIMARY;
+ let instance = wgpu::Instance::new(desc);
+ let adapter = pollster::block_on(
+ instance.request_adapter(&wgpu::RequestAdapterOptions { power_preference: wgpu::PowerPreference::HighPerformance, ..Default::default() }),
+ )
+ .ok()?;
+ let info = adapter.get_info();
+ if info.device_type == wgpu::DeviceType::Cpu {
+ log::info!("gpu: only a software adapter ({}), using the CPU pipeline", info.name);
+ return None;
+ }
+ let limits = adapter.limits();
+ if limits.max_storage_buffers_per_shader_stage < 10 {
+ log::info!("gpu: {} has too few storage buffers per stage", info.name);
+ return None;
+ }
+ let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor {
+ label: Some("lightcraft"),
+ required_limits: limits.clone(),
+ ..Default::default()
+ }))
+ .ok()?;
+ // Errors outside an error scope (e.g. out of memory during a render): log, and stop using
+ // the GPU (the render in flight is redone on the CPU).
+ device.on_uncaptured_error(std::sync::Arc::new(|e| {
+ log::error!("gpu: {e}");
+ crate::mark_broken();
+ }));
+ let scope = device.push_error_scope(wgpu::ErrorFilter::Validation);
+ let consts = constants();
+ let mut kernels = HashMap::new();
+ for m in MODULES {
+ let src = module_source(m, &consts);
+ let module = device.create_shader_module(wgpu::ShaderModuleDescriptor { label: None, source: wgpu::ShaderSource::Wgsl(src.into()) });
+ let mut entries = vec![storage_entry(0, true)];
+ entries.extend(m.bindings.iter().enumerate().map(|(i, (_, w, _))| storage_entry(i as u32 + 1, !w)));
+ let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: None, entries: &entries });
+ let pl = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+ label: None,
+ bind_group_layouts: &[Some(&layout)],
+ immediate_size: 0,
+ });
+ for e in m.entries {
+ let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+ label: Some(e),
+ layout: Some(&pl),
+ module: &module,
+ entry_point: Some(e),
+ compilation_options: Default::default(),
+ cache: None,
+ });
+ kernels.insert(*e, Kernel { pipeline, layout: layout.clone(), nbuf: m.bindings.len() });
+ }
+ }
+ if let Some(e) = pollster::block_on(scope.pop()) {
+ log::error!("gpu: kernels failed to build on {}: {e}", info.name);
+ return None;
+ }
+ let dummy = device.create_buffer(&wgpu::BufferDescriptor {
+ label: Some("dummy"),
+ size: 16,
+ usage: wgpu::BufferUsages::STORAGE,
+ mapped_at_creation: false,
+ });
+ 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, free: Default::default(), max_buffer })
+ }
+
+ /// 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 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: Some(buf), len }
+ }
+
+ /// Submit command buffers, then make the buffers this thread dropped before reusable.
+ pub fn submit(&self, cmds: impl IntoIterator- ) {
+ 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.
+ pub fn upload(&self, data: &[T]) -> Buf {
+ let bytes: &[u8] = bytemuck::cast_slice(data);
+ let len = bytes.len() / 4;
+ if bytes.is_empty() {
+ return self.buffer(1);
+ }
+ let buf = self.device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+ label: None,
+ contents: bytes,
+ usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::COPY_DST,
+ });
+ Buf { buf: Some(buf), len }
+ }
+
+ /// Whether a buffer of `len` 32-bit values fits the device limits.
+ pub fn fits(&self, len: usize) -> bool {
+ (len as u64) * 4 <= self.max_buffer
+ }
+
+ pub fn encoder(&self) -> wgpu::CommandEncoder {
+ self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None })
+ }
+
+ /// Record kernel `name` with parameters `p` over buffers `bufs` (None = unused binding).
+ pub fn run(&self, enc: &mut wgpu::CommandEncoder, name: &str, p: &[u32], bufs: &[Option<&Buf>], groups: [u32; 3]) {
+ let k = self.kernels.get(name).unwrap_or_else(|| panic!("unknown kernel {name}"));
+ assert_eq!(bufs.len(), k.nbuf, "{name}: binding count");
+ let params = self.device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+ label: None,
+ contents: bytemuck::cast_slice(if p.is_empty() { &[0u32] } else { p }),
+ usage: wgpu::BufferUsages::STORAGE,
+ });
+ 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.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 });
+ let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some(name), timestamp_writes: None });
+ pass.set_pipeline(&k.pipeline);
+ pass.set_bind_group(0, &bg, &[]);
+ pass.dispatch_workgroups(groups[0], groups[1], groups[2]);
+ }
+
+ /// Submit `enc`, then read back the first `len` 32-bit values of `src`.
+ pub fn finish_and_read(&self, mut enc: wgpu::CommandEncoder, src: &Buf, len: usize) -> Vec {
+ assert!(len <= src.len, "readback of {len} values from a buffer of {}", src.len);
+ let size = (len * 4) as u64;
+ let staging = self.device.create_buffer(&wgpu::BufferDescriptor {
+ label: Some("readback"),
+ size: size.max(4),
+ usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
+ mapped_at_creation: false,
+ });
+ 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| {
+ let _ = tx.send(r);
+ });
+ let _ = self.device.poll(wgpu::PollType::wait_indefinitely());
+ match rx.recv() {
+ Ok(Ok(())) => {}
+ other => panic!("gpu readback failed: {other:?}"),
+ }
+ let data = slice.get_mapped_range().expect("mapped");
+ let out: Vec = bytemuck::cast_slice(&data).to_vec();
+ drop(data);
+ staging.unmap();
+ out
+ }
+}
+
+fn storage_entry(binding: u32, read_only: bool) -> wgpu::BindGroupLayoutEntry {
+ wgpu::BindGroupLayoutEntry {
+ binding,
+ visibility: wgpu::ShaderStages::COMPUTE,
+ ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only }, has_dynamic_offset: false, min_binding_size: None },
+ count: None,
+ }
+}
+
+/// Workgroups covering a `w × h` grid with `wg`-sized groups.
+pub fn groups2(w: usize, h: usize, wg: [usize; 2]) -> [u32; 3] {
+ [w.div_ceil(wg[0]).max(1) as u32, h.div_ceil(wg[1]).max(1) as u32, 1]
+}
+
+/// Workgroups for a 1-D kernel of `n` threads (256 per group) on a 2-D grid (see `lin_index`).
+pub fn groups1(n: usize) -> [u32; 3] {
+ let g = n.div_ceil(256).max(1);
+ let x = g.min(32768);
+ [x as u32, g.div_ceil(x) as u32, 1]
+}
diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs
new file mode 100644
index 00000000..bdfd5b3e
--- /dev/null
+++ b/crates/gpu/src/lib.rs
@@ -0,0 +1,119 @@
+//! The LightCraft develop pipeline on the GPU (wgpu compute, WGSL kernels).
+//!
+//! The CPU pipeline (`lightcraft-pipeline`) is the reference: every kernel here is a port of a CPU
+//! stage, both read the same resolved parameters ([`lightcraft_pipeline::Plan`],
+//! [`lightcraft_pipeline::finish::FinishParams`]), and the equivalence tests (`tests/`) render the
+//! same settings on both and bound the difference in 8-bit sRGB. Stages without a kernel run on the
+//! CPU inside the same render (per-stage hybrid); see `docs/gpu-pipeline.md`.
+//!
+//! Use [`render`]: it returns `None` when the GPU is unavailable, disabled (`LIGHTCRAFT_GPU=0` or
+//! [`set_enabled`]), or the render does not fit the device — callers then render on the CPU.
+//! The browser build has no GPU path yet (WebGPU device creation is asynchronous): everything here
+//! compiles to the CPU fallback on wasm32.
+#![forbid(unsafe_code)]
+
+use std::sync::Arc;
+use std::sync::atomic::{AtomicBool, Ordering};
+
+use lightcraft_develop::DevelopSettings;
+use lightcraft_pipeline::{RenderRequest, Rendered, SourceInfo, StageCache};
+use lightcraft_raster::Rgb32f;
+
+#[cfg(not(target_arch = "wasm32"))]
+mod ctx;
+#[cfg(not(target_arch = "wasm32"))]
+mod params;
+#[cfg(not(target_arch = "wasm32"))]
+mod render;
+
+#[cfg(not(target_arch = "wasm32"))]
+pub use render::GpuStages;
+
+static ENABLED: AtomicBool = AtomicBool::new(true);
+/// Set when a GPU render failed: the process stays on the CPU from then on.
+static BROKEN: AtomicBool = AtomicBool::new(false);
+
+/// Stop using the GPU for the rest of the process (after a device error).
+#[cfg_attr(target_arch = "wasm32", allow(dead_code))]
+fn mark_broken() {
+ BROKEN.store(true, Ordering::Relaxed);
+}
+
+/// Allow or forbid GPU rendering at runtime (a preference). `LIGHTCRAFT_GPU=0` forbids it for the
+/// whole process regardless.
+pub fn set_enabled(on: bool) {
+ ENABLED.store(on, Ordering::Relaxed);
+}
+
+/// Whether GPU rendering is allowed (environment, preference, no earlier failure).
+pub fn enabled() -> bool {
+ ENABLED.load(Ordering::Relaxed) && !BROKEN.load(Ordering::Relaxed) && !env_disabled()
+}
+
+fn env_disabled() -> bool {
+ static OFF: std::sync::OnceLock = std::sync::OnceLock::new();
+ *OFF.get_or_init(|| std::env::var("LIGHTCRAFT_GPU").is_ok_and(|v| matches!(v.trim(), "0" | "off" | "false" | "no")))
+}
+
+#[cfg(not(target_arch = "wasm32"))]
+pub(crate) fn device() -> Option<&'static ctx::Gpu> {
+ static GPU: std::sync::OnceLock> = std::sync::OnceLock::new();
+ if env_disabled() {
+ return None;
+ }
+ GPU.get_or_init(|| std::panic::catch_unwind(ctx::Gpu::new).ok().flatten()).as_ref()
+}
+
+/// Whether a usable GPU adapter exists and GPU rendering is enabled (creates the device on first
+/// use).
+pub fn available() -> bool {
+ #[cfg(not(target_arch = "wasm32"))]
+ {
+ enabled() && device().is_some()
+ }
+ #[cfg(target_arch = "wasm32")]
+ {
+ false
+ }
+}
+
+/// The adapter's name and backend, e.g. "Apple M4 Pro (Metal)".
+pub fn adapter_name() -> Option {
+ #[cfg(not(target_arch = "wasm32"))]
+ {
+ device().map(|g| format!("{} ({:?})", g.info.name, g.info.backend))
+ }
+ #[cfg(target_arch = "wasm32")]
+ {
+ None
+ }
+}
+
+/// Render `src` with `s` on the GPU, reusing the device-resident stages kept with `stages` (the
+/// view's CPU stage cache). `None`: render on the CPU instead.
+pub fn render(src: &Arc, info: &SourceInfo, s: &DevelopSettings, req: &RenderRequest, stages: Option<&StageCache>) -> Option {
+ #[cfg(not(target_arch = "wasm32"))]
+ {
+ if !enabled() {
+ return None;
+ }
+ let gpu = device()?;
+ let ext = stages.map(|c| c.extension::());
+ let r = std::panic::catch_unwind(std::panic::AssertUnwindSafe(|| render::render(gpu, src, info, s, req, ext.as_deref())));
+ match r {
+ // a device error during the render: its result is not trustworthy
+ Ok(_) if BROKEN.load(Ordering::Relaxed) => None,
+ Ok(r) => r,
+ Err(_) => {
+ log::error!("gpu: render failed; using the CPU pipeline from now on");
+ BROKEN.store(true, Ordering::Relaxed);
+ None
+ }
+ }
+ }
+ #[cfg(target_arch = "wasm32")]
+ {
+ let _ = (src, info, s, req, stages);
+ None
+ }
+}
diff --git a/crates/gpu/src/params.rs b/crates/gpu/src/params.rs
new file mode 100644
index 00000000..64d722d7
--- /dev/null
+++ b/crates/gpu/src/params.rs
@@ -0,0 +1,211 @@
+//! Parameter blocks: the per-pixel stage's [`FinishParams`] packed into the 32-bit words the
+//! `finish` kernel reads (`F_*` constants, generated from [`FIELDS`]), plus its auxiliary tables.
+
+use lightcraft_develop::VignetteStyle;
+use lightcraft_pipeline::finish::{FinishParams, MASK_TERMS, srgb_lut};
+
+/// Entries per tone-curve table (`lightcraft_pipeline::finish` builds them at this size).
+pub const CURVE_N: u32 = 1024;
+
+/// Field name and number of 32-bit words.
+const FIELDS: &[(&str, usize)] = &[
+ ("W", 1),
+ ("H", 1),
+ ("NMASK", 1),
+ ("MASK_OFF", 1),
+ ("SRGB_OFF", 1),
+ ("CURVE_OFF", 1),
+ ("CURVES", 1),
+ ("GAIN", 1),
+ ("EV", 1),
+ ("AIR", 1),
+ ("AIR_PRE", 1),
+ ("HL", 1),
+ ("SH", 1),
+ ("CLAR", 1),
+ ("TEX", 1),
+ ("DEHAZE", 1),
+ ("SHARPEN", 1),
+ ("SHARPEN_MASK", 1),
+ ("HAS_CLAR", 1),
+ ("HAS_TEX", 1),
+ ("HAS_DARK", 1),
+ ("OPS_IDENTITY", 1),
+ ("VIBRANCE", 1),
+ ("SATURATION", 1),
+ ("MIXER", 1),
+ ("MIX_HUE", 8),
+ ("MIX_SAT", 8),
+ ("MIX_LUM", 8),
+ ("BW", 1),
+ ("BW_MIX", 8),
+ ("GRADING", 1),
+ ("WHEELS", 12),
+ ("BLENDING", 1),
+ ("BALANCE", 1),
+ ("SKIN", 1),
+ ("BANDH", 8),
+ ("VIG", 1),
+ ("VIG_AMOUNT", 1),
+ ("VIG_START", 1),
+ ("VIG_WIDTH", 1),
+ ("VIG_ASPECT_MIX", 1),
+ ("VIG_POWER", 1),
+ ("VIG_HL", 1),
+ ("VIG_STYLE", 1),
+ ("GRAIN", 1),
+ ("GRAIN_AMT", 1),
+ ("GRAIN_SC", 1),
+ ("GRAIN_ROUGH", 1),
+ ("GRAIN_SEED", 1),
+ ("GRAIN_AFF", 6),
+];
+
+/// `(name, index)` of every field (for the WGSL constants).
+pub fn finish_fields() -> Vec<(&'static str, usize)> {
+ let mut i = 0;
+ FIELDS
+ .iter()
+ .map(|(n, len)| {
+ let r = (*n, i);
+ i += len;
+ r
+ })
+ .collect()
+}
+
+fn index(name: &str) -> usize {
+ let mut i = 0;
+ for (n, len) in FIELDS {
+ if *n == name {
+ return i;
+ }
+ i += len;
+ }
+ panic!("unknown finish field {name}")
+}
+
+/// A parameter block being filled by name.
+pub struct Block(pub Vec);
+
+impl Block {
+ fn new() -> Block {
+ Block(vec![0; FIELDS.iter().map(|f| f.1).sum()])
+ }
+ fn f(&mut self, name: &str, v: f32) {
+ let i = index(name);
+ self.0[i] = v.to_bits();
+ }
+ fn u(&mut self, name: &str, v: u32) {
+ let i = index(name);
+ self.0[i] = v;
+ }
+ fn b(&mut self, name: &str, v: bool) {
+ self.u(name, v as u32);
+ }
+ fn fs(&mut self, name: &str, v: &[f32]) {
+ let i = index(name);
+ for (k, x) in v.iter().enumerate() {
+ self.0[i + k] = x.to_bits();
+ }
+ }
+}
+
+/// Which optional planes are bound.
+pub struct Present {
+ pub clarity: bool,
+ pub texture: bool,
+ pub dark: bool,
+}
+
+/// The `finish` kernel's parameter block and auxiliary table (tone LUT | sRGB LUT | curve LUTs |
+/// mask terms) for `fp` with `masks` (their local adjustments' terms).
+pub fn finish_block(fp: &FinishParams, masks: &[[f32; MASK_TERMS]], present: &Present) -> (Vec, Vec) {
+ let mut aux: Vec = fp.tone.lut().to_vec();
+ let srgb_off = aux.len();
+ aux.extend_from_slice(&srgb_lut()[..]);
+ let curve_off = aux.len();
+ if let Some(c) = &fp.curves {
+ for l in c {
+ assert_eq!(l.v.len(), CURVE_N as usize);
+ aux.extend_from_slice(&l.v);
+ }
+ }
+ let mask_off = aux.len();
+ for m in masks {
+ aux.extend_from_slice(m);
+ }
+
+ let mut p = Block::new();
+ p.u("W", fp.w as u32);
+ p.u("H", fp.h as u32);
+ p.u("NMASK", masks.len() as u32);
+ p.u("MASK_OFF", mask_off as u32);
+ p.u("SRGB_OFF", srgb_off as u32);
+ p.u("CURVE_OFF", curve_off as u32);
+ p.b("CURVES", fp.curves.is_some());
+ p.f("GAIN", fp.gain);
+ p.f("EV", fp.ev);
+ p.f("AIR", fp.air);
+ p.f("AIR_PRE", fp.air_pre);
+ p.f("HL", fp.hl);
+ p.f("SH", fp.sh);
+ p.f("CLAR", fp.clar);
+ p.f("TEX", fp.tex);
+ p.f("DEHAZE", fp.dehaze);
+ p.f("SHARPEN", fp.sharpen);
+ p.f("SHARPEN_MASK", fp.sharpen_mask);
+ p.b("HAS_CLAR", present.clarity);
+ p.b("HAS_TEX", present.texture);
+ p.b("HAS_DARK", present.dark);
+
+ let ops = &fp.ops;
+ p.b("OPS_IDENTITY", ops.is_identity());
+ p.f("VIBRANCE", ops.vibrance);
+ p.f("SATURATION", ops.saturation);
+ p.b("MIXER", ops.mixer);
+ p.fs("MIX_HUE", &ops.hue);
+ p.fs("MIX_SAT", &ops.sat);
+ p.fs("MIX_LUM", &ops.lum);
+ p.b("BW", ops.bw.is_some());
+ p.fs("BW_MIX", &ops.bw.unwrap_or([0.0; 8]));
+ if let Some((wheels, blending, balance)) = &ops.grading {
+ p.b("GRADING", true);
+ let w: Vec = wheels.iter().flat_map(|k| [k.a, k.b, k.lum]).collect();
+ p.fs("WHEELS", &w);
+ p.f("BLENDING", *blending);
+ p.f("BALANCE", *balance);
+ }
+ p.f("SKIN", ops.skin);
+ p.fs("BANDH", lightcraft_pipeline::colorops::band_hues());
+
+ if let Some(v) = &fp.vig {
+ p.b("VIG", true);
+ p.f("VIG_AMOUNT", v.amount);
+ p.f("VIG_START", v.start);
+ p.f("VIG_WIDTH", v.width);
+ p.f("VIG_ASPECT_MIX", v.aspect_mix);
+ p.f("VIG_POWER", v.power);
+ p.f("VIG_HL", v.highlights);
+ p.u(
+ "VIG_STYLE",
+ match v.style {
+ VignetteStyle::HighlightPriority => 1,
+ VignetteStyle::PaintOverlay => 2,
+ VignetteStyle::ColorPriority => 0,
+ },
+ );
+ }
+ if let Some((amt, cell, rough, seed)) = fp.grain {
+ p.b("GRAIN", true);
+ p.f("GRAIN_AMT", amt);
+ p.f("GRAIN_SC", fp.px_per_long as f32 / cell);
+ p.f("GRAIN_ROUGH", rough);
+ p.u("GRAIN_SEED", seed);
+ let long = fp.ow.max(fp.oh);
+ let [a, b, c, d, e, f] = fp.out_to_norm.0;
+ let (kx, ky) = (fp.ow / long, fp.oh / long);
+ p.fs("GRAIN_AFF", &[(a * kx) as f32, (b * ky) as f32, (c * kx) as f32, (d * ky) as f32, (e * kx) as f32, (f * ky) as f32]);
+ }
+ (p.0, aux)
+}
diff --git a/crates/gpu/src/render.rs b/crates/gpu/src/render.rs
new file mode 100644
index 00000000..72c51f5f
--- /dev/null
+++ b/crates/gpu/src/render.rs
@@ -0,0 +1,773 @@
+//! The GPU renderer: the CPU pipeline's stages, evaluated on the device where a kernel exists and
+//! on the CPU otherwise (per-stage hybrid). Stage results stay on the device and are cached per
+//! view exactly like [`lightcraft_pipeline::StageCache`] (same keys, see [`lightcraft_pipeline::Plan`]).
+
+use std::sync::{Arc, Mutex};
+
+use lightcraft_develop::DevelopSettings;
+use lightcraft_geom::Orientation;
+use lightcraft_pipeline::finish::{FinishParams, MASK_TERMS, mask_terms};
+use lightcraft_pipeline::geometry::SampleMode;
+use lightcraft_pipeline::{Plan, RenderRequest, Rendered, SourceInfo, local};
+use lightcraft_raster::resample::Filter;
+use lightcraft_raster::{Histogram, Plane, Rgb32f, Rgba8};
+
+use crate::ctx::{Buf, Gpu, groups1, groups2};
+use crate::params::{Present, finish_block};
+
+/// Device-resident stages of one view, kept next to the CPU [`lightcraft_pipeline::StageCache`]
+/// (as its extension). Holds the last two output sizes, like the CPU cache.
+#[derive(Default)]
+pub struct GpuStages {
+ entries: Mutex>,
+ /// The uploaded source (a view renders one photo at several sizes).
+ source: Mutex, Arc)>>,
+}
+
+const CAPACITY: usize = 2;
+
+#[derive(Clone)]
+struct Entry {
+ src: Arc,
+ geo: u64,
+ sampled: Arc,
+ lin: Option<(u64, Arc)>,
+ planes: Planes,
+}
+
+/// Spatial planes of one linear image (`key` = its `lin_key`), each tagged with its radius.
+#[derive(Clone, Default)]
+struct Planes {
+ key: u64,
+ log_l: Option>,
+ base: Option<(u32, Arc)>,
+ clarity: Option<(u32, Arc)>,
+ texture: Option<(u32, Arc)>,
+ dark: Option<(u32, Arc, f32)>,
+}
+
+impl GpuStages {
+ fn get(&self, src: &Arc, geo: u64) -> Option {
+ self.lock().iter().find(|e| e.geo == geo && Arc::ptr_eq(&e.src, src)).cloned()
+ }
+
+ fn put(&self, e: Entry) {
+ let mut v = self.lock();
+ v.retain(|o| !(o.geo == e.geo && Arc::ptr_eq(&o.src, &e.src)));
+ v.push(e);
+ while v.len() > CAPACITY {
+ v.remove(0);
+ }
+ }
+
+ fn source(&self, src: &Arc, upload: impl FnOnce() -> Buf) -> Arc {
+ let mut g = self.source.lock().unwrap_or_else(|e| e.into_inner());
+ match &*g {
+ Some((s, b)) if Arc::ptr_eq(s, src) => b.clone(),
+ _ => {
+ *g = None; // free the previous photo's buffer first
+ let b = Arc::new(upload());
+ *g = Some((src.clone(), b.clone()));
+ b
+ }
+ }
+ }
+
+ fn lock(&self) -> std::sync::MutexGuard<'_, Vec> {
+ self.entries.lock().unwrap_or_else(|e| e.into_inner())
+ }
+
+ /// Number of cached output sizes.
+ pub fn len(&self) -> usize {
+ self.lock().len()
+ }
+
+ pub fn is_empty(&self) -> bool {
+ self.len() == 0
+ }
+}
+
+/// A render in progress: the device and the commands recorded so far. Reading a result back
+/// submits everything recorded before it.
+pub(crate) struct Cx<'a> {
+ pub gpu: &'a Gpu,
+ enc: Option,
+}
+
+impl<'a> Cx<'a> {
+ pub fn new(gpu: &'a Gpu) -> Cx<'a> {
+ Cx { gpu, enc: Some(gpu.encoder()) }
+ }
+
+ /// Record kernel `name` (see [`Gpu::run`]).
+ pub fn run(&mut self, name: &str, p: &[u32], bufs: &[Option<&Buf>], groups: [u32; 3]) {
+ let enc = self.enc.get_or_insert_with(|| self.gpu.encoder());
+ self.gpu.run(enc, name, p, bufs, groups);
+ }
+
+ /// Submit what is recorded and read back `len` values of `b`.
+ pub fn read(&mut self, b: &Buf, len: usize) -> Vec {
+ let enc = self.enc.take().unwrap_or_else(|| self.gpu.encoder());
+ 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.raw(), 0, out.raw(), 0, (b.len * 4) as u64);
+ out
+ }
+
+ pub fn read_rgb(&mut self, b: &Buf, w: usize, h: usize) -> Rgb32f {
+ Rgb32f { width: w, height: h, data: self.read(b, w * h * 3) }
+ }
+
+ pub fn read_plane(&mut self, b: &Buf, w: usize, h: usize) -> Plane {
+ Plane { width: w, height: h, data: self.read(b, w * h) }
+ }
+}
+
+pub(crate) fn rgb_words(img: &Rgb32f) -> &[f32] {
+ bytemuck::cast_slice(&img.data)
+}
+
+// ---------------------------------------------------------------------------------------------
+// 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
+/// 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
+/// (`lightcraft_raster::blur::gaussian`).
+pub(crate) fn gaussian(cx: &mut Cx<'_>, src: &Buf, w: usize, h: usize, nc: usize, sigma: f32) -> Buf {
+ if sigma <= 0.3 || w * h == 0 {
+ 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();
+ 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];
+ 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.
+fn taps(src: usize, dst: usize, filter: Filter) -> Vec {
+ let w = lightcraft_raster::resample::weights(src, dst, filter);
+ let mut t = Vec::with_capacity(dst * 3);
+ let mut weights = Vec::new();
+ let base = dst * 3;
+ for (lo, ws) in &w {
+ t.extend_from_slice(&[*lo as u32, ws.len() as u32, (base + weights.len()) as u32]);
+ weights.extend(ws.iter().map(|v| v.to_bits()));
+ }
+ t.extend(weights);
+ t
+}
+
+/// `lightcraft_raster::resample::resize` of an `nc`-channel image.
+pub(crate) fn resize(cx: &mut Cx<'_>, src: &Buf, (sw, sh): (usize, usize), (dw, dh): (usize, usize), nc: usize, filter: Filter) -> Buf {
+ let (dw, dh) = (dw.max(1), dh.max(1));
+ if (sw, sh) == (dw, dh) {
+ return cx.copy(src);
+ }
+ let tx = cx.gpu.upload(&taps(sw, dw, filter));
+ let tmp = cx.gpu.buffer(dw * sh * nc);
+ cx.run("resize_h", &[sw as u32, sh as u32, dw as u32, nc as u32], &[Some(src), Some(&tx), Some(&tmp)], groups2(dw, sh, [16, 16]));
+ let ty = cx.gpu.upload(&taps(sh, dh, filter));
+ let out = cx.gpu.buffer(dw * dh * nc);
+ cx.run("resize_v", &[dw as u32, sh as u32, dh as u32, nc as u32], &[Some(&tmp), Some(&ty), Some(&out)], groups2(dw, dh, [16, 16]));
+ out
+}
+
+/// An element-wise kernel of the `map` module over `n` items.
+pub(crate) fn map(cx: &mut Cx<'_>, k: &str, n: usize, extra: &[u32], ins: [Option<&Buf>; 3], out: &Buf) {
+ let mut p = vec![n as u32];
+ p.extend_from_slice(extra);
+ cx.run(k, &p, &[ins[0], ins[1], ins[2], Some(out)], groups1(n));
+}
+
+/// Self-guided filter of a plane (`local::guided`).
+fn guided(cx: &mut Cx<'_>, p: &Buf, w: usize, h: usize, sigma: f32, eps: f32) -> Buf {
+ let ab = guided_coeffs(cx, p, w, h, sigma, eps);
+ let q = cx.gpu.buffer(w * h);
+ map(cx, "guided_apply", w * h, &[], [Some(p), Some(&ab), None], &q);
+ q
+}
+
+/// Blurred (a, b) coefficients, interleaved.
+fn guided_coeffs(cx: &mut Cx<'_>, p: &Buf, w: usize, h: usize, sigma: f32, eps: f32) -> Buf {
+ let n = w * h;
+ let pp = cx.gpu.buffer(2 * n);
+ map(cx, "guided_pre", n, &[], [Some(p), None, None], &pp);
+ let m = gaussian(cx, &pp, w, h, 2, sigma);
+ let ab = cx.gpu.buffer(2 * n);
+ map(cx, "guided_ab", n, &[eps.to_bits()], [Some(&m), None, None], &ab);
+ gaussian(cx, &ab, w, h, 2, sigma)
+}
+
+/// Fast guided filter (`local::guided_fast`): coefficients on a subsampled grid.
+fn guided_fast(cx: &mut Cx<'_>, p: &Buf, w: usize, h: usize, sigma: f32, eps: f32) -> Buf {
+ let s = local::guided_fast_step(sigma);
+ if s <= 1 {
+ return guided(cx, p, w, h, sigma, eps);
+ }
+ let (lw, lh) = (w.div_ceil(s).max(1), h.div_ceil(s).max(1));
+ let lo = resize(cx, p, (w, h), (lw, lh), 1, Filter::Box);
+ let ab = guided_coeffs(cx, &lo, lw, lh, sigma / s as f32, eps);
+ let up = resize(cx, &ab, (lw, lh), (w, h), 2, Filter::Bilinear);
+ let q = cx.gpu.buffer(w * h);
+ map(cx, "guided_apply", w * h, &[], [Some(p), Some(&up), None], &q);
+ q
+}
+
+// ---------------------------------------------------------------------------------------------
+// Geometry (`Frame::sample`)
+
+/// Integer pixel map `(x, y) → (m0·x + m1·y + m2, m3·x + m4·y + m5)` from an oriented image back
+/// to the source of `w × h` (`Image::oriented`: optional horizontal flip, then quarter turns).
+pub(crate) fn orient_map(o: Orientation, w: usize, h: usize) -> [i32; 6] {
+ let (flip, turns) = o.to_parts();
+ let (w, h) = (w as i32, h as i32);
+ let mut ops: Vec [i32; 6]> = Vec::new();
+ if flip {
+ ops.push(|w, _| [-1, 0, w - 1, 0, 1, 0]);
+ }
+ let cw: fn(i32, i32) -> [i32; 6] = |_, h| [0, 1, 0, -1, 0, h - 1];
+ let half: fn(i32, i32) -> [i32; 6] = |w, h| [-1, 0, w - 1, 0, -1, h - 1];
+ match turns {
+ 1 => ops.push(cw),
+ 2 => ops.push(half),
+ 3 => {
+ ops.push(half);
+ ops.push(cw);
+ }
+ _ => {}
+ }
+ // compose: T ← T ∘ op, tracking each op's input size
+ let mut t = [1, 0, 0, 0, 1, 0];
+ let (mut cw_, mut ch_) = (w, h);
+ for op in ops {
+ let o = op(cw_, ch_);
+ t = [
+ t[0] * o[0] + t[1] * o[3],
+ t[0] * o[1] + t[1] * o[4],
+ t[0] * o[2] + t[1] * o[5] + t[2],
+ t[3] * o[0] + t[4] * o[3],
+ t[3] * o[1] + t[4] * o[4],
+ t[3] * o[2] + t[4] * o[5] + t[5],
+ ];
+ if o[0] == 0 {
+ (cw_, ch_) = (ch_, cw_);
+ }
+ }
+ t
+}
+
+fn affine_bits(a: &lightcraft_geom::Affine) -> [u32; 6] {
+ a.0.map(|v| (v as f32).to_bits())
+}
+
+/// The `sample_warp` kernel's parameters (layout documented in `geom.wgsl`).
+fn warp_params(
+ wp: &lightcraft_pipeline::optics::Warp,
+ base: (usize, usize),
+ out: (usize, usize),
+ o2t: &lightcraft_geom::Affine,
+ sx: f64,
+ sy: f64,
+) -> Vec {
+ let f = |v: f64| (v as f32).to_bits();
+ let mut p = vec![base.0 as u32, base.1 as u32, out.0 as u32, out.1 as u32];
+ p.extend(affine_bits(o2t));
+ p.extend([f(sx), f(sy), f(wp.w), f(wp.h)]);
+ let persp = wp.persp != lightcraft_geom::Homography::IDENTITY;
+ p.push(persp as u32);
+ p.extend(wp.persp_inv.0.map(f));
+ p.push(f(wp.k1));
+ p.extend(wp.ca.map(f));
+ p.push(f(wp.lens_dist));
+ let warp = wp.lens.and_then(|l| l.warp);
+ p.push(warp.is_some() as u32);
+ let w = warp.unwrap_or_default();
+ p.extend(w.planes.iter().flatten().map(|v| f(*v)));
+ p.extend([f(w.center.x), f(w.center.y), f(w.radius)]);
+ p.extend([f(wp.vig_stops), f(wp.vig_power), f(wp.lens_vig)]);
+ let vig = wp.lens.and_then(|l| l.vignette);
+ p.push(vig.is_some() as u32);
+ let v = vig.unwrap_or_default();
+ p.extend(v.k.map(f));
+ p.extend([f(v.center.x), f(v.center.y), f(v.radius)]);
+ p.push(wp.per_channel() as u32);
+ p.push(wp.has_gain() as u32);
+ debug_assert_eq!(p.len(), 65);
+ p
+}
+
+/// Resample the source into the output frame (`Frame::sample`). May return the source buffer
+/// itself when the frame is the identity.
+fn sample(cx: &mut Cx<'_>, src: &Arc, src_buf: Arc, plan: &Plan<'_>) -> Arc {
+ let (w, h) = (plan.w, plan.h);
+ let fr = &plan.frame;
+ let sp = fr.sample_plan(src.width, src.height, w, h);
+ let oriented = if fr.orient == Orientation::Normal {
+ src_buf
+ } else {
+ let out = cx.gpu.buffer(sp.ow * sp.oh * 3);
+ let mut p = vec![src.width as u32, src.height as u32, sp.ow as u32, sp.oh as u32];
+ p.extend(orient_map(fr.orient, src.width, src.height).map(|v| v as u32));
+ cx.run("orient", &p, &[Some(&src_buf), Some(&out)], groups2(sp.ow, sp.oh, [16, 16]));
+ Arc::new(out)
+ };
+ let (base, (bw, bh)) = match sp.prefilter {
+ Some(d) => (Arc::new(resize(cx, &oriented, (sp.ow, sp.oh), d, 3, Filter::Mitchell)), d),
+ None => (oriented, (sp.ow, sp.oh)),
+ };
+ match &sp.mode {
+ SampleMode::Copy => base,
+ SampleMode::Affine(xf) => {
+ let out = cx.gpu.buffer(w * h * 3);
+ let mut p = vec![bw as u32, bh as u32, w as u32, h as u32];
+ p.extend(affine_bits(xf));
+ cx.run("sample_affine", &p, &[Some(&base), Some(&out)], groups2(w, h, [16, 16]));
+ Arc::new(out)
+ }
+ SampleMode::Warp(o2t) => {
+ let wp = fr.warp.as_ref().expect("warp mode has a warp");
+ let out = cx.gpu.buffer(w * h * 3);
+ let p = warp_params(wp, (bw, bh), (w, h), o2t, sp.sx, sp.sy);
+ cx.run("sample_warp", &p, &[Some(&base), Some(&out)], groups2(w, h, [16, 16]));
+ Arc::new(out)
+ }
+ }
+}
+
+// ---------------------------------------------------------------------------------------------
+// The render
+
+/// CPU copies made along the way (for stages that still run on the CPU).
+#[derive(Default)]
+struct Host {
+ sampled: Option,
+ lin: Option>,
+ log_l: Option>,
+}
+
+fn profiling() -> bool {
+ static ON: std::sync::OnceLock = std::sync::OnceLock::new();
+ *ON.get_or_init(|| std::env::var_os("LIGHTCRAFT_PROFILE").is_some())
+}
+
+/// Render on `gpu`, reusing `stages` (if given). `None` when the render does not fit the device.
+pub fn render(
+ gpu: &Gpu,
+ src: &Arc,
+ info: &SourceInfo,
+ s: &DevelopSettings,
+ req: &RenderRequest,
+ stages: Option<&GpuStages>,
+) -> Option {
+ let mut t = profiling().then(std::time::Instant::now);
+ // Under `LIGHTCRAFT_PROFILE` each stage is submitted and waited for, so the timings are real.
+ let lap = |what: &str, t: &mut Option, cx: &mut Cx<'_>| {
+ if let Some(t) = t {
+ cx.sync();
+ eprintln!(" gpu {what}: {:.1} ms", t.elapsed().as_secs_f64() * 1e3);
+ *t = std::time::Instant::now();
+ }
+ };
+ let plan = lightcraft_pipeline::plan(src, info, s, req);
+ let (w, h) = (plan.w, plan.h);
+ let n = w * h;
+ if !gpu.fits(n * 3) {
+ return None;
+ }
+ let s = &*plan.settings;
+ 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 {
+ Some(e) => e.sampled.clone(),
+ None if gpu.fits(src.data.len() * 3) => {
+ let upload = || gpu.upload(rgb_words(src));
+ let src_buf = match stages {
+ Some(c) => c.source(src, upload),
+ None => Arc::new(upload()),
+ };
+ sample(&mut cx, src, src_buf, &plan)
+ }
+ None => {
+ let img = plan.frame.sample(src, w, h);
+ let b = Arc::new(gpu.upload(rgb_words(&img)));
+ host.sampled = Some(img);
+ b
+ }
+ };
+ 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, &mut cx);
+
+ // 3. spatial planes
+ let mut planes = match cached.map(|e| e.planes) {
+ Some(p) if p.key == plan.lin_key => p,
+ _ => Planes { key: plan.lin_key, ..Default::default() },
+ };
+ let prep = prepare(&mut cx, &lin, &plan, req, &mut planes);
+ 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, &mut cx);
+
+ // 5. per-pixel stage
+ let fp = FinishParams::new(s, &plan.frame, info, w, h, plan.px_per_long, prep.air);
+ let present = Present { clarity: prep.clarity.is_some(), texture: prep.texture.is_some(), dark: prep.dark.is_some() };
+ let (p, aux) = finish_block(&fp, &terms, &present);
+ let aux = gpu.upload(&aux);
+ let out = gpu.buffer(n);
+ cx.run(
+ "main",
+ &p,
+ &[
+ Some(&lin),
+ Some(&prep.log_l),
+ Some(&prep.base),
+ prep.clarity.as_deref(),
+ prep.texture.as_deref(),
+ prep.dark.as_deref(),
+ masks.as_ref(),
+ Some(&aux),
+ Some(&out),
+ ],
+ groups2(w, h, [16, 16]),
+ );
+ let data: Vec<[u8; 4]> = cx.read(&out, n);
+ let image = Rgba8 { width: w, height: h, data };
+ lap("finish + readback", &mut t, &mut cx);
+ let histogram = Histogram::of_srgb8(&image);
+ lap("histogram", &mut t, &mut cx);
+ Some(Rendered { image, histogram })
+}
+
+/// The white-balanced, retouched, denoised image.
+fn linear(cx: &mut Cx<'_>, sampled: &Buf, info: &SourceInfo, plan: &Plan<'_>, host: &mut Host) -> Buf {
+ let (w, h) = (plan.w, plan.h);
+ let n = w * h;
+ let s = &*plan.settings;
+ let img = if lightcraft_pipeline::lin_needs_cpu(s) {
+ // defringe / spot removal: CPU
+ let mut img = match host.sampled.take() {
+ Some(i) => i,
+ None => cx.read_rgb(sampled, w, h),
+ };
+ lightcraft_pipeline::lin_cpu(&mut img, info, plan);
+ cx.gpu.upload(rgb_words(&img))
+ } else {
+ let m = local::wb_matrix_for(info, s);
+ let mut p = vec![m.is_some() as u32];
+ p.extend(m.unwrap_or_default().iter().flatten().map(|v| v.to_bits()));
+ let out = cx.gpu.buffer(n * 3);
+ map(cx, "wb_k", n, &p, [Some(sampled), None, None], &out);
+ out
+ };
+ denoise(cx, img, plan)
+}
+
+/// Noise reduction (`local::denoise`).
+fn denoise(cx: &mut Cx<'_>, img: Buf, plan: &Plan<'_>) -> Buf {
+ let (w, h) = (plan.w, plan.h);
+ let n = w * h;
+ let (lum, col) = local::nr_params(&plan.settings, plan.src_long, w.max(h));
+ let mut img = img;
+ if let Some(nr) = lum {
+ let l = cx.gpu.buffer(n);
+ map(cx, "log_lum_k", n, &[], [Some(&img), None, None], &l);
+ let f = guided(cx, &l, w, h, nr.sigma, nr.eps);
+ let out = cx.gpu.buffer(n * 3);
+ map(cx, "nr_lum", n, &[nr.k.to_bits()], [Some(&img), Some(&l), Some(&f)], &out);
+ img = out;
+ }
+ if let Some(nr) = col {
+ let chroma = cx.gpu.buffer(n * 3);
+ map(cx, "chroma_k", n, &[], [Some(&img), None, None], &chroma);
+ let b = gaussian(cx, &chroma, w, h, 3, nr.sigma);
+ let out = cx.gpu.buffer(n * 3);
+ map(cx, "nr_col", n, &[nr.t.to_bits()], [Some(&img), Some(&chroma), Some(&b)], &out);
+ img = out;
+ }
+ img
+}
+
+/// The planes the per-pixel stage reads.
+struct Prep {
+ log_l: Arc,
+ base: Arc,
+ clarity: Option>,
+ texture: Option>,
+ dark: Option>,
+ air: f32,
+}
+
+fn plane_at(slot: &mut Option<(u32, Arc)>, sigma: f32, f: impl FnOnce() -> Buf) -> Arc {
+ match slot {
+ Some((k, p)) if *k == sigma.to_bits() => p.clone(),
+ _ => {
+ let p = Arc::new(f());
+ *slot = Some((sigma.to_bits(), p.clone()));
+ p
+ }
+ }
+}
+
+/// The spatial planes (`local::prepare`), reusing those in `planes`.
+fn prepare(cx: &mut Cx<'_>, lin: &Buf, plan: &Plan<'_>, req: &RenderRequest, planes: &mut Planes) -> Prep {
+ let (w, h) = (plan.w, plan.h);
+ let n = w * h;
+ let sig = local::plane_sigmas(&plan.settings, plan.px_per_long, req.quality);
+ let log_l = match &planes.log_l {
+ Some(b) => b.clone(),
+ None => {
+ let l = cx.gpu.buffer(n);
+ map(cx, "log_lum_k", n, &[], [Some(lin), None, None], &l);
+ let b = Arc::new(l);
+ planes.log_l = Some(b.clone());
+ b
+ }
+ };
+ let base = match sig.base {
+ Some(sg) => plane_at(&mut planes.base, sg, || guided_fast(cx, &log_l, w, h, sg, local::BASE_EPS)),
+ None => log_l.clone(),
+ };
+ let clarity = sig.clarity.map(|sg| plane_at(&mut planes.clarity, sg, || guided_fast(cx, &log_l, w, h, sg, local::CLARITY_EPS)));
+ let texture = sig.texture.map(|sg| plane_at(&mut planes.texture, sg, || gaussian(cx, &log_l, w, h, 1, sg)));
+ let (dark, air) = match sig.dark {
+ None => (None, 1.0),
+ Some(sg) => match &planes.dark {
+ Some((k, d, air)) if *k == sg.to_bits() => (Some(d.clone()), *air),
+ _ => {
+ let d0 = cx.gpu.buffer(n);
+ map(cx, "dark_k", n, &[], [Some(lin), None, None], &d0);
+ let d = gaussian(cx, &d0, w, h, 1, sg);
+ // airlight: a percentile of every 7th value, as on the CPU
+ let m = n.div_ceil(local::AIRLIGHT_STEP);
+ let sub = cx.gpu.buffer(m);
+ map(cx, "subsample", m, &[local::AIRLIGHT_STEP as u32], [Some(&d), None, None], &sub);
+ let air = local::airlight_of(cx.read(&sub, m));
+ let b = Arc::new(d);
+ planes.dark = Some((sg.to_bits(), b.clone(), air));
+ (Some(b), air)
+ }
+ },
+ };
+ Prep { log_l, base, clarity, texture, dark, air }
+}
+
+/// Most brush dabs the mask kernel evaluates per pixel (more: the CPU rasterizes the brush).
+const MAX_DABS: usize = 4096;
+
+/// Evaluate the masks (`lightcraft_pipeline::masks::evaluate`): their alpha planes (concatenated)
+/// and their adjustment terms. Shapes without a kernel (Sky, Subject, …) run on the CPU.
+fn masks(cx: &mut Cx<'_>, lin: &Buf, prep: &Prep, plan: &Plan<'_>, host: &mut Host) -> (Option, Vec<[f32; MASK_TERMS]>) {
+ use lightcraft_develop::{MaskOp, MaskShape};
+ let s = &*plan.settings;
+ let list: Vec<_> = s.masks.iter().filter(|m| m.visible && !m.components.is_empty()).collect();
+ if list.is_empty() {
+ return (None, Vec::new());
+ }
+ let (w, h) = (plan.w, plan.h);
+ let n = w * h;
+ let frame = &plan.frame;
+ let ev = s.light.exposure as f32;
+ let long = frame.ow.max(frame.oh);
+ 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.zeroed(n * list.len());
+ let comp_plane = cx.gpu.buffer(n);
+ for (mi, m) in list.iter().enumerate() {
+ let mut first = true;
+ for comp in &m.components {
+ let op = match comp.op {
+ _ if first && comp.op != MaskOp::Intersect => 0u32,
+ MaskOp::Add => 1,
+ MaskOp::Subtract => 2,
+ MaskOp::Intersect => 3,
+ };
+ first = false;
+ let mut p = vec![w as u32, h as u32, 0, comp.invert as u32];
+ p.extend_from_slice(&aff);
+ let f32s = |v: &[f64]| v.iter().map(|x| (*x as f32).to_bits()).collect::>();
+ let mut aux: Vec = Vec::new();
+ let kind = match &comp.shape {
+ MaskShape::Linear { start, end } => {
+ let (a, b) = (frame.norm_to_long(*start), frame.norm_to_long(*end));
+ let d = b - a;
+ p.extend(f32s(&[a.x, a.y, d.x, d.y, d.dot(d).max(1e-12)]));
+ Some(0)
+ }
+ MaskShape::Radial { center, rx, ry, angle, feather, invert } => {
+ let c = frame.norm_to_long(*center);
+ let (sn, co) = (-angle.to_radians()).sin_cos();
+ p.extend(f32s(&[c.x, c.y, co, sn, rx.max(1e-6), ry.max(1e-6), (*feather / 100.0).clamp(0.0, 1.0)]));
+ p.push(*invert as u32);
+ Some(1)
+ }
+ MaskShape::LuminanceRange { lo, hi, lo_feather, hi_feather } => {
+ p.extend(f32s(&[*lo, *hi, lo_feather.max(1e-3), hi_feather.max(1e-3)]));
+ p.push(ev.to_bits());
+ Some(2)
+ }
+ MaskShape::ColorRange { samples, refine } => {
+ let tol = 0.04 + 0.16 * (*refine as f32 / 100.0);
+ p.extend([tol.to_bits(), ev.exp2().to_bits(), samples.len() as u32]);
+ aux.extend(samples.iter().flat_map(|s| s.map(|v| v as f32)));
+ Some(3)
+ }
+ MaskShape::Brush { strokes } => {
+ let dabs: Vec<_> = strokes.iter().map(|st| lightcraft_pipeline::masks::brush_dabs(st, frame, w, h)).collect();
+ if dabs.iter().map(|d| d.dabs.len()).sum::() <= MAX_DABS {
+ p.push(dabs.len() as u32);
+ let mut off = 11 * dabs.len();
+ for d in &dabs {
+ let (mut x0, mut y0, mut x1, mut y1) = (f64::MAX, f64::MAX, f64::MIN, f64::MIN);
+ for q in &d.dabs {
+ (x0, y0, x1, y1) = (x0.min(q.x - d.r), y0.min(q.y - d.r), x1.max(q.x + d.r), y1.max(q.y + d.r));
+ }
+ aux.extend([off as f32, d.dabs.len() as f32, d.r as f32, d.hard as f32, d.flow, d.density, d.erase as u8 as f32]);
+ aux.extend([x0, y0, x1, y1].map(|v| v as f32));
+ off += 2 * d.dabs.len();
+ }
+ for d in &dabs {
+ aux.extend(d.dabs.iter().flat_map(|q| [q.x as f32, q.y as f32]));
+ }
+ Some(4)
+ } else {
+ None
+ }
+ }
+ _ => None,
+ };
+ let plane = match kind {
+ Some(k) => {
+ p[2] = k;
+ let aux = (!aux.is_empty()).then(|| cx.gpu.upload(&aux));
+ cx.run("shape", &p, &[Some(lin), Some(&prep.log_l), aux.as_ref(), Some(&comp_plane), Some(&alpha)], groups2(w, h, [16, 16]));
+ None
+ }
+ None => {
+ // no kernel: evaluate on the CPU
+ let img = host.lin.get_or_insert_with(|| Arc::new(cx.read_rgb(lin, w, h))).clone();
+ let l = host.log_l.get_or_insert_with(|| Arc::new(cx.read_plane(&prep.log_l, w, h))).clone();
+ let mut v = lightcraft_pipeline::masks::shape_alpha(&comp.shape, frame, w, h, &img, &l, ev);
+ if comp.invert {
+ v.data.iter_mut().for_each(|x| *x = 1.0 - *x);
+ }
+ Some(cx.gpu.upload(&v.data))
+ }
+ };
+ let cbuf = plane.as_ref().unwrap_or(&comp_plane);
+ let k = [n as u32, 0, mi as u32, op];
+ cx.run("combine", &k, &[None, None, None, Some(cbuf), Some(&alpha)], groups1(n));
+ }
+ let amt = (m.adjust.amount / 100.0) as f32;
+ cx.run(
+ "finalize",
+ &[n as u32, 0, mi as u32, m.invert as u32, amt.to_bits()],
+ &[None, None, None, Some(&comp_plane), Some(&alpha)],
+ groups1(n),
+ );
+ }
+ let terms = list.iter().map(|m| mask_terms(&m.adjust)).collect();
+ (Some(alpha), terms)
+}
+
+#[cfg(test)]
+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::*;
+ let img = Rgb32f::from_fn(5, 3, |x, y| [x as f32, y as f32, 0.0]);
+ for o in [Normal, Rotate90, Rotate180, Rotate270, FlipH, Transverse, FlipV, Transpose] {
+ let r = img.oriented(o);
+ let m = orient_map(o, 5, 3);
+ for y in 0..r.height {
+ for x in 0..r.width {
+ let (x, y) = (x as i32, y as i32);
+ let (sx, sy) = (m[0] * x + m[1] * y + m[2], m[3] * x + m[4] * y + m[5]);
+ assert_eq!(r.get(x as usize, y as usize), img.get(sx as usize, sy as usize), "{o:?} at {x},{y}");
+ }
+ }
+ }
+ }
+}
diff --git a/crates/gpu/src/wgsl/blur.wgsl b/crates/gpu/src/wgsl/blur.wgsl
new file mode 100644
index 00000000..576aff09
--- /dev/null
+++ b/crates/gpu/src/wgsl/blur.wgsl
@@ -0,0 +1,80 @@
+// 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.
+
+// Pixel `i` (all nc ≤ 3 channels at once: one contiguous read per pixel).
+fn ld(i: u32, nc: u32) -> vec3 {
+ let j = i * nc;
+ if (nc == 1u) {
+ return vec3(src[j], 0.0, 0.0);
+ }
+ if (nc == 2u) {
+ return vec3(src[j], src[j + 1u], 0.0);
+ }
+ return vec3(src[j], src[j + 1u], src[j + 2u]);
+}
+
+fn st(i: u32, nc: u32, v: vec3) {
+ 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) {
+ let w = pu(0u);
+ let h = pu(1u);
+ let nc = pu(2u);
+ let r = i32(pu(3u));
+ let ch = pu(4u);
+ let y = g.y;
+ let x0 = g.x * ch;
+ if (y >= h || x0 >= w) {
+ 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);
+ }
+}
+
+// Vertical: one thread per CH consecutive rows of a column (the CPU re-primes every BAND rows the
+// same way).
+@compute @workgroup_size(64, 4)
+fn box_v(@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 y0 = g.y * ch;
+ if (x >= w || y0 >= h) {
+ return;
+ }
+ let last = i32(h) - 1;
+ let inv = 1.0 / f32(2 * r + 1);
+ let y1 = min(y0 + ch, h);
+ var acc = vec3(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);
+ }
+}
diff --git a/crates/gpu/src/wgsl/common.wgsl b/crates/gpu/src/wgsl/common.wgsl
new file mode 100644
index 00000000..409972e4
--- /dev/null
+++ b/crates/gpu/src/wgsl/common.wgsl
@@ -0,0 +1,67 @@
+// Shared helpers, prepended to every kernel (after the generated constants and bindings).
+// Each function mirrors its CPU twin in `lightcraft-pipeline` / `lightcraft-color`; keep them in step.
+
+fn pu(i: u32) -> u32 {
+ return P[i];
+}
+
+fn pf(i: u32) -> f32 {
+ return bitcast(P[i]);
+}
+
+// Linear index of a 1-D kernel dispatched on a 2-D grid of 256-wide workgroups.
+fn lin_index(g: vec3, n: vec3) -> u32 {
+ return g.x + g.y * n.x * 256u;
+}
+
+fn sstep(e0: f32, e1: f32, x: f32) -> f32 {
+ let t = clamp((x - e0) / (e1 - e0), 0.0, 1.0);
+ return t * t * (3.0 - 2.0 * t);
+}
+
+fn rem_euclid(a: f32, b: f32) -> f32 {
+ let r = a % b;
+ return select(r, r + b, r < 0.0);
+}
+
+const PI: f32 = 3.14159265358979323846;
+const TAU: f32 = 6.28318530717958647692;
+const GREY: f32 = 0.18;
+
+fn wrap_angle(a: f32) -> f32 {
+ return rem_euclid(a + PI, TAU) - PI;
+}
+
+// Rec.2020 luminance (`lightcraft_color::luminance_2020`).
+fn lum2020(c: vec3) -> f32 {
+ return c.x * 0.2627 + c.y * 0.6780 + c.z * 0.0593;
+}
+
+// `lightcraft_pipeline::local::log_lum`.
+fn log_lum(c: vec3) -> f32 {
+ return log2(max(lum2020(c), 1e-7) / GREY);
+}
+
+fn mul3(m: array, 3>, v: vec3) -> vec3 {
+ return vec3(
+ m[0].x * v.x + m[0].y * v.y + m[0].z * v.z,
+ m[1].x * v.x + m[1].y * v.y + m[1].z * v.z,
+ m[2].x * v.x + m[2].y * v.y + m[2].z * v.z,
+ );
+}
+
+fn cbrt(x: f32) -> f32 {
+ return sign(x) * pow(abs(x), 1.0 / 3.0);
+}
+
+// Linear Rec.2020 → OkLab.
+fn oklab(c: vec3) -> vec3 {
+ let lms = mul3(OK_TO_LMS, c);
+ return mul3(OK_TO_LAB, vec3(cbrt(lms.x), cbrt(lms.y), cbrt(lms.z)));
+}
+
+// OkLab → linear Rec.2020.
+fn oklab_inv(lab: vec3) -> vec3 {
+ let l = mul3(OK_FROM_LAB, lab);
+ return mul3(OK_FROM_LMS, l * l * l);
+}
diff --git a/crates/gpu/src/wgsl/finish.wgsl b/crates/gpu/src/wgsl/finish.wgsl
new file mode 100644
index 00000000..0548ab5d
--- /dev/null
+++ b/crates/gpu/src/wgsl/finish.wgsl
@@ -0,0 +1,358 @@
+// The per-pixel stage: a straight port of `lightcraft_pipeline::finish` (keep in step with it).
+// Bindings: img (rgb, pre-exposure), log_l, base, clar, tex, dark, masks (NMASK planes), aux
+// (tone LUT | sRGB LUT | curve LUTs | mask terms), out (packed RGBA8).
+
+fn tone_apply(y: f32) -> f32 {
+ if (y <= 0.0) {
+ return 0.0;
+ }
+ let ev = log2(y / GREY);
+ let f = clamp((ev - TONE_MIN_EV) / (TONE_MAX_EV - TONE_MIN_EV), 0.0, 1.0) * f32(TONE_N - 1u);
+ let i = min(u32(f), TONE_N - 2u);
+ let t = f - f32(i);
+ let v = aux[i] + (aux[i + 1u] - aux[i]) * t;
+ if (ev < TONE_MIN_EV) {
+ return v * (y / (GREY * TONE_MIN_GAIN));
+ }
+ return v;
+}
+
+fn encode_srgb(v: f32) -> f32 {
+ let o = pu(F_SRGB_OFF);
+ let f = clamp(v, 0.0, 1.0) * f32(SRGB_N);
+ let i = min(u32(f), SRGB_N - 1u);
+ let t = f - f32(i);
+ return aux[o + i] + (aux[o + i + 1u] - aux[o + i]) * t;
+}
+
+fn curve(ch: u32, x: f32) -> f32 {
+ let o = pu(F_CURVE_OFF) + ch * CURVE_N;
+ let f = clamp(x, 0.0, 1.0) * f32(CURVE_N - 1u);
+ let i = min(u32(f), CURVE_N - 2u);
+ let t = f - f32(i);
+ return aux[o + i] + (aux[o + i + 1u] - aux[o + i]) * t;
+}
+
+fn band_weights(h: f32) -> array {
+ var w: array;
+ for (var i = 0u; i < 8u; i++) {
+ let a = pf(F_BANDH + i);
+ let b = pf(F_BANDH + (i + 1u) % 8u);
+ let span = rem_euclid(wrap_angle(b - a), TAU);
+ let d = rem_euclid(wrap_angle(h - a), TAU);
+ if (d <= span) {
+ let t = d / span;
+ let s = 0.5 - 0.5 * cos(t * PI);
+ w[i] += 1.0 - s;
+ w[(i + 1u) % 8u] += s;
+ break;
+ }
+ }
+ return w;
+}
+
+// `ColorOps::apply`.
+fn color_ops(rgb: vec3, local_sat: f32, local_hue: f32) -> vec3 {
+ if (pu(F_OPS_IDENTITY) != 0u && local_sat == 0.0 && local_hue == 0.0) {
+ return rgb;
+ }
+ let lab0 = oklab(rgb);
+ var l = lab0.x;
+ var c = sqrt(lab0.y * lab0.y + lab0.z * lab0.z);
+ var h = atan2(lab0.z, lab0.y);
+ if (pu(F_MIXER) != 0u) {
+ let w = band_weights(h);
+ var dh = 0.0;
+ var ds = 0.0;
+ var dl = 0.0;
+ for (var i = 0u; i < 8u; i++) {
+ dh += w[i] * pf(F_MIX_HUE + i);
+ ds += w[i] * pf(F_MIX_SAT + i);
+ dl += w[i] * pf(F_MIX_LUM + i);
+ }
+ let chroma_w = min(c / 0.12, 1.0);
+ h += dh * chroma_w;
+ c *= max(1.0 + ds, 0.0);
+ l += dl * chroma_w * sqrt(max(l, 0.05));
+ }
+ let vib = pf(F_VIBRANCE);
+ if (vib != 0.0) {
+ let low = 1.0 - clamp(c / 0.22, 0.0, 1.0);
+ var skin = 1.0;
+ if (vib > 0.0) {
+ let q = wrap_angle(h - pf(F_SKIN)) / 0.35;
+ skin = 1.0 - 0.6 * exp(-(q * q));
+ }
+ c *= max(1.0 + vib * low * low * skin * 1.2, 0.0);
+ }
+ let sat = pf(F_SATURATION);
+ if (sat != 0.0 || local_sat != 0.0) {
+ c *= max(1.0 + sat + local_sat, 0.0);
+ }
+ h += local_hue;
+ if (pu(F_BW) != 0u) {
+ let w = band_weights(h);
+ var mix = 0.0;
+ for (var i = 0u; i < 8u; i++) {
+ mix += w[i] * pf(F_BW_MIX + i);
+ }
+ l = max(l + mix * min(c / 0.2, 1.0) * 0.25, 0.0);
+ c = 0.0;
+ }
+ var lab = vec3(l, c * cos(h), c * sin(h));
+ if (pu(F_GRADING) != 0u) {
+ let m = 0.5 - pf(F_BALANCE) * 0.25;
+ let width = 0.15 + pf(F_BLENDING) * 0.5;
+ let ws = 1.0 - sstep(m - width, m + width * 0.25, lab.x);
+ let wh = sstep(m - width * 0.25, m + width, lab.x);
+ let wm = max(1.0 - ws - wh, 0.0);
+ let wts = array(ws, wm, wh, 1.0);
+ for (var k = 0u; k < 4u; k++) {
+ let wt = wts[k];
+ lab.y += pf(F_WHEELS + 3u * k) * wt;
+ lab.z += pf(F_WHEELS + 3u * k + 1u) * wt;
+ lab.x += pf(F_WHEELS + 3u * k + 2u) * wt;
+ }
+ }
+ return oklab_inv(lab);
+}
+
+fn ghash(i: i32, j: i32, seed: u32) -> f32 {
+ var v = (bitcast(i) * GRAIN_H0) ^ (bitcast(j) * GRAIN_H1) ^ (seed * GRAIN_H2);
+ v ^= v >> 13u;
+ v = v * GRAIN_H3;
+ v ^= v >> 15u;
+ return f32(v & 0xffffu) / 32768.0 - 1.0;
+}
+
+fn grain_noise(x: f32, y: f32, seed: u32) -> f32 {
+ let x0 = floor(x);
+ let y0 = floor(y);
+ let fx = x - x0;
+ let fy = y - y0;
+ let i = i32(x0);
+ let j = i32(y0);
+ let u = fx * fx * (3.0 - 2.0 * fx);
+ let v = fy * fy * (3.0 - 2.0 * fy);
+ let a = ghash(i, j, seed) + (ghash(i + 1, j, seed) - ghash(i, j, seed)) * u;
+ let b = ghash(i, j + 1, seed) + (ghash(i + 1, j + 1, seed) - ghash(i, j + 1, seed)) * u;
+ return a + (b - a) * v;
+}
+
+fn enc8(v: f32) -> u32 {
+ return u32(clamp(v, 0.0, 1.0) * 255.0 + 0.5);
+}
+
+@compute @workgroup_size(16, 16)
+fn main(@builtin(global_invocation_id) gid: vec3) {
+ let w = pu(F_W);
+ let h = pu(F_H);
+ if (gid.x >= w || gid.y >= h) {
+ return;
+ }
+ let x = gid.x;
+ let y = gid.y;
+ let i = y * w + x;
+ let n = w * h;
+ let gain = pf(F_GAIN);
+ let ev = pf(F_EV);
+ var c = vec3(img[3u * i], img[3u * i + 1u], img[3u * i + 2u]);
+ if (gain != 1.0) {
+ c = c * gain;
+ }
+ let l_pre = log_l[i];
+ let l0 = l_pre + ev;
+
+ // --- local (mask) contributions
+ var lt: array;
+ var tint_on = false;
+ var tint_dir = vec2(0.0, 0.0);
+ var tint_amt = 0.0;
+ let nm = pu(F_NMASK);
+ for (var m = 0u; m < nm; m++) {
+ let a = masks[m * n + i];
+ if (a <= 0.0) {
+ continue;
+ }
+ let t = pu(F_MASK_OFF) + m * MASK_TERMS;
+ for (var k = 0u; k < 14u; k++) {
+ lt[k] += a * aux[t + k];
+ }
+ if (aux[t + 14u] > 0.0) {
+ tint_on = true;
+ tint_dir = vec2(aux[t + 15u], aux[t + 16u]);
+ tint_amt = a * aux[t + 17u];
+ }
+ }
+ let l_exp = lt[0];
+ let l_temp = lt[1];
+ let l_tint = lt[2];
+
+ // --- dehaze (scene linear)
+ let dz = pf(F_DEHAZE) + lt[10];
+ if (dz != 0.0 && pu(F_HAS_DARK) != 0u) {
+ let d = clamp(dark[i] / pf(F_AIR_PRE), 0.0, 1.0);
+ let air = pf(F_AIR);
+ if (dz > 0.0) {
+ let t = max(1.0 - 0.95 * min(dz, 1.0) * d, 0.12);
+ c = max((c - air * (1.0 - t)) / t, vec3(0.0));
+ } else {
+ let k = min(-dz, 1.0) * 0.7 * (0.35 + 0.65 * d);
+ c = c + (air * 0.9 - c) * k;
+ }
+ }
+
+ // --- local exposure / temp / tint
+ if (l_exp != 0.0) {
+ c = c * exp2(l_exp);
+ }
+ if (l_temp != 0.0 || l_tint != 0.0) {
+ let y0 = lum2020(c);
+ c = vec3(c.x * (1.0 + 0.3 * l_temp), c.y * (1.0 - 0.22 * l_tint), c.z * max(1.0 - 0.3 * l_temp, 0.0));
+ let y1 = max(lum2020(c), 1e-9);
+ c = c * y0 / y1;
+ }
+
+ // --- local tone in log luminance
+ var l1 = l0;
+ if (dz != 0.0 || l_exp != 0.0) {
+ l1 = log_lum(c);
+ }
+ let shift = l1 - l0;
+ let base = base_p[i] + ev + shift;
+ var delta = 0.0;
+ let hh = pf(F_HL) + lt[4];
+ let ss = pf(F_SH) + lt[5];
+ if (hh != 0.0 || ss != 0.0) {
+ let ws = 1.0 - sstep(-4.8, 0.3, base);
+ let wh = sstep(-1.0, 2.8, base);
+ delta += ss * 1.7 * ws * sqrt(ws) + hh * 1.7 * wh;
+ }
+ if (lt[6] != 0.0) {
+ delta += lt[6] * 0.8 * sstep(0.5, 3.0, l1);
+ }
+ if (lt[7] != 0.0) {
+ delta += lt[7] * 0.8 * (1.0 - sstep(-6.0, -1.5, l1));
+ }
+ if (lt[3] != 0.0) {
+ delta += lt[3] * 0.14 * clamp(l1, -6.0, 4.0);
+ }
+ let cl = pf(F_CLAR) + lt[9];
+ if (cl != 0.0 && pu(F_HAS_CLAR) != 0u) {
+ let det = clamp(l_pre - clar[i], -2.5, 2.5);
+ let q = base / 3.2;
+ let mid = exp(-(q * q));
+ delta += cl * 0.85 * det * (0.35 + 0.65 * mid);
+ }
+ let tx = pf(F_TEX) + lt[8];
+ let sp = lt[13] * 0.6 + pf(F_SHARPEN);
+ if ((tx != 0.0 || sp != 0.0) && pu(F_HAS_TEX) != 0u) {
+ let det = l_pre - tex[i];
+ let tame = 1.0 - 0.6 * sstep(0.4, 1.6, abs(det));
+ delta += tx * 1.1 * clamp(det, -1.0, 1.0) * tame;
+ if (sp != 0.0) {
+ let sm = pf(F_SHARPEN_MASK);
+ var mk = 1.0;
+ if (sm > 0.0) {
+ mk = sstep(sm * 0.25, sm * 0.25 + 0.15, abs(det));
+ }
+ delta += sp * 1.3 * clamp(det, -0.8, 0.8) * mk;
+ }
+ }
+ if (delta != 0.0) {
+ c = c * exp2(delta);
+ }
+
+ // --- tone map on luminance, highlight desaturation
+ let yl = lum2020(c);
+ let o = tone_apply(yl);
+ var d = vec3(0.0);
+ if (yl > 1e-9) {
+ d = c * o / yl;
+ }
+ let mx = max(d.x, max(d.y, d.z));
+ if (mx > 1.0) {
+ let t = clamp((mx - 1.0) / max(mx - o, 1e-6), 0.0, 1.0);
+ d = d + (o - d) * t;
+ }
+
+ // --- colour
+ d = color_ops(d, lt[11], lt[12]);
+ if (tint_on) {
+ let lab = oklab(d);
+ d = oklab_inv(vec3(lab.x, lab.y + tint_dir.x * 0.08 * tint_amt, lab.z + tint_dir.y * 0.08 * tint_amt));
+ }
+
+ // --- vignette (display linear, post-crop)
+ if (pu(F_VIG) != 0u) {
+ let fw = f32(w);
+ let fh = f32(h);
+ let aspect = fw / fh;
+ let amount = pf(F_VIG_AMOUNT);
+ let mixa = pf(F_VIG_ASPECT_MIX);
+ let power = pf(F_VIG_POWER);
+ let u = (f32(x) + 0.5) / fw * 2.0 - 1.0;
+ let vv = (f32(y) + 0.5) / fh * 2.0 - 1.0;
+ let sx = 1.0 + (aspect - 1.0) * mixa;
+ let sy = 1.0 + (1.0 / aspect - 1.0) * mixa;
+ let ax = abs(u * max(sx, 1.0) / max(sx, sy));
+ let ay = abs(vv * max(sy, 1.0) / max(sx, sy));
+ let dist = pow(pow(ax, power) + pow(ay, power), 1.0 / power);
+ let start = pf(F_VIG_START);
+ let t = sstep(start, start + pf(F_VIG_WIDTH), dist);
+ if (t > 0.0) {
+ let lum = clamp(lum2020(d), 0.0, 1.0);
+ if (amount < 0.0) {
+ var f = 1.0 + amount * t;
+ let style = pu(F_VIG_STYLE);
+ if (style == 1u) {
+ f += (1.0 - f) * pf(F_VIG_HL) * sstep(0.4, 1.0, lum);
+ }
+ if (style == 2u) {
+ d = d * (1.0 - (-amount) * t) + 0.0;
+ } else {
+ d = d * f;
+ }
+ } else {
+ d = d + (1.0 - d) * amount * t * 0.85;
+ }
+ }
+ }
+
+ // --- gamut map to sRGB (desaturate towards luminance until in range)
+ var r = mul3(TO_SRGB, d);
+ let yy = clamp(0.2126 * r.x + 0.7152 * r.y + 0.0722 * r.z, 0.0, 1.0);
+ var tg = 1.0;
+ for (var k = 0u; k < 3u; k++) {
+ let cc = r[k];
+ if (cc < 0.0) {
+ tg = min(tg, yy / max(yy - cc, 1e-9));
+ } else if (cc > 1.0) {
+ tg = min(tg, (1.0 - yy) / max(cc - yy, 1e-9));
+ }
+ }
+ if (tg < 1.0) {
+ r = yy + (r - yy) * tg;
+ }
+
+ // --- encode, curves, grain
+ var e = vec3(encode_srgb(r.x), encode_srgb(r.y), encode_srgb(r.z));
+ if (pu(F_CURVES) != 0u) {
+ e = vec3(curve(0u, e.x), curve(1u, e.y), curve(2u, e.z));
+ }
+ if (pu(F_GRAIN) != 0u) {
+ let px = f32(x) + 0.5;
+ let py = f32(y) + 0.5;
+ let gx = pf(F_GRAIN_AFF) * px + pf(F_GRAIN_AFF + 2u) * py + pf(F_GRAIN_AFF + 4u);
+ let gy = pf(F_GRAIN_AFF + 1u) * px + pf(F_GRAIN_AFF + 3u) * py + pf(F_GRAIN_AFF + 5u);
+ let sc = pf(F_GRAIN_SC);
+ let rough = pf(F_GRAIN_ROUGH);
+ let seed = pu(F_GRAIN_SEED);
+ var g = grain_noise(gx * sc, gy * sc, seed);
+ g = g * (1.0 - rough * 0.5) + grain_noise(gx * sc * 2.3, gy * sc * 2.3, seed ^ 0x55u) * rough * 0.7;
+ let lum = 0.2126 * e.x + 0.7152 * e.y + 0.0722 * e.z;
+ let k = pf(F_GRAIN_AMT) * g * (0.35 + 2.6 * lum * (1.0 - lum));
+ e = e + k;
+ }
+ out[i] = enc8(e.x) | (enc8(e.y) << 8u) | (enc8(e.z) << 16u) | (255u << 24u);
+}
diff --git a/crates/gpu/src/wgsl/geom.wgsl b/crates/gpu/src/wgsl/geom.wgsl
new file mode 100644
index 00000000..973f386d
--- /dev/null
+++ b/crates/gpu/src/wgsl/geom.wgsl
@@ -0,0 +1,195 @@
+// Geometry (`lightcraft_pipeline::geometry::Frame::sample`): orientation, and the single bilinear
+// resample through the crop/straighten/flip affine or the full lens + perspective warp.
+// Bindings: src (rgb), dst (rgb).
+
+fn src_px(w: u32, h: u32, x: i32, y: i32) -> vec3 {
+ let cx = u32(clamp(x, 0, i32(w) - 1));
+ let cy = u32(clamp(y, 0, i32(h) - 1));
+ let j = 3u * (cy * w + cx);
+ return vec3(src[j], src[j + 1u], src[j + 2u]);
+}
+
+// `Rgb32f::sample_bilinear` (pixel centres at +0.5, clamped edges).
+fn bilinear(w: u32, h: u32, x: f32, y: f32) -> vec3 {
+ let fx = x - 0.5;
+ let fy = y - 0.5;
+ let x0 = floor(fx);
+ let y0 = floor(fy);
+ let tx = fx - x0;
+ let ty = fy - y0;
+ let ix = i32(x0);
+ let iy = i32(y0);
+ let a = src_px(w, h, ix, iy);
+ let b = src_px(w, h, ix + 1, iy);
+ let c = src_px(w, h, ix, iy + 1);
+ let d = src_px(w, h, ix + 1, iy + 1);
+ let top = a + (b - a) * tx;
+ let bot = c + (d - c) * tx;
+ return top + (bot - top) * ty;
+}
+
+fn put(i: u32, v: vec3) {
+ dst[3u * i] = v.x;
+ dst[3u * i + 1u] = v.y;
+ dst[3u * i + 2u] = v.z;
+}
+
+// Orientation: dst (ow × oh) pixel (x, y) = src pixel (P[4]·x + P[5]·y + P[6], P[7]·x + P[8]·y + P[9])
+// (integers). P: sw, sh, ow, oh, map[6].
+@compute @workgroup_size(16, 16)
+fn orient(@builtin(global_invocation_id) g: vec3) {
+ let sw = pu(0u);
+ let ow = pu(2u);
+ let oh = pu(3u);
+ if (g.x >= ow || g.y >= oh) {
+ return;
+ }
+ let x = i32(g.x);
+ let y = i32(g.y);
+ let m = array(bitcast(pu(4u)), bitcast(pu(5u)), bitcast(pu(6u)), bitcast(pu(7u)), bitcast(pu(8u)), bitcast(pu(9u)));
+ let sx = u32(m[0] * x + m[1] * y + m[2]);
+ let sy = u32(m[3] * x + m[4] * y + m[5]);
+ let j = 3u * (sy * sw + sx);
+ put(g.y * ow + g.x, vec3(src[j], src[j + 1u], src[j + 2u]));
+}
+
+// Bilinear resample through an affine (output px → base px). P: bw, bh, w, h, affine[6].
+@compute @workgroup_size(16, 16)
+fn sample_affine(@builtin(global_invocation_id) g: vec3) {
+ let w = pu(2u);
+ let h = pu(3u);
+ if (g.x >= w || g.y >= h) {
+ return;
+ }
+ let px = f32(g.x) + 0.5;
+ let py = f32(g.y) + 0.5;
+ let x = pf(4u) * px + pf(6u) * py + pf(8u);
+ let y = pf(5u) * px + pf(7u) * py + pf(9u);
+ put(g.y * w + g.x, bilinear(pu(0u), pu(1u), x, y));
+}
+
+// --- the lens / perspective warp (`lightcraft_pipeline::optics::Warp`), in oriented-source px.
+// P layout (see `geom_warp_params`): 0 bw, 1 bh, 2 w, 3 h, 4..10 output → transformed affine,
+// 10 sx, 11 sy, 12 W, 13 H, 14 persp?, 15..24 persp_inv, 24 k1, 25..28 ca, 28 lens_dist,
+// 29 warp?, 30..48 warp planes (3 × 6), 48 warp centre x, 49 y, 50 warp radius, 51 vig_stops,
+// 52 vig_power, 53 lens_vig, 54 vignette?, 55..60 vignette k, 60 vignette centre x, 61 y,
+// 62 vignette radius, 63 per-channel?, 64 gain?
+const WP_N: u32 = 65u;
+
+fn centre() -> vec2 {
+ return vec2(pf(12u) / 2.0, pf(13u) / 2.0);
+}
+
+fn to_corrected(p: vec2) -> vec2 {
+ if (pu(14u) == 0u) {
+ return p;
+ }
+ let c = centre();
+ let l = max(pf(12u), pf(13u)) / 2.0;
+ let q = (p - c) / l;
+ var wq = pf(21u) * q.x + pf(22u) * q.y + pf(23u);
+ if (abs(wq) < 1e-30) {
+ wq = 1e-30;
+ }
+ let x = (pf(15u) * q.x + pf(16u) * q.y + pf(17u)) / wq;
+ let y = (pf(18u) * q.x + pf(19u) * q.y + pf(20u)) / wq;
+ return c + vec2(x, y) * l;
+}
+
+fn embedded_warp(p: vec2, ch: u32) -> vec2 {
+ let w = pf(12u);
+ let h = pf(13u);
+ let cc = vec2(pf(48u) * w, pf(49u) * h);
+ let m = max(pf(50u) * max(w, h), 1e-9);
+ let o = 30u + 6u * min(ch, 2u);
+ let d = (p - cc) / m;
+ let r2 = d.x * d.x + d.y * d.y;
+ let f = pf(o) + r2 * (pf(o + 1u) + r2 * (pf(o + 2u) + r2 * pf(o + 3u)));
+ let kt0 = pf(o + 4u);
+ let kt1 = pf(o + 5u);
+ let sx = d.x * f + kt0 * 2.0 * d.x * d.y + kt1 * (r2 + 2.0 * d.x * d.x);
+ let sy = d.y * f + kt1 * 2.0 * d.x * d.y + kt0 * (r2 + 2.0 * d.y * d.y);
+ return cc + m * vec2(sx, sy);
+}
+
+fn corrected_to_source(p: vec2, ch: u32) -> vec2 {
+ let c = centre();
+ let hd = max(sqrt(pf(12u) * pf(12u) + pf(13u) * pf(13u)) / 2.0, 1e-9);
+ var q = p;
+ let k1 = pf(24u);
+ if (k1 != 0.0) {
+ let d = q - c;
+ let r2 = (d.x * d.x + d.y * d.y) / (hd * hd);
+ q = c + d * (1.0 + k1 * r2);
+ }
+ let ld = pf(28u);
+ if (ld != 0.0 && pu(29u) != 0u) {
+ let s = embedded_warp(q, ch);
+ q = q + (s - q) * ld;
+ }
+ let k = pf(25u + ch);
+ if (k != 0.0) {
+ q = c + (q - c) * (1.0 + k);
+ }
+ return q;
+}
+
+fn warp_gain(s: vec2) -> f32 {
+ var g = 1.0;
+ let stops = pf(51u);
+ let w = pf(12u);
+ let h = pf(13u);
+ if (stops != 0.0) {
+ let hd = max(sqrt(w * w + h * h) / 2.0, 1e-9);
+ let r = min(length(s - centre()) / hd, 1.5);
+ g *= exp2(stops * pow(r, pf(52u)));
+ }
+ let lv = pf(53u);
+ if (lv != 0.0 && pu(54u) != 0u) {
+ let cc = vec2(pf(60u) * w, pf(61u) * h);
+ let m = max(pf(62u) * max(w, h), 1e-9);
+ let d = s - cc;
+ let r2 = (d.x * d.x + d.y * d.y) / (m * m);
+ var vg = 1.0;
+ var rp = r2;
+ for (var k = 0u; k < 5u; k++) {
+ vg += pf(55u + k) * rp;
+ rp *= r2;
+ }
+ g *= 1.0 + (vg - 1.0) * lv;
+ }
+ return g;
+}
+
+@compute @workgroup_size(16, 16)
+fn sample_warp(@builtin(global_invocation_id) g: vec3) {
+ let w = pu(2u);
+ let h = pu(3u);
+ if (g.x >= w || g.y >= h) {
+ return;
+ }
+ let bw = pu(0u);
+ let bh = pu(1u);
+ let px = f32(g.x) + 0.5;
+ let py = f32(g.y) + 0.5;
+ let t = vec2(pf(4u) * px + pf(6u) * py + pf(8u), pf(5u) * px + pf(7u) * py + pf(9u));
+ let c = to_corrected(t);
+ let s = corrected_to_source(c, 1u);
+ let sx = pf(10u);
+ let sy = pf(11u);
+ if (!(s.x >= -0.5 && s.y >= -0.5 && s.x <= pf(12u) + 0.5 && s.y <= pf(13u) + 0.5)) {
+ put(g.y * w + g.x, vec3(BLANK_R, BLANK_G, BLANK_B));
+ return;
+ }
+ var v = bilinear(bw, bh, s.x * sx, s.y * sy);
+ if (pu(63u) != 0u) {
+ let qr = corrected_to_source(c, 0u);
+ v.x = bilinear(bw, bh, qr.x * sx, qr.y * sy).x;
+ let qb = corrected_to_source(c, 2u);
+ v.z = bilinear(bw, bh, qb.x * sx, qb.y * sy).z;
+ }
+ if (pu(64u) != 0u) {
+ v = v * warp_gain(s);
+ }
+ put(g.y * w + g.x, v);
+}
diff --git a/crates/gpu/src/wgsl/map.wgsl b/crates/gpu/src/wgsl/map.wgsl
new file mode 100644
index 00000000..c0a4419a
--- /dev/null
+++ b/crates/gpu/src/wgsl/map.wgsl
@@ -0,0 +1,137 @@
+// Element-wise kernels (one thread per pixel, 1-D on a 2-D grid). Bindings: a, b, c (inputs), dst.
+// P[0] = pixel count; further parameters per kernel. Each mirrors a CPU expression in
+// `lightcraft_pipeline::local` (named in the comment).
+
+fn rgb_a(i: u32) -> vec3 {
+ return vec3(a[3u * i], a[3u * i + 1u], a[3u * i + 2u]);
+}
+
+fn put_rgb(i: u32, v: vec3) {
+ dst[3u * i] = v.x;
+ dst[3u * i + 1u] = v.y;
+ dst[3u * i + 2u] = v.z;
+}
+
+// `log_lum` of an RGB image.
+@compute @workgroup_size(256)
+fn log_lum_k(@builtin(global_invocation_id) g: vec3, @builtin(num_workgroups) nw: vec3) {
+ let i = lin_index(g, nw);
+ if (i >= pu(0u)) {
+ return;
+ }
+ dst[i] = log_lum(rgb_a(i));
+}
+
+// `dark_of`: the dark channel (dehaze).
+@compute @workgroup_size(256)
+fn dark_k(@builtin(global_invocation_id) g: vec3, @builtin(num_workgroups) nw: vec3) {
+ let i = lin_index(g, nw);
+ if (i >= pu(0u)) {
+ return;
+ }
+ let c = rgb_a(i);
+ dst[i] = min(min(c.x, c.y), c.z);
+}
+
+// Guided filter input pair (p, p²).
+@compute @workgroup_size(256)
+fn guided_pre(@builtin(global_invocation_id) g: vec3, @builtin(num_workgroups) nw: vec3) {
+ let i = lin_index(g, nw);
+ if (i >= pu(0u)) {
+ return;
+ }
+ let p = a[i];
+ dst[2u * i] = p;
+ dst[2u * i + 1u] = p * p;
+}
+
+// Guided filter coefficients (a, b) from blurred (mean, corr). P[1] = eps.
+@compute @workgroup_size(256)
+fn guided_ab(@builtin(global_invocation_id) g: vec3, @builtin(num_workgroups) nw: vec3) {
+ let i = lin_index(g, nw);
+ if (i >= pu(0u)) {
+ return;
+ }
+ let m = a[2u * i];
+ let c = a[2u * i + 1u];
+ let v = max(c - m * m, 0.0);
+ let k = v / (v + pf(1u));
+ dst[2u * i] = k;
+ dst[2u * i + 1u] = m - k * m;
+}
+
+// `q = a·p + b` (a: p, b: interleaved blurred coefficients).
+@compute @workgroup_size(256)
+fn guided_apply(@builtin(global_invocation_id) g: vec3, @builtin(num_workgroups) nw: vec3) {
+ let i = lin_index(g, nw);
+ if (i >= pu(0u)) {
+ return;
+ }
+ dst[i] = b[2u * i] * a[i] + b[2u * i + 1u];
+}
+
+// White balance: 3×3 matrix (P[1] = has matrix, P[2..11] row-major), clamped at 0.
+@compute @workgroup_size(256)
+fn wb_k(@builtin(global_invocation_id) g: vec3, @builtin(num_workgroups) nw: vec3) {
+ let i = lin_index(g, nw);
+ if (i >= pu(0u)) {
+ return;
+ }
+ var c = rgb_a(i);
+ if (pu(1u) != 0u) {
+ c = vec3(
+ pf(2u) * c.x + pf(3u) * c.y + pf(4u) * c.z,
+ pf(5u) * c.x + pf(6u) * c.y + pf(7u) * c.z,
+ pf(8u) * c.x + pf(9u) * c.y + pf(10u) * c.z,
+ );
+ }
+ put_rgb(i, max(c * 1.0, vec3(0.0)));
+}
+
+// Luminance NR: scale by 2^((f − l)·k) (a: image, b: log luminance l, c: filtered f). P[1] = k.
+@compute @workgroup_size(256)
+fn nr_lum(@builtin(global_invocation_id) g: vec3, @builtin(num_workgroups) nw: vec3) {
+ let i = lin_index(g, nw);
+ if (i >= pu(0u)) {
+ return;
+ }
+ let d = (c[i] - b[i]) * pf(1u);
+ put_rgb(i, rgb_a(i) * exp2(d));
+}
+
+// Chromaticity rgb / Y.
+@compute @workgroup_size(256)
+fn chroma_k(@builtin(global_invocation_id) g: vec3, @builtin(num_workgroups) nw: vec3) {
+ let i = lin_index(g, nw);
+ if (i >= pu(0u)) {
+ return;
+ }
+ let c0 = rgb_a(i);
+ let y = max(lum2020(c0), 1e-6);
+ put_rgb(i, vec3(c0.x / y, c0.y / y, c0.z / y));
+}
+
+// Colour NR: blend chromaticity towards its blur and re-apply luminance (a: image, b: chroma,
+// c: blurred chroma). P[1] = t.
+@compute @workgroup_size(256)
+fn nr_col(@builtin(global_invocation_id) g: vec3, @builtin(num_workgroups) nw: vec3) {
+ let i = lin_index(g, nw);
+ if (i >= pu(0u)) {
+ return;
+ }
+ let yl = lum2020(rgb_a(i));
+ let t = pf(1u);
+ let c0 = vec3(b[3u * i], b[3u * i + 1u], b[3u * i + 2u]);
+ let cb = vec3(c[3u * i], c[3u * i + 1u], c[3u * i + 2u]);
+ put_rgb(i, max((c0 + (cb - c0) * t) * yl, vec3(0.0)));
+}
+
+// Every P[1]-th value of `a` (airlight sampling). P[0] = output count.
+@compute @workgroup_size(256)
+fn subsample(@builtin(global_invocation_id) g: vec3, @builtin(num_workgroups) nw: vec3) {
+ let i = lin_index(g, nw);
+ if (i >= pu(0u)) {
+ return;
+ }
+ dst[i] = a[i * pu(1u)];
+}
diff --git a/crates/gpu/src/wgsl/mask.wgsl b/crates/gpu/src/wgsl/mask.wgsl
new file mode 100644
index 00000000..2075e186
--- /dev/null
+++ b/crates/gpu/src/wgsl/mask.wgsl
@@ -0,0 +1,149 @@
+// Mask evaluation (`lightcraft_pipeline::masks`): one component shape → `c`, then combined into
+// the mask's alpha plane. Bindings: img (rgb, pre-exposure), log_l, aux (shape data), c, alpha.
+
+// Position of output pixel (x, y) in long-edge units: P[4..10] = affine (a, b, c, d, e, f).
+fn pos(x: u32, y: u32) -> vec2 {
+ let px = f32(x) + 0.5;
+ let py = f32(y) + 0.5;
+ return vec2(pf(4u) * px + pf(6u) * py + pf(8u), pf(5u) * px + pf(7u) * py + pf(9u));
+}
+
+// P: w, h, kind, invert, affine[6], then per kind:
+// 0 linear: a.xy, d.xy, |d|² 1 radial: centre.xy, cos, sin, rx, ry, feather, invert
+// 2 luminance range: lo, hi, lo feather, hi feather, ev
+// 3 colour range: tol, gain, sample count (OkLab triples in aux)
+// 4 brush: stroke count (stroke records + dabs in aux)
+@compute @workgroup_size(16, 16)
+fn shape(@builtin(global_invocation_id) g: vec3) {
+ let w = pu(0u);
+ let h = pu(1u);
+ if (g.x >= w || g.y >= h) {
+ return;
+ }
+ let i = g.y * w + g.x;
+ let kind = pu(2u);
+ var v = 0.0;
+ if (kind == 0u) {
+ let p = pos(g.x, g.y);
+ let a = vec2(pf(10u), pf(11u));
+ let d = vec2(pf(12u), pf(13u));
+ let t = dot(p - a, d) / pf(14u);
+ v = 1.0 - sstep(0.0, 1.0, t);
+ } else if (kind == 1u) {
+ let p = pos(g.x, g.y);
+ let d = p - vec2(pf(10u), pf(11u));
+ let co = pf(12u);
+ let s = pf(13u);
+ let x = d.x * co - d.y * s;
+ let y = d.x * s + d.y * co;
+ let qx = x / pf(14u);
+ let qy = y / pf(15u);
+ let r = sqrt(qx * qx + qy * qy);
+ let f = pf(16u);
+ v = 1.0 - sstep(1.0 - max(f, 0.01), 1.0, r);
+ if (pu(17u) != 0u) {
+ v = 1.0 - v;
+ }
+ } else if (kind == 2u) {
+ let lo = pf(10u);
+ let hi = pf(11u);
+ let y = clamp((log_l[i] + pf(14u) + 8.0) / 12.0, 0.0, 1.0);
+ v = sstep(lo - pf(12u), lo, y) * (1.0 - sstep(hi, hi + pf(13u), y));
+ } else if (kind == 3u) {
+ let tol = pf(10u);
+ let gain = pf(11u);
+ let c = vec3(img[3u * i], img[3u * i + 1u], img[3u * i + 2u]) * gain;
+ let lab = oklab(c / (1.0 + c));
+ var best = 3.402823e38;
+ for (var k = 0u; k < pu(12u); k++) {
+ let s = vec3(aux[3u * k], aux[3u * k + 1u], aux[3u * k + 2u]);
+ let da = lab.y - s.y;
+ let db = lab.z - s.z;
+ let dl = lab.x - s.x;
+ best = min(best, sqrt(da * da + db * db + 0.25 * (dl * dl)));
+ }
+ v = 1.0 - sstep(tol * 0.5, tol, best);
+ } else if (kind == 4u) {
+ // stroke record: dab offset, dab count, r, hard, flow, density, erase, bbox x0 y0 x1 y1
+ let px = f32(g.x) + 0.5;
+ let py = f32(g.y) + 0.5;
+ for (var k = 0u; k < pu(10u); k++) {
+ let o = 11u * k;
+ if (px < aux[o + 7u] || py < aux[o + 8u] || px > aux[o + 9u] || py > aux[o + 10u]) {
+ continue;
+ }
+ let off = u32(aux[o]);
+ let cnt = u32(aux[o + 1u]);
+ let r = aux[o + 2u];
+ let hard = aux[o + 3u];
+ let flow = aux[o + 4u];
+ var keep = 1.0;
+ for (var j = 0u; j < cnt; j++) {
+ let dx = px - aux[off + 2u * j];
+ let dy = py - aux[off + 2u * j + 1u];
+ let dd = sqrt(dx * dx + dy * dy);
+ if (dd > r) {
+ continue;
+ }
+ var a = 1.0;
+ if (dd > hard) {
+ a = 1.0 - sstep(hard, r, dd);
+ }
+ keep *= 1.0 - a * flow;
+ }
+ let sa = min(1.0 - keep, aux[o + 5u]);
+ if (aux[o + 6u] != 0.0) {
+ v = v * (1.0 - sa);
+ } else {
+ v = max(v, sa);
+ }
+ }
+ }
+ if (pu(3u) != 0u) {
+ v = 1.0 - v;
+ }
+ c[i] = v;
+}
+
+// Combine `c` into mask plane P[2]: P[3] = 0 copy, 1 add (max), 2 subtract, 3 intersect.
+@compute @workgroup_size(256)
+fn combine(@builtin(global_invocation_id) g: vec3, @builtin(num_workgroups) nw: vec3) {
+ let i = lin_index(g, nw);
+ let n = pu(0u);
+ if (i >= n) {
+ return;
+ }
+ let j = pu(2u) * n + i;
+ let op = pu(3u);
+ let a = alpha[j];
+ let v = c[i];
+ if (op == 0u) {
+ alpha[j] = v;
+ } else if (op == 1u) {
+ alpha[j] = max(a, v);
+ } else if (op == 2u) {
+ alpha[j] = a * (1.0 - v);
+ } else {
+ alpha[j] = a * v;
+ }
+}
+
+// Finish mask plane P[2]: P[3] = invert, P[4] = amount (applied when ≠ 1).
+@compute @workgroup_size(256)
+fn finalize(@builtin(global_invocation_id) g: vec3, @builtin(num_workgroups) nw: vec3) {
+ let i = lin_index(g, nw);
+ let n = pu(0u);
+ if (i >= n) {
+ return;
+ }
+ let j = pu(2u) * n + i;
+ var a = alpha[j];
+ if (pu(3u) != 0u) {
+ a = 1.0 - a;
+ }
+ let amt = pf(4u);
+ if (abs(amt - 1.0) > 1e-6) {
+ a = a * amt;
+ }
+ alpha[j] = a;
+}
diff --git a/crates/gpu/src/wgsl/resize.wgsl b/crates/gpu/src/wgsl/resize.wgsl
new file mode 100644
index 00000000..a766e53a
--- /dev/null
+++ b/crates/gpu/src/wgsl/resize.wgsl
@@ -0,0 +1,51 @@
+// Separable resample with precomputed taps (`lightcraft_raster::resample::weights`, uploaded):
+// table[3i..3i+3] = (first source index, tap count, weight offset), weights stored as f32 bits.
+// Bindings: src, table (u32), dst.
+
+// Horizontal. P: sw, h, dw, nc.
+@compute @workgroup_size(16, 16)
+fn resize_h(@builtin(global_invocation_id) g: vec3) {
+ let sw = pu(0u);
+ let h = pu(1u);
+ let dw = pu(2u);
+ let nc = pu(3u);
+ let x = g.x;
+ let y = g.y;
+ if (x >= dw || y >= h) {
+ return;
+ }
+ let lo = table[3u * x];
+ let cnt = table[3u * x + 1u];
+ let off = table[3u * x + 2u];
+ for (var c = 0u; c < nc; c++) {
+ var acc = 0.0;
+ for (var k = 0u; k < cnt; k++) {
+ acc = acc + src[(y * sw + lo + k) * nc + c] * bitcast(table[off + k]);
+ }
+ dst[(y * dw + x) * nc + c] = acc;
+ }
+}
+
+// Vertical. P: w, sh, dh, nc.
+@compute @workgroup_size(16, 16)
+fn resize_v(@builtin(global_invocation_id) g: vec3) {
+ let w = pu(0u);
+ let sh = pu(1u);
+ let dh = pu(2u);
+ let nc = pu(3u);
+ let x = g.x;
+ let y = g.y;
+ if (x >= w || y >= dh) {
+ return;
+ }
+ let lo = table[3u * y];
+ let cnt = table[3u * y + 1u];
+ let off = table[3u * y + 2u];
+ for (var c = 0u; c < nc; c++) {
+ var acc = 0.0;
+ for (var k = 0u; k < cnt; k++) {
+ acc = acc + src[((lo + k) * w + x) * nc + c] * bitcast(table[off + k]);
+ }
+ dst[(y * w + x) * nc + c] = acc;
+ }
+}
diff --git a/crates/gpu/tests/equivalence.rs b/crates/gpu/tests/equivalence.rs
new file mode 100644
index 00000000..474d3404
--- /dev/null
+++ b/crates/gpu/tests/equivalence.rs
@@ -0,0 +1,359 @@
+//! CPU ↔ GPU equivalence: the same settings rendered by the CPU pipeline (the reference) and by
+//! `lightcraft_gpu::render`, compared in 8-bit sRGB.
+//!
+//! Bounds (documented in `docs/gpu-pipeline.md`): mean |Δ| < 0.5 LSB and max |Δ| ≤ 3 LSB per
+//! channel for every case. Skips (passes with a note) when no GPU adapter exists, e.g. on CI.
+
+use std::sync::Arc;
+
+use lightcraft_develop::{BrushStroke, DevelopSettings, Mask, MaskComponent, MaskOp, MaskShape, Spot, Treatment, VignetteStyle, WbMode, Wheel};
+use lightcraft_geom::{Orientation, Point, Rect};
+use lightcraft_pipeline::{Quality, RenderRequest, SourceInfo, StageCache, render};
+use lightcraft_raster::{Rgb32f, Rgba8};
+
+const MEAN_LSB: f64 = 0.5;
+const MAX_LSB: u8 = 3;
+
+fn gpu() -> bool {
+ let ok = lightcraft_gpu::available();
+ if !ok {
+ eprintln!("skipped: no GPU adapter");
+ }
+ ok
+}
+
+fn scene(i: usize, w: usize, h: usize) -> Arc {
+ Arc::new(lightcraft_scenes::demo_library()[i].render(w, h))
+}
+
+/// (mean |Δ|, max |Δ|, share of channels with |Δ| > 1).
+fn diff(a: &Rgba8, b: &Rgba8) -> (f64, u8, f64) {
+ assert_eq!((a.width, a.height), (b.width, b.height));
+ let (mut sum, mut max, mut over) = (0u64, 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]);
+ sum += d as u64;
+ max = max.max(d);
+ over += (d > 1) as u64;
+ }
+ }
+ let n = (a.data.len() * 3) as f64;
+ (sum as f64 / n, max, over as f64 / n)
+}
+
+fn check(name: &str, src: &Arc, info: &SourceInfo, s: &DevelopSettings, req: &RenderRequest) -> (f64, u8) {
+ let cpu = render(src, info, s, req).image;
+ let gpu = lightcraft_gpu::render(src, info, s, req, None).expect("gpu render");
+ let (mean, max, over) = diff(&cpu, &gpu.image);
+ eprintln!("{name:<28} {}x{} mean {mean:.4} max {max} >1: {:.4}%", cpu.width, cpu.height, over * 100.0);
+ assert!(mean < MEAN_LSB && max <= MAX_LSB, "{name}: mean {mean:.4} LSB, max {max} LSB");
+ (mean, max)
+}
+
+type Edit = fn(&mut DevelopSettings);
+
+fn typical(s: &mut DevelopSettings) {
+ s.light.exposure = 0.3;
+ s.light.highlights = -40.0;
+ s.light.shadows = 30.0;
+ s.effects.clarity = 15.0;
+ s.effects.texture = 10.0;
+ s.effects.dehaze = 10.0;
+ s.detail.nr_luminance = 30.0;
+ s.detail.nr_color = 25.0;
+ s.detail.sharpen_amount = 40.0;
+}
+
+fn cases() -> Vec<(&'static str, Edit)> {
+ vec![
+ ("default", |_| {}),
+ ("typical", typical),
+ ("exposure/contrast/whites", |s| {
+ s.light.exposure = -0.7;
+ s.light.contrast = 45.0;
+ s.light.whites = 30.0;
+ s.light.blacks = -25.0;
+ }),
+ ("vibrance/saturation", |s| {
+ s.color.vibrance = 40.0;
+ s.color.saturation = -20.0;
+ }),
+ ("colour mixer", |s| {
+ s.mixer.blue.sat = -60.0;
+ s.mixer.orange.hue = 30.0;
+ s.mixer.green.lum = 40.0;
+ }),
+ ("b&w mix", |s| {
+ s.treatment = Treatment::Bw;
+ s.bw_mix.blue = 60.0;
+ s.bw_mix.red = -30.0;
+ }),
+ ("colour grading", |s| {
+ s.grading.shadows = Wheel { hue: 220.0, sat: 40.0, lum: -10.0 };
+ s.grading.highlights = Wheel { hue: 40.0, sat: 30.0, lum: 10.0 };
+ s.grading.global = Wheel { hue: 300.0, sat: 10.0, lum: 0.0 };
+ s.grading.blending = 70.0;
+ s.grading.balance = 20.0;
+ }),
+ ("tone curves", |s| {
+ s.curve.highlights = -30.0;
+ s.curve.shadows = 25.0;
+ s.curve.master = vec![Point::new(0.0, 0.05), Point::new(0.5, 0.55), Point::new(1.0, 0.95)];
+ s.curve.blue = vec![Point::new(0.0, 0.0), Point::new(0.5, 0.45), Point::new(1.0, 1.0)];
+ }),
+ ("vignette (highlight)", |s| {
+ s.vignette.amount = -60.0;
+ s.vignette.highlights = 50.0;
+ s.vignette.roundness = -40.0;
+ }),
+ ("vignette (paint, +)", |s| {
+ s.vignette.amount = 40.0;
+ s.vignette.style = VignetteStyle::PaintOverlay;
+ }),
+ ("vignette (paint, -)", |s| {
+ s.vignette.amount = -40.0;
+ s.vignette.style = VignetteStyle::PaintOverlay;
+ }),
+ ("grain", |s| {
+ s.grain.amount = 50.0;
+ s.grain.size = 40.0;
+ s.grain.roughness = 60.0;
+ }),
+ ("dehaze -", |s| s.effects.dehaze = -50.0),
+ ("dehaze +", |s| s.effects.dehaze = 60.0),
+ ("clarity/texture -", |s| {
+ s.effects.clarity = -60.0;
+ s.effects.texture = -40.0;
+ }),
+ ("sharpen masking", |s| {
+ s.detail.sharpen_amount = 90.0;
+ s.detail.sharpen_masking = 60.0;
+ }),
+ ("white balance", |s| {
+ s.wb.mode = WbMode::Custom;
+ s.wb.temp = 8200.0;
+ s.wb.tint = 15.0;
+ }),
+ ("noise reduction", |s| {
+ s.detail.nr_luminance = 70.0;
+ s.detail.nr_detail = 30.0;
+ s.detail.nr_color = 60.0;
+ s.detail.nr_color_smoothness = 50.0;
+ }),
+ ("crop + straighten + flip", |s| {
+ s.crop.geometry.rect = Rect::new(0.1, 0.05, 0.85, 0.9);
+ s.crop.geometry.angle = 7.5;
+ s.crop.flip_h = true;
+ }),
+ ("orientation", |s| s.orientation = Orientation::Rotate90),
+ ("lens + perspective", |s| {
+ s.optics.distortion = 30.0;
+ s.optics.vignetting = 40.0;
+ s.optics.ca_red = 50.0;
+ s.optics.ca_blue = -40.0;
+ s.geometry.vertical = 20.0;
+ }),
+ ("masks (gradients)", |s| {
+ s.masks = vec![
+ Mask {
+ components: vec![MaskComponent {
+ op: MaskOp::Add,
+ invert: false,
+ shape: MaskShape::Linear { start: Point::new(0.5, 0.0), end: Point::new(0.5, 0.6) },
+ }],
+ adjust: lightcraft_develop::LocalAdjustments { exposure: -0.8, temp: -30.0, saturation: 20.0, ..Default::default() },
+ ..Default::default()
+ },
+ Mask {
+ components: vec![
+ MaskComponent {
+ op: MaskOp::Add,
+ invert: false,
+ shape: MaskShape::Radial { center: Point::new(0.4, 0.6), rx: 0.25, ry: 0.15, angle: 20.0, feather: 60.0, invert: false },
+ },
+ MaskComponent {
+ op: MaskOp::Intersect,
+ invert: true,
+ shape: MaskShape::Linear { start: Point::new(0.0, 0.0), end: Point::new(1.0, 1.0) },
+ },
+ ],
+ adjust: lightcraft_develop::LocalAdjustments {
+ shadows: 40.0,
+ clarity: 30.0,
+ contrast: 20.0,
+ color_hue: 30.0,
+ color_sat: 50.0,
+ hue: 10.0,
+ ..Default::default()
+ },
+ ..Default::default()
+ },
+ ];
+ }),
+ ("masks (brush + ranges)", |s| {
+ let stroke = BrushStroke {
+ points: vec![Point::new(0.1, 0.3), Point::new(0.5, 0.4), Point::new(0.8, 0.35)],
+ size: 0.04,
+ flow: 60.0,
+ ..Default::default()
+ };
+ let erase = BrushStroke { points: vec![Point::new(0.5, 0.4)], size: 0.03, erase: true, ..Default::default() };
+ s.masks = vec![
+ Mask {
+ components: vec![MaskComponent { op: MaskOp::Add, invert: false, shape: MaskShape::Brush { strokes: vec![stroke, erase] } }],
+ adjust: lightcraft_develop::LocalAdjustments { exposure: 0.7, whites: 20.0, blacks: -20.0, dehaze: 30.0, ..Default::default() },
+ ..Default::default()
+ },
+ Mask {
+ components: vec![MaskComponent {
+ op: MaskOp::Add,
+ invert: false,
+ shape: MaskShape::LuminanceRange { lo: 0.5, hi: 0.8, lo_feather: 0.1, hi_feather: 0.1 },
+ }],
+ adjust: lightcraft_develop::LocalAdjustments { highlights: -50.0, texture: 30.0, sharpness: 40.0, ..Default::default() },
+ ..Default::default()
+ },
+ Mask {
+ components: vec![MaskComponent {
+ op: MaskOp::Add,
+ invert: false,
+ shape: MaskShape::ColorRange { samples: vec![[0.6, -0.05, -0.08]], refine: 50.0 },
+ }],
+ adjust: lightcraft_develop::LocalAdjustments { saturation: 40.0, tint: 20.0, ..Default::default() },
+ invert: true,
+ ..Default::default()
+ },
+ ];
+ }),
+ ("masks (cpu shapes, ops, amount)", |s| {
+ s.masks = vec![
+ Mask {
+ components: vec![
+ MaskComponent { op: MaskOp::Add, invert: false, shape: MaskShape::Sky },
+ MaskComponent {
+ op: MaskOp::Subtract,
+ invert: false,
+ shape: MaskShape::Radial { center: Point::new(0.7, 0.2), rx: 0.1, ry: 0.1, angle: 0.0, feather: 30.0, invert: false },
+ },
+ ],
+ adjust: lightcraft_develop::LocalAdjustments { exposure: -0.5, dehaze: 40.0, amount: 70.0, ..Default::default() },
+ ..Default::default()
+ },
+ Mask {
+ components: vec![MaskComponent { op: MaskOp::Intersect, invert: false, shape: MaskShape::Subject }],
+ adjust: lightcraft_develop::LocalAdjustments { exposure: 0.5, ..Default::default() },
+ ..Default::default()
+ },
+ Mask {
+ components: vec![MaskComponent { op: MaskOp::Add, invert: true, shape: MaskShape::Background }],
+ adjust: lightcraft_develop::LocalAdjustments { saturation: -60.0, ..Default::default() },
+ invert: true,
+ ..Default::default()
+ },
+ ];
+ }),
+ ("spots + defringe (cpu stage)", |s| {
+ s.spots = vec![Spot { points: vec![Point::new(0.3, 0.3)], size: 0.03, source_offset: Some(Point::new(0.1, 0.0)), ..Default::default() }];
+ s.optics.defringe_purple_amount = 5.0;
+ }),
+ ]
+}
+
+#[test]
+fn gpu_matches_cpu() {
+ if !gpu() {
+ return;
+ }
+ let src = scene(0, 960, 640);
+ let raw = SourceInfo { raw: true, ..Default::default() };
+ let req = RenderRequest::fit(720, 720);
+ let mut worst = (0.0f64, 0u8);
+ for (name, edit) in cases() {
+ let mut s = DevelopSettings::default();
+ edit(&mut s);
+ let (m, x) = check(name, &src, &raw, &s, &req);
+ worst = (worst.0.max(m), worst.1.max(x));
+ }
+ eprintln!("worst: mean {:.4} LSB, max {} LSB", worst.0, worst.1);
+}
+
+#[test]
+fn rendered_sources_and_other_scenes() {
+ if !gpu() {
+ return;
+ }
+ // A display-referred source (JPEG-like: display tone map) and other scenes / sizes / draft.
+ let jpeg = SourceInfo::default();
+ for (i, (w, h)) in [(3usize, (800usize, 533usize)), (5, (640, 960)), (7, (1200, 800))] {
+ let src = scene(i, w, h);
+ let mut s = DevelopSettings::default();
+ typical(&mut s);
+ s.color.vibrance = 25.0;
+ check(&format!("scene {i} display-referred"), &src, &jpeg, &s, &RenderRequest::fit(700, 700));
+ let draft = RenderRequest { quality: Quality::Draft, ..RenderRequest::fit(500, 500) };
+ check(&format!("scene {i} draft"), &src, &SourceInfo { raw: true, ..Default::default() }, &s, &draft);
+ check(&format!("scene {i} full size"), &src, &SourceInfo { raw: true, ..Default::default() }, &s, &RenderRequest::fit(w, h));
+ }
+}
+
+#[test]
+fn geometry_variants() {
+ if !gpu() {
+ return;
+ }
+ use lightcraft_develop::{EmbeddedLens, EmbeddedVignette, EmbeddedWarp};
+ let src = scene(2, 900, 600);
+ let raw = SourceInfo { raw: true, ..Default::default() };
+ use Orientation::*;
+ for o in [Normal, Rotate90, Rotate180, Rotate270, FlipH, Transverse, FlipV, Transpose] {
+ for (crop, req) in [(false, RenderRequest::fit(900, 900)), (true, RenderRequest::fit(500, 500))] {
+ let mut s = DevelopSettings { orientation: o, ..Default::default() };
+ if crop {
+ s.crop.geometry.rect = Rect::new(0.2, 0.1, 0.9, 0.8);
+ s.crop.geometry.angle = -4.0;
+ s.crop.flip_v = true;
+ }
+ check(&format!("{o:?} crop={crop}"), &src, &raw, &s, &req);
+ }
+ }
+ // embedded DNG lens corrections (per-plane warp + vignette) with manual CA
+ let lens = EmbeddedLens {
+ warp: Some(EmbeddedWarp {
+ planes: [[1.0, -0.03, 0.01, 0.0, 0.001, -0.002], [1.0, -0.028, 0.01, 0.0, 0.001, -0.002], [1.0, -0.026, 0.01, 0.0, 0.001, -0.002]],
+ center: Point::new(0.52, 0.48),
+ radius: 0.6,
+ }),
+ vignette: Some(EmbeddedVignette { k: [0.4, -0.1, 0.02, 0.0, 0.0], center: Point::new(0.5, 0.5), radius: 0.6 }),
+ };
+ let info = SourceInfo { lens: Some(lens), ..raw };
+ let mut s = DevelopSettings::default();
+ s.optics.lens_profile = true;
+ s.optics.ca_red = 30.0;
+ s.geometry.horizontal = -15.0;
+ check("embedded lens + perspective", &src, &info, &s, &RenderRequest::fit(700, 700));
+ s.orientation = Rotate270;
+ check("embedded lens rotated", &src, &info, &s, &RenderRequest::fit(700, 700));
+}
+
+#[test]
+fn cached_renders_match_uncached() {
+ if !gpu() {
+ return;
+ }
+ // Slider drags reuse device-resident stages: the result must equal a fresh render.
+ let src = scene(1, 900, 600);
+ let info = SourceInfo { raw: true, ..Default::default() };
+ let cache = StageCache::default();
+ let req = RenderRequest::fit(640, 640);
+ let mut s = DevelopSettings::default();
+ typical(&mut s);
+ for k in 0..4 {
+ s.light.exposure = 0.1 * k as f64;
+ s.effects.clarity = 10.0 + 5.0 * (k / 2) as f64;
+ s.detail.nr_luminance = 20.0 + 10.0 * (k % 2) as f64;
+ let warm = lightcraft_gpu::render(&src, &info, &s, &req, Some(&cache)).expect("gpu");
+ let fresh = lightcraft_gpu::render(&src, &info, &s, &req, None).expect("gpu");
+ assert_eq!(warm.image, fresh.image, "step {k}");
+ }
+}
diff --git a/crates/pipeline/src/colorops.rs b/crates/pipeline/src/colorops.rs
index b271e560..5d378d19 100644
--- a/crates/pipeline/src/colorops.rs
+++ b/crates/pipeline/src/colorops.rs
@@ -16,7 +16,8 @@ pub fn oklch_hue_of_srgb_hue(deg: f64) -> f32 {
lab_to_lch(oklab_from_2020(m))[2]
}
-fn band_hues() -> &'static [f32; 8] {
+/// OkLCh hue (radians) of each colour-mixer band centre.
+pub fn band_hues() -> &'static [f32; 8] {
static H: OnceLock<[f32; 8]> = OnceLock::new();
H.get_or_init(|| MIXER_HUES.map(oklch_hue_of_srgb_hue))
}
@@ -47,23 +48,28 @@ pub fn band_weights(h: f32) -> [f32; 8] {
w
}
-#[derive(Clone, Copy)]
-struct WheelK {
- a: f32,
- b: f32,
- lum: f32,
+/// A colour-grading wheel as an OkLab offset (a, b) and a lightness shift.
+#[derive(Clone, Copy, Debug)]
+pub struct WheelK {
+ pub a: f32,
+ pub b: f32,
+ pub lum: f32,
}
+/// The colour tools' parameters, resolved once per render (fields are read by the GPU kernel).
+#[derive(Clone, Debug)]
pub struct ColorOps {
- vibrance: f32,
- saturation: f32,
- hue: [f32; 8],
- sat: [f32; 8],
- lum: [f32; 8],
- mixer: bool,
- bw: Option<[f32; 8]>,
- grading: Option<([WheelK; 4], f32, f32)>,
- skin: f32,
+ pub vibrance: f32,
+ pub saturation: f32,
+ pub hue: [f32; 8],
+ pub sat: [f32; 8],
+ pub lum: [f32; 8],
+ pub mixer: bool,
+ pub bw: Option<[f32; 8]>,
+ /// Wheels (shadows, midtones, highlights, global), blending, balance.
+ pub grading: Option<([WheelK; 4], f32, f32)>,
+ /// OkLCh hue of skin tones (protected by vibrance).
+ pub skin: f32,
}
fn wheel(w: &lightcraft_develop::Wheel) -> WheelK {
diff --git a/crates/pipeline/src/finish.rs b/crates/pipeline/src/finish.rs
index 196a75b2..5a79a2de 100644
--- a/crates/pipeline/src/finish.rs
+++ b/crates/pipeline/src/finish.rs
@@ -3,7 +3,7 @@
use lightcraft_color::spline::{Lut1, MonotoneCurve};
use lightcraft_color::transfer::linear_to_srgb;
use lightcraft_color::{REC2020, SRGB, luminance_2020};
-use lightcraft_develop::{DevelopSettings, ToneCurve, VignetteStyle};
+use lightcraft_develop::{DevelopSettings, LocalAdjustments, ToneCurve, VignetteStyle};
use lightcraft_geom::Point;
use lightcraft_raster::Rgba8;
@@ -55,14 +55,16 @@ fn curve_luts(c: &ToneCurve) -> Option<[Lut1; 3]> {
Some(std::array::from_fn(|i| if chans[i] { MonotoneCurve::new(&to_pts(per[i])).to_lut(N).compose(&base) } else { base.clone() }))
}
-struct Vig {
- amount: f32,
- start: f32,
- width: f32,
- aspect_mix: f32,
- power: f32,
- highlights: f32,
- style: VignetteStyle,
+/// Vignette (post-crop) parameters.
+#[derive(Clone, Copy, Debug)]
+pub struct Vig {
+ pub amount: f32,
+ pub start: f32,
+ pub width: f32,
+ pub aspect_mix: f32,
+ pub power: f32,
+ pub highlights: f32,
+ pub style: VignetteStyle,
}
fn vignette(s: &DevelopSettings) -> Option {
@@ -81,14 +83,17 @@ fn vignette(s: &DevelopSettings) -> Option {
})
}
+/// Hash constants of [`grain_noise`] (shared with the GPU kernel).
+pub const GRAIN_HASH: [u32; 4] = [0x8da6_b343, 0xd816_3841, 0xcb1a_b31f, 0x5bd1_e995];
+
#[inline]
fn grain_noise(x: f32, y: f32, seed: u32) -> f32 {
let (x0, y0) = (x.floor(), y.floor());
let (fx, fy) = (x - x0, y - y0);
let h = |i: i32, j: i32| {
- let mut v = (i as u32).wrapping_mul(0x8da6_b343) ^ (j as u32).wrapping_mul(0xd816_3841) ^ seed.wrapping_mul(0xcb1a_b31f);
+ let mut v = (i as u32).wrapping_mul(GRAIN_HASH[0]) ^ (j as u32).wrapping_mul(GRAIN_HASH[1]) ^ seed.wrapping_mul(GRAIN_HASH[2]);
v ^= v >> 13;
- v = v.wrapping_mul(0x5bd1_e995);
+ v = v.wrapping_mul(GRAIN_HASH[3]);
v ^= v >> 15;
(v & 0xffff) as f32 / 32768.0 - 1.0
};
@@ -99,33 +104,130 @@ fn grain_noise(x: f32, y: f32, seed: u32) -> f32 {
a + (b - a) * v
}
-pub fn finish(p: &Prepared, s: &DevelopSettings, frame: &Frame, info: &SourceInfo) -> Rgba8 {
+/// Number of per-mask terms in [`mask_terms`].
+pub const MASK_TERMS: usize = 18;
+
+/// A mask's local adjustments as the per-pixel stage uses them (alpha-weighted sums), in order:
+/// exposure, temp, tint, contrast, highlights, shadows, whites, blacks, texture, clarity, dehaze,
+/// saturation, hue, sharpness, then the colour overlay: on (0/1), cos and sin of its OkLCh hue, and
+/// its strength.
+pub fn mask_terms(j: &LocalAdjustments) -> [f32; MASK_TERMS] {
+ let (on, cos, sin, amt) = if j.color_sat > 0.0 {
+ let hue = crate::colorops::oklch_hue_of_srgb_hue(j.color_hue);
+ (1.0, hue.cos(), hue.sin(), (j.color_sat / 100.0) as f32)
+ } else {
+ (0.0, 0.0, 0.0, 0.0)
+ };
+ [
+ j.exposure as f32,
+ (j.temp / 100.0) as f32,
+ (j.tint / 100.0) as f32,
+ (j.contrast / 100.0) as f32,
+ (j.highlights / 100.0) as f32,
+ (j.shadows / 100.0) as f32,
+ (j.whites / 100.0) as f32,
+ (j.blacks / 100.0) as f32,
+ (j.texture / 100.0) as f32,
+ (j.clarity / 100.0) as f32,
+ (j.dehaze / 100.0) as f32,
+ (j.saturation / 100.0) as f32,
+ (j.hue / 100.0) as f32 * 0.6,
+ (j.sharpness / 100.0) as f32,
+ on,
+ cos,
+ sin,
+ amt,
+ ]
+}
+
+/// Everything the per-pixel stage computes once per render: the CPU loop below and the GPU kernel
+/// (`lightcraft-gpu`) both read their parameters from here, so the two cannot drift apart.
+pub struct FinishParams {
+ pub tone: ToneMap,
+ pub ops: ColorOps,
+ /// Tone curves (parametric ∘ point, per channel) on encoded values, 1024 entries each.
+ pub curves: Option<[Lut1; 3]>,
+ pub vig: Option,
+ pub to_srgb: [[f32; 3]; 3],
+ pub hl: f32,
+ pub sh: f32,
+ pub clar: f32,
+ pub tex: f32,
+ pub dehaze: f32,
+ pub sharpen: f32,
+ pub sharpen_mask: f32,
+ /// Airlight after and before exposure, exposure gain and EV (see [`crate::Prepared`]).
+ pub air: f32,
+ pub air_pre: f32,
+ pub gain: f32,
+ pub ev: f32,
+ /// Grain: amount, cell size (px), roughness, seed.
+ pub grain: Option<(f32, f32, f32, u32)>,
+ /// Output px → normalized oriented coordinates.
+ pub out_to_norm: lightcraft_geom::Affine,
+ pub ow: f64,
+ pub oh: f64,
+ pub w: usize,
+ pub h: usize,
+ pub px_per_long: f64,
+}
+
+impl FinishParams {
+ /// Parameters for a `w × h` render; `ev`/`air_pre` as in [`crate::Prepared`].
+ pub fn new(s: &DevelopSettings, frame: &Frame, info: &SourceInfo, w: usize, h: usize, px_per_long: f64, air_pre: f32) -> FinishParams {
+ let effects = s.section_enabled("effects");
+ let (clar, tex, dehaze) = if effects {
+ ((s.effects.clarity / 100.0) as f32, (s.effects.texture / 100.0) as f32, (s.effects.dehaze / 100.0) as f32)
+ } else {
+ (0.0, 0.0, 0.0)
+ };
+ let ev = s.light.exposure as f32;
+ let gain = 2f32.powf(ev);
+ let grain = (s.grain.amount > 0.0 && effects).then(|| {
+ let cell = (0.0006 + (s.grain.size / 100.0) as f32 * 0.0024) * px_per_long as f32;
+ ((s.grain.amount / 100.0) as f32 * 0.13, cell.max(0.6), (s.grain.roughness / 100.0) as f32, s.grain.seed)
+ });
+ FinishParams {
+ tone: if info.raw {
+ ToneMap::new(s.light.contrast, s.light.whites, s.light.blacks)
+ } else {
+ ToneMap::display(s.light.contrast, s.light.whites, s.light.blacks)
+ },
+ ops: ColorOps::new(s),
+ curves: curve_luts(&s.curve),
+ vig: if effects { vignette(s) } else { None },
+ to_srgb: REC2020.to_space(&SRGB).to_f32(),
+ hl: (s.light.highlights / 100.0) as f32,
+ sh: (s.light.shadows / 100.0) as f32,
+ clar,
+ tex,
+ dehaze,
+ sharpen: (s.detail.sharpen_amount / 150.0) as f32,
+ sharpen_mask: (s.detail.sharpen_masking / 100.0) as f32,
+ air: air_pre * gain,
+ air_pre,
+ gain,
+ ev,
+ grain,
+ out_to_norm: frame.out_to_norm(w, h),
+ ow: frame.ow,
+ oh: frame.oh,
+ w,
+ h,
+ px_per_long,
+ }
+ }
+}
+
+pub(crate) fn finish(p: &Prepared, s: &DevelopSettings, frame: &Frame, info: &SourceInfo) -> Rgba8 {
let (w, h) = (p.img.width, p.img.height);
- let tone = if info.raw {
- ToneMap::new(s.light.contrast, s.light.whites, s.light.blacks)
- } else {
- ToneMap::display(s.light.contrast, s.light.whites, s.light.blacks)
- };
- let ops = ColorOps::new(s);
- let curves = curve_luts(&s.curve);
- let vig = if s.section_enabled("effects") { vignette(s) } else { None };
- let to_srgb: [[f32; 3]; 3] = REC2020.to_space(&SRGB).to_f32();
- let hl = (s.light.highlights / 100.0) as f32;
- let sh = (s.light.shadows / 100.0) as f32;
- let (clar, tex, dehaze) = if s.section_enabled("effects") {
- ((s.effects.clarity / 100.0) as f32, (s.effects.texture / 100.0) as f32, (s.effects.dehaze / 100.0) as f32)
- } else {
- (0.0, 0.0, 0.0)
- };
- let sharpen = (s.detail.sharpen_amount / 150.0) as f32;
- let sharpen_mask = (s.detail.sharpen_masking / 100.0) as f32;
- // Planes are pre-exposure: scale the airlight, shift log luminance (see `Prepared`).
- let (air, air_pre, gain, ev) = (p.air * p.gain, p.air, p.gain, p.ev);
- let grain = (s.grain.amount > 0.0 && s.section_enabled("effects")).then(|| {
- let cell = (0.0006 + (s.grain.size / 100.0) as f32 * 0.0024) * p.px_per_long as f32;
- ((s.grain.amount / 100.0) as f32 * 0.13, cell.max(0.6), (s.grain.roughness / 100.0) as f32, s.grain.seed)
- });
- let out_to_norm = frame.out_to_norm(w, h);
+ let fp = FinishParams::new(s, frame, info, w, h, p.px_per_long, p.air);
+ let FinishParams { tone, ops, curves, vig, to_srgb, hl, sh, clar, tex, dehaze, sharpen, sharpen_mask, air, air_pre, gain, ev, grain, .. } = &fp;
+ let (hl, sh, clar, tex, dehaze, sharpen, sharpen_mask, air, air_pre, gain, ev) =
+ (*hl, *sh, *clar, *tex, *dehaze, *sharpen, *sharpen_mask, *air, *air_pre, *gain, *ev);
+ let terms: Vec<[f32; MASK_TERMS]> = p.masks.iter().map(|m| mask_terms(&m.adjust)).collect();
+ let out_to_norm = fp.out_to_norm;
+ let long = fp.ow.max(fp.oh);
let aspect = w as f32 / h as f32;
let srgb = srgb_lut();
@@ -138,35 +240,22 @@ pub fn finish(p: &Prepared, s: &DevelopSettings, frame: &Frame, info: &SourceInf
let l_pre = p.log_l.data[i];
let l0 = l_pre + ev;
- // --- local (mask) contributions
- let (mut l_exp, mut l_temp, mut l_tint, mut l_con, mut l_hl, mut l_sh, mut l_wh, mut l_bl) = (0.0f32, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0);
- let (mut l_tex, mut l_clar, mut l_dehaze, mut l_sat, mut l_hue, mut l_sharp) = (0.0f32, 0.0, 0.0, 0.0, 0.0, 0.0);
+ // --- local (mask) contributions: alpha-weighted sums of the masks' terms
+ let mut l = [0.0f32; 14];
let mut tint_col: Option<([f32; 3], f32)> = None;
- for m in &p.masks {
+ for (m, t) in p.masks.iter().zip(&terms) {
let a = m.alpha.data[i];
if a <= 0.0 {
continue;
}
- let j = &m.adjust;
- l_exp += a * j.exposure as f32;
- l_temp += a * (j.temp / 100.0) as f32;
- l_tint += a * (j.tint / 100.0) as f32;
- l_con += a * (j.contrast / 100.0) as f32;
- l_hl += a * (j.highlights / 100.0) as f32;
- l_sh += a * (j.shadows / 100.0) as f32;
- l_wh += a * (j.whites / 100.0) as f32;
- l_bl += a * (j.blacks / 100.0) as f32;
- l_tex += a * (j.texture / 100.0) as f32;
- l_clar += a * (j.clarity / 100.0) as f32;
- l_dehaze += a * (j.dehaze / 100.0) as f32;
- l_sat += a * (j.saturation / 100.0) as f32;
- l_hue += a * (j.hue / 100.0) as f32 * 0.6;
- l_sharp += a * (j.sharpness / 100.0) as f32;
- if j.color_sat > 0.0 {
- let hue = crate::colorops::oklch_hue_of_srgb_hue(j.color_hue);
- tint_col = Some(([hue.cos(), hue.sin(), 0.0], a * (j.color_sat / 100.0) as f32));
+ for k in 0..14 {
+ l[k] += a * t[k];
+ }
+ if t[14] > 0.0 {
+ tint_col = Some(([t[15], t[16], 0.0], a * t[17]));
}
}
+ let [l_exp, l_temp, l_tint, l_con, l_hl, l_sh, l_wh, l_bl, l_tex, l_clar, l_dehaze, l_sat, l_hue, l_sharp] = l;
// --- dehaze (scene linear)
let dz = dehaze + l_dehaze;
@@ -259,7 +348,7 @@ pub fn finish(p: &Prepared, s: &DevelopSettings, frame: &Frame, info: &SourceInf
}
// --- vignette (display linear, post-crop)
- if let Some(v) = &vig {
+ if let Some(v) = vig {
let u = (x as f32 + 0.5) / w as f32 * 2.0 - 1.0;
let vv = (y as f32 + 0.5) / h as f32 * 2.0 - 1.0;
let sx = 1.0 + (aspect - 1.0) * v.aspect_mix;
@@ -286,7 +375,7 @@ pub fn finish(p: &Prepared, s: &DevelopSettings, frame: &Frame, info: &SourceInf
}
// --- gamut map to sRGB (desaturate towards luminance until in range)
- let m = &to_srgb;
+ let m = to_srgb;
let mut r = [
m[0][0] * d[0] + m[0][1] * d[1] + m[0][2] * d[2],
m[1][0] * d[0] + m[1][1] * d[1] + m[1][2] * d[2],
@@ -307,13 +396,13 @@ pub fn finish(p: &Prepared, s: &DevelopSettings, frame: &Frame, info: &SourceInf
// --- encode, curves, grain
let mut e = r.map(|v| encode_srgb(srgb, v));
- if let Some(l) = &curves {
+ if let Some(l) = curves {
e = [l[0].eval(e[0]), l[1].eval(e[1]), l[2].eval(e[2])];
}
- if let Some((amt, cell, rough, seed)) = grain {
+ if let Some((amt, cell, rough, seed)) = *grain {
let n = out_to_norm.apply(Point::new(x as f64 + 0.5, y as f64 + 0.5));
- let (gx, gy) = ((n.x * frame.ow) as f32 / frame.ow.max(frame.oh) as f32, (n.y * frame.oh) as f32 / frame.ow.max(frame.oh) as f32);
- let sc = p.px_per_long as f32 / cell;
+ let (gx, gy) = ((n.x * fp.ow) as f32 / long as f32, (n.y * fp.oh) as f32 / long as f32);
+ let sc = fp.px_per_long as f32 / cell;
let mut g = grain_noise(gx * sc, gy * sc, seed);
g = g * (1.0 - rough * 0.5) + grain_noise(gx * sc * 2.3, gy * sc * 2.3, seed ^ 0x55) * rough * 0.7;
let lum = 0.2126 * e[0] + 0.7152 * e[1] + 0.0722 * e[2];
@@ -326,9 +415,10 @@ pub fn finish(p: &Prepared, s: &DevelopSettings, frame: &Frame, info: &SourceInf
out
}
-const SRGB_LUT_N: usize = 4096;
+pub const SRGB_LUT_N: usize = 4096;
-fn srgb_lut() -> &'static [f32; SRGB_LUT_N + 1] {
+/// Linear → sRGB-encoded table (`SRGB_LUT_N + 1` entries over 0..1), interpolated linearly.
+pub fn srgb_lut() -> &'static [f32; SRGB_LUT_N + 1] {
static LUT: std::sync::OnceLock> = std::sync::OnceLock::new();
LUT.get_or_init(|| {
let mut t = Box::new([0.0f32; SRGB_LUT_N + 1]);
diff --git a/crates/pipeline/src/geometry.rs b/crates/pipeline/src/geometry.rs
index 5742c8ab..a8ed0cd2 100644
--- a/crates/pipeline/src/geometry.rs
+++ b/crates/pipeline/src/geometry.rs
@@ -177,37 +177,48 @@ impl Frame {
Point::new(p.x * self.ow / l, p.y * self.oh / l)
}
+ /// How [`Frame::sample`] resamples a (EXIF-oriented) source of `src_w × src_h` into `w × h`.
+ pub fn sample_plan(&self, src_w: usize, src_h: usize, w: usize, h: usize) -> SamplePlan {
+ let (ow, oh) = if self.orient.swaps_axes() { (src_h, src_w) } else { (src_w, src_h) };
+ let crop_px = self.crop.rect_px(self.ow, self.oh);
+ let k = crop_px.width() / w as f64;
+ let prefilter = (k > 1.25).then(|| (((ow as f64 / k).round() as usize).max(1), ((oh as f64 / k).round() as usize).max(1)));
+ let (bw, bh) = prefilter.unwrap_or((ow, oh));
+ let (sx, sy) = (bw as f64 / ow as f64, bh as f64 / oh as f64);
+ let mode = if self.warp.is_some() {
+ SampleMode::Warp(self.out_to_oriented(w, h))
+ } else if self.crop.is_identity() && !self.flip_h && !self.flip_v && bw == w && bh == h {
+ SampleMode::Copy
+ } else {
+ SampleMode::Affine(Affine::scale(sx, sy) * self.out_to_oriented(w, h))
+ };
+ SamplePlan { ow, oh, prefilter, sx, sy, mode }
+ }
+
/// Sample the source into a `w × h` output buffer: orientation, crop, rotation, flips — one
/// bilinear resample from a pre-filtered (area-downscaled) copy, so minification never aliases.
pub fn sample(&self, src: &Rgb32f, w: usize, h: usize) -> Rgb32f {
let oriented = if self.orient == Orientation::Normal { None } else { Some(src.oriented(self.orient)) };
let o = oriented.as_ref().unwrap_or(src);
- let crop_px = self.crop.rect_px(self.ow, self.oh);
- let k = crop_px.width() / w as f64;
- let (base, s) = if k > 1.25 {
- let nw = ((o.width as f64 / k).round() as usize).max(1);
- let nh = ((o.height as f64 / k).round() as usize).max(1);
- (resize(o, nw, nh, Filter::Mitchell), nw as f64 / o.width as f64)
- } else {
- (o.clone(), 1.0)
+ let plan = self.sample_plan(src.width, src.height, w, h);
+ let base = match plan.prefilter {
+ Some((nw, nh)) => std::borrow::Cow::Owned(resize(o, nw, nh, Filter::Mitchell)),
+ None => std::borrow::Cow::Borrowed(o),
};
- let (sx, sy) = (s, base.height as f64 / o.height as f64);
- if let Some(wp) = &self.warp {
- return sample_warped(&base, sx, sy, wp, self.out_to_oriented(w, h), w, h);
- }
- let xf = Affine::scale(sx, sy) * self.out_to_oriented(w, h);
- let identity_like = self.crop.is_identity() && !self.flip_h && !self.flip_v && base.width == w && base.height == h;
- if identity_like {
- return base;
- }
- let mut out = Rgb32f::new(w, h);
- par_rows(&mut out.data, w, |y, row| {
- for (x, px) in row.iter_mut().enumerate() {
- let p = xf.apply(Point::new(x as f64 + 0.5, y as f64 + 0.5));
- *px = base.sample_bilinear(p.x as f32, p.y as f32);
+ match plan.mode {
+ SampleMode::Copy => base.into_owned(),
+ SampleMode::Warp(o2t) => sample_warped(&base, plan.sx, plan.sy, self.warp.as_ref().expect("warp"), o2t, w, h),
+ SampleMode::Affine(xf) => {
+ let mut out = Rgb32f::new(w, h);
+ par_rows(&mut out.data, w, |y, row| {
+ for (x, px) in row.iter_mut().enumerate() {
+ let p = xf.apply(Point::new(x as f64 + 0.5, y as f64 + 0.5));
+ *px = base.sample_bilinear(p.x as f32, p.y as f32);
+ }
+ });
+ out
}
- });
- out
+ }
}
/// The crop rectangle as a quad in normalized oriented coordinates (for overlays).
@@ -221,6 +232,30 @@ impl Frame {
}
}
+/// How a render resamples its source (see [`Frame::sample_plan`]).
+#[derive(Clone, Debug)]
+pub struct SamplePlan {
+ /// Oriented source size.
+ pub ow: usize,
+ pub oh: usize,
+ /// Size of the Mitchell prefilter (area downscale before the bilinear resample), if any.
+ pub prefilter: Option<(usize, usize)>,
+ /// Oriented-source px → prefiltered px.
+ pub sx: f64,
+ pub sy: f64,
+ pub mode: SampleMode,
+}
+
+#[derive(Clone, Debug)]
+pub enum SampleMode {
+ /// The (prefiltered) oriented source is the output.
+ Copy,
+ /// Bilinear through this affine (output px → prefiltered px).
+ Affine(Affine),
+ /// Through the warp; the affine maps output px → transformed px.
+ Warp(Affine),
+}
+
/// The perspective homography of the settings (Upright, then the manual Transform sliders), lens-corrected →
/// transformed, in centred coordinates.
pub fn perspective(s: &DevelopSettings, ow: f64, oh: f64) -> Homography {
diff --git a/crates/pipeline/src/lib.rs b/crates/pipeline/src/lib.rs
index ac62a327..3ccee51b 100644
--- a/crates/pipeline/src/lib.rs
+++ b/crates/pipeline/src/lib.rs
@@ -23,19 +23,21 @@
#![forbid(unsafe_code)]
pub mod auto;
-mod colorops;
-mod finish;
+pub mod colorops;
+pub mod finish;
pub mod geometry;
-mod local;
+pub mod local;
pub mod masks;
pub mod optics;
pub mod profiles;
pub mod spots;
-mod tone;
+pub mod tone;
pub mod transform;
pub mod upright;
use lightcraft_develop::{DevelopSettings, Treatment};
+use std::any::Any;
+use std::borrow::Cow;
use std::sync::{Arc, Mutex};
use lightcraft_raster::{Histogram, Plane, Rgb32f, Rgba8, par_rows};
@@ -105,9 +107,6 @@ pub(crate) struct Prepared {
/// Airlight of `dark` (before exposure).
pub air: f32,
pub masks: Vec,
- /// Exposure in EV and as a linear gain.
- pub ev: f32,
- pub gain: f32,
/// Output pixels per unit of the source long edge.
pub px_per_long: f64,
}
@@ -137,11 +136,13 @@ pub fn frame_for(src: &Rgb32f, info: &SourceInfo, s: &DevelopSettings, apply_cro
pub struct StageCache {
entries: Mutex>,
capacity: usize,
+ /// Another renderer's per-view state (the GPU renderer keeps its device-resident stages here).
+ ext: Mutex>>,
}
impl Default for StageCache {
fn default() -> Self {
- StageCache { entries: Mutex::new(Vec::new()), capacity: 2 }
+ StageCache { entries: Mutex::new(Vec::new()), capacity: 2, ext: Mutex::new(None) }
}
}
@@ -149,6 +150,21 @@ impl StageCache {
/// Drop everything (e.g. when memory is needed).
pub fn clear(&self) {
self.lock().clear();
+ *self.ext.lock().unwrap_or_else(|e| e.into_inner()) = None;
+ }
+
+ /// Per-view state of type `T` kept alongside this cache (created on first use). Lets another
+ /// renderer of the same view (e.g. `lightcraft-gpu`) keep its own stage cache with this one.
+ pub fn extension(&self) -> Arc {
+ let mut g = self.ext.lock().unwrap_or_else(|e| e.into_inner());
+ if let Some(e) = g.as_ref()
+ && let Ok(t) = e.clone().downcast::()
+ {
+ return t;
+ }
+ let t = Arc::new(T::default());
+ *g = Some(t.clone());
+ t
}
/// Number of cached output sizes.
@@ -192,6 +208,66 @@ fn hash_of(parts: impl std::hash::Hash) -> u64 {
BuildHasherDefault::::default().hash_one(parts)
}
+/// What a render resolves to before any pixel work: the effective settings (profile and Upright
+/// applied), the geometric frame, the output size and the stage-cache keys. Shared by the CPU
+/// renderer and the GPU renderer (`lightcraft-gpu`), so both key their caches identically.
+pub struct Plan<'a> {
+ pub settings: Cow<'a, DevelopSettings>,
+ pub frame: geometry::Frame,
+ pub w: usize,
+ pub h: usize,
+ /// Output pixels per unit of the oriented source's long edge.
+ pub px_per_long: f64,
+ /// Long edge of the source buffer (px).
+ pub src_long: usize,
+ /// Key of the resampled source (together with the source buffer's identity).
+ pub geo: u64,
+ /// Key of the white-balanced, retouched, denoised image (and of the planes computed from it).
+ pub lin_key: u64,
+}
+
+/// Resolve `s` against `src` for `req` (see [`Plan`]).
+pub fn plan<'a>(src: &Rgb32f, info: &SourceInfo, s: &'a DevelopSettings, req: &RenderRequest) -> Plan<'a> {
+ let settings: Cow<'a, DevelopSettings> = match profiles::effective(s) {
+ Cow::Borrowed(b) => upright::resolve(src, info, b),
+ Cow::Owned(o) => Cow::Owned(upright::resolve(src, info, &o).into_owned()),
+ };
+ let s = &*settings;
+ let frame = frame_for(src, info, s, req.apply_crop);
+ let (w, h) = frame.fit(req.max_w, req.max_h);
+ let px_per_long = frame.px_per_long(w);
+ let src_long = src.width.max(src.height);
+ let geo = hash_of((format!("{frame:?}"), w, h));
+ let (wb_t, wb_tint) = local::effective_wb(info, s);
+ let d = &s.detail;
+ let lin_key = hash_of((
+ geo,
+ [wb_t, wb_tint, info.as_shot_temp, info.as_shot_tint].map(f64::to_bits),
+ format!("{:?}", s.spots),
+ // defringe runs in this stage
+ format!("{:?}", s.optics),
+ [d.nr_luminance, d.nr_detail, d.nr_color, d.nr_color_detail, d.nr_color_smoothness].map(f64::to_bits),
+ src_long,
+ ));
+ Plan { settings, frame, w, h, px_per_long, src_long, geo, lin_key }
+}
+
+/// Whether the scene-linear stage needs work only the CPU does (defringe, spot removal).
+pub fn lin_needs_cpu(s: &DevelopSettings) -> bool {
+ let o = &s.optics;
+ let defringe = s.section_enabled("optics") && (o.defringe_purple_amount > 0.0 || o.defringe_green_amount > 0.0);
+ defringe || !s.spots.is_empty()
+}
+
+/// The white-balanced, defringed, retouched image (before noise reduction): the CPU part of the
+/// scene-linear stage, in place.
+pub fn lin_cpu(img: &mut Rgb32f, info: &SourceInfo, p: &Plan<'_>) {
+ let s = &*p.settings;
+ local::white_balance(img, info, s);
+ optics::defringe(img, s, p.px_per_long / optics::DEFRINGE_REF_LONG);
+ spots::apply(img, &s.spots, &p.frame, p.px_per_long);
+}
+
/// Render `src` with settings `s`.
pub fn render(src: &Rgb32f, info: &SourceInfo, s: &DevelopSettings, req: &RenderRequest) -> Rendered {
render_impl(Src::Borrowed(src), info, s, req, None)
@@ -217,22 +293,18 @@ fn render_impl(src: Src<'_>, info: &SourceInfo, s: &DevelopSettings, req: &Rende
}
};
let mut t = profiling().then(std::time::Instant::now);
- let s = &*profiles::effective(s);
let src_img: &Rgb32f = match &src {
Src::Borrowed(r) => r,
Src::Shared(a) => a,
};
- let s = &*upright::resolve(src_img, info, s);
- let frame = frame_for(src_img, info, s, req.apply_crop);
- let (w, h) = frame.fit(req.max_w, req.max_h);
- let px_per_long = frame.px_per_long(w);
- let src_long = src_img.width.max(src_img.height);
+ let plan = plan(src_img, info, s, req);
+ let Plan { ref frame, w, h, px_per_long, src_long, geo, lin_key, .. } = plan;
+ let s = &*plan.settings;
let shared = match (&src, cache) {
(Src::Shared(a), Some(c)) => Some((*a, c)),
_ => None,
};
- let geo = hash_of((format!("{frame:?}"), w, h));
let cached = shared.and_then(|(a, c)| c.get(a, geo));
let sampled = match &cached {
Some(e) => e.sampled.clone(),
@@ -240,25 +312,12 @@ fn render_impl(src: Src<'_>, info: &SourceInfo, s: &DevelopSettings, req: &Rende
};
lap("sample", &mut t);
- let (wb_t, wb_tint) = local::effective_wb(info, s);
- let d = &s.detail;
- let lin_key = hash_of((
- geo,
- [wb_t, wb_tint, info.as_shot_temp, info.as_shot_tint].map(f64::to_bits),
- format!("{:?}", s.spots),
- // defringe runs in this stage
- format!("{:?}", s.optics),
- [d.nr_luminance, d.nr_detail, d.nr_color, d.nr_color_detail, d.nr_color_smoothness].map(f64::to_bits),
- src_long,
- ));
let lin = match cached.as_ref().and_then(|e| e.lin.clone()).filter(|(k, _)| *k == lin_key) {
Some((_, img)) => img,
None => {
// Without a cache the resampled buffer is ours: work on it in place.
let mut img = if shared.is_some() { (*sampled).clone() } else { Arc::unwrap_or_clone(sampled.clone()) };
- local::white_balance(&mut img, info, s);
- optics::defringe(&mut img, s, px_per_long / optics::DEFRINGE_REF_LONG);
- spots::apply(&mut img, &s.spots, &frame, px_per_long);
+ lin_cpu(&mut img, info, &plan);
local::denoise(&mut img, s, src_long, w.max(h));
Arc::new(img)
}
@@ -268,12 +327,12 @@ fn render_impl(src: Src<'_>, info: &SourceInfo, s: &DevelopSettings, req: &Rende
Some(p) if p.key == lin_key => p,
_ => local::Planes { key: lin_key, ..Default::default() },
};
- let prep = local::prepare(lin.clone(), s, &frame, px_per_long, req.quality, &mut planes);
+ let prep = local::prepare(lin.clone(), s, frame, px_per_long, req.quality, &mut planes);
lap("prepare", &mut t);
if let Some((a, c)) = shared {
c.put(CacheEntry { src: a.clone(), geo, sampled, lin: Some((lin_key, lin)), planes });
}
- let image = finish::finish(&prep, s, &frame, info);
+ let image = finish::finish(&prep, s, frame, info);
lap("finish", &mut t);
let histogram = Histogram::of_srgb8(&image);
lap("histogram", &mut t);
diff --git a/crates/pipeline/src/local.rs b/crates/pipeline/src/local.rs
index 27bb1d0d..8efb4126 100644
--- a/crates/pipeline/src/local.rs
+++ b/crates/pipeline/src/local.rs
@@ -22,19 +22,24 @@ pub fn white_balance(img: &mut Rgb32f, info: &SourceInfo, s: &DevelopSettings) {
wb_gain(img, info, s, 1.0);
}
-fn wb_gain(img: &mut Rgb32f, info: &SourceInfo, s: &DevelopSettings, gain: f32) {
+/// The white-balance matrix (linear Rec.2020, luminance-preserving) for the settings, or `None`
+/// when the as-shot white is kept.
+pub fn wb_matrix_for(info: &SourceInfo, s: &DevelopSettings) -> Option<[[f32; 3]; 3]> {
let (t, tint) = effective_wb(info, s);
- let m = if (t - info.as_shot_temp).abs() < 1e-6 && (tint - info.as_shot_tint).abs() < 1e-6 {
- None
- } else {
- let set = wb_matrix(&REC2020, temp_tint_to_xy(t, tint));
- let shot = wb_matrix(&REC2020, temp_tint_to_xy(info.as_shot_temp, info.as_shot_tint));
- let m = set.mul(&shot.inverse().unwrap_or(lightcraft_color::Mat3::IDENTITY));
- // Normalize so neutral luminance is preserved (WB shouldn't change exposure).
- let g = m.apply([1.0, 1.0, 1.0]);
- let y = g[0] * 0.2627 + g[1] * 0.6780 + g[2] * 0.0593;
- Some(m.mul(&lightcraft_color::Mat3::diag(1.0 / y, 1.0 / y, 1.0 / y)).to_f32())
- };
+ if (t - info.as_shot_temp).abs() < 1e-6 && (tint - info.as_shot_tint).abs() < 1e-6 {
+ return None;
+ }
+ let set = wb_matrix(&REC2020, temp_tint_to_xy(t, tint));
+ let shot = wb_matrix(&REC2020, temp_tint_to_xy(info.as_shot_temp, info.as_shot_tint));
+ let m = set.mul(&shot.inverse().unwrap_or(lightcraft_color::Mat3::IDENTITY));
+ // Normalize so neutral luminance is preserved (WB shouldn't change exposure).
+ let g = m.apply([1.0, 1.0, 1.0]);
+ let y = g[0] * 0.2627 + g[1] * 0.6780 + g[2] * 0.0593;
+ Some(m.mul(&lightcraft_color::Mat3::diag(1.0 / y, 1.0 / y, 1.0 / y)).to_f32())
+}
+
+fn wb_gain(img: &mut Rgb32f, info: &SourceInfo, s: &DevelopSettings, gain: f32) {
+ let m = wb_matrix_for(info, s);
let w = img.width;
for_rows(&mut img.data, w, |_, row| {
for p in row.iter_mut() {
@@ -96,7 +101,7 @@ fn guided_apply(p: &Plane, a: &Plane, b: &Plane) -> Plane {
/// large windows at a fraction of the cost.
pub fn guided_fast(p: &Plane, sigma: f32, eps: f32) -> Plane {
use lightcraft_raster::resample::{Filter, resize};
- let s = (sigma / 3.0).floor().clamp(1.0, 16.0) as usize;
+ let s = guided_fast_step(sigma);
if s <= 1 {
return guided(p, sigma, eps);
}
@@ -107,28 +112,57 @@ pub fn guided_fast(p: &Plane, sigma: f32, eps: f32) -> Plane {
guided_apply(p, &up(&ma), &up(&mb))
}
+/// Subsampling step of [`guided_fast`] for `sigma` (1 = no subsampling: the plain guided filter).
+pub fn guided_fast_step(sigma: f32) -> usize {
+ (sigma / 3.0).floor().clamp(1.0, 16.0) as usize
+}
+
+/// Luminance noise reduction: guided filter of log luminance (`sigma` px, `eps`), blended by `k`.
+#[derive(Clone, Copy, Debug, PartialEq)]
+pub struct NrLum {
+ pub sigma: f32,
+ pub eps: f32,
+ pub k: f32,
+}
+
+/// Colour noise reduction: chromaticity blurred by `sigma` px, mixed in by `t`.
+#[derive(Clone, Copy, Debug, PartialEq)]
+pub struct NrColor {
+ pub sigma: f32,
+ pub t: f32,
+}
+
+/// Noise-reduction parameters at an output long edge of `out_long` px (see [`denoise`]).
+pub fn nr_params(s: &DevelopSettings, src_long: usize, out_long: usize) -> (Option, Option) {
+ let lum = (s.detail.nr_luminance / 100.0) as f32;
+ let col = (s.detail.nr_color / 100.0) as f32;
+ let scale = (out_long as f32 / src_long.max(1) as f32).clamp(0.05, 1.0);
+ let l = (lum > 0.0).then(|| {
+ let detail = (s.detail.nr_detail / 100.0) as f32;
+ NrLum { sigma: (1.0 + 2.5 * lum) * scale.max(0.4), eps: 0.002 + lum * lum * 0.25 * (1.0 - 0.7 * detail), k: lum.sqrt() }
+ });
+ let c = (col > 0.0).then(|| {
+ let sigma = (1.5 + 6.0 * col) * scale.max(0.35) * (1.0 + (s.detail.nr_color_smoothness / 100.0) as f32);
+ let keep = (s.detail.nr_color_detail / 100.0) as f32 * 0.5;
+ NrColor { sigma, t: col * (1.0 - keep) }
+ });
+ (l, c)
+}
+
/// Noise reduction at output resolution: luminance via an edge-aware self-guided filter on
/// log-luminance, colour by blurring chromaticity (rgb / Y) and re-applying the original luminance.
/// Radii scale with how much the source was downsampled (preview noise is already averaged out).
pub fn denoise(img: &mut Rgb32f, s: &DevelopSettings, src_long: usize, out_long: usize) {
- let lum = (s.detail.nr_luminance / 100.0) as f32;
- let col = (s.detail.nr_color / 100.0) as f32;
- if lum <= 0.0 && col <= 0.0 {
- return;
- }
- let scale = (out_long as f32 / src_long.max(1) as f32).clamp(0.05, 1.0);
+ let (lum, col) = nr_params(s, src_long, out_long);
let w = img.width;
- if lum > 0.0 {
+ if let Some(nr) = lum {
let _t = crate::profiling().then(std::time::Instant::now);
let l = img.map(log_lum);
- let detail = (s.detail.nr_detail / 100.0) as f32;
- let eps = 0.002 + lum * lum * 0.25 * (1.0 - 0.7 * detail);
- let f = guided(&l, (1.0 + 2.5 * lum) * scale.max(0.4), eps);
- let k = lum.sqrt();
+ let f = guided(&l, nr.sigma, nr.eps);
for_rows(&mut img.data, w, |y, row| {
for (x, p) in row.iter_mut().enumerate() {
let i = y * w + x;
- let d = (f.data[i] - l.data[i]) * k;
+ let d = (f.data[i] - l.data[i]) * nr.k;
let g = d.exp2();
*p = p.map(|v| v * g);
}
@@ -137,23 +171,20 @@ pub fn denoise(img: &mut Rgb32f, s: &DevelopSettings, src_long: usize, out_long:
eprintln!(" nr luminance: {:.1} ms", t.elapsed().as_secs_f64() * 1e3);
}
}
- if col > 0.0 {
+ if let Some(nr) = col {
let _t = crate::profiling().then(std::time::Instant::now);
let chroma = img.map(|c| {
let y = luminance_2020(c).max(1e-6);
[c[0] / y, c[1] / y, c[2] / y]
});
- let sigma = (1.5 + 6.0 * col) * scale.max(0.35) * (1.0 + (s.detail.nr_color_smoothness / 100.0) as f32);
- let b = gaussian(&chroma, sigma);
- let keep = (s.detail.nr_color_detail / 100.0) as f32 * 0.5;
+ let b = gaussian(&chroma, nr.sigma);
for_rows(&mut img.data, w, |y, row| {
for (x, p) in row.iter_mut().enumerate() {
let i = y * w + x;
let yl = luminance_2020(*p);
let c0 = chroma.data[i];
let cb = b.data[i];
- let t = col * (1.0 - keep);
- *p = [0, 1, 2].map(|k| ((c0[k] + (cb[k] - c0[k]) * t) * yl).max(0.0));
+ *p = [0, 1, 2].map(|k| ((c0[k] + (cb[k] - c0[k]) * nr.t) * yl).max(0.0));
}
});
if let Some(t) = _t {
@@ -191,32 +222,60 @@ fn plane_at(slot: &mut Option<(u32, Arc)>, sigma: f32, f: impl FnOnce() -
}
}
-/// The spatial planes the per-pixel stage needs for `img` (white-balanced, before exposure),
-/// reusing whatever `planes` already holds for it. Missing planes are computed side by side (each
-/// one alone scales poorly: the guided filters work on small subsampled grids).
-pub(crate) fn prepare(img: Arc, s: &DevelopSettings, frame: &Frame, px_per_long: f64, q: Quality, planes: &mut Planes) -> Prepared {
- let log_l = planes.log_l.get_or_insert_with(|| Arc::new(timed("log_l", || img.map(log_lum)))).clone();
+/// Guided-filter epsilons (EV²) of the highlights/shadows base and the clarity band.
+pub const BASE_EPS: f32 = 0.35;
+pub const CLARITY_EPS: f32 = 0.8;
+
+/// Dark channel of a pixel (dehaze).
+#[inline]
+pub fn dark_of(c: [f32; 3]) -> f32 {
+ c[0].min(c[1]).min(c[2])
+}
+
+/// Radii (px) of the spatial planes the settings need (`None` = not needed).
+#[derive(Clone, Copy, Debug, Default, PartialEq)]
+pub struct PlaneSigmas {
+ /// Edge-aware base for highlights/shadows (fast guided filter, [`BASE_EPS`]).
+ pub base: Option,
+ /// Clarity band (fast guided filter, [`CLARITY_EPS`]).
+ pub clarity: Option,
+ /// Texture / sharpening band (Gaussian).
+ pub texture: Option,
+ /// Dehaze dark channel (Gaussian).
+ pub dark: Option,
+}
+
+pub fn plane_sigmas(s: &DevelopSettings, px_per_long: f64, q: Quality) -> PlaneSigmas {
let ppl = px_per_long as f32;
let local_any = |f: fn(&lightcraft_develop::LocalAdjustments) -> f64| s.masks.iter().any(|m| f(&m.adjust) != 0.0);
let tone_active =
s.light.highlights != 0.0 || s.light.shadows != 0.0 || s.masks.iter().any(|m| m.adjust.highlights != 0.0 || m.adjust.shadows != 0.0);
// Edge-aware base at ~1.5% of the long edge (EV² epsilon: edges of > ~0.6 EV are preserved).
- let base_sigma = tone_active.then(|| {
+ let base = tone_active.then(|| {
let sigma = (0.015 * ppl).max(1.0);
if q == Quality::Draft { sigma.min(24.0) } else { sigma }
});
- let clarity_sigma = (s.effects.clarity != 0.0 || local_any(|a| a.clarity)).then(|| (0.012 * ppl).max(1.0));
- let texture_sigma = (s.effects.texture != 0.0 || s.detail.sharpen_amount != 0.0 || local_any(|a| a.texture) || local_any(|a| a.sharpness))
+ let clarity = (s.effects.clarity != 0.0 || local_any(|a| a.clarity)).then(|| (0.012 * ppl).max(1.0));
+ let texture = (s.effects.texture != 0.0 || s.detail.sharpen_amount != 0.0 || local_any(|a| a.texture) || local_any(|a| a.sharpness))
.then(|| (0.0018 * ppl).max(0.6));
- let dark_sigma = (s.effects.dehaze != 0.0 || local_any(|a| a.dehaze)).then(|| (0.02 * ppl).max(1.0));
+ let dark = (s.effects.dehaze != 0.0 || local_any(|a| a.dehaze)).then(|| (0.02 * ppl).max(1.0));
+ PlaneSigmas { base, clarity, texture, dark }
+}
+
+/// The spatial planes the per-pixel stage needs for `img` (white-balanced, before exposure),
+/// reusing whatever `planes` already holds for it. Missing planes are computed side by side (each
+/// one alone scales poorly: the guided filters work on small subsampled grids).
+pub(crate) fn prepare(img: Arc, s: &DevelopSettings, frame: &Frame, px_per_long: f64, q: Quality, planes: &mut Planes) -> Prepared {
+ let log_l = planes.log_l.get_or_insert_with(|| Arc::new(timed("log_l", || img.map(log_lum)))).clone();
+ let PlaneSigmas { base: base_sigma, clarity: clarity_sigma, texture: texture_sigma, dark: dark_sigma } = plane_sigmas(s, px_per_long, q);
let Planes { base: sb, clarity: sc, texture: st, dark: sd, .. } = planes;
let l = &log_l;
let (base, (clarity_blur, (texture_blur, dark))) = par_join(
- || base_sigma.map(|sg| plane_at(sb, sg, || timed("base", || guided_fast(l, sg, 0.35)))),
+ || base_sigma.map(|sg| plane_at(sb, sg, || timed("base", || guided_fast(l, sg, BASE_EPS)))),
|| {
par_join(
- || clarity_sigma.map(|sg| plane_at(sc, sg, || timed("clarity", || guided_fast(l, sg, 0.8)))),
+ || clarity_sigma.map(|sg| plane_at(sc, sg, || timed("clarity", || guided_fast(l, sg, CLARITY_EPS)))),
|| {
par_join(
|| texture_sigma.map(|sg| plane_at(st, sg, || timed("texture", || gaussian(l, sg)))),
@@ -224,7 +283,7 @@ pub(crate) fn prepare(img: Arc, s: &DevelopSettings, frame: &Frame, px_p
dark_sigma.map(|sg| match sd {
Some((k, d, air)) if *k == sg.to_bits() => (d.clone(), *air),
_ => {
- let d = timed("dehaze", || Arc::new(gaussian(&img.map(|c| c[0].min(c[1]).min(c[2])), sg)));
+ let d = timed("dehaze", || Arc::new(gaussian(&img.map(dark_of), sg)));
let air = airlight(&d);
*sd = Some((sg.to_bits(), d.clone(), air));
(d, air)
@@ -242,14 +301,20 @@ pub(crate) fn prepare(img: Arc, s: &DevelopSettings, frame: &Frame, px_p
None => (None, 1.0),
};
let ev = s.light.exposure as f32;
- let gain = 2f32.powf(ev);
let masks = timed("masks", || masks::evaluate(&s.masks, frame, img.width, img.height, &img, &log_l, ev));
- Prepared { img, log_l, base, clarity_blur, texture_blur, dark, air, masks, ev, gain, px_per_long }
+ Prepared { img, log_l, base, clarity_blur, texture_blur, dark, air, masks, px_per_long }
}
+/// The airlight is estimated from every `AIRLIGHT_STEP`-th value of the dark channel.
+pub const AIRLIGHT_STEP: usize = 7;
+
/// Airlight estimate: bright end of the dark channel.
pub fn airlight(dark: &Plane) -> f32 {
- let mut v: Vec = dark.data.iter().step_by(7).copied().collect();
+ airlight_of(dark.data.iter().step_by(AIRLIGHT_STEP).copied().collect())
+}
+
+/// Airlight from the sampled dark-channel values (see [`airlight`]): their 99.5th percentile.
+pub fn airlight_of(mut v: Vec) -> f32 {
if v.is_empty() {
return 1.0;
}
diff --git a/crates/pipeline/src/masks.rs b/crates/pipeline/src/masks.rs
index bd27f90f..8f1796a7 100644
--- a/crates/pipeline/src/masks.rs
+++ b/crates/pipeline/src/masks.rs
@@ -172,29 +172,44 @@ fn tonemap_for_select(c: [f32; 3]) -> [f32; 3] {
c.map(|v| v / (1.0 + v))
}
-/// Stamp brush strokes into `out` (max-combine; erase strokes subtract).
-fn rasterize_brush(strokes: &[BrushStroke], frame: &Frame, w: usize, h: usize, out: &mut Plane) {
+/// A brush stroke resolved to output pixels: dab centres (spacing r/4 along the path), radius,
+/// hard-core radius, flow and density (0..1), and whether it erases.
+pub struct BrushDabs {
+ pub dabs: Vec,
+ pub r: f64,
+ pub hard: f64,
+ pub flow: f32,
+ pub density: f32,
+ pub erase: bool,
+}
+
+/// The dabs of stroke `s` in a `w × h` output.
+pub fn brush_dabs(s: &BrushStroke, frame: &Frame, w: usize, h: usize) -> BrushDabs {
let to_out = frame.norm_to_out(w, h);
let ppl = frame.px_per_long(w);
- for s in strokes {
- let r = (s.size * ppl).max(0.5);
- let hard = r * (1.0 - (s.feather / 100.0).clamp(0.0, 1.0));
- let flow = (s.flow / 100.0).clamp(0.0, 1.0) as f32;
- let dens = (s.density / 100.0).clamp(0.0, 1.0) as f32;
- // Densify the path so dabs overlap (spacing r/4).
- let mut dabs: Vec = Vec::new();
- for (i, p) in s.points.iter().enumerate() {
- let q = to_out.apply(*p);
- if i > 0 {
- let prev = to_out.apply(s.points[i - 1]);
- let dist = prev.dist(q);
- let n = (dist / (r / 4.0).max(0.5)).ceil() as usize;
- for k in 1..n {
- dabs.push(prev.lerp(q, k as f64 / n as f64));
- }
+ let r = (s.size * ppl).max(0.5);
+ let hard = r * (1.0 - (s.feather / 100.0).clamp(0.0, 1.0));
+ // Densify the path so dabs overlap (spacing r/4).
+ let mut dabs: Vec = Vec::new();
+ for (i, p) in s.points.iter().enumerate() {
+ let q = to_out.apply(*p);
+ if i > 0 {
+ let prev = to_out.apply(s.points[i - 1]);
+ let dist = prev.dist(q);
+ let n = (dist / (r / 4.0).max(0.5)).ceil() as usize;
+ for k in 1..n {
+ dabs.push(prev.lerp(q, k as f64 / n as f64));
}
- dabs.push(q);
}
+ dabs.push(q);
+ }
+ BrushDabs { dabs, r, hard, flow: (s.flow / 100.0).clamp(0.0, 1.0) as f32, density: (s.density / 100.0).clamp(0.0, 1.0) as f32, erase: s.erase }
+}
+
+/// Stamp brush strokes into `out` (max-combine; erase strokes subtract).
+fn rasterize_brush(strokes: &[BrushStroke], frame: &Frame, w: usize, h: usize, out: &mut Plane) {
+ for s in strokes {
+ let BrushDabs { dabs, r, hard, flow, density: dens, erase } = brush_dabs(s, frame, w, h);
let mut stroke_alpha = Plane::new(w, h);
for d in &dabs {
let (x0, x1) = (((d.x - r).floor().max(0.0)) as usize, ((d.x + r).ceil().max(0.0) as usize).min(w));
@@ -213,7 +228,7 @@ fn rasterize_brush(strokes: &[BrushStroke], frame: &Frame, w: usize, h: usize, o
}
}
for (o, a) in out.data.iter_mut().zip(&stroke_alpha.data) {
- *o = if s.erase { *o * (1.0 - a) } else { o.max(*a) };
+ *o = if erase { *o * (1.0 - a) } else { o.max(*a) };
}
}
}
diff --git a/crates/pipeline/src/tone.rs b/crates/pipeline/src/tone.rs
index 2419cf19..63d23abc 100644
--- a/crates/pipeline/src/tone.rs
+++ b/crates/pipeline/src/tone.rs
@@ -9,9 +9,10 @@
//! S-curve adjustments in a gamma-2.2 perceptual domain and a short shoulder above 0.95.
pub const GREY: f32 = 0.18;
-const LUT_MIN_EV: f32 = -14.0;
-const LUT_MAX_EV: f32 = 10.0;
-const LUT_N: usize = 4096;
+/// The tone LUT spans `LUT_MIN_EV..LUT_MAX_EV` around grey in `LUT_N` steps.
+pub const LUT_MIN_EV: f32 = -14.0;
+pub const LUT_MAX_EV: f32 = 10.0;
+pub const LUT_N: usize = 4096;
#[derive(Clone, Debug)]
pub struct ToneMap {
@@ -86,6 +87,11 @@ impl ToneMap {
ToneMap { lut }
}
+ /// The table (`LUT_N` entries, see [`ToneMap::apply`]).
+ pub fn lut(&self) -> &[f32] {
+ &self.lut
+ }
+
/// Scene luminance → display-linear luminance.
#[inline]
pub fn apply(&self, y: f32) -> f32 {
diff --git a/crates/raster/src/blur.rs b/crates/raster/src/blur.rs
index ac8eeda6..8deba2d1 100644
--- a/crates/raster/src/blur.rs
+++ b/crates/raster/src/blur.rs
@@ -4,7 +4,7 @@ use crate::resample::Pixel;
use crate::{Image, par_rows};
/// Box radii approximating a Gaussian of `sigma` with 3 passes (Kovesi / Wells).
-fn box_radii(sigma: f32) -> [usize; 3] {
+pub fn box_radii(sigma: f32) -> [usize; 3] {
let n = 3.0f32;
let w_ideal = (12.0 * sigma * sigma / n + 1.0).sqrt();
let mut wl = w_ideal.floor() as i32;
@@ -40,7 +40,7 @@ fn box_row(row: &mut [T], tmp: &mut Vec, r: usize) {
}
/// Rows per parallel band in the vertical pass (each band re-primes its running sums).
-const BAND: usize = 32;
+pub const BAND: usize = 32;
/// One vertical box pass `src → dst` (clamped edges): a running sum per column, vectorized across
/// the row, parallel over horizontal bands. Cache-friendly — no transposes.
diff --git a/crates/raster/src/resample.rs b/crates/raster/src/resample.rs
index b6c37b79..a5759583 100644
--- a/crates/raster/src/resample.rs
+++ b/crates/raster/src/resample.rs
@@ -54,8 +54,9 @@ impl Filter {
}
}
-/// Contributions for each output coordinate: (first source index, weights).
-fn weights(src: usize, dst: usize, filter: Filter) -> Vec<(usize, Vec)> {
+/// Contributions for each output coordinate: (first source index, weights) — the exact taps
+/// [`resize`] uses (GPU kernels upload these tables).
+pub fn weights(src: usize, dst: usize, filter: Filter) -> Vec<(usize, Vec)> {
let scale = src as f32 / dst as f32;
let fscale = scale.max(1.0);
let support = filter.support() * fscale;
diff --git a/crates/ui-egui/src/control.rs b/crates/ui-egui/src/control.rs
index 3b6885c4..8d6cbfae 100644
--- a/crates/ui-egui/src/control.rs
+++ b/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::>(),
"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,
})
}
diff --git a/crates/ui-egui/src/panels/right.rs b/crates/ui-egui/src/panels/right.rs
index e797ecfb..45e65291 100644
--- a/crates/ui-egui/src/panels/right.rs
+++ b/crates/ui-egui/src/panels/right.rs
@@ -117,40 +117,39 @@ fn crop(app: &mut LightcraftApp, ui: &mut egui::Ui, id: PhotoId) {
super::edit::apply_slider_out(app, &spec, out, |app, v| app.run("crop.straighten", json!({"angle": v})));
}
padded(ui, |ui| {
- ui.horizontal_wrapped(|ui| {
- ui.label("Overlay");
- for (label, o) in [
- ("Thirds", crate::state::CropOverlay::Thirds),
- ("Grid", crate::state::CropOverlay::Grid),
- ("Golden", crate::state::CropOverlay::Golden),
- ("None", crate::state::CropOverlay::None),
- ] {
- if text_button(ui, &format!("cropOverlay-{}", label.to_lowercase()), label, app.ui.crop_overlay == o).clicked() {
- app.ui.crop_overlay = o;
- }
+ ui.label("Overlay");
+ {
+ use crate::state::CropOverlay as O;
+ let opts = [(O::Thirds, "Thirds", "thirds"), (O::Grid, "Grid", "grid"), (O::Golden, "Golden", "golden"), (O::None, "None", "none")];
+ let items: Vec<(&str, &str)> = opts.iter().map(|(_, l, k)| (*l, *k)).collect();
+ let active = opts.iter().position(|(o, _, _)| *o == app.ui.crop_overlay);
+ if let Some(i) = crate::widgets::segmented(ui, "cropOverlay", &items, active, 4) {
+ app.ui.crop_overlay = opts[i].0;
}
- });
+ }
});
divider(ui);
header(ui, "Geometry");
padded(ui, |ui| {
ui.label("Upright");
- ui.horizontal_wrapped(|ui| {
+ {
use lightcraft_develop::Upright;
- for (label, mode, key) in [
+ let modes = [
("Off", Upright::Off, "off"),
("Auto", Upright::Auto, "auto"),
("Guided", Upright::Guided, "guided"),
("Level", Upright::Level, "level"),
("Vertical", Upright::Vertical, "vertical"),
("Full", Upright::Full, "full"),
- ] {
- if text_button(ui, &format!("upright-{key}"), label, d.geometry.upright == mode).clicked() {
- let _ = app.run("geometry.upright", json!({"mode": key}));
- app.ui.tool = if mode == Upright::Guided { "guidedUpright".into() } else { String::new() };
- }
+ ];
+ let items: Vec<(&str, &str)> = modes.iter().map(|(l, _, k)| (*l, *k)).collect();
+ let active = modes.iter().position(|(_, m, _)| *m == d.geometry.upright);
+ if let Some(i) = crate::widgets::segmented(ui, "upright", &items, active, 3) {
+ let (_, mode, key) = modes[i];
+ let _ = app.run("geometry.upright", json!({"mode": key}));
+ app.ui.tool = if mode == Upright::Guided { "guidedUpright".into() } else { String::new() };
}
- });
+ }
let guided = d.geometry.upright == lightcraft_develop::Upright::Guided;
ui.horizontal_wrapped(|ui| {
if !guided && d.geometry.upright != lightcraft_develop::Upright::Off && text_button(ui, "uprightUpdate", "Update", false).clicked() {
diff --git a/crates/ui-egui/src/widgets.rs b/crates/ui-egui/src/widgets.rs
index be4fa561..482b906f 100644
--- a/crates/ui-egui/src/widgets.rs
+++ b/crates/ui-egui/src/widgets.rs
@@ -267,6 +267,41 @@ pub fn text_button(ui: &mut Ui, id: &str, label: &str, active: bool) -> Response
resp
}
+/// A segmented control: `labels` share the available width equally, `per_row` per row. Each segment is
+/// addressable as `button:{id}-{key}` (keys parallel to labels). Returns the clicked index.
+pub fn segmented(ui: &mut Ui, id: &str, items: &[(&str, &str)], active: Option, per_row: usize) -> Option {
+ let t = Tokens::get(ui.ctx());
+ let per_row = per_row.max(1);
+ let gap = 4.0;
+ let w = ui.available_width();
+ let seg_w = ((w - gap * (per_row as f32 - 1.0)) / per_row as f32).max(24.0);
+ let mut clicked = None;
+ for (row_i, row) in items.chunks(per_row).enumerate() {
+ ui.horizontal(|ui| {
+ ui.spacing_mut().item_spacing.x = gap;
+ for (j, (label, key)) in row.iter().enumerate() {
+ let i = row_i * per_row + j;
+ let (r, resp) = ui.allocate_exact_size(vec2(seg_w, 24.0), Sense::click());
+ register(ui.ctx(), format!("button:{id}-{key}"), r);
+ let fill = if active == Some(i) {
+ t.pressed
+ } else if resp.hovered() {
+ t.hover
+ } else {
+ t.button
+ };
+ let p = ui.painter();
+ p.rect(r, CornerRadius::same(4), fill, Stroke::new(1.0, t.button_border), StrokeKind::Inside);
+ p.text(r.center(), Align2::CENTER_CENTER, *label, t.semibold(11.5), t.text);
+ if resp.clicked() {
+ clicked = Some(i);
+ }
+ }
+ });
+ }
+ clicked
+}
+
/// An icon-only button. `active` draws the selected background (tool strip).
pub fn icon_button(ui: &mut Ui, id: &str, icon: Icon, size: egui::Vec2, active: bool, enabled: bool, tooltip: &str) -> Response {
let t = Tokens::get(ui.ctx());
diff --git a/docs/gpu-pipeline.md b/docs/gpu-pipeline.md
index 94b99179..3dc9299d 100644
--- a/docs/gpu-pipeline.md
+++ b/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` 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.