Merge main (GPU pipeline) into the web branch

This commit is contained in:
Brandon Thomas 2026-10-01 09:02:23 -04:00
commit 8e5d6cb521
39 changed files with 4227 additions and 329 deletions

17
Cargo.lock generated
View File

@ -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"

View File

@ -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 }

View File

@ -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 🚧 |
<br>

View File

@ -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::<f64>() / a.len() as f64;
10.0 * (255.0f64.powi(2) / mse.max(1e-9)).log10()
}
fn best<T>(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))
);
}
}

View File

@ -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<u8>)> = 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 {

View File

@ -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<u8>,
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<u8> {
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<f32>,
}
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<u8>, 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<u8>, 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<u8>)],
) -> Vec<u8> {
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<Vec<f32>> = (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<usize> = (0..my.div_ceil(rows_per_band)).collect();
let encode_band = |band: &usize| -> Vec<u8> {
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<Vec<u8>> = bands.par_iter().map(encode_band).collect();
#[cfg(not(feature = "parallel"))]
let coded: Vec<Vec<u8>> = bands.iter().map(encode_band).collect();
// headers
let mut out = Vec::with_capacity(coded.iter().map(Vec::len).sum::<usize>() + 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<u8> {
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<u8> {
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::<f64>() / 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::<usize>(), 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<u8> = (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);
}
}

View File

@ -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;

View File

@ -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] {

View File

@ -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 }

View File

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

View File

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

View File

@ -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 {

View File

@ -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<Rgb32f>, info: &SourceInfo, s: &DevelopSettings, req: &RenderRequest, stages: Option<&StageCache>, gpu: bool) -> Rendered {
if gpu && let Some(r) = lightcraft_gpu::render(src, info, s, req, stages) {
return r;
}
match stages {
Some(st) => lightcraft_pipeline::render_cached(src, info, s, req, st),
None => lightcraft_pipeline::render(src, info, s, req),
}
}
/// What identifies a photo's pixels for caching: its content hash, else its source.
pub fn content_key(p: &Photo) -> String {
match (&p.content_hash, &p.source) {

27
crates/gpu/Cargo.toml Normal file
View File

@ -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 }

339
crates/gpu/src/ctx.rs Normal file
View File

@ -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<wgpu::Buffer>,
pub len: usize,
}
impl Buf {
pub fn raw(&self) -> &wgpu::Buffer {
self.buf.as_ref().expect("live buffer")
}
}
impl Drop for Buf {
fn drop(&mut self) {
if let Some(b) = self.buf.take() {
RETIRED.with(|r| r.borrow_mut().push(b));
}
}
}
thread_local! {
/// Buffers dropped on this thread since its last submit.
static RETIRED: std::cell::RefCell<Vec<wgpu::Buffer>> = const { std::cell::RefCell::new(Vec::new()) };
}
/// Most bytes kept in the free pool.
const POOL_BYTES: u64 = 2 << 30;
pub struct Gpu {
pub device: wgpu::Device,
pub queue: wgpu::Queue,
pub info: wgpu::AdapterInfo,
kernels: HashMap<&'static str, Kernel>,
dummy: wgpu::Buffer,
/// Recycled buffers (see [`Buf`]).
free: std::sync::Mutex<Vec<wgpu::Buffer>>,
/// Largest storage buffer the device accepts (bytes).
pub max_buffer: u64,
}
/// 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<String> = m.iter().map(|r| format!("vec3<f32>({:?}, {:?}, {:?})", r[0], r[1], r[2])).collect();
s += &format!("const {name} = array<vec3<f32>, 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<storage, read> P: array<u32>;\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<storage, {access}> {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<Gpu> {
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<Item = wgpu::CommandBuffer>) {
self.queue.submit(cmds);
let retired = RETIRED.with(|r| std::mem::take(&mut *r.borrow_mut()));
let mut f = self.free.lock().unwrap_or_else(|e| e.into_inner());
f.extend(retired);
let mut total: u64 = f.iter().map(|b| b.size()).sum();
while total > POOL_BYTES && !f.is_empty() {
total -= f.remove(0).size();
}
}
/// Upload 32-bit values.
pub fn upload<T: bytemuck::Pod>(&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<T: bytemuck::Pod>(&self, mut enc: wgpu::CommandEncoder, src: &Buf, len: usize) -> Vec<T> {
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<T> = 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]
}

119
crates/gpu/src/lib.rs Normal file
View File

@ -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<bool> = 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<Option<ctx::Gpu>> = 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<String> {
#[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<Rgb32f>, info: &SourceInfo, s: &DevelopSettings, req: &RenderRequest, stages: Option<&StageCache>) -> Option<Rendered> {
#[cfg(not(target_arch = "wasm32"))]
{
if !enabled() {
return None;
}
let gpu = device()?;
let ext = stages.map(|c| c.extension::<GpuStages>());
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
}
}

211
crates/gpu/src/params.rs Normal file
View File

@ -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<u32>);
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<u32>, Vec<f32>) {
let mut aux: Vec<f32> = 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<f32> = 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)
}

