Merge branch 'worktree-watershed-plateaux'

core/dr-gpu/src/segment.rs

@@ -43,6 +43,14 @@ pub struct SegmentOptions {
     pub blur_radius: i32,
     pub w_luma: f32,
     pub w_chroma: f32,
+    /// How far the lower-completion carries a distance inward from a
+    /// plateau's rim, in breadth-first steps.
+    ///
+    /// Bounds the widest plateau that resolves fully. Beyond it, the interior
+    /// keeps the behaviour it had before the pass existed — a fan of diagonal
+    /// chains — so this trades dispatches against the size of flat area the
+    /// watershed handles cleanly, and never against correctness elsewhere.
+    pub plateau_iterations: u32,
 }
 
 impl Default for SegmentOptions {
@@ -58,6 +66,12 @@ impl Default for SegmentOptions {
             // boundaries anyone would draw, so it counts for less — but not
             // zero, or a red flower on green leaves has no edge at all.
             w_chroma: 0.5,
+            // **Zero: the pass is off.** It is implemented, dispatched
+            // correctly and measurably changes nothing — see the ignored test
+            // below and §12 of docs/segmentation.md. Until that is understood,
+            // running it would buy 64 dispatches per segmentation and no
+            // improvement, so the default declines to pay.
+            plateau_iterations: 0,
         }
     }
 }
@@ -87,6 +101,8 @@ pub struct SegmentPass {
     features: Stage,
     blur: Stage,
     gradient: Stage,
+    plateau_init: Stage,
+    plateau_step: Stage,
     flow: Stage,
     jump: Stage,
 }
@@ -144,10 +160,30 @@ impl SegmentPass {
             Stage { layout, pipeline }
         };
 
+        // Three bindings rather than two: these read the gradient *and* a
+        // distance field, and write a second one.
+        let triple_stage = |a: u32, b: u32, c: u32, entry: &str, label: &str| {
+            let layout = ctx
+                .device
+                .create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                    label: Some(label),
+                    entries: &[
+                        uniform_entry(0),
+                        storage_entry(a, true),
+                        storage_entry(b, true),
+                        storage_entry(c, false),
+                    ],
+                });
+            let pipeline = compute(ctx, &module, &layout, entry);
+            Stage { layout, pipeline }
+        };
+
         let blur = buffer_stage(3, 4, "blur", "watershed-blur-bgl");
         let gradient = buffer_stage(5, 6, "gradient", "watershed-gradient-bgl");
-        let flow = buffer_stage(7, 8, "flow", "watershed-flow-bgl");
+        let plateau_init = buffer_stage(7, 8, "plateau_init", "watershed-pinit-bgl");
-        let jump = buffer_stage(9, 10, "jump", "watershed-jump-bgl");
+        let plateau_step = triple_stage(9, 10, 11, "plateau_step", "watershed-pstep-bgl");
+        let flow = triple_stage(12, 13, 14, "flow", "watershed-flow-bgl");
+        let jump = buffer_stage(15, 16, "jump", "watershed-jump-bgl");
 
         if let Some(err) = pollster::block_on(scope.pop()) {
             return Err(GpuError::ShaderCompilation(err.to_string()));
@@ -158,6 +194,8 @@ impl SegmentPass {
             features,
             blur,
             gradient,
+            plateau_init,
+            plateau_step,
             flow,
             jump,
         })
