From 0233df4bf279f65c1cd67281b81767c4902a9c0f Mon Sep 17 00:00:00 2001 From: Duncan Tourolle Date: Mon, 17 Aug 2026 09:55:40 +0200 Subject: [PATCH] See what the highlights are doing: a live histogram (FR-DSP-7) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Exposure, blacks and whites were set by eye. Nothing said a highlight had blown — the canvas shows white where a channel is at 250 and white where it is at 255, and the difference is the whole question. **Counted on the GPU, not on the readback.** There is a full frame sitting in CPU memory on every canvas update right now — `AdjustPass::read_output`, the bridge spike S1 removes — and walking it would have been thirty lines and no shader. FR-DSP-7 states the mechanism and 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 founded on the bridge would be correct today and deleted by S1, and would meanwhile be the reason the bridge could not go. What crosses the bus here is 4104 bytes whatever the image size. The reduction tallies into workgroup memory first and merges once per workgroup. A photograph is not noise: a clear sky puts tens of thousands of adjacent pixels in one bin, and contending for that single global atomic serialises the dispatch. **On the settled frame only.** `render_now` already knows whether a gesture is still moving — `draft` is the flag `redraw` derives from `was_coalesced` — so the dispatch and its transfer happen once when the slider stops rather than on each of the forty frames a drag emits. Nothing is lost: a histogram flickering past under a finger is not a reading anyone takes. FR-DSP-7 requires exactly this, that it not extend the FR-DSP-3 frame budget. Luma is weighted in 8.8 fixed point — 54, 183, 19, summing to 256 exactly — rather than in floats. Not thrift: it makes the shader's arithmetic reproducible bit for bit, which is what lets the test below be an `assert_eq` against a CPU count rather than a tolerance. ARCH §6.13's line about integer state, applied where it happens to also be free. **What the numbers were checked against.** A flat frame must put all 4096 pixels in one bin and one only. A 256-wide ramp must occupy every level with exactly the same count, which is what catches an off-by-one in the quantisation — a `floor` where a rounding was needed shifts the whole photograph one bin left and looks like nothing at all. And a 101x37 frame of seeded pseudo-random pixels — deliberately not a multiple of the 16x16 workgroup, so the edge tiles run off the image — is compared slot for slot against a second, obvious CPU implementation. Exact equality, no tolerance. The CPU version is a deliberate reimplementation rather than shared code: the bugs worth catching here are ones shared code would commit identically on both sides. Above that, the presentation arithmetic is unit-tested headless, because it is where a wrong answer is invisible. A histogram of the wrong shape looks exactly as plausible as one of the right shape. So: 64 columns because it divides 256 and an uneven fold draws an even ramp as a comb; the peak excludes the end columns, or a night scene scaled against its own black spike is a flat line with no information in it; heights are clamped into the plot; and "0%" is kept distinct from "<0.1%" and from "—", since an indicator reading "clipped" over a figure reading "none" is a panel contradicting itself. Clipping counts a *pixel* with any channel at an extreme, not a channel. Any, because a blown red has no gradation left in it however much green and blue still hold — and it is the saturated highlight, the sunset and the red jersey, that clips first and recovers worst. Per pixel, because counting channels can report 200% of a frame clipped, and a percentage above 100 is a readout nobody trusts again. Two affordances for it, which NFR-A11Y-3 asks for: a bar standing at the end of the plot the tones are piling against, and a figure saying how much. Either alone reads. The panel sits directly under the capture metadata and above every control, because it is what the controls are judged against. It is hand-built rather than generated, and ARCH §4.3a is untroubled: a histogram is not an operation — no parameters, changes nothing, answers a question rather than asking one — and nothing in it reads a parameter out of a descriptor. Three plot colours and a neutral luma trace join the palette. That is the swatch's exception rather than a second one: a per-channel histogram has to say which channel, and no achromatic treatment distinguishes red from blue, so the hue is data exactly as the image beside it is. Held well back from full strength for the reason the theme preamble gives. The bounded, non-parking map wait moves out of `AdjustPass` into `readback::await_mapping`, shared with the histogram's transfer. Thirty lines of load-bearing reasoning about frozen interfaces and lost devices, and two copies of it would have drifted. The histogram describes the frame on the canvas, so it is in the output colour space FR-DSP-7 asks for, and when zoomed it describes the visible region — a photographer inspecting a highlight at 4x is asking about that highlight. A device that cannot build the reduction loses the histogram and keeps the photograph. Still to do for FR-DSP-7: the pixel colour readout under the cursor. 324 tests pass, clippy and fmt clean. Co-Authored-By: Claude Opus 5 --- core/dr-gpu/src/adjust.rs | 49 +- core/dr-gpu/src/histogram.rs | 591 +++++++++++++++++++++++++ core/dr-gpu/src/lib.rs | 5 + core/dr-gpu/src/readback.rs | 56 +++ core/dr-gpu/src/shaders/histogram.wgsl | 112 +++++ ui/dr-ui/src/develop.rs | 120 ++++- ui/dr-ui/src/histogram.rs | 330 ++++++++++++++ ui/dr-ui/src/lib.rs | 28 ++ ui/dr-ui/style.yaml | 25 ++ ui/dr-ui/ui/app.slint | 23 +- ui/dr-ui/ui/histogram.slint | 248 +++++++++++ 11 files changed, 1541 insertions(+), 46 deletions(-) create mode 100644 core/dr-gpu/src/histogram.rs create mode 100644 core/dr-gpu/src/readback.rs create mode 100644 core/dr-gpu/src/shaders/histogram.wgsl create mode 100644 ui/dr-ui/src/histogram.rs create mode 100644 ui/dr-ui/ui/histogram.slint 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; + } + } + } +}