See what the highlights are doing: a live histogram (FR-DSP-7)

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 <noreply@anthropic.com>

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<f32>;
+@group(0) @binding(1) var<storage, read_write> bins: array<atomic<u32>>;
+@group(0) @binding(2) var<uniform> dims: Dims;
+
+var<workgroup> tile: array<atomic<u32>, 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<u32>,
+    @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>(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);
+        }
+    }
+}