773
crates/gpu/src/render.rs Normal file
View File

@ -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<Vec<Entry>>,
/// The uploaded source (a view renders one photo at several sizes).
source: Mutex<Option<(Arc<Rgb32f>, Arc<Buf>)>>,
}
const CAPACITY: usize = 2;
#[derive(Clone)]
struct Entry {
src: Arc<Rgb32f>,
geo: u64,
sampled: Arc<Buf>,
lin: Option<(u64, Arc<Buf>)>,
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<Arc<Buf>>,
base: Option<(u32, Arc<Buf>)>,
clarity: Option<(u32, Arc<Buf>)>,
texture: Option<(u32, Arc<Buf>)>,
dark: Option<(u32, Arc<Buf>, f32)>,
}
impl GpuStages {
fn get(&self, src: &Arc<Rgb32f>, geo: u64) -> Option<Entry> {
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<Rgb32f>, upload: impl FnOnce() -> Buf) -> Arc<Buf> {
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<Entry>> {
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<wgpu::CommandEncoder>,
}
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<T: bytemuck::Pod>(&mut self, b: &Buf, len: usize) -> Vec<T> {
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<u32> {
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<fn(i32, i32) -> [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<u32> {
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<Rgb32f>, src_buf: Arc<Buf>, plan: &Plan<'_>) -> Arc<Buf> {
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<Rgb32f>,
lin: Option<Arc<Rgb32f>>,
log_l: Option<Arc<Plane>>,
}
fn profiling() -> bool {
static ON: std::sync::OnceLock<bool> = 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<Rgb32f>,
info: &SourceInfo,
s: &DevelopSettings,
req: &RenderRequest,
stages: Option<&GpuStages>,
) -> Option<Rendered> {
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<std::time::Instant>, 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<Buf>,
base: Arc<Buf>,
clarity: Option<Arc<Buf>>,
texture: Option<Arc<Buf>>,
dark: Option<Arc<Buf>>,
air: f32,
}
fn plane_at(slot: &mut Option<(u32, Arc<Buf>)>, sigma: f32, f: impl FnOnce() -> Buf) -> Arc<Buf> {
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<Buf>, 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::<Vec<u32>>();
let mut aux: Vec<f32> = 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::<usize>() <= 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}");
}
}
}
}
}

View File

