diff --git a/core/dr-gpu/src/adjust.rs b/core/dr-gpu/src/adjust.rs index 39c3c22..efb3329 100644 --- a/core/dr-gpu/src/adjust.rs +++ b/core/dr-gpu/src/adjust.rs @@ -19,6 +19,7 @@ use std::collections::HashMap; use dr_pipeline::ComposedShader; use wgpu::util::DeviceExt; +use crate::readback::await_mapping; use crate::{DemosaicedImage, GpuContext, GpuError}; /// Leading floats the composer reserves before any operation's own uniforms: @@ -31,18 +32,6 @@ use crate::{DemosaicedImage, GpuContext, GpuError}; /// reads them. const RESERVED_FIELDS: usize = dr_pipeline::RESERVED_UNIFORM_FIELDS; -/// How many non-blocking polls a readback gets before it is called failed. -/// -/// A bound rather than a spin forever: if the device is lost the map callback -/// never arrives, and an unbounded loop would hang the interface rather than -/// surfacing the error. Set far above any plausible completion — the copy this -/// waits on is milliseconds — so it is reached only when something is wrong. -/// -/// Ungated along with `export_pixels`: an export reads pixels back in a -/// shipping build, and the bound that stops a lost device hanging the app -/// applies at least as much there as it does to the display bridge. -const READBACK_POLL_LIMIT: u32 = 100_000; - /// Runs composed operation chains against demosaiced images. pub struct AdjustPass { ctx: GpuContext, @@ -399,38 +388,10 @@ impl AdjustPass { let _ = tx.send(r); }); - // **Polled without blocking, then checked.** - // - // `Maintain::Wait` parks the calling thread until the GPU has finished, - // and this is called from the UI thread — so that park was a frozen - // interface for the duration of the copy (~7 ms at 4K, per the note - // above). `Poll` drives the same callbacks without sleeping, so the - // loop below stays interruptible and the mapping still completes. - // - // The bounded spin matters: a lost device would otherwise never - // deliver the callback and this would hang the app instead of - // reporting an error. - let mut mapped = None; - for _ in 0..READBACK_POLL_LIMIT { - // A poll error is a lost device, which is exactly the case the - // bounded spin exists to escape — returning here reports it - // immediately rather than spinning out the full limit first. - self.ctx - .device - .poll(wgpu::PollType::Poll) - .map_err(|e| GpuError::Readback(e.to_string()))?; - match rx.try_recv() { - Ok(r) => { - mapped = Some(r); - break; - } - Err(std::sync::mpsc::TryRecvError::Empty) => continue, - Err(e) => return Err(GpuError::Readback(e.to_string())), - } - } - mapped - .ok_or_else(|| GpuError::Readback("readback did not complete".into()))? - .map_err(|e| GpuError::Readback(e.to_string()))?; + // Polled rather than parked, and bounded rather than spun forever — + // see `readback::await_mapping`, which the histogram's own transfer + // shares for exactly the same reasons. + await_mapping(&self.ctx, &rx)?; let data = slice.get_mapped_range(); let mut out = Vec::with_capacity((unpadded * h) as usize); diff --git a/core/dr-gpu/src/histogram.rs b/core/dr-gpu/src/histogram.rs new file mode 100644 index 0000000..716248d --- /dev/null +++ b/core/dr-gpu/src/histogram.rs @@ -0,0 +1,591 @@ +//! TRACES: FR-DSP-7 +//! Counting the display frame, on the device that drew it. +//! +//! # Why a compute reduction and not a CPU pass over the readback +//! +//! There is, today, a whole frame already sitting in CPU memory every time the +//! canvas updates — `AdjustPass::read_output`, the temporary bridge that spike +//! S1 removes. Walking it to build a histogram would have been perhaps thirty +//! lines and no shader at all, and it was the obvious thing to reach for. +//! +//! It was rejected for two reasons, in this order. +//! +//! FR-DSP-7 states the mechanism, not just the feature: "these derive from a +//! GPU-side reduction into a small buffer. Per-frame CPU readback of image data +//! is prohibited." A histogram built on the bridge would be correct today and +//! *deleted* by S1 — it would be new code whose only foundation is the one +//! thing the architecture is committed to removing, and the histogram would +//! then be the reason the bridge could not go. +//! +//! And the cost does not scale the way the shortcut implies. Counting 2 MP on +//! one CPU thread is several milliseconds of the settle frame; the reduction +//! below is a fraction of one, and what crosses the bus is 4104 bytes +//! regardless of the image. The simpler-looking option is simpler only while +//! the frame happens to be lying there. +//! +//! # What it costs and when it runs +//! +//! One dispatch plus a 4 KB buffer copy and a mapping — a device sync point. +//! FR-DSP-7 requires that this not extend the FR-DSP-3 frame budget, so the +//! interface runs it on the *settled* frame only, never on the draft frames a +//! drag produces. The histogram of an image being dragged past is not read +//! anyway; the one that arrives when the slider stops is. + +use wgpu::util::DeviceExt; + +use crate::readback::await_mapping; +use crate::{GpuContext, GpuError}; + +/// Levels per channel. 256, so a bin *is* an output code value and no +/// re-bucketing stands between the count and what the display shows. +pub const BINS: usize = 256; + +/// Four channel histograms plus the two clip counters, as the shader lays them +/// out. Kept next to the shader's own constants because the two must agree. +const CHANNELS: usize = 4; +const CLIPPED_HIGH: usize = CHANNELS * BINS; +const CLIPPED_LOW: usize = CHANNELS * BINS + 1; +const SLOTS: usize = CHANNELS * BINS + 2; + +/// TRACES: FR-DSP-7 +/// A counted frame: how many pixels sit at each output level. +/// +/// Counts, not proportions. Turning these into something drawable — folding +/// 256 bins into the columns a 280px panel can show, choosing a peak to scale +/// against — is presentation, and belongs to whoever is drawing (ARCH §4.3a). +/// What this crate owes is the numbers. +#[derive(Clone, Debug, PartialEq, Eq)] +pub struct Histogram { + red: [u32; BINS], + green: [u32; BINS], + blue: [u32; BINS], + luma: [u32; BINS], + clipped_highlights: u32, + clipped_shadows: u32, + pixels: u32, +} + +impl Histogram { + /// Counts per output level, darkest first. + pub fn red(&self) -> &[u32; BINS] { + &self.red + } + + pub fn green(&self) -> &[u32; BINS] { + &self.green + } + + pub fn blue(&self) -> &[u32; BINS] { + &self.blue + } + + /// Rec.709 luma of the encoded values — the axis a photographer reads + /// exposure off. See the shader for why it is weighted in fixed point. + pub fn luma(&self) -> &[u32; BINS] { + &self.luma + } + + /// Pixels with **any** channel at 255, and with any channel at 0. + /// + /// Any rather than all, because a single blown channel is detail that is + /// already gone: a red that has hit the ceiling has no gradation left in it + /// however much green and blue still hold. + pub fn clipped_highlights(&self) -> u32 { + self.clipped_highlights + } + + pub fn clipped_shadows(&self) -> u32 { + self.clipped_shadows + } + + /// Pixels counted. The denominator for the two figures above. + pub fn pixels(&self) -> u32 { + self.pixels + } + + /// Rebuild from the flat slot array the shader writes. + /// + /// `pixels` is summed from the red channel rather than taken from the image + /// dimensions: every pixel lands in exactly one red bin, so the sum *is* + /// the count, and deriving it that way makes a dropped or double-counted + /// texel show up as a wrong denominator instead of hiding. + fn from_slots(slots: &[u32]) -> Result { + if slots.len() < SLOTS { + return Err(GpuError::Readback(format!( + "histogram readback was {} slots, expected {SLOTS}", + slots.len() + ))); + } + let channel = |i: usize| -> [u32; BINS] { + let mut out = [0u32; BINS]; + out.copy_from_slice(&slots[i * BINS..(i + 1) * BINS]); + out + }; + let red = channel(0); + Ok(Self { + pixels: red.iter().sum(), + red, + green: channel(1), + blue: channel(2), + luma: channel(3), + clipped_highlights: slots[CLIPPED_HIGH], + clipped_shadows: slots[CLIPPED_LOW], + }) + } +} + +/// The dispatch's view of the frame. Padded to 16 bytes for std140. +#[repr(C)] +#[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] +struct Dims { + width: u32, + height: u32, + pad_0: u32, + pad_1: u32, +} + +/// TRACES: FR-DSP-7 +/// Counts a rendered frame into [`Histogram`]. +/// +/// Holds its buffers for the life of the session. They are a fixed 4104 bytes +/// whatever the image size — the one property that makes this affordable — so +/// there is nothing to reallocate when the viewport changes, unlike the display +/// target beside it. +pub struct HistogramPass { + ctx: GpuContext, + pipeline: wgpu::ComputePipeline, + bind_group_layout: wgpu::BindGroupLayout, + /// Where the shader accumulates. Cleared before each dispatch. + bins: wgpu::Buffer, + /// Mappable destination; a storage buffer cannot also be `MAP_READ`. + staging: wgpu::Buffer, + dims: wgpu::Buffer, +} + +impl HistogramPass { + pub fn new(ctx: &GpuContext) -> Result { + // A validation error here is a bug in the shader beside this file, not + // anything a user did — surfaced as a `Result` rather than left to + // wgpu's default handler, which panics. + let scope = ctx.device.push_error_scope(wgpu::ErrorFilter::Validation); + + let module = ctx + .device + .create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("histogram"), + source: wgpu::ShaderSource::Wgsl(include_str!("shaders/histogram.wgsl").into()), + }); + + let bind_group_layout = + ctx.device + .create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("histogram-bgl"), + entries: &[ + // The rendered frame, sampled with `textureLoad` — the + // same texture the compositor shows, so what is counted + // is what is on screen. + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: true }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 2, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + ], + }); + + let layout = ctx + .device + .create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("histogram-layout"), + bind_group_layouts: &[Some(&bind_group_layout)], + immediate_size: 0, + }); + + let pipeline = ctx + .device + .create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("histogram-pipeline"), + layout: Some(&layout), + module: &module, + entry_point: Some("main"), + compilation_options: Default::default(), + cache: None, + }); + + if let Some(err) = pollster::block_on(scope.pop()) { + return Err(GpuError::ShaderCompilation(err.to_string())); + } + + let bytes = (SLOTS * std::mem::size_of::()) as u64; + let bins = ctx.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("histogram-bins"), + size: bytes, + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::COPY_SRC + | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + let staging = ctx.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("histogram-staging"), + size: bytes, + usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ, + mapped_at_creation: false, + }); + let dims = ctx + .device + .create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("histogram-dims"), + contents: bytemuck::bytes_of(&Dims { + width: 0, + height: 0, + pad_0: 0, + pad_1: 0, + }), + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + }); + + Ok(Self { + ctx: ctx.clone(), + pipeline, + bind_group_layout, + bins, + staging, + dims, + }) + } + + /// TRACES: FR-DSP-7 + /// Count one rendered frame. + /// + /// `texture` must carry `TEXTURE_BINDING`, which `AdjustPass`'s output + /// does because the compositor samples it. + pub fn compute(&self, texture: &wgpu::Texture) -> Result { + let (width, height) = (texture.width(), texture.height()); + self.ctx.queue.write_buffer( + &self.dims, + 0, + bytemuck::bytes_of(&Dims { + width, + height, + pad_0: 0, + pad_1: 0, + }), + ); + + let view = texture.create_view(&Default::default()); + let bind_group = self + .ctx + .device + .create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("histogram-bg"), + layout: &self.bind_group_layout, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::TextureView(&view), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: self.bins.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: self.dims.as_entire_binding(), + }, + ], + }); + + let mut enc = self + .ctx + .device + .create_command_encoder(&wgpu::CommandEncoderDescriptor { + label: Some("histogram-encoder"), + }); + // The accumulator is reused between frames, so it carries the previous + // frame's counts until this line. Forgetting it does not fail — it + // quietly integrates every frame since the image opened, which looks + // like a histogram that will not respond to the exposure slider. + enc.clear_buffer(&self.bins, 0, None); + { + let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("histogram-pass"), + timestamp_writes: None, + }); + pass.set_pipeline(&self.pipeline); + pass.set_bind_group(0, &bind_group, &[]); + pass.dispatch_workgroups(width.div_ceil(16), height.div_ceil(16), 1); + } + enc.copy_buffer_to_buffer(&self.bins, 0, &self.staging, 0, self.staging.size()); + self.ctx.queue.submit(Some(enc.finish())); + + let slice = self.staging.slice(..); + let (tx, rx) = std::sync::mpsc::channel(); + slice.map_async(wgpu::MapMode::Read, move |r| { + let _ = tx.send(r); + }); + await_mapping(&self.ctx, &rx)?; + + let data = slice.get_mapped_range(); + let slots: Vec = bytemuck::cast_slice::(&data).to_vec(); + drop(data); + self.staging.unmap(); + + Histogram::from_slots(&slots) + } +} + +#[cfg(test)] +mod tests { + use super::*; + + fn ctx() -> Option { + match pollster::block_on(GpuContext::new_headless()) { + Ok(c) => Some(c), + Err(e) => { + eprintln!("skipping: no GPU adapter ({e})"); + None + } + } + } + + /// Upload `rgba` as a texture the pass can read, the way `AdjustPass` + /// hands its output over. + fn texture(ctx: &GpuContext, rgba: &[u8], width: u32, height: u32) -> wgpu::Texture { + let tex = ctx.device.create_texture(&wgpu::TextureDescriptor { + label: Some("histogram-test-source"), + size: wgpu::Extent3d { + width, + height, + depth_or_array_layers: 1, + }, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu::TextureDimension::D2, + format: wgpu::TextureFormat::Rgba8Unorm, + usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::COPY_DST, + view_formats: &[], + }); + ctx.queue.write_texture( + wgpu::TexelCopyTextureInfo { + texture: &tex, + mip_level: 0, + origin: wgpu::Origin3d::ZERO, + aspect: wgpu::TextureAspect::All, + }, + rgba, + wgpu::TexelCopyBufferLayout { + offset: 0, + bytes_per_row: Some(width * 4), + rows_per_image: Some(height), + }, + wgpu::Extent3d { + width, + height, + depth_or_array_layers: 1, + }, + ); + ctx.queue.submit(std::iter::empty()); + tex + } + + /// The reference the shader is checked against: the same bucketing, written + /// the obvious way on the CPU. + /// + /// Deliberately a *second* implementation rather than shared code. The bugs + /// this is here to catch — a workgroup tile that is never merged, an edge + /// tile counted twice, a luma weighting that carries past 255 — are all + /// bugs a shared implementation would commit identically on both sides and + /// so could not detect. + fn expected(rgba: &[u8]) -> Histogram { + let mut slots = vec![0u32; SLOTS]; + for px in rgba.chunks_exact(4) { + let (r, g, b) = (px[0] as usize, px[1] as usize, px[2] as usize); + let y = (54 * r + 183 * g + 19 * b) >> 8; + slots[r] += 1; + slots[BINS + g] += 1; + slots[2 * BINS + b] += 1; + slots[3 * BINS + y] += 1; + if px[0] == 255 || px[1] == 255 || px[2] == 255 { + slots[CLIPPED_HIGH] += 1; + } + if px[0] == 0 || px[1] == 0 || px[2] == 0 { + slots[CLIPPED_LOW] += 1; + } + } + Histogram::from_slots(&slots).expect("slot count") + } + + #[test] + fn a_flat_frame_puts_every_pixel_in_one_bin() { + // The arithmetic at its most checkable: 64x64 pixels of one value must + // produce exactly 4096 in exactly one bin and nothing anywhere else. + // A tile that failed to merge, or merged twice, changes this number — + // and a histogram that is merely "roughly right" is a histogram nobody + // can set a black point from. + let Some(ctx) = ctx() else { return }; + let pass = HistogramPass::new(&ctx).expect("pass"); + + let rgba: Vec = std::iter::repeat_n([90u8, 140, 200, 255], 64 * 64) + .flatten() + .collect(); + let tex = texture(&ctx, &rgba, 64, 64); + let hist = pass.compute(&tex).expect("compute"); + + assert_eq!(hist.pixels(), 4096); + assert_eq!(hist.red()[90], 4096); + assert_eq!(hist.green()[140], 4096); + assert_eq!(hist.blue()[200], 4096); + assert_eq!( + hist.red().iter().filter(|c| **c > 0).count(), + 1, + "one value can only occupy one bin" + ); + // (54*90 + 183*140 + 19*200) >> 8 = 34280 >> 8 = 133. + assert_eq!(hist.luma()[133], 4096); + assert_eq!(hist.clipped_highlights(), 0); + assert_eq!(hist.clipped_shadows(), 0); + } + + #[test] + fn every_level_is_reachable_and_lands_where_it_belongs() { + // A ramp covering all 256 codes, four pixels each. This is the test + // that would catch an off-by-one in the quantisation — a `floor` where + // a rounding was needed shifts the whole ramp down one bin and leaves + // 255 empty, which on a real photograph looks like nothing at all. + let Some(ctx) = ctx() else { return }; + let pass = HistogramPass::new(&ctx).expect("pass"); + + let (w, h) = (256u32, 4u32); + let mut rgba = Vec::with_capacity((w * h * 4) as usize); + for _ in 0..h { + for x in 0..w { + let v = x as u8; + rgba.extend_from_slice(&[v, v, v, 255]); + } + } + let tex = texture(&ctx, &rgba, w, h); + let hist = pass.compute(&tex).expect("compute"); + + assert_eq!(hist.pixels(), w * h); + for level in 0..BINS { + assert_eq!( + hist.red()[level], + h, + "level {level} should hold exactly {h} pixels" + ); + // Neutral, so luma must land on the same bin as the channels do. + assert_eq!(hist.luma()[level], h, "luma drifted at level {level}"); + } + } + + #[test] + fn the_shader_agrees_with_a_cpu_count_of_the_same_frame() { + // The cross-check, on a frame with no structure for a wrong dispatch to + // hide behind: a size that is not a multiple of the 16x16 workgroup, so + // the edge tiles run off the image, and pseudo-random content so every + // bin is occupied unevenly. Exact equality — the reduction is integer + // throughout precisely so this can be an `assert_eq`, not a tolerance. + let Some(ctx) = ctx() else { return }; + let pass = HistogramPass::new(&ctx).expect("pass"); + + let (w, h) = (101u32, 37u32); + let mut rgba = Vec::with_capacity((w * h * 4) as usize); + let mut state = 0x2545_F491_4F6C_DD1Du64; + for _ in 0..w * h { + for _ in 0..3 { + // xorshift64*, so the frame is identical on every machine and a + // failure can be reproduced rather than merely observed. + state ^= state >> 12; + state ^= state << 25; + state ^= state >> 27; + rgba.push((state.wrapping_mul(0x2545_F491_4F6C_DD1D) >> 56) as u8); + } + rgba.push(255); + } + + let tex = texture(&ctx, &rgba, w, h); + let hist = pass.compute(&tex).expect("compute"); + + assert_eq!(hist.pixels(), w * h, "an edge tile was dropped or doubled"); + assert_eq!(hist, expected(&rgba)); + } + + #[test] + fn clipping_is_counted_per_pixel_and_not_per_channel() { + // The distinction the indicator rests on. A pixel with two channels at + // the ceiling is *one* clipped pixel; counting channels would report + // 200% of a frame clipped, and a percentage that can exceed 100 is a + // readout nobody will trust again. + let Some(ctx) = ctx() else { return }; + let pass = HistogramPass::new(&ctx).expect("pass"); + + let rgba: Vec = [ + // Two channels blown, one pixel clipped. + [255u8, 255, 10, 255], + // One channel blown — still clipped, which is the point of "any". + [255, 10, 10, 255], + // Clean. + [10, 10, 10, 255], + // Black in one channel only: a clipped shadow. + [0, 10, 10, 255], + ] + .concat(); + let tex = texture(&ctx, &rgba, 4, 1); + let hist = pass.compute(&tex).expect("compute"); + + assert_eq!(hist.pixels(), 4); + assert_eq!(hist.clipped_highlights(), 2); + assert_eq!(hist.clipped_shadows(), 1); + } + + #[test] + fn a_second_frame_replaces_the_first_rather_than_adding_to_it() { + // The accumulator is reused, so a missing clear integrates every frame + // since the session opened. The symptom is subtle and awful: the + // histogram keeps its shape and simply stops responding to the sliders, + // because each frame's contribution shrinks against the running total. + let Some(ctx) = ctx() else { return }; + let pass = HistogramPass::new(&ctx).expect("pass"); + + let dark: Vec = std::iter::repeat_n([40u8, 40, 40, 255], 16 * 16) + .flatten() + .collect(); + let bright: Vec = std::iter::repeat_n([210u8, 210, 210, 255], 16 * 16) + .flatten() + .collect(); + + let first = pass.compute(&texture(&ctx, &dark, 16, 16)).expect("first"); + assert_eq!(first.red()[40], 256); + + let second = pass + .compute(&texture(&ctx, &bright, 16, 16)) + .expect("second"); + assert_eq!(second.pixels(), 256, "the previous frame was still counted"); + assert_eq!(second.red()[40], 0); + assert_eq!(second.red()[210], 256); + } +} diff --git a/core/dr-gpu/src/lib.rs b/core/dr-gpu/src/lib.rs index fcc204a..be66e39 100644 --- a/core/dr-gpu/src/lib.rs +++ b/core/dr-gpu/src/lib.rs @@ -15,10 +15,15 @@ mod adjust; mod demosaic; mod error; pub mod hierarchy; +mod histogram; +mod readback; mod segment; pub use adjust::AdjustPass; pub use demosaic::{DemosaicedImage, Demosaicer}; pub use error::GpuError; +// Renamed on the way out: `BINS` says enough inside `histogram`, and nothing +// at all at a crate root shared with demosaic and segmentation. +pub use histogram::{Histogram, HistogramPass, BINS as HISTOGRAM_BINS}; pub use segment::{SegmentOptions, SegmentPass, Segmentation}; /// Owns the wgpu device and queue. diff --git a/core/dr-gpu/src/readback.rs b/core/dr-gpu/src/readback.rs new file mode 100644 index 0000000..af40d17 --- /dev/null +++ b/core/dr-gpu/src/readback.rs @@ -0,0 +1,56 @@ +//! Waiting for a buffer mapping without parking the interface. +//! +//! Shared by every transfer off the device — the display bridge, an export, and +//! the histogram's 4 KB of bin counts. It was written once inside `AdjustPass` +//! and the reasoning below is the whole of why it is shaped this way; copying +//! thirty lines of that reasoning into a second caller would have left two +//! copies to keep true of each other. + +use crate::{GpuContext, GpuError}; + +/// How many non-blocking polls a readback gets before it is called failed. +/// +/// A bound rather than a spin forever: if the device is lost the map callback +/// never arrives, and an unbounded loop would hang the interface rather than +/// surfacing the error. Set far above any plausible completion — the copies +/// this waits on are milliseconds — so it is reached only when something is +/// wrong. +const POLL_LIMIT: u32 = 100_000; + +/// Drive the device until a `map_async` callback lands. +/// +/// **Polled without blocking, then checked.** +/// +/// `PollType::Wait` parks the calling thread until the GPU has finished, and +/// these transfers are called from the UI thread — so that park was a frozen +/// interface for the duration of the copy (~7 ms at 4K for a whole frame). +/// `Poll` drives the same callbacks without sleeping, so the loop below stays +/// interruptible and the mapping still completes. +/// +/// The bounded spin matters: a lost device would otherwise never deliver the +/// callback and this would hang the app instead of reporting an error. +pub(crate) fn await_mapping( + ctx: &GpuContext, + rx: &std::sync::mpsc::Receiver>, +) -> Result<(), GpuError> { + let mut mapped = None; + for _ in 0..POLL_LIMIT { + // A poll error is a lost device, which is exactly the case the bounded + // spin exists to escape — returning here reports it immediately rather + // than spinning out the full limit first. + ctx.device + .poll(wgpu::PollType::Poll) + .map_err(|e| GpuError::Readback(e.to_string()))?; + match rx.try_recv() { + Ok(r) => { + mapped = Some(r); + break; + } + Err(std::sync::mpsc::TryRecvError::Empty) => continue, + Err(e) => return Err(GpuError::Readback(e.to_string())), + } + } + mapped + .ok_or_else(|| GpuError::Readback("readback did not complete".into()))? + .map_err(|e| GpuError::Readback(e.to_string())) +} diff --git a/core/dr-gpu/src/shaders/histogram.wgsl b/core/dr-gpu/src/shaders/histogram.wgsl new file mode 100644 index 0000000..eaf83bb --- /dev/null +++ b/core/dr-gpu/src/shaders/histogram.wgsl @@ -0,0 +1,112 @@ +// TRACES: FR-DSP-7 +// A reduction of the display frame into bin counts, run where the pixels are. +// +// FR-DSP-7 does not merely permit this shape, it names it: "these derive from a +// GPU-side reduction into a small buffer", because the alternative — dragging +// the frame back across the bus to count it — is the per-frame round-trip +// ARCH §6.1 forbids and the bottleneck darktable documents. What leaves the +// device here is 4104 bytes whatever the image size. +// +// **Counted per workgroup first, then merged.** A photograph is not noise: a +// clear sky puts tens of thousands of adjacent pixels in one bin, and having +// every invocation contend for that single global atomic serialises the whole +// dispatch. Each workgroup therefore tallies its own 256 pixels into workgroup +// memory — where the atomic is cheap and the contention is between 256 threads +// rather than two million — and contributes one add per non-empty bin at the +// end. Four kilobytes of workgroup storage against a 16 KB floor. + +const BINS: u32 = 256u; +// Two counters past the four channels: how many pixels clip at each end. They +// live in the same buffer because they are gathered from the same texel read +// and would otherwise need a second reduction to answer a question the first +// one already had the data for. +const CLIPPED_HIGH: u32 = 4u * BINS; +const CLIPPED_LOW: u32 = 4u * BINS + 1u; +const SLOTS: u32 = 4u * BINS + 2u; +// 16x16. Stated as a constant because the clear and merge loops below stride by +// it, and a workgroup size that disagreed would leave slots uncleared. +const THREADS: u32 = 256u; + +struct Dims { + width: u32, + height: u32, + // std140 rounds a uniform block up to 16 bytes; named rather than left + // implicit so the Rust side's padding is visibly the same shape. + pad_0: u32, + pad_1: u32, +} + +@group(0) @binding(0) var source: texture_2d; +@group(0) @binding(1) var bins: array>; +@group(0) @binding(2) var dims: Dims; + +var tile: array, SLOTS>; + +/// The 8-bit level a sampled texel came from. +/// +/// The source is `Rgba8Unorm`, so the sampler hands back exactly n/255 and this +/// recovers n. `floor(x + 0.5)` rather than `round`, which ties to even in WGSL +/// and away from zero in Rust — a difference invisible except on an exact tie, +/// which is precisely the kind of disagreement that makes a cross-check against +/// a CPU reference fail once in a thousand runs and look like flakiness. +fn level(v: f32) -> u32 { + return u32(clamp(floor(v * 255.0 + 0.5), 0.0, 255.0)); +} + +@compute @workgroup_size(16, 16, 1) +fn main( + @builtin(global_invocation_id) gid: vec3, + @builtin(local_invocation_index) lid: u32, +) { + for (var i = lid; i < SLOTS; i = i + THREADS) { + atomicStore(&tile[i], 0u); + } + workgroupBarrier(); + + // Guarded rather than dispatched exactly: the workgroup is 16x16 and an + // image is not, so the last row and column of workgroups run off the edge. + if (gid.x < dims.width && gid.y < dims.height) { + let texel = textureLoad(source, vec2(i32(gid.x), i32(gid.y)), 0); + let r = level(texel.r); + let g = level(texel.g); + let b = level(texel.b); + + // Rec.709 luma in 8.8 fixed point. 54 + 183 + 19 is exactly 256, so the + // weights sum to unity and the shift can never carry past 255. + // + // Integer rather than float on purpose (ARCH §6.13): a float weighting + // is reproducible only to within the vendor's rounding, and the whole + // value of the CPU cross-check in the tests is that it is exact. + // + // Weighted on the *encoded* values, not on linear light. That is what + // every histogram a photographer has read is: the axis is the output + // level, so a mid-grey has to sit in the middle of it. + let y = (54u * r + 183u * g + 19u * b) >> 8u; + + atomicAdd(&tile[r], 1u); + atomicAdd(&tile[BINS + g], 1u); + atomicAdd(&tile[2u * BINS + b], 1u); + atomicAdd(&tile[3u * BINS + y], 1u); + + // Any channel, not all three: a blown red channel is detail that is + // gone, whatever green and blue still hold. Counting only neutral white + // would stay silent on exactly the saturated highlight — a sunset, a + // red jersey — that clips first and recovers worst. + if (r == 255u || g == 255u || b == 255u) { + atomicAdd(&tile[CLIPPED_HIGH], 1u); + } + if (r == 0u || g == 0u || b == 0u) { + atomicAdd(&tile[CLIPPED_LOW], 1u); + } + } + workgroupBarrier(); + + for (var i = lid; i < SLOTS; i = i + THREADS) { + let count = atomicLoad(&tile[i]); + // Most bins of most workgroups are empty — a 16x16 tile can touch 256 + // of 1026 slots at the very most, and usually far fewer. + if (count != 0u) { + atomicAdd(&bins[i], count); + } + } +} diff --git a/ui/dr-ui/src/develop.rs b/ui/dr-ui/src/develop.rs index f03f4d3..6005912 100644 --- a/ui/dr-ui/src/develop.rs +++ b/ui/dr-ui/src/develop.rs @@ -10,7 +10,7 @@ //! new operation appears in the panel with no change here (FR-DEV-3c). use dr_decode::RawImage; -use dr_gpu::{AdjustPass, DemosaicedImage, Demosaicer, GpuContext}; +use dr_gpu::{AdjustPass, DemosaicedImage, Demosaicer, GpuContext, Histogram, HistogramPass}; use dr_pipeline::ops::curve; use dr_pipeline::{ CropRect, EditGraph, OpCapability, OpId, ParamId, ParamKind, Presentation, Preset, Scope, Unit, @@ -25,6 +25,12 @@ pub struct DevelopSession { graph: EditGraph, demosaiced: DemosaicedImage, adjust: AdjustPass, + /// TRACES: FR-DSP-7 + /// Optional, because a session that cannot count its frames is still a + /// session that can develop them. If the reduction fails to build — an + /// old driver, a device without the storage-buffer atomics it needs — the + /// photographer loses the histogram and keeps the photograph. + histogram: Option, } impl DevelopSession { @@ -77,6 +83,9 @@ impl DevelopSession { graph, demosaiced, adjust: AdjustPass::new(ctx), + histogram: HistogramPass::new(ctx) + .inspect_err(|e| log::warn!("no histogram on this device: {e}")) + .ok(), } } @@ -533,6 +542,34 @@ impl DevelopSession { Ok(slint::Image::from_rgba8(buffer)) } + /// TRACES: FR-DSP-7 + /// Count the frame that is currently on the canvas. + /// + /// **Reads the frame [`Self::render`] last produced rather than rendering + /// its own.** The histogram has to describe what the photographer is + /// looking at, and rendering a second time to count it would both cost a + /// second pass and open the possibility of the two disagreeing. + /// + /// That the frame is the *displayed* one has two consequences worth being + /// explicit about. It is in the output colour space, which is what + /// FR-DSP-7 asks for — the levels counted are the levels the display will + /// show, so a clipped bin means a highlight that is actually gone rather + /// than one the transform might still recover. And when the view is zoomed + /// or cropped it describes the visible region, not the whole file: a + /// photographer inspecting a highlight at 4× is asking about *that* + /// highlight, and a histogram of the parts of the frame off screen would + /// be answering a question nobody asked. + /// + /// `None` where nothing has been rendered yet, or where the device could + /// not build the reduction. + pub fn histogram(&self) -> Option { + let pass = self.histogram.as_ref()?; + let frame = self.adjust.output()?; + pass.compute(frame) + .inspect_err(|e| log::warn!("histogram failed: {e}")) + .ok() + } + /// Render the *whole* frame for the crop overlay to be drawn over. /// /// Crop mode cannot use [`Self::render`]: that applies the crop, so the @@ -1871,6 +1908,87 @@ mod tests { } } + /// A frame black on the left half and white on the right, at `size` + /// square. Both ends of the histogram are occupied and both clipping + /// counters are non-zero, and cropping to one half leaves exactly one of + /// them so. + fn split_frame(size: u32) -> Vec { + let mut rgba = Vec::with_capacity((size * size * 4) as usize); + for _ in 0..size { + for x in 0..size { + let v = if x < size / 2 { 0u8 } else { 255 }; + rgba.extend_from_slice(&[v, v, v, 255]); + } + } + rgba + } + + /// TRACES: FR-DSP-7 + #[test] + fn the_histogram_counts_the_frame_that_is_actually_on_the_canvas() { + // The wiring, end to end and against exact numbers: a 64x64 frame that + // is half black and half white must come back as 2048 pixels at level + // 0, 2048 at 255, and both clipping counters at 2048. + // + // Asserted at the session rather than at the pass because the mistake + // this catches is not arithmetic — `dr_gpu` has its own tests for that + // — it is counting the *wrong texture*. Reading a stale target, or the + // demosaiced source instead of the adjusted output, produces a + // perfectly well-formed histogram of an image the photographer is not + // looking at, which is the one failure mode that cannot be seen. + let Ok(ctx) = pollster::block_on(dr_gpu::GpuContext::new_headless()) else { + log::warn!("no GPU adapter; skipping"); + return; + }; + + let rgba = split_frame(64); + let mut session = + DevelopSession::open_rgb(&ctx, &rgba, 64, 64, dr_types::Orientation::NORMAL) + .expect("session"); + session.render(64, 64).expect("render"); + + let hist = session.histogram().expect("a rendered session must count"); + assert_eq!(hist.pixels(), 64 * 64); + assert_eq!(hist.red()[0], 2048, "the black half"); + assert_eq!(hist.red()[255], 2048, "the white half"); + assert_eq!(hist.clipped_shadows(), 2048); + assert_eq!(hist.clipped_highlights(), 2048); + } + + /// TRACES: FR-DSP-7 + #[test] + fn the_histogram_follows_the_edit_rather_than_the_file() { + // The property that makes it *live*. A histogram computed once from the + // source would pass the test above and be useless — the whole reason + // FR-DSP-7 exists is to show what an adjustment is doing, so cropping + // away the white half must leave a histogram with no white in it and + // no highlight clipping to report. + let Ok(ctx) = pollster::block_on(dr_gpu::GpuContext::new_headless()) else { + log::warn!("no GPU adapter; skipping"); + return; + }; + + let rgba = split_frame(64); + let mut session = + DevelopSession::open_rgb(&ctx, &rgba, 64, 64, dr_types::Orientation::NORMAL) + .expect("session"); + + session.set_crop(CropRect { + x: 0.0, + y: 0.0, + width: 0.5, + height: 1.0, + }); + session.render(64, 64).expect("render"); + + let hist = session.histogram().expect("histogram"); + assert_eq!(hist.pixels(), 32 * 64, "the crop halved the frame"); + assert_eq!(hist.red()[0], 32 * 64); + assert_eq!(hist.red()[255], 0, "the white half was cropped away"); + assert_eq!(hist.clipped_highlights(), 0); + assert_eq!(hist.clipped_shadows(), 32 * 64); + } + #[test] fn routing_indices_map_back_to_the_right_parameter() { // A wrong index would silently move the wrong slider's value, which diff --git a/ui/dr-ui/src/histogram.rs b/ui/dr-ui/src/histogram.rs new file mode 100644 index 0000000..b70c5d2 --- /dev/null +++ b/ui/dr-ui/src/histogram.rs @@ -0,0 +1,330 @@ +//! TRACES: FR-DSP-7 +//! Turning bin counts into something a 280-pixel column can be read from. +//! +//! `dr_gpu` counts; this decides what the counting looks like. The split is +//! ARCH §4.3a's: how many columns a panel can show, which peak to scale +//! against and when a handful of specular pixels is worth an alarm are all +//! questions about *this* interface, and none of them belong beside the shader. +//! +//! Free-standing functions over plain numbers, deliberately. Every judgement +//! here is one that shows up as a wrong picture rather than as a crash — a +//! histogram flattened by a black-point spike, a clipping figure that reads +//! 0.0% when a quarter of the sky is gone — and none of them can be checked by +//! looking at a running application, because a plausible wrong shape and the +//! right shape look equally plausible. So they are all reachable without a GPU, +//! a window or a photograph. + +use dr_gpu::{Histogram, HISTOGRAM_BINS}; + +use crate::HistogramView; + +/// Columns the plot draws. +/// +/// **A divisor of [`HISTOGRAM_BINS`], and that is not a detail.** 96 columns +/// over 256 bins folds two levels into some columns and three into others, so a +/// perfectly even ramp is drawn as a comb — structure the photograph does not +/// have. Four levels per column, always. +/// +/// 64 rather than 128 because the plot is about 256px wide in the develop +/// column: four pixels per column is a bar that can be seen, where two is a +/// hairline. It also halves the repeater, and this panel is instantiated for +/// the life of the window rather than built when it is looked at. +pub(crate) const COLUMNS: usize = 64; + +/// Fraction of the frame that must clip before the indicator lights. +/// +/// Not zero. Almost every photograph has a few pixels at an extreme — a +/// specular glint on chrome, a sensor hot pixel, the dark corner of a vignette +/// — and an indicator that fires on all of them is one a photographer stops +/// reading within a day. A thousandth of the frame is roughly where clipping +/// stops being an artefact and starts being a decision. +const CLIP_VISIBLE: f32 = 0.001; + +/// The whole panel's state, for one counted frame. +pub(crate) fn view(hist: &Histogram) -> HistogramView { + HistogramView { + // Zero pixels means nothing has been rendered yet, not an empty + // photograph — the panel draws its frame and no data rather than a flat + // line, which would be a claim about an image that does not exist. + available: hist.pixels() > 0, + luma: model(&scaled(hist.luma())), + red: model(&scaled(hist.red())), + green: model(&scaled(hist.green())), + blue: model(&scaled(hist.blue())), + highlights_clipped: fraction(hist.clipped_highlights(), hist.pixels()) >= CLIP_VISIBLE, + shadows_clipped: fraction(hist.clipped_shadows(), hist.pixels()) >= CLIP_VISIBLE, + highlights_label: percentage(hist.clipped_highlights(), hist.pixels()).into(), + shadows_label: percentage(hist.clipped_shadows(), hist.pixels()).into(), + } +} + +/// An empty panel — no image open, or a frame that could not be counted. +pub(crate) fn empty() -> HistogramView { + HistogramView { + available: false, + luma: model(&[0.0; COLUMNS]), + red: model(&[0.0; COLUMNS]), + green: model(&[0.0; COLUMNS]), + blue: model(&[0.0; COLUMNS]), + highlights_clipped: false, + shadows_clipped: false, + // Not "0%". Nothing has been counted, and claiming no clipping about a + // frame that does not exist is the one thing an instrument must not do. + highlights_label: UNKNOWN.into(), + shadows_label: UNKNOWN.into(), + } +} + +/// What a figure reads before there is a frame behind it. +const UNKNOWN: &str = "—"; + +fn model(heights: &[f32; COLUMNS]) -> slint::ModelRc { + slint::ModelRc::new(slint::VecModel::from(heights.to_vec())) +} + +/// Fold 256 levels into [`COLUMNS`] columns and scale to 0..1. +pub(crate) fn scaled(bins: &[u32; HISTOGRAM_BINS]) -> [f32; COLUMNS] { + let folded = fold(bins); + let peak = peak(&folded); + if peak == 0 { + return [0.0; COLUMNS]; + } + let mut out = [0.0f32; COLUMNS]; + for (o, count) in out.iter_mut().zip(folded.iter()) { + // Clamped because `peak` deliberately ignores the end columns, so the + // spike at a clipped white is taller than the scale it is drawn on. + *o = (*count as f32 / peak as f32).min(1.0); + } + out +} + +/// Sum adjacent levels into the columns the plot has room for. +fn fold(bins: &[u32; HISTOGRAM_BINS]) -> [u32; COLUMNS] { + const PER_COLUMN: usize = HISTOGRAM_BINS / COLUMNS; + let mut out = [0u32; COLUMNS]; + for (level, count) in bins.iter().enumerate() { + // Saturating: a 24 MP frame of one colour is 24 million in one bin, + // which fits, but two folded bins of a hypothetical larger frame need + // not — and a histogram that wrapped to near-zero at the exact moment + // the image went flat would be worse than useless. + out[level / PER_COLUMN] = out[level / PER_COLUMN].saturating_add(*count); + } + out +} + +/// The count the plot's full height stands for. +/// +/// **The end columns are excluded, and this is the single most consequential +/// decision in the file.** A photograph shot against a black backdrop puts a +/// third of its pixels in level 0; scaled against that, every tone the +/// photographer is actually working with is drawn two pixels tall and the +/// histogram says nothing. The same happens at the top with a blown sky. Both +/// spikes are exactly what the clipping indicators report separately and in +/// figures, so nothing is hidden by leaving them off the scale — the plot stops +/// being dominated by the one fact it was already stating twice. +/// +/// Falls back to the true maximum when the ends are all there is, so a frame +/// that really is entirely black still draws something rather than nothing. +fn peak(folded: &[u32; COLUMNS]) -> u32 { + let interior = folded[1..COLUMNS - 1].iter().copied().max().unwrap_or(0); + if interior > 0 { + interior + } else { + folded.iter().copied().max().unwrap_or(0) + } +} + +fn fraction(clipped: u32, pixels: u32) -> f32 { + if pixels == 0 { + 0.0 + } else { + clipped as f32 / pixels as f32 + } +} + +/// How much of the frame is gone, as a figure rather than a colour. +/// +/// NFR-A11Y-3 asks that no status be carried by hue alone, and this is the +/// text half of that: the marker beside it says *whether*, and this says how +/// much — which is the more useful half anyway, since the choice between +/// pulling a stop back and leaving a specular highlight alone is a choice about +/// magnitude. +/// +/// Distinguishes "none" from "not none but under a tenth of a percent": those +/// are different answers, and rounding the second to `0.0%` would tell a +/// photographer their highlights were safe when the indicator beside it is lit. +fn percentage(clipped: u32, pixels: u32) -> String { + if pixels == 0 || clipped == 0 { + return "0%".into(); + } + let pct = 100.0 * fraction(clipped, pixels); + if pct < 0.1 { + "<0.1%".into() + } else { + format!("{pct:.1}%") + } +} + +#[cfg(test)] +mod tests { + use super::*; + + /// A histogram with `count` pixels at each named level and nothing else. + /// + /// Built through the GPU pass's own constructor path would need a device; + /// these are assertions about arithmetic, so they take the counts directly. + fn bins(levels: &[(usize, u32)]) -> [u32; HISTOGRAM_BINS] { + let mut out = [0u32; HISTOGRAM_BINS]; + for (level, count) in levels { + out[*level] = *count; + } + out + } + + #[test] + fn every_column_covers_the_same_number_of_levels() { + // The comb bug. With a column count that does not divide 256, some + // columns gather three levels and some two, so a perfectly even ramp is + // drawn with every third bar 50% taller — which reads as structure in + // the image that is not there. Asserted on the fold rather than on the + // constant, because the constant being wrong is only a problem through + // what the fold then does with it. + assert_eq!(HISTOGRAM_BINS % COLUMNS, 0); + + let flat = [7u32; HISTOGRAM_BINS]; + let folded = fold(&flat); + let first = folded[0]; + assert!(first > 0); + assert!( + folded.iter().all(|c| *c == first), + "an even distribution must fold to even columns, got {folded:?}" + ); + } + + #[test] + fn a_level_lands_in_the_column_that_covers_it() { + // Routing, level by level. An off-by-one here draws the whole + // photograph one column left of where it belongs, which is invisible on + // any image and wrong on all of them. + let per = HISTOGRAM_BINS / COLUMNS; + for level in [0usize, 1, per, per + 1, HISTOGRAM_BINS - 1] { + let folded = fold(&bins(&[(level, 5)])); + let expected = level / per; + assert_eq!( + folded[expected], 5, + "level {level} missed column {expected}" + ); + assert_eq!( + folded.iter().sum::(), + 5, + "level {level} was counted in more than one column" + ); + } + } + + #[test] + fn a_clipped_spike_does_not_flatten_everything_else() { + // The reason `peak` ignores the ends. A frame that is nine-tenths pure + // black — a studio shot on a black backdrop, or any night scene — must + // still show the tones the photographer is working on. Scaled against + // the black spike they would be a hundredth of the plot's height, which + // is a flat line. + let mut levels = vec![(0usize, 900_000u32)]; + levels.push((128, 1000)); + levels.push((129, 500)); + let heights = scaled(&bins(&levels)); + + let mid = heights[128 / (HISTOGRAM_BINS / COLUMNS)]; + assert!( + mid > 0.9, + "the tallest interior column should reach the top, got {mid}" + ); + // And the spike is still drawn, at the ceiling rather than off it. + assert_eq!(heights[0], 1.0); + } + + #[test] + fn an_entirely_black_frame_still_draws_its_spike() { + // The fallback in `peak`. With every pixel at level 0 the interior is + // empty, and dividing by that peak would either panic or produce a + // plot with nothing in it — which says "no data" about a frame that has + // a great deal of data, all of it bad news. + let heights = scaled(&bins(&[(0, 4096)])); + assert_eq!(heights[0], 1.0); + assert!(heights[1..].iter().all(|h| *h == 0.0)); + } + + #[test] + fn an_empty_frame_scales_to_nothing_rather_than_dividing_by_zero() { + let heights = scaled(&[0u32; HISTOGRAM_BINS]); + assert!(heights.iter().all(|h| *h == 0.0)); + } + + #[test] + fn heights_never_leave_the_plot() { + // Everything downstream multiplies these by the plot's height, so a + // value above 1 paints outside the box and one below 0 paints upward + // out of the panel. Checked across the shapes most likely to break it: + // all the weight at one end, and all of it in one interior column. + for levels in [ + vec![(0usize, 100u32)], + vec![(255, 100)], + vec![(0, 100), (255, 100)], + vec![(64, 100)], + vec![(0, 1_000_000), (64, 1)], + ] { + for h in scaled(&bins(&levels)) { + assert!((0.0..=1.0).contains(&h), "height {h} left the plot"); + } + } + } + + #[test] + fn folding_cannot_overflow_on_a_large_frame() { + // `u32` counts summed two at a time. A saturating add rather than a + // wrapping one, because the wrap would land near zero — a histogram + // that emptied itself at the exact moment the image became flat. + let folded = fold(&[u32::MAX; HISTOGRAM_BINS]); + assert_eq!(folded[0], u32::MAX); + } + + #[test] + fn a_clipping_figure_distinguishes_none_from_nearly_none() { + // `0.0%` and `<0.1%` are different answers, and the second is the one + // that appears beside a lit marker. Rounding it to the first would have + // the panel contradict itself: an indicator saying "clipped" over a + // figure saying "none". + assert_eq!(percentage(0, 1000), "0%"); + assert_eq!(percentage(1, 1_000_000), "<0.1%"); + assert_eq!(percentage(12, 1000), "1.2%"); + assert_eq!(percentage(1000, 1000), "100.0%"); + // No frame yet: not a claim that nothing is clipped so much as nothing + // to claim, and "0%" is the honest reading of zero pixels counted. + assert_eq!(percentage(0, 0), "0%"); + } + + #[test] + fn an_uncounted_frame_says_it_does_not_know_rather_than_says_zero() { + // The panel is emptied when an image is closed and again when one fails + // to render, and it stays empty until the first settled frame. Reading + // "Highlights 0%" through that gap is a positive claim about a + // photograph nobody has counted — the failure mode of an instrument + // that is worse than no instrument. + let blank = empty(); + assert!(!blank.available); + assert_eq!(blank.highlights_label, UNKNOWN); + assert_eq!(blank.shadows_label, UNKNOWN); + assert!(!blank.highlights_clipped && !blank.shadows_clipped); + } + + #[test] + fn the_indicator_ignores_a_handful_of_specular_pixels() { + // The threshold, from both sides. Below it the marker stays dark while + // the figure still reports what it found — an indicator that fired on + // every glint would be one nobody reads, and one that hid the number + // would be one nobody could check. + let just_under = (CLIP_VISIBLE * 1_000_000.0) as u32 - 1; + assert!(fraction(just_under, 1_000_000) < CLIP_VISIBLE); + assert!(fraction(just_under + 1, 1_000_000) >= CLIP_VISIBLE); + } +} diff --git a/ui/dr-ui/src/lib.rs b/ui/dr-ui/src/lib.rs index d7acf76..3ee9156 100644 --- a/ui/dr-ui/src/lib.rs +++ b/ui/dr-ui/src/lib.rs @@ -19,6 +19,7 @@ mod collections_ui; mod derived_sync; mod develop; mod export; +mod histogram; mod labels; mod library; mod library_ui; @@ -256,6 +257,11 @@ fn reset_view_state(window: &AppWindow) { window.set_flip_h(false); window.set_flip_v(false); window.set_framing_modified(false); + // TRACES: FR-DSP-7 + // Emptied rather than left standing: the previous photograph's histogram + // beside the next one's filename is a confident, precise lie, and the gap + // before the new frame settles is exactly long enough to read it. + window.set_histogram(histogram::empty()); } /// Push the framing back to the geometry panel. @@ -1002,10 +1008,32 @@ pub fn run(paths: Vec) -> Result<()> { // duration of every gesture. let (vw, vh) = *viewport.borrow(); window.set_magnified(s.magnifies_source(vw, vh)); + + // TRACES: FR-DSP-7 + // **Counted on the settled frame and no other.** + // + // FR-DSP-7 requires the histogram not extend the FR-DSP-3 + // frame budget, and `draft` is already exactly the flag + // that says a gesture is still moving — so this reduction + // and its 4 KB transfer happen once when the slider stops + // rather than on every one of the forty frames a drag + // emits. Nothing is lost by it: a histogram of a + // half-resolution frame flickering past under a finger is + // not a reading anyone takes. + if !draft { + window.set_histogram( + s.histogram() + .as_ref() + .map_or_else(histogram::empty, histogram::view), + ); + } } Err(e) => { log::warn!("render failed: {e}"); window.set_load_error(e.into()); + // No frame, so nothing to describe. The stale plot would + // otherwise sit beside the error message looking current. + window.set_histogram(histogram::empty()); } } }) diff --git a/ui/dr-ui/style.yaml b/ui/dr-ui/style.yaml index e52db69..8f09731 100644 --- a/ui/dr-ui/style.yaml +++ b/ui/dr-ui/style.yaml @@ -117,6 +117,31 @@ colors: state, and it is worth the one exception to say that instantly. Muted rather than saturated so it does not shift perception of a nearby image. + _plot: + section: plot series + note: | + The histogram's four traces (FR-DSP-7). Hue here is the same exception the + swatch takes and not a second one: a per-channel histogram has to say + *which channel*, and no achromatic treatment can distinguish red from blue + — so the colour is data, exactly as the image beside it is. + + Held well back from full strength, and darker than `ink`, for the reason + the theme preamble gives: three saturated traces sitting a few centimetres + from the photograph would compete with it and shift how its colours read. + These are legible against `ground` at a glance and no louder than that. + + plot-red: "#C4626A" + plot-green: "#6BA867" + plot-blue: "#5F8CCB" + + plot-luma: + value: "#4E5257" + doc: | + The luminance trace, drawn filled and behind the three channels. Neutral + because luminance is not a channel — it is the axis exposure is read off, + and a hue would imply it were one series among four rather than the one + the other three are decomposing. + lengths: _gaps: { break: true } gap-sm: 6 diff --git a/ui/dr-ui/ui/app.slint b/ui/dr-ui/ui/app.slint index 1be7c1e..eb9bc88 100644 --- a/ui/dr-ui/ui/app.slint +++ b/ui/dr-ui/ui/app.slint @@ -4,9 +4,10 @@ import { LaunchScreen } from "launch.slint"; import { LibraryGrid, LibraryCell, TimelineBar } from "library.slint"; import { Button, PanelHeading, Label, Value, Caption, Panel, EmptyState, ProgressBar, ActivityRow } from "widgets.slint"; import { CollectionsPanel, CollectionRow } from "collections.slint"; +import { HistogramPanel, HistogramView } from "histogram.slint"; import { SettingsPage } from "settings.slint"; -export { LibraryCell, TimelineBar, CollectionRow, ActivityRow } +export { LibraryCell, TimelineBar, CollectionRow, ActivityRow, HistogramView } // Status strip — surfaces the GPU backend and adapter, which matters during // v0.1 because assumption A1 is exactly "does this compositing path work on @@ -199,6 +200,11 @@ export component AppWindow inherits Window { in property total: 0; in property load-error: ""; + /// TRACES: FR-DSP-7 + /// The counted frame. Set only when the render has settled — a histogram + /// of a draft frame is a histogram of an image nobody is reading. + in property histogram; + // --- zoom, pan and crop (FR-DEV-4) --- // // Zoom is a *viewing* state, not an edit: it changes the resolution the @@ -1449,6 +1455,21 @@ in property panel-visible: true; background: Theme.rule; } + // Directly under the capture metadata and above every + // control, because it is the thing the controls are + // judged against: exposure, blacks and whites are all + // set by watching this move (FR-DSP-7). An instrument + // below the sliders it reports on would have the + // photographer looking away from it to use it. + HistogramPanel { + data: root.histogram; + } + + Rectangle { + height: 1px; + background: Theme.rule; + } + // Framing above the colour work, matching how the edit is // made rather than how it is applied: the frame is decided // by eye first and the pipeline runs it last (see diff --git a/ui/dr-ui/ui/histogram.slint b/ui/dr-ui/ui/histogram.slint new file mode 100644 index 0000000..d333368 --- /dev/null +++ b/ui/dr-ui/ui/histogram.slint @@ -0,0 +1,248 @@ +// TRACES: FR-DSP-7 +// The live histogram, and what it says about clipping. +// +// **Why this is hand-built when the rest of the column is generated.** The +// develop panel is built from operation capabilities, and a histogram is not an +// operation: it has no parameters, changes nothing about the photograph, and +// answers a question rather than asking one. It is an *instrument* — the same +// kind of thing as the zoom readout — so it is written, and ARCH §4.3a is +// untroubled by it: nothing here reads a parameter out of a descriptor. +// +// **Everything numeric was decided in Rust.** The heights arriving here are +// already 0..1 against a chosen scale, and the clipping figures are already +// strings. That is not tidiness — it is the only way the arithmetic is +// testable. A peak chosen in Slint could only be checked by looking at it, and +// a histogram that is the wrong shape looks exactly as plausible as one that is +// right (see `src/histogram.rs`). + +import { Theme } from "theme.slint"; +import { PanelHeading, Caption, Panel } from "widgets.slint"; + +// One counted frame, ready to draw. +// +// Heights are 0..1 with y up, darkest level first — the same convention the +// tone curve's samples use, so the two plots in this column cannot end up +// disagreeing about which way is up. +export struct HistogramView { + // Whether a frame has been counted at all. Distinct from an all-zero + // histogram, which would be a claim about an image rather than the absence + // of one. + available: bool, + + luma: [float], + red: [float], + green: [float], + blue: [float], + + // Past the threshold worth reporting — see `CLIP_VISIBLE` in Rust. + highlights-clipped: bool, + shadows-clipped: bool, + + // How much, as text. NFR-A11Y-3: the marker beside these says *whether* + // by hue and position, and these say *how much* in a form that survives + // being unable to tell the two apart. + highlights-label: string, + shadows-label: string, +} + +// One series, as a run of columns. +// +// Two shapes from one component because the two differ only in where the column +// starts: the luminance trace is filled from the baseline, the three channels +// are a line. Splitting them into two components would duplicate the column +// arithmetic, which is the part that has to stay identical or the four series +// would no longer be plotted against the same axis. +// +// One thin Rectangle per column: Slint has no polyline primitive, and this is +// the same construction `CurveEditor` uses a few files away. +component Trace inherits Rectangle { + /// 0..1, y up. + in property <[float]> heights; + in property ink; + /// Filled from the baseline rather than drawn as a line. + in property filled: false; + in property shade: 1.0; + + background: transparent; + + for h[i] in root.heights: Rectangle { + // The next column's height, so a line segment can span the gap between + // them. Without this a steep edge draws as a dotted stair rather than a + // rising line — the same fix `CurveEditor` makes for the same reason. + property next: i + 1 < root.heights.length + ? root.heights[i + 1] : h; + + x: parent.width * i / max(root.heights.length, 1); + width: parent.width / max(root.heights.length, 1) + 1px; + + y: root.filled + ? parent.height * (1.0 - h) + : parent.height * (1.0 - max(h, self.next)); + height: root.filled + ? parent.height * h + // A floor, so a flat stretch of the trace is still a line and not a + // gap in one. + : max(parent.height * abs(self.next - h), 1.5px); + + background: root.ink; + opacity: root.shade; + } +} + +// A clipping readout: a lit marker and a figure. +// +// **Two affordances for one fact, and NFR-A11Y-3 is why.** No status in this +// application is carried by hue alone, and clipping is the status most worth +// getting right — it is the one the histogram exists to warn about. So the +// marker appears and disappears (a shape, not a tint) and the figure states the +// magnitude in words. Either alone is enough to read it. +component ClipReadout inherits HorizontalLayout { + in property title; + in property figure; + in property lit; + + spacing: Theme.gap-sm; + + Rectangle { + width: 6px; + height: 6px; + y: (parent.height - self.height) / 2; + border-radius: 1px; + // `warn-ink`, the palette's one sanctioned hue: a blown highlight is a + // caution about the photograph, which is exactly the kind of thing that + // token exists to say instantly. + background: Theme.warn-ink; + visible: root.lit; + } + + Caption { text: root.title; } + Caption { text: root.figure; warn: root.lit; } +} + +// TRACES: FR-DSP-7 +export component HistogramPanel inherits Rectangle { + in property data; + + background: transparent; + height: panel.preferred-height; + + panel := Panel { + // Flat, like every panel in this column: the rule between it and its + // neighbour belongs to the column that stacks them. + flat: true; + width: 100%; + + PanelHeading { text: "HISTOGRAM"; } + + Rectangle { + // Tall enough to read a shape off and no taller. The column is the + // photographer's instrument panel, and every pixel this takes is a + // slider pushed below the fold. + height: 84px; + background: Theme.ground; + border-width: 1px; + border-color: Theme.rule; + // The traces are positioned by fraction of the plot and a value at + // the ceiling lands exactly on the border; clipping keeps the top + // row of pixels inside the box rather than over its edge. + clip: true; + + // Quarter gridlines, so a tone can be placed on the axis without + // counting. Same construction and same weight as the tone curve's, + // because the two plots sit in one column and any difference + // between them would read as meaning something. + for i in [1, 2, 3]: Rectangle { + x: parent.width * i / 4; + width: 1px; + background: Theme.rule; + opacity: 0.5; + } + + // Luminance first, so it sits behind: it is the envelope the three + // channels decompose, and a filled area drawn over a line hides it. + Trace { + width: 100%; + height: 100%; + heights: root.data.luma; + ink: Theme.plot-luma; + filled: true; + } + + // The channels over it, held back so three overlapping traces stay + // legible where they cross — which on a neutral subject is + // everywhere. + Trace { + width: 100%; + height: 100%; + heights: root.data.red; + ink: Theme.plot-red; + shade: 0.85; + } + Trace { + width: 100%; + height: 100%; + heights: root.data.green; + ink: Theme.plot-green; + shade: 0.85; + } + Trace { + width: 100%; + height: 100%; + heights: root.data.blue; + ink: Theme.plot-blue; + shade: 0.85; + } + + // The clipping markers, drawn on the plot's own ends. + // + // Here as well as in the figures below because this is where the + // eye already is: a bar standing at the edge the tones are piling + // against says which end has gone in the same glance that reads the + // shape, and the figures underneath say how much. + Rectangle { + x: 0; + width: 2px; + height: 100%; + background: Theme.warn-ink; + visible: root.data.shadows-clipped; + } + Rectangle { + x: parent.width - self.width; + width: 2px; + height: 100%; + background: Theme.warn-ink; + visible: root.data.highlights-clipped; + } + + // Nothing rendered yet. Said rather than shown as an empty plot, + // which would be a claim that the photograph has no tones in it. + Caption { + text: "no frame yet"; + width: 100%; + horizontal-alignment: center; + y: (parent.height - self.height) / 2; + visible: !root.data.available; + } + } + + // Always drawn, never hidden. With no frame counted the figures read + // "—", which is a different statement from "0%" and the true one: the + // panel does not yet know. Hiding the row instead would move the + // controls below it up and back down as each image opens. + HorizontalLayout { + ClipReadout { + title: "Shadows"; + figure: root.data.shadows-label; + lit: root.data.shadows-clipped; + } + + Rectangle { horizontal-stretch: 1; } + + ClipReadout { + title: "Highlights"; + figure: root.data.highlights-label; + lit: root.data.highlights-clipped; + } + } + } +}