@@ -197,6 +235,8 @@ impl SegmentPass {
         let feat_a = self.buffer("watershed-feat-a", n * 16, false);
         let feat_b = self.buffer("watershed-feat-b", n * 16, false);
         let gradient = self.buffer("watershed-gradient", n * 4, true);
+        let dist_a = self.buffer("watershed-dist-a", n * 4, false);
+        let dist_b = self.buffer("watershed-dist-b", n * 4, false);
         let parent_a = self.buffer("watershed-parent-a", n * 4, true);
         let parent_b = self.buffer("watershed-parent-b", n * 4, true);
 
@@ -233,9 +273,39 @@ impl SegmentPass {
 
         let blur_bg = self.bind(&self.blur.layout, &params, 3, &feat_a, 4, &feat_b);
         let gradient_bg = self.bind(&self.gradient.layout, &params, 5, &feat_b, 6, &gradient);
-        let flow_bg = self.bind(&self.flow.layout, &params, 7, &gradient, 8, &parent_a);
+        let pinit_bg = self.bind(&self.plateau_init.layout, &params, 7, &gradient, 8, &dist_a);
-        let jump_ab = self.bind(&self.jump.layout, &params, 9, &parent_a, 10, &parent_b);
+        let pstep_ab = self.bind3(
-        let jump_ba = self.bind(&self.jump.layout, &params, 9, &parent_b, 10, &parent_a);
+            &self.plateau_step.layout,
+            &params,
+            (9, &gradient),
+            (10, &dist_a),
+            (11, &dist_b),
+        );
+        let pstep_ba = self.bind3(
+            &self.plateau_step.layout,
+            &params,
+            (9, &gradient),
+            (10, &dist_b),
+            (11, &dist_a),
+        );
+
+        // An odd number of plateau steps leaves the distance field in B.
+        let plateau_steps = opts.plateau_iterations;
+        let final_dist = if plateau_steps % 2 == 0 {
+            &dist_a
+        } else {
+            &dist_b
+        };
+
+        let flow_bg = self.bind3(
+            &self.flow.layout,
+            &params,
+            (12, &gradient),
+            (13, final_dist),
+            (14, &parent_a),
+        );
+        let jump_ab = self.bind(&self.jump.layout, &params, 15, &parent_a, 16, &parent_b);
+        let jump_ba = self.bind(&self.jump.layout, &params, 15, &parent_b, 16, &parent_a);
 
         // Pointer jumping halves every path per pass, so log2 of the pixel
         // count bounds it — that is the longest possible descent chain. A
@@ -253,13 +323,24 @@ impl SegmentPass {
                 (&self.features.pipeline, &features_bg),
                 (&self.blur.pipeline, &blur_bg),
                 (&self.gradient.pipeline, &gradient_bg),
-                (&self.flow.pipeline, &flow_bg),
+                (&self.plateau_init.pipeline, &pinit_bg),
             ] {
                 pass.set_pipeline(pipeline);
                 pass.set_bind_group(0, bg, &[]);
                 pass.dispatch_workgroups(groups.0, groups.1, 1);
             }
 
+            pass.set_pipeline(&self.plateau_step.pipeline);
+            for i in 0..plateau_steps {
+                let bg = if i % 2 == 0 { &pstep_ab } else { &pstep_ba };
+                pass.set_bind_group(0, bg, &[]);
+                pass.dispatch_workgroups(groups.0, groups.1, 1);
+            }
+
+            pass.set_pipeline(&self.flow.pipeline);
+            pass.set_bind_group(0, &flow_bg, &[]);
+            pass.dispatch_workgroups(groups.0, groups.1, 1);
+
             pass.set_pipeline(&self.jump.pipeline);
             for i in 0..jumps {
                 let bg = if i % 2 == 0 { &jump_ab } else { &jump_ba };
@@ -295,6 +376,40 @@ impl SegmentPass {
         })
     }
 
+    fn bind3(
+        &self,
+        layout: &wgpu::BindGroupLayout,
+        params: &wgpu::Buffer,
+        a: (u32, &wgpu::Buffer),
+        b: (u32, &wgpu::Buffer),
+        c: (u32, &wgpu::Buffer),
+    ) -> wgpu::BindGroup {
+        self.ctx
+            .device
+            .create_bind_group(&wgpu::BindGroupDescriptor {
+                label: Some("watershed-bg3"),
+                layout,
+                entries: &[
+                    wgpu::BindGroupEntry {
+                        binding: 0,
+                        resource: params.as_entire_binding(),
+                    },
+                    wgpu::BindGroupEntry {
+                        binding: a.0,
+                        resource: a.1.as_entire_binding(),
+                    },
+                    wgpu::BindGroupEntry {
+                        binding: b.0,
+                        resource: b.1.as_entire_binding(),
+                    },
+                    wgpu::BindGroupEntry {
+                        binding: c.0,
+                        resource: c.1.as_entire_binding(),
+                    },
+                ],
+            })
+    }
+
     fn bind(
         &self,
         layout: &wgpu::BindGroupLayout,
@@ -556,6 +671,186 @@ mod tests {
         );
     }
 
+    /// A flat disc on flat ground: two plateaux and one boundary between
+    /// them. Nothing here has a downhill direction except at the rim.
+    /// A linear ramp between two flat fields.
+    ///
+    /// The watershed runs on gradient *magnitude*, and that changes which
+    /// images contain a plateau worth resolving. A flat region of the picture
+    /// has gradient zero — the global minimum — and a plateau at the minimum
+    /// has no descending exit at all, which makes it a single basin by
+    /// definition with nothing for lower-completion to do. The plateaux that
+    /// do have an exit are regions of constant *non-zero* gradient: linear
+    /// ramps. So that is what this builds.
+    fn ramp(w: u32, h: u32) -> Vec<u8> {
+        let mut px = vec![0u8; (w * h * 4) as usize];
+        let (lo, hi) = (w / 4, w - w / 4);
+        for y in 0..h {
+            for x in 0..w {
+                let v = if x < lo {
+                    40u8
+                } else if x >= hi {
+                    210u8
+                } else {
+                    // Constant slope, so the gradient is constant and
+                    // non-zero across the whole band.
+                    (40.0 + (x - lo) as f32 * (170.0 / (hi - lo) as f32)) as u8
+                };
+                let i = ((y * w + x) * 4) as usize;
+                px[i] = v;
+                px[i + 1] = v;
+                px[i + 2] = v;
+                px[i + 3] = 255;
+            }
+        }
+        px
+    }
+
+    /// A terraced disc: dark ground, a mid-level annulus, a bright centre.
+    ///
+    /// The annulus is the point. It is a wide plateau that *has* a descending
+    /// exit — the ground outside it — which is the only situation
+    /// lower-completion is defined for. `disc` below has only flat regions at
+    /// the gradient's global minimum, and a plateau with no exit at all is a
+    /// minimum: one basin by definition, with nothing to resolve.
+    fn terrace(w: u32, h: u32) -> Vec<u8> {
+        let mut px = vec![0u8; (w * h * 4) as usize];
+        let (cx, cy) = (w as f32 / 2.0, h as f32 / 2.0);
+        for y in 0..h {
+            for x in 0..w {
+                let d = ((x as f32 - cx).powi(2) + (y as f32 - cy).powi(2)).sqrt();
+                let v = if d < w as f32 * 0.16 {
+                    200
+                } else if d < w as f32 * 0.40 {
+                    130
+                } else {
+                    60
+                };
+                let i = ((y * w + x) * 4) as usize;
+                px[i] = v;
+                px[i + 1] = v;
+                px[i + 2] = v;
+                px[i + 3] = 255;
+            }
+        }
+        px
+    }
+
+    fn disc(w: u32, h: u32) -> Vec<u8> {
+        let mut px = vec![0u8; (w * h * 4) as usize];
+        let (cx, cy) = (w as f32 / 2.0, h as f32 / 2.0);
+        for y in 0..h {
+            for x in 0..w {
+                let d = ((x as f32 - cx).powi(2) + (y as f32 - cy).powi(2)).sqrt();
+                let v = if d < w as f32 * 0.3 { 200 } else { 60 };
+                let i = ((y * w + x) * 4) as usize;
+                px[i] = v;
+                px[i + 1] = v;
+                px[i + 2] = v;
+                px[i + 3] = 255;
+            }
+        }
+        px
+    }
+
+    #[test]
+    #[ignore = "the plateau pass is a measured no-op; see docs/segmentation.md §12"]
+    fn lower_completion_drains_a_plateau_instead_of_shattering_it() {
+        // F1, asserted rather than eyeballed, and asserted at the level where
+        // it matters.
+        //
+        // Two claims, because they are different claims. First: carrying the
+        // distance inward genuinely reduces fragmentation — a plateau with an
+        // exit now drains to it instead of fanning into diagonal chains.
+        // Second, and the one a user would notice: whatever fragments survive
+        // are separated by zero-height saddles, so the hierarchy merges them
+        // at its very first steps and the plateau reads as one region.
+        //
+        // The second claim is what makes the first one's *residue* tolerable.
+        // A perfectly flat regional minimum — the inside of a uniform disc,
+        // with no exit anywhere — cannot be drained by a distance that has
+        // nowhere to descend to, and collapsing it fully would need connected
+        // component labelling rather than a local rule. It is not worth it:
+        // see docs/segmentation.md §12.
+        let Some(ctx) = ctx() else { return };
+        let (w, h) = (96u32, 96u32);
+        let src = DemosaicedImage::from_rgba8(&ctx, &ramp(w, h), w, h).expect("source");
+        let pass = SegmentPass::new(&ctx).expect("segment pass");
+
+        let labels_in = |labels: &[u32], inside: bool| {
+            let (cx, cy) = (w as f32 / 2.0, h as f32 / 2.0);
+            let mut seen = std::collections::HashSet::new();
+            for y in 0..h {
+                for x in 0..w {
+                    let d = ((x as f32 - cx).powi(2) + (y as f32 - cy).powi(2)).sqrt();
+                    let take = if inside {
+                        d < w as f32 * 0.20
+                    } else {
+                        d > w as f32 * 0.42
+                    };
+                    if take {
+                        seen.insert(labels[(y * w + x) as usize]);
+                    }
+                }
+            }
+            seen
+        };
+
+        let field = |iterations: u32| {
+            pass.run(
+                &src,
+                SegmentOptions {
+                    plateau_iterations: iterations,
+                    ..Default::default()
+                },
+            )
+            .expect("run")
+            .read_field()
+            .expect("field")
+        };
+
+        let shallow = field(1);
+        let deep = field(64);
+
+        // Before anything else: does the pass change the labelling at all? If
+        // the distance field were never populated — a binding astray, a level
+        // test that never matches — every downstream claim would be excused
+        // by a no-op rather than tested. This is the one assertion that
+        // cannot pass vacuously.
+        let differs = shallow
+            .labels
+            .iter()
+            .zip(deep.labels.iter())
+            .filter(|(a, b)| a != b)
+            .count();
+        assert!(
+            differs > 0,
+            "the plateau distance changed no pixel's basin, so the pass is a \
+             no-op: {} pixels, {} differ",
+            shallow.labels.len(),
+            differs
+        );
+
+        // Claim one: fewer basins, because plateaux with an exit now use it.
+        assert!(
+            deep.region_count < shallow.region_count,
+            "carrying the distance inward should reduce fragmentation: \
+             {} basins against {}",
+            deep.region_count,
+            shallow.region_count
+        );
+
+        // Claim two: what survives costs nothing, because the hierarchy
+        // dissolves it immediately.
+        let tree = MergeTree::build(&deep);
+        let grouped = deep.apply(&tree.cut_to(2));
+        let inside = labels_in(&grouped, true);
+        let outside = labels_in(&grouped, false);
+        assert_eq!(inside.len(), 1, "the disc should read as one region");
+        assert_eq!(outside.len(), 1, "the ground should read as one region");
+        assert_ne!(inside, outside, "and they must not be the same region");
+    }
+
     #[test]
     fn the_same_image_segments_identically_twice() {
         // M5 on one device — the weaker half of the determinism question, but

core/dr-gpu/src/shaders/watershed.wgsl

@@ -1,12 +1,14 @@
 // Watershed segmentation — the passes behind arm A of S15 (docs/segmentation.md).
 //
-// Five entry points forming one chain:
+// Seven entry points forming one chain:
 //
-//   features  source texture -> perceptual triple, box-downscaled to proxy size
+//   features      source texture -> perceptual triple, downscaled to proxy size
-//   blur      pre-smoothing, without which every noise grain becomes a basin
+//   blur          pre-smoothing, without which every grain becomes a basin
-//   gradient  Sobel magnitude — the surface the watershed floods
+//   gradient      Sobel magnitude — the surface the watershed floods
-//   flow      each pixel points downhill to its steepest neighbour
+//   plateau_init  seed the distance field at every real descent
-//   jump      pointer-jumping, until every pixel points at its basin root
+//   plateau_step  carry it inward, so flat ground drains toward its exit
+//   flow          each pixel points downhill to its steepest neighbour
+//   jump          pointer-jumping, until every pixel points at its basin root
 //
 // Everything after `features` works in storage buffers rather than textures.
 // That is deliberate: the flow and jump passes need read-write access to the
@@ -184,19 +186,162 @@ fn gradient(@builtin(global_invocation_id) gid: vec3<u32>) {
     grad_out[gid.y * u.width + gid.x] = sqrt(dot(wx, wx) + dot(wy, wy));
 }
 
+// ---------------------------------------------------------- lower-complete
+//
+// A watershed needs every non-minimum pixel to have a lower neighbour. A real
+// gradient does not oblige: a flat wall, a clipped sky or the inside of a
+// uniform object is a **plateau**, where every neighbour is exactly equal and
+// there is no downhill direction to follow.
+//
+// Left alone, the tie-break in `flow` sends every plateau pixel to its
+// lowest-indexed neighbour, which is up and to the left. Each pixel therefore
+// walks diagonally until it falls off the plateau, and one flat region becomes
+// a fan of diagonal chains rather than one basin — visible as hatching across
+// what should be a single area (docs/segmentation.md §12, F1).
+//
+// The fix is the standard lower-completion: give each plateau pixel its
+// geodesic distance to the nearest pixel that *does* have a lower neighbour,
+// then let `flow` order on (gradient, distance). Water on a plateau now runs
+// toward the plateau's exit, which is what it would physically do.
+//
+// A plateau with no exit at all is a genuine regional minimum — the inside of
+// a uniform disc, say. Those pixels keep `PLATEAU_UNRESOLVED`, tie with each
+// other, and fall through to the index tie-break, which collapses the whole
+// connected plateau onto its lowest-indexed pixel. One basin, which is the
+// right answer for a regional minimum.
+
+const PLATEAU_UNRESOLVED: u32 = 0xffffffffu;
+
+// How close two gradients must be to count as the same level.
+//
+// **Exact equality does not work here, and that is not a rounding nicety.**
+// The gradient is a float computed from 8-bit samples, so a region the eye
+// and the algorithm both consider flat still has neighbours differing in the
+// sixth decimal. With `==`, the breadth-first step never advances past its
+// seeds and the whole pass is a no-op; with `<`, nearly every pixel finds
+// some marginally lower neighbour and is seeded at zero, which is the same
+// no-op wearing a different hat. Both were measured before this constant
+// existed.
+//
+// Sized against the gradient's own scale: features are normalised to 0..1, so
+// a Sobel magnitude runs to a few units, and 1e-4 is far below any step a
+// real edge produces while sitting comfortably above f32 noise from a blur.
+const LEVEL_EPS: f32 = 1e-4;
+
+// Whether `b` lies below `a` by more than the level tolerance.
+fn strictly_below(b: f32, a: f32) -> bool {
+    return b < a - LEVEL_EPS;
+}
+
+// Whether two gradients belong to the same plateau.
+fn same_level(a: f32, b: f32) -> bool {
+    return abs(a - b) <= LEVEL_EPS;
+}
+
+@group(0) @binding(7) var<storage, read> pinit_grad: array<f32>;
+@group(0) @binding(8) var<storage, read_write> pinit_out: array<u32>;
+
+// Seed the distance field: zero where a real descent exists, unresolved on a
+// plateau.
+@compute @workgroup_size(8, 8, 1)
+fn plateau_init(@builtin(global_invocation_id) gid: vec3<u32>) {
+    if (gid.x >= u.width || gid.y >= u.height) {
+        return;
+    }
+    let idx = gid.y * u.width + gid.x;
+    let here = pinit_grad[idx];
+
+    for (var dy = -1; dy <= 1; dy = dy + 1) {
+        for (var dx = -1; dx <= 1; dx = dx + 1) {
+            if (dx == 0 && dy == 0) {
+                continue;
+            }
+            let nx = i32(gid.x) + dx;
+            let ny = i32(gid.y) + dy;
+            if (nx < 0 || ny < 0 || nx >= i32(u.width) || ny >= i32(u.height)) {
+                continue;
+            }
+            if (strictly_below(pinit_grad[u32(ny) * u.width + u32(nx)], here)) {
+                pinit_out[idx] = 0u;
+                return;
+            }
+        }
+    }
+
+    pinit_out[idx] = PLATEAU_UNRESOLVED;
+}
+
+@group(0) @binding(9) var<storage, read> pstep_grad: array<f32>;
+@group(0) @binding(10) var<storage, read> pstep_in: array<u32>;
+@group(0) @binding(11) var<storage, read_write> pstep_out: array<u32>;
+
+// One breadth-first step inward from the plateau's rim.
+//
+// Iterated by the host a fixed number of times rather than to convergence: a
+// convergence test costs a readback per pass, and the count only has to cover
+// the widest plateau in the frame. Pixels still unresolved when the budget
+// runs out keep `PLATEAU_UNRESOLVED` and behave exactly as they did before
+// this pass existed — the degradation is graceful, not a wrong answer.
+@compute @workgroup_size(8, 8, 1)
+fn plateau_step(@builtin(global_invocation_id) gid: vec3<u32>) {
+    if (gid.x >= u.width || gid.y >= u.height) {
+        return;
+    }
+    let idx = gid.y * u.width + gid.x;
+    let current = pstep_in[idx];
+    if (current != PLATEAU_UNRESOLVED) {
+        pstep_out[idx] = current;
+        return;
+    }
+
+    let here = pstep_grad[idx];
+    var best = PLATEAU_UNRESOLVED;
+    for (var dy = -1; dy <= 1; dy = dy + 1) {
+        for (var dx = -1; dx <= 1; dx = dx + 1) {
+            if (dx == 0 && dy == 0) {
+                continue;
+            }
+            let nx = i32(gid.x) + dx;
+            let ny = i32(gid.y) + dy;
+            if (nx < 0 || ny < 0 || nx >= i32(u.width) || ny >= i32(u.height)) {
+                continue;
+            }
+            let ni = u32(ny) * u.width + u32(nx);
+            // Only within the same plateau: a neighbour at a different height
+            // is across a boundary, and its distance says nothing about the
+            // way out of this one.
+            if (!same_level(pstep_grad[ni], here)) {
+                continue;
+            }
+            let nd = pstep_in[ni];
+            if (nd != PLATEAU_UNRESOLVED && nd < best) {
+                best = nd;
+            }
+        }
+    }
+
+    if (best == PLATEAU_UNRESOLVED) {
+        pstep_out[idx] = PLATEAU_UNRESOLVED;
+    } else {
+        pstep_out[idx] = best + 1u;
+    }
+}
+
 // -------------------------------------------------------------------- flow
 
-@group(0) @binding(7) var<storage, read> flow_grad: array<f32>;
+@group(0) @binding(12) var<storage, read> flow_grad: array<f32>;
-@group(0) @binding(8) var<storage, read_write> flow_out: array<u32>;
+@group(0) @binding(13) var<storage, read> flow_dist: array<u32>;
+@group(0) @binding(14) var<storage, read_write> flow_out: array<u32>;
 
 // Each pixel points at the steepest-descent neighbour among its 8, or at
 // itself if it is a local minimum — a basin seed.
 //
-// **The tie-break is load-bearing, twice over.** Comparing on (value, index)
+// **The ordering is load-bearing, twice over.** Comparing on (gradient,
-// rather than value alone gives a strict total order, so the pointer graph
+// plateau distance, index) rather than gradient alone gives a strict total
-// descends monotonically and cannot contain a cycle — plateaux, which are
+// order, so the pointer graph descends monotonically and cannot contain a
-// everywhere in a smoothed image, would otherwise make two equal pixels point
+// cycle — plateaux, which are everywhere in a smoothed image, would otherwise
-// at each other and hang the pointer-jumping below.
+// make two equal pixels point at each other and hang the pointer-jumping
+// below.
 //
 // It is also what makes the result reproducible. S15's M5 asks whether a
 // label field is stable enough across GPU vendors to be a cache key
@@ -210,6 +355,7 @@ fn flow(@builtin(global_invocation_id) gid: vec3<u32>) {
     let idx = gid.y * u.width + gid.x;
 
     var best_val = flow_grad[idx];
+    var best_dist = flow_dist[idx];
     var best_idx = idx;
 
     for (var dy = -1; dy <= 1; dy = dy + 1) {
@@ -224,8 +370,32 @@ fn flow(@builtin(global_invocation_id) gid: vec3<u32>) {
             }
             let ni = u32(ny) * u.width + u32(nx);
             let nv = flow_grad[ni];
-            if (nv < best_val || (nv == best_val && ni < best_idx)) {
+            let nd = flow_dist[ni];
+
+            // Lexicographic on (gradient, plateau distance, index) rather
+            // than one fused scalar. Folding the distance into the gradient
+            // as a small epsilon would need a scale factor that is small
+            // enough never to cross a real gradient step and large enough to
+            // survive f32 — a tuning problem with a silent failure mode,
+            // where three explicit keys have neither.
+            // The same tolerance the plateau passes use, and for the same
+            // reason: with exact equality this tie never fires on real data,
+            // so the distance carried inward above would be computed and then
+            // never consulted — the pass measurably did nothing.
+            var better = false;
+            if (strictly_below(nv, best_val)) {
+                better = true;
+            } else if (same_level(nv, best_val)) {
+                if (nd < best_dist) {
+                    better = true;
+                } else if (nd == best_dist && ni < best_idx) {
+                    better = true;
+                }
+            }
+
+            if (better) {
                 best_val = nv;
+                best_dist = nd;
                 best_idx = ni;
             }
         }
@@ -236,8 +406,8 @@ fn flow(@builtin(global_invocation_id) gid: vec3<u32>) {
 
 // -------------------------------------------------------------------- jump
 
-@group(0) @binding(9) var<storage, read> jump_in: array<u32>;
+@group(0) @binding(15) var<storage, read> jump_in: array<u32>;
-@group(0) @binding(10) var<storage, read_write> jump_out: array<u32>;
+@group(0) @binding(16) var<storage, read_write> jump_out: array<u32>;
 
 // Pointer jumping: parent = parent[parent].
 //

docs/segmentation.md

@@ -303,6 +303,30 @@ up an artefact of the flow pass is fragile. The principled fix is a **lower-comp
 extra pass giving plateau pixels a gradient toward their nearest descending exit. Standard, cheap,
 and worth doing before the corpus work.
 
+*Attempted, and parked.* The pass exists — `plateau_init` seeds every pixel
+that has a strictly lower neighbour, `plateau_step` carries a breadth-first
+distance inward within a level set, and `flow` takes that distance as the
+second key of a lexicographic tie-break. Bindings, ping-pong and dispatch were
+all checked and are right. It is nonetheless a **measured no-op**: with a test
+comparing the labelling at one iteration against sixty-four, *zero* of 9216
+pixels change basin. That test is committed and ignored rather than deleted,
+because it is the thing that turned "we think this works" into a fact.
+
+Three explanations were tried and none of them was it. Exact float equality is
+certainly wrong — a gradient computed from 8-bit samples is never exactly
+equal across a region the eye calls flat — and a `LEVEL_EPS` tolerance now
+replaces `==` and `<` in all three comparisons; it did not change the outcome.
+Nor did the test image: a flat disc, a terraced disc and a constant-slope ramp
+all behave identically. Worth knowing for whoever picks this up: on a
+gradient-*magnitude* watershed, every flat region of the picture sits at
+gradient zero, which is the global minimum, and a plateau with no descending
+exit is a minimum — one basin by definition, with nothing for lower-completion
+to resolve. The plateaux that do have an exit are regions of constant non-zero
+gradient, which are rarer in a photograph than F1's phrasing suggests.
+
+`plateau_iterations` therefore defaults to **0**. The pass is off, costs
+nothing, and F1 stands open.
+
 **F2 — `cut_to(N)` is a visualisation, not the interaction.** A global cut by region count spends its
 budget wherever the saddles happen to be densest: at blur 5 the soft-edged disc's interior held a
 cluster of near-equal saddles and ate the budget, fragmenting at a level where everything else was