@ -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<f32> {
let j = i * nc;
if (nc == 1u) {
return vec3<f32>(src[j], 0.0, 0.0);
}
if (nc == 2u) {
return vec3<f32>(src[j], src[j + 1u], 0.0);
}
return vec3<f32>(src[j], src[j + 1u], src[j + 2u]);
}
fn st(i: u32, nc: u32, v: vec3<f32>) {
let j = i * nc;
dst[j] = v.x;
if (nc > 1u) {
dst[j + 1u] = v.y;
}
if (nc > 2u) {
dst[j + 2u] = v.z;
}
}
// Horizontal: one thread per CH consecutive pixels of a row.
@compute @workgroup_size(64, 4)
fn box_h(@builtin(global_invocation_id) g: vec3<u32>) {
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<f32>(0.0);
for (var k = -r; k <= r; k++) {
acc += ld(row + u32(clamp(i32(x0) + k, 0, last)), nc);
}
for (var x = x0; x < x1; x++) {
st(row + x, nc, acc * inv);
acc = acc + ld(row + u32(min(i32(x) + r + 1, last)), nc) - ld(row + u32(max(i32(x) - r, 0)), nc);
}
}
// 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<u32>) {
let w = pu(0u);
let h = pu(1u);
let nc = pu(2u);
let r = i32(pu(3u));
let ch = pu(4u);
let x = g.x;
let 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<f32>(0.0);
for (var k = -r; k <= r; k++) {
acc += ld(u32(clamp(i32(y0) + k, 0, last)) * w + x, nc);
}
for (var y = y0; y < y1; y++) {
st(y * w + x, nc, acc * inv);
acc = acc + ld(u32(min(i32(y) + r + 1, last)) * w + x, nc) - ld(u32(max(i32(y) - r, 0)) * w + x, nc);
}
}

View File

