Let a detail pass carry a list, not only a kernel

Every neighbourhood pass so far has been a convolution, whose whole
description fits in the uniform block because its structure fixes how many
numbers it needs. Spot removal is not that shape: sixty-four repairs and
one repair are the same shader with a different buffer behind it.

So a pass may declare `storage`, which arrives at binding 3 as
`array<vec4<f32>>` with `arrayLength` in scope. The alternative — packing
the list into uniforms — needs a fixed maximum paid for on every frame, a
composer that can emit vec4 fields because a uniform array's stride is 16
whatever it holds, and it gives the next operation that wants a table
nothing to build on.

The property worth having is what stays out of the generated source: the
count is in the buffer, so placing the tenth spot uploads 512 bytes and
reuses the compiled pipeline, exactly as moving a slider does for the
fused pass. `changing_the_list_does_not_recompile` is that, asserted.

One bind group entry rather than two more layouts, and one placeholder
buffer allocated in `new` rather than sixteen bytes per pass per frame —
a zero-length storage buffer cannot be bound, and per-frame allocation is
what this module's documentation exists to refuse.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

core/dr-gpu/src/detail.rs

@@ -149,6 +149,8 @@ pub(crate) struct DetailRunner {
     /// Compiled pipelines by pass structure hash.
     cache: HashMap<u64, wgpu::ComputePipeline>,
     pool: Intermediates,
+    /// See [`placeholder_instances`].
+    no_instances: wgpu::Buffer,
 }
 
 struct Layout {
@@ -156,6 +158,22 @@ struct Layout {
     pipeline: wgpu::PipelineLayout,
 }
 
+/// What binding 3 holds for a pass that declared no instance list.
+///
+/// One zeroed element, allocated once. Zero-length storage buffers cannot be
+/// bound, and the passes that read this binding are exactly the ones that
+/// uploaded something of their own, so nothing ever reads the placeholder's
+/// contents — it exists to keep one bind group layout serving both kinds of
+/// pass.
+fn placeholder_instances(ctx: &GpuContext) -> wgpu::Buffer {
+    ctx.device
+        .create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: Some("detail-instances-placeholder"),
+            contents: bytemuck::cast_slice(&[[0.0f32; 4]]),
+            usage: wgpu::BufferUsages::STORAGE,
+        })
+}
+
 impl DetailRunner {
     pub(crate) fn new(ctx: &GpuContext) -> Self {
         Self {
@@ -164,6 +182,7 @@ impl DetailRunner {
             to_output: Layout::new(ctx, crate::AdjustPass::FORMAT, "detail-output"),
             cache: HashMap::new(),
             pool: Intermediates::new(),
+            no_instances: placeholder_instances(ctx),
         }
     }
 
@@ -230,6 +249,24 @@ impl DetailRunner {
                     usage: wgpu::BufferUsages::UNIFORM,
                 });
 
+            // TRACES: FR-DEV-8
+            // The instance list, uploaded only by the passes that have one. A
+            // kernel pass — which is every pass that is a convolution — is
+            // handed the placeholder allocated once in `new`, because a storage
+            // buffer of length zero is not bindable and allocating a fresh
+            // sixteen bytes per pass per frame is the per-frame allocation this
+            // module's documentation exists to refuse.
+            let instances = (!pass.storage.is_empty()).then(|| {
+                self.ctx
+                    .device
+                    .create_buffer_init(&wgpu::util::BufferInitDescriptor {
+                        label: Some("detail-instances"),
+                        contents: bytemuck::cast_slice(pass.storage.as_slice()),
+                        usage: wgpu::BufferUsages::STORAGE,
+                    })
+            });
+            let instances = instances.as_ref().unwrap_or(&self.no_instances);
+
             let bind_group = self
                 .ctx
                 .device
@@ -249,6 +286,10 @@ impl DetailRunner {
                             binding: 2,
                             resource: wgpu::BindingResource::TextureView(destination),
                         },
+                        wgpu::BindGroupEntry {
+                            binding: 3,
+                            resource: instances.as_entire_binding(),
+                        },
                     ],
                 });
 
@@ -338,6 +379,20 @@ impl DetailRunner {
     }
 }
 
