diff --git a/core/dr-gpu/src/segment.rs b/core/dr-gpu/src/segment.rs index e3a5953..fb4e1c6 100644 --- a/core/dr-gpu/src/segment.rs +++ b/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 jump = buffer_stage(9, 10, "jump", "watershed-jump-bgl"); + let plateau_init = buffer_stage(7, 8, "plateau_init", "watershed-pinit-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, ¶ms, 3, &feat_a, 4, &feat_b); let gradient_bg = self.bind(&self.gradient.layout, ¶ms, 5, &feat_b, 6, &gradient); - let flow_bg = self.bind(&self.flow.layout, ¶ms, 7, &gradient, 8, &parent_a); - let jump_ab = self.bind(&self.jump.layout, ¶ms, 9, &parent_a, 10, &parent_b); - let jump_ba = self.bind(&self.jump.layout, ¶ms, 9, &parent_b, 10, &parent_a); + let pinit_bg = self.bind(&self.plateau_init.layout, ¶ms, 7, &gradient, 8, &dist_a); + let pstep_ab = self.bind3( + &self.plateau_step.layout, + ¶ms, + (9, &gradient), + (10, &dist_a), + (11, &dist_b), + ); + let pstep_ba = self.bind3( + &self.plateau_step.layout, + ¶ms, + (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, + ¶ms, + (12, &gradient), + (13, final_dist), + (14, &parent_a), + ); + let jump_ab = self.bind(&self.jump.layout, ¶ms, 15, &parent_a, 16, &parent_b); + let jump_ba = self.bind(&self.jump.layout, ¶ms, 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 { + 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 { + 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 { + 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 diff --git a/core/dr-gpu/src/shaders/watershed.wgsl b/core/dr-gpu/src/shaders/watershed.wgsl index c9e8e22..cbd0bdb 100644 --- a/core/dr-gpu/src/shaders/watershed.wgsl +++ b/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 -// blur pre-smoothing, without which every noise grain becomes a basin -// gradient Sobel magnitude — the surface the watershed floods -// flow each pixel points downhill to its steepest neighbour -// jump pointer-jumping, until every pixel points at its basin root +// features source texture -> perceptual triple, downscaled to proxy size +// blur pre-smoothing, without which every grain becomes a basin +// gradient Sobel magnitude — the surface the watershed floods +// plateau_init seed the distance field at every real descent +// 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) { 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 pinit_grad: array; +@group(0) @binding(8) var pinit_out: array; + +// 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) { + 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 pstep_grad: array; +@group(0) @binding(10) var pstep_in: array; +@group(0) @binding(11) var pstep_out: array; + +// 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) { + 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 flow_grad: array; -@group(0) @binding(8) var flow_out: array; +@group(0) @binding(12) var flow_grad: array; +@group(0) @binding(13) var flow_dist: array; +@group(0) @binding(14) var flow_out: array; // 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) -// rather than value alone gives a strict total order, so the pointer graph -// descends monotonically and cannot contain a cycle — plateaux, which are -// everywhere in a smoothed image, would otherwise make two equal pixels point -// at each other and hang the pointer-jumping below. +// **The ordering is load-bearing, twice over.** Comparing on (gradient, +// plateau distance, index) rather than gradient alone gives a strict total +// order, so the pointer graph descends monotonically and cannot contain a +// cycle — plateaux, which are everywhere in a smoothed image, would otherwise +// 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) { 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) { } 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) { // -------------------------------------------------------------------- jump -@group(0) @binding(9) var jump_in: array; -@group(0) @binding(10) var jump_out: array; +@group(0) @binding(15) var jump_in: array; +@group(0) @binding(16) var jump_out: array; // Pointer jumping: parent = parent[parent]. // diff --git a/docs/segmentation.md b/docs/segmentation.md index eeca7c1..28725eb 100644 --- a/docs/segmentation.md +++ b/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