WIP: clarity and texture

Checkpoint committed by the coordinator, not by the authoring agent: the
session hit its API limit mid-task and left this work uncommitted. Committed
so it survives, NOT because it is finished - expect failing tests and
half-applied changes. The agent resumes from here.

core/dr-pipeline/src/detail.rs

@@ -315,6 +315,37 @@ pub struct DetailPass {
     ///   them in the shader. Prefer computing lengths on the CPU in
     ///   [`DetailStage::passes`], where the units are named methods rather
     ///   than an untyped float.
+    /// - `aux: f32` and `tap_aux(coord, offset) -> f32` — **one scalar per
+    ///   pixel that survives to the next pass**, pre-loaded with what the
+    ///   previous pass left there and written back out unless the body
+    ///   assigns it.
+    ///
+    /// # Why `aux` exists
+    ///
+    /// The ping-pong hands each pass exactly one texture: what the pass before
+    /// it wrote. That is enough for a chain of filters — a separable blur is
+    /// two of them — and it is *not* enough for an unsharp mask, which is the
+    /// shape of sharpening, clarity, texture and dehaze alike. An unsharp mask
+    /// needs the blur **and** the original in the same place at the same time,
+    /// and once the first pass has written its blur the original is gone.
+    ///
+    /// Three channels cannot carry both. Even restricted to the case where the
+    /// operation only moves luminance — so the colour is a luminance and two
+    /// chromaticity degrees of freedom — the combining pass needs four
+    /// numbers: the original luminance, two of chromaticity, and the blurred
+    /// luminance. Four does not fit in three, and no encoding makes it fit.
+    ///
+    /// The intermediate is `rgba16float` and its alpha was being written as a
+    /// constant `1.0` and read by nobody, so the fourth number goes there. A
+    /// blur pass leaves `c` alone and puts its result in `aux`; the pass after
+    /// it therefore receives the untouched original *and* the blur, and can
+    /// subtract one from the other. An operation with no use for the lane says
+    /// nothing and hands on what it was given.
+    ///
+    /// The last pass in the chain writes the display texture, whose alpha is
+    /// opacity rather than scratch space, so `aux` is readable there and not
+    /// written. That is exactly the right way round: the combining pass is the
+    /// one that reads it.
     ///
     /// Uniforms are addressed by the bare names declared in [`Self::uniforms`],
     /// exactly as a fused fragment addresses its own; the composer rewrites
@@ -536,7 +567,11 @@ fn compose_one(
             "rgba16float",
             "    // Another linear intermediate: no clip and no encode, because\n\
              \x20   // the pass after this one still has to read real values.\n\
-             \x20   textureStore(output, coord, vec4<f32>(c, 1.0));"
+             \x20   //\n\
+             \x20   // `aux` rides in alpha. A pass that never touches it hands on\n\
+             \x20   // whatever it was given, so the lane costs an operation that\n\
+             \x20   // does not want it exactly one copy of a value it already read.\n\
+             \x20   textureStore(output, coord, vec4<f32>(c, aux));"
                 .to_string(),
         )
     };
@@ -582,6 +617,12 @@ fn tap(coord: vec2<i32>, offset: vec2<i32>) -> vec3<f32> {{
     return textureLoad(source, clamp(coord + offset, vec2<i32>(0), last), 0).rgb;
 }}
 
+// The same neighbour's scratch lane — see `aux` in the body below.
+fn tap_aux(coord: vec2<i32>, offset: vec2<i32>) -> f32 {{
+    let last = vec2<i32>(textureDimensions(source)) - vec2<i32>(1);
+    return textureLoad(source, clamp(coord + offset, vec2<i32>(0), last), 0).a;
+}}
+
 {helper_src}{encode_fn}
 @compute @workgroup_size(8, 8, 1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {{
@@ -596,6 +637,10 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {{
     let render_scale = u.detail_base.z;
 
     var c = tap(coord, vec2<i32>(0));
+    // One scalar per pixel that survives the hand-off from one pass to the
+    // next, alongside the colour. See `DetailPass::wgsl` for what it is for
+    // and why three channels were not enough.
+    var aux = tap_aux(coord, vec2<i32>(0));
 
     {{
 {indented}
@@ -880,6 +925,44 @@ mod tests {
         }
     }
 
+    #[test]
+    fn a_pass_can_hand_a_scalar_to_the_next_one_alongside_the_colour() {
+        // What makes an unsharp mask — sharpening, clarity, texture, dehaze —
+        // expressible at all in a chain that hands each pass exactly one
+        // texture. The blur goes in `aux` and the colour rides through
+        // untouched, so the pass that combines them receives both; without the
+        // lane, four numbers would have to fit in three channels and the
+        // operation could only ever be a blur.
+        //
+        // A pass that says nothing about `aux` hands on what it was given,
+        // which is why the box blur below needs no knowledge of it.
+        let ops = with_blur(0.05);
+        let composed = compose_detail(
+            &ops,
+            RenderScale::full((512, 512)),
+            dr_types::ColourSpace::Srgb,
+        );
+
+        for pass in &composed.passes {
+            assert!(
+                pass.source.contains("fn tap_aux(") && pass.source.contains("var aux = tap_aux("),
+                "{} cannot read the scratch lane",
+                pass.label
+            );
+        }
+        assert!(
+            composed.passes[0]
+                .source
+                .contains("textureStore(output, coord, vec4<f32>(c, aux));"),
+            "an intermediate must carry the lane to the pass after it"
+        );
+        // The last pass writes the display texture, whose alpha is opacity and
+        // not scratch space. Readable there, not written — which is the right
+        // way round, because the combining pass is the one that reads it.
+        assert!(composed.passes[1].writes_output);
+        assert!(!composed.passes[1].source.contains("vec4<f32>(c, aux)"));
+    }
+
     #[test]
     fn an_edit_with_no_detail_operation_composes_no_passes() {
         // The property that keeps the cost of this stage at zero for the