+/// A read-only storage buffer entry, as `mask.rs` declares its strokes.
+fn storage_entry(binding: u32) -> wgpu::BindGroupLayoutEntry {
+    wgpu::BindGroupLayoutEntry {
+        binding,
+        visibility: wgpu::ShaderStages::COMPUTE,
+        ty: wgpu::BindingType::Buffer {
+            ty: wgpu::BufferBindingType::Storage { read_only: true },
+            has_dynamic_offset: false,
+            min_binding_size: None,
+        },
+        count: None,
+    }
+}
+
 impl Layout {
     fn new(ctx: &GpuContext, format: wgpu::TextureFormat, label: &str) -> Self {
         let bind_group = ctx
@@ -376,6 +431,14 @@ impl Layout {
                         },
                         count: None,
                     },
+                    // TRACES: FR-DEV-8
+                    // The instance list, for a pass whose work is a list rather
+                    // than a kernel (`DetailPass::storage`). Every other pass
+                    // gets `Intermediates`' placeholder here — one entry on both
+                    // layouts rather than two more layouts, since a convolution
+                    // that never reads the buffer costs nothing for it being
+                    // bound.
+                    storage_entry(3),
                 ],
             });
 

core/dr-gpu/tests/detail_instances.rs

@@ -0,0 +1,169 @@
+//! TRACES: FR-DEV-8
+//! The instance binding: a detail pass whose work is a list, not a kernel.
+//!
+//! Spot removal needs a detail pass to read a variable number of records —
+//! sixty-four repairs and one repair are the same shader with a different
+//! buffer behind it. That is binding 3, and this file proves the three things
+//! about it that a picture would not tell you clearly:
+//!
+//! - the data uploaded is the data the shader reads, in order;
+//! - a pass that declares no list still runs, bound to the placeholder;
+//! - the same shader with a *different* list does not recompile, which is what
+//!   keeps placing a spot as cheap as moving a slider.
+//!
+//! The passes here are synthetic on purpose. `spot_removal.rs` asserts the
+//! repair; this asserts the plumbing, so a failure in one does not have to be
+//! read to work out which of the two broke.
+
+use dr_gpu::{AdjustPass, DemosaicedImage, GpuContext};
+use dr_pipeline::detail::{ComposedDetail, ComposedDetailPass};
+use dr_pipeline::{Affects, EditGraph};
+use dr_types::ColourSpace;
+
+const SIZE: u32 = 8;
+
+fn ctx() -> Option<GpuContext> {
+    match pollster::block_on(GpuContext::new_headless()) {
+        Ok(c) => Some(c),
+        Err(e) => {
+            eprintln!("skipping: no GPU adapter ({e})");
+            None
+        }
+    }
+}
+
+/// A flat mid-grey frame, so anything the pass adds is the whole answer.
+fn grey(ctx: &GpuContext) -> DemosaicedImage {
+    let data: Vec<u8> = (0..SIZE * SIZE).flat_map(|_| [0u8, 0, 0, 255]).collect();
+    DemosaicedImage::from_rgba8(ctx, &data, SIZE, SIZE).expect("upload")
+}
+
+/// A pass that sums the instance list into the red channel and writes the
+/// output. Deliberately trivial: the value on screen is then a direct readout
+/// of what arrived in the buffer.
+fn summing_pass(storage: Vec<[f32; 4]>, structure: u64) -> ComposedDetailPass {
+    let source = "
+@group(0) @binding(0) var source: texture_2d<f32>;
+struct Params { detail_base: vec4<f32> }
+@group(0) @binding(1) var<uniform> u: Params;
+@group(0) @binding(2) var output: texture_storage_2d<rgba8unorm, write>;
+@group(0) @binding(3) var<storage, read> instances: array<vec4<f32>>;
+
+@compute @workgroup_size(8, 8, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let dims = textureDimensions(output);
+    if (gid.x >= dims.x || gid.y >= dims.y) { return; }
+
+    // Weighted by index, so a buffer read back to front fails this rather
+    // than passing by symmetry.
+    var total = 0.0;
+    let n = arrayLength(&instances);
+    for (var i = 0u; i < n; i = i + 1u) {
+        total = total + instances[i].x * f32(i + 1u);
+    }
+    textureStore(output, vec2<i32>(gid.xy), vec4<f32>(total, f32(n) / 255.0, 0.0, 1.0));
+}
+"
+    .to_string();
+
+    ComposedDetailPass {
+        label: "test/instances".to_string(),
+        source,
+        uniforms: vec![SIZE as f32, SIZE as f32, 1.0, 0.0],
+        storage,
+        radius: 0,
+        writes_output: true,
+        // Any distinct number: the hash is a cache key, and these tests are
+        // what decide whether two chains share a pipeline.
+        structure_hash: structure,
+    }
+}
+
+fn render(pass: &mut AdjustPass, source: &DemosaicedImage, chain: &ComposedDetail) -> Vec<u8> {
+    // The fused half has to be composed knowing a detail stage follows it, or
+    // it encodes its own output and the chain would quantise twice — a mismatch
+    // `render_detailed` refuses outright. The probe is the graph that says so;
+    // its own passes are not used, since the chain here is hand-built.
+    let mut graph = EditGraph::with_detail_probe();
+    graph.set_param(
+        dr_pipeline::descriptor::OpId("detail_probe"),
+        dr_pipeline::descriptor::ParamId("radius"),
+        0.05,
+    );
+    let shader = graph.compose_for(ColourSpace::Srgb);
+    let key = graph.invalidation().through(Affects::Colour);
+    pass.render_detailed(source, &shader, SIZE, SIZE, None, chain, key)
+        .expect("render");
+    pass.export_pixels().expect("readback").0
+}
+
+/// The list arrives whole, in order, and the shader can tell how long it is.
+#[test]
+fn a_pass_reads_the_list_it_was_given() {
+    let Some(ctx) = ctx() else { return };
+    let source = grey(&ctx);
+    let mut pass = AdjustPass::new(&ctx);
+
+    // 0.1·1 + 0.2·2 + 0.3·3 = 1.4, which clips to 1.0 — so instead: values
+    // chosen to land at a quarter, unambiguously distinguishable from both the
+    // "read nothing" answer of 0 and the "read them unweighted" answer of 0.15.
+    let chain = ComposedDetail {
+        passes: vec![summing_pass(
+            vec![[0.05, 0.0, 0.0, 0.0], [0.1, 0.0, 0.0, 0.0]],
+            1,
+        )],
+    };
+
+    let pixels = render(&mut pass, &source, &chain);
+    let (red, green) = (pixels[0], pixels[1]);
+
+    // 0.05·1 + 0.1·2 = 0.25, written straight to an rgba8 target.
+    assert!(
+        red.abs_diff((0.25 * 255.0) as u8) <= 1,
+        "the shader summed {red}, not the list it was handed"
+    );
+    assert_eq!(green, 2, "arrayLength saw both entries");
+}
+
+/// A convolution declares no list and must still run: it is bound to the
+/// placeholder rather than to nothing, because a zero-length storage buffer
+/// cannot be bound at all and a second bind group layout for the difference
+/// would be two layouts to keep in step.
+#[test]
+fn a_pass_with_no_list_still_runs() {
+    let Some(ctx) = ctx() else { return };
+    let source = grey(&ctx);
+    let mut pass = AdjustPass::new(&ctx);
+
+    let chain = ComposedDetail {
+        passes: vec![summing_pass(Vec::new(), 2)],
+    };
+
+    let pixels = render(&mut pass, &source, &chain);
+    assert_eq!(pixels[0], 0, "the placeholder is zeroed");
+    assert_eq!(pixels[1], 1, "and is exactly one element long");
+}
+
+/// The property that makes placing the tenth spot as cheap as moving a slider:
+/// the list is in the buffer, not in the source, so the pipeline is compiled
+/// once however many entries arrive.
+#[test]
+fn changing_the_list_does_not_recompile() {
+    let Some(ctx) = ctx() else { return };
+    let source = grey(&ctx);
+    let mut pass = AdjustPass::new(&ctx);
+
+    for count in 1..=6 {
+        let list = (0..count).map(|_| [0.01, 0.0, 0.0, 0.0]).collect();
+        let chain = ComposedDetail {
+            passes: vec![summing_pass(list, 3)],
+        };
+        render(&mut pass, &source, &chain);
+    }
+
+    assert_eq!(
+        pass.cached_detail_pipelines(),
+        1,
+        "six different lists, one compiled pipeline"
+    );
+}