@ -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<f32>(P[i]);
}
// Linear index of a 1-D kernel dispatched on a 2-D grid of 256-wide workgroups.
fn lin_index(g: vec3<u32>, n: vec3<u32>) -> 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>) -> 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>) -> f32 {
return log2(max(lum2020(c), 1e-7) / GREY);
}
fn mul3(m: array<vec3<f32>, 3>, v: vec3<f32>) -> vec3<f32> {
return vec3<f32>(
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<f32>) -> vec3<f32> {
let lms = mul3(OK_TO_LMS, c);
return mul3(OK_TO_LAB, vec3<f32>(cbrt(lms.x), cbrt(lms.y), cbrt(lms.z)));
}
// OkLab → linear Rec.2020.
fn oklab_inv(lab: vec3<f32>) -> vec3<f32> {
let l = mul3(OK_FROM_LAB, lab);
return mul3(OK_FROM_LMS, l * l * l);
}

View File

@ -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<f32, 8> {
var w: array<f32, 8>;
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<f32>, local_sat: f32, local_hue: f32) -> vec3<f32> {
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<f32>(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<f32, 4>(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<u32>(i) * GRAIN_H0) ^ (bitcast<u32>(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<u32>) {
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<f32>(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<f32, 14>;
var tint_on = false;
var tint_dir = vec2<f32>(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<f32>(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<f32>(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<f32>(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<f32>(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<f32>(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<f32>(encode_srgb(r.x), encode_srgb(r.y), encode_srgb(r.z));
if (pu(F_CURVES) != 0u) {
e = vec3<f32>(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);
}

View File

@ -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<f32> {
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<f32>(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<f32> {
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<f32>) {
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<u32>) {
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<i32, 6>(bitcast<i32>(pu(4u)), bitcast<i32>(pu(5u)), bitcast<i32>(pu(6u)), bitcast<i32>(pu(7u)), bitcast<i32>(pu(8u)), bitcast<i32>(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<f32>(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<u32>) {
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<f32> {
return vec2<f32>(pf(12u) / 2.0, pf(13u) / 2.0);
}
fn to_corrected(p: vec2<f32>) -> vec2<f32> {
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<f32>(x, y) * l;
}
fn embedded_warp(p: vec2<f32>, ch: u32) -> vec2<f32> {
let w = pf(12u);
let h = pf(13u);
let cc = vec2<f32>(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<f32>(sx, sy);
}
fn corrected_to_source(p: vec2<f32>, ch: u32) -> vec2<f32> {
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>) -> 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<f32>(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<u32>) {
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<f32>(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<f32>(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);
}

View File

@ -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<f32> {
return vec3<f32>(a[3u * i], a[3u * i + 1u], a[3u * i + 2u]);
}
fn put_rgb(i: u32, v: vec3<f32>) {
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<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
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<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
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<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
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<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
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<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
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<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
let i = lin_index(g, nw);
if (i >= pu(0u)) {
return;
}
var c = rgb_a(i);
if (pu(1u) != 0u) {
c = vec3<f32>(
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<f32>(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<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
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<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
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<f32>(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<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
let i = lin_index(g, nw);
if (i >= pu(0u)) {
return;
}
let yl = lum2020(rgb_a(i));
let t = pf(1u);
let c0 = vec3<f32>(b[3u * i], b[3u * i + 1u], b[3u * i + 2u]);
let cb = vec3<f32>(c[3u * i], c[3u * i + 1u], c[3u * i + 2u]);
put_rgb(i, max((c0 + (cb - c0) * t) * yl, vec3<f32>(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<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
let i = lin_index(g, nw);
if (i >= pu(0u)) {
return;
}
dst[i] = a[i * pu(1u)];
}

View File

@ -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<f32> {
let px = f32(x) + 0.5;
let py = f32(y) + 0.5;
return vec2<f32>(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<u32>) {
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<f32>(pf(10u), pf(11u));
let d = vec2<f32>(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<f32>(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<f32>(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<f32>(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<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
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<u32>, @builtin(num_workgroups) nw: vec3<u32>) {
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;
}

View File

@ -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<u32>) {
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<f32>(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<u32>) {
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<f32>(table[off + k]);
}
dst[(y * w + x) * nc + c] = acc;
}
}

View File

@ -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<Rgb32f> {
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<Rgb32f>, 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}");
}
}

View File

@ -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 {

View File

@ -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<Vig> {
@ -81,14 +83,17 @@ fn vignette(s: &DevelopSettings) -> Option<Vig> {
})
}
/// 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<Vig>,
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<Box<[f32; SRGB_LUT_N + 1]>> = std::sync::OnceLock::new();
LUT.get_or_init(|| {
let mut t = Box::new([0.0f32; SRGB_LUT_N + 1]);

View File

@ -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 {

View File

@ -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<masks::Evaluated>,
/// 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<Vec<CacheEntry>>,
capacity: usize,
/// Another renderer's per-view state (the GPU renderer keeps its device-resident stages here).
ext: Mutex<Option<Arc<dyn Any + Send + Sync>>>,
}
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<T: Any + Send + Sync + Default>(&self) -> Arc<T> {
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::<T>()
{
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::<DefaultHasher>::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);

View File

@ -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<NrLum>, Option<NrColor>) {
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<Plane>)>, 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<Rgb32f>, 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<f32>,
/// Clarity band (fast guided filter, [`CLARITY_EPS`]).
pub clarity: Option<f32>,
/// Texture / sharpening band (Gaussian).
pub texture: Option<f32>,
/// Dehaze dark channel (Gaussian).
pub dark: Option<f32>,
}
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<Rgb32f>, 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<Rgb32f>, 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<Rgb32f>, 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<f32> = 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>) -> f32 {
if v.is_empty() {
return 1.0;
}

View File

@ -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<Point>,
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<Point> = 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<Point> = 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) };
}
}
}

View File

@ -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 {

View File

@ -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<T: Pixel>(row: &mut [T], tmp: &mut Vec<T>, 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.

View File

@ -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<f32>)> {
/// 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<f32>)> {
let scale = src as f32 / dst as f32;
let fscale = scale.max(1.0);
let support = filter.support() * fscale;

View File

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

View File

@ -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() {

View File

@ -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<usize>, per_row: usize) -> Option<usize> {
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());

View File

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