Merge branch 'histogram'

# Conflicts:
#	ui/dr-ui/src/lib.rs

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);
+        }
+    }
+}