core/dr-pipeline/src/detail.rs

@@ -359,6 +359,29 @@ pub struct DetailPass {
 
     /// Uniform values this pass's body reads.
     pub uniforms: Vec<Uniform>,
+
+    /// TRACES: FR-DEV-8
+    /// Per-instance data, for a pass whose work is a *list* rather than a
+    /// kernel.
+    ///
+    /// Reaches the body as `instances: array<vec4<f32>>`, with
+    /// `instance_count` in scope as a `u32`. Empty for every pass that is a
+    /// convolution, which is every pass that existed before spot removal.
+    ///
+    /// # Why not the uniform block
+    ///
+    /// Because the uniform block is fixed by the pass's *structure*, and this
+    /// is not: sixty-four repairs and one repair are the same shader with a
+    /// different buffer behind it. Packing the list into uniforms would need a
+    /// fixed maximum, paid for on every frame whether the photograph carries
+    /// one spot or none, and it would need the composer to emit `vec4` fields —
+    /// a WGSL uniform array has a stride of 16 whatever it holds.
+    ///
+    /// The property that matters more: with the list in storage the generated
+    /// source does not mention how many there are, so placing the tenth spot
+    /// uploads 512 bytes and reuses the compiled pipeline, exactly as moving a
+    /// slider does for the fused pass.
+    pub storage: Vec<[f32; 4]>,
 }
 
 /// TRACES: FR-DEV-3 | FR-DEV-8
@@ -396,6 +419,9 @@ pub struct ComposedDetailPass {
     pub source: String,
     /// Uniform values in the order the generated struct declares them.
     pub uniforms: Vec<f32>,
+    /// TRACES: FR-DEV-8
+    /// The instance list, if this pass declared one. See [`DetailPass::storage`].
+    pub storage: Vec<[f32; 4]>,
     /// See [`DetailPass::radius`].
     pub radius: u32,
     /// Whether this pass writes the display/export texture rather than another
@@ -509,6 +535,7 @@ pub fn compose_detail(
                     radius: 0,
                     wgsl: String::new(),
                     uniforms: Vec::new(),
+                    storage: Vec::new(),
                 },
                 0,
                 scale,
@@ -651,6 +678,10 @@ struct Params {{
 @group(0) @binding(0) var source: texture_2d<f32>;
 @group(0) @binding(1) var<uniform> u: Params;
 @group(0) @binding(2) var output: texture_storage_2d<{store_format}, write>;
+// A pass whose work is a list rather than a kernel reads it here; every other
+// pass leaves this bound to a single empty element and never looks at it. See
+// `DetailPass::storage` for why the list is not in the uniform block.
+@group(0) @binding(3) var<storage, read> instances: array<vec4<f32>>;
 
 // A neighbour, clamped to the edge of the image.
 //
@@ -681,6 +712,10 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {{
     // What this render is, relative to the export it has to match.
     let render_dims = u.detail_base.xy;
     let render_scale = u.detail_base.z;
+    // How many entries `instances` actually holds, read from the buffer itself
+    // rather than from a uniform so the two cannot disagree. A pass that
+    // declared no list is bound to a one-element placeholder and never asks.
+    let instance_count = arrayLength(&instances);
 
     var c = tap(coord, vec2<i32>(0));
     // One scalar per pixel that survives the hand-off from one pass to the
@@ -708,6 +743,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {{
         label,
         source,
         uniforms: uniform_values,
+        storage: pass.storage.clone(),
         radius: pass.radius,
         writes_output,
         structure_hash,

core/dr-pipeline/src/detail/probe.rs

@@ -141,6 +141,8 @@ impl DetailStage for BoxBlur {
             .map(|(axis, _)| DetailPass {
                 label: if axis == 0 { "horizontal" } else { "vertical" },
                 radius: r,
+                // A convolution, not a list: nothing to bind at binding 3.
+                storage: Vec::new(),
                 uniforms: vec![
                     Uniform {
                         name: "radius",

core/dr-pipeline/src/ops/capture_sharpen.rs

@@ -358,6 +358,8 @@ impl DetailStage for CaptureSharpen {
             .map(|label| DetailPass {
                 label,
                 radius: extent,
+                // A convolution, not a list: nothing to bind at binding 3.
+                storage: Vec::new(),
                 uniforms: vec![
                     Uniform {
                         // −100…100 as a gain around zero. A hundred percent is
@@ -410,6 +412,8 @@ fn nothing_to_sharpen() -> DetailPass {
         label: "unresolved",
         // Reads only the pixel it writes, so a tile needs no halo at all.
         radius: 0,
+        // A convolution, not a list: nothing to bind at binding 3.
+        storage: Vec::new(),
         uniforms: Vec::new(),
         wgsl: "// The chosen radius is finer than one pixel of this render, so the detail
 // it would act on is not in this texture — it was lost to the downscale

core/dr-pipeline/src/ops/local_contrast.rs

@@ -475,12 +475,16 @@ impl<B: Band> DetailStage for LocalContrast<B> {
             DetailPass {
                 label: "base",
                 radius,
+                // A convolution, not a list: nothing to bind at binding 3.
+                storage: Vec::new(),
                 uniforms: shape,
                 wgsl: BASE_X.to_string(),
             },
             DetailPass {
                 label: "combine",
                 radius,
+                // A convolution, not a list: nothing to bind at binding 3.
+                storage: Vec::new(),
                 uniforms: combine,
                 wgsl: combine_body(B::RECIPE.midtone_taper),
             },

core/dr-pipeline/src/ops/noise_reduction.rs

@@ -436,6 +436,8 @@ impl DetailStage for NoiseReduction {
             passes.push(DetailPass {
                 label: "luminance",
                 radius: luma,
+                // A convolution, not a list: nothing to bind at binding 3.
+                storage: Vec::new(),
                 uniforms: vec![
                     Uniform {
                         name: "radius",
@@ -474,6 +476,8 @@ impl DetailStage for NoiseReduction {
                         "chroma-vertical"
                     },
                     radius: chroma,
+                    // A convolution, not a list: nothing to bind at binding 3.
+                    storage: Vec::new(),
                     uniforms: vec![
                         Uniform {
                             name: "radius",