refinement_stepper: compatibility views are kept on purpose (comment only)

The views and the PeriodicPatchStepper alias stay for the existing tests
and docs (about 130 references); the comment no longer says REMOVE AFTER.
No code change.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

README.md

@@ -112,8 +112,9 @@ from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
 cfg = LBMConfig(nx=32, ny=24, viscosity=.1)
 body = SimObject(0, geometry=CircleGeometry(16.0, 12.0, 2.15), center=(16.0, 12.0), radius=2.15,
                  surface_actuation=SurfaceActuation(0.7, 0.85, 0.002))  # angle, half-width, speed
-with CudaContext(0) as ctx, FixedPatchStepper(cfg, (8, 6, 24, 18), context=ctx,
-                                              fine_objects=[body], fine_block=(16, 20)) as run:
+with CudaContext(0) as ctx, FixedPatchStepper(cfg, context=ctx, refinements=[(8, 6, 24, 18)],
+                                              level_blocks=[None, (16, 20)],
+                                              level_objects=[(), (body,)]) as run:
     rest = np.asarray(compute_feq_d2q9(1.0, 0.0, 0.0), np.float32)[:, None, None]
     # (9, ny, nx) postcollision populations per level; non-leaf cells are discarded.
     run.initialize(np.broadcast_to(rest, (9, *run.patch.coarse_shape)),
@@ -124,8 +125,8 @@ with CudaContext(0) as ctx, FixedPatchStepper(cfg, (8, 6, 24, 18), context=ctx,
 
 Scope and limits:
 - The examples below use D2Q9, SRT, FP32, double buffer, unshifted, no LES or
-  forcing. Fixed 2:1 hierarchies can use level lists of arbitrary depth; the
-  legacy two/three-level forms are shown below.
+  forcing. Fixed 2:1 hierarchies use level lists of arbitrary depth; two- and
+  three-level examples below use that same constructor.
 - `coarse_boundary="periodic"` (default) or `"slot_reservoir"` (bottom slot, top
   reservoir; patch needs `y0 >= 4`, `y1 <= ny-4`). Periodic mode ignores the wall/slot config fields.
 - Object centers/radii use the coarse **cell-edge** frame (cell `I` spans `[I, I+1)`),
@@ -133,50 +134,51 @@ Scope and limits:
 - Bodies must be stationary circles whose whole boundary stencil lies in the fine
   active region, clear of the coarse/fine transfer cells; violations are rejected
   before any GPU allocation.
-- `fine_block=(by, bx)` stores the fine level in fixed-size blocks; `by`/`bx` must
+- `level_blocks=[None, (by, bx)]` stores the fine level in fixed-size blocks; `by`/`bx` must
   divide the WHOLE fine allocation, i.e. `2*(y1-y0)+8` by `2*(x1-x0)+8` including the
   four-cell ghost band. It gives results identical to dense storage and is
   infrastructure for future block AMR, not a speedup.
 - Several regions: pass a sequence of `(x0, y0, x1, y1)` rectangles as
-  `bounds`, with both `fine_block` and `base_block`. Transfers are built once
+  `refinements[0]`, with both levels blocked (`level_blocks`). Transfers are built once
   from their union; adjacent, overlapping and concave unions are accepted when
   their actual dependencies close. There is no fixed rectangle-gap rule.
   The fine level is then one level-wide grid `(2*ny, 2*nx)`:
   `run.patch.fine_shape` gives it, `initialize`/`download` use it, and only the
   grid-aligned tiles intersecting the union plus its four-cell ghost band are stored.
   Block sizes must divide the level-wide grid; rectangle edges need not align
-  with tiles. Order of rectangles does not matter. `fine_objects`/`jets` work
+  with tiles. Order of rectangles does not matter. L1 bodies/Jets (`level_objects[1]`/`level_jets[1]`) work
   as for one patch: each circle (same coarse cell-edge frame) must lie wholly
   inside the union's active region with
   its boundary stencil clear of transfer cells; a jet may not select links of two
   different bodies, even across regions. Wall-touching patches remain unsupported.
-  For nesting, use the level lists or legacy child forms below.
+  For nesting, add further `refinements` entries as shown below.
 
   ```python
   cfg = LBMConfig(nx=32, ny=16, viscosity=.1)
   rest = np.asarray(compute_feq_d2q9(1.0, 0.0, 0.0), np.float32)[:, None, None]
   with CudaContext(0) as ctx, FixedPatchStepper(
-          cfg, ((4, 4, 12, 12), (16, 4, 24, 12)), context=ctx,
-          fine_block=(4, 4), base_block=(4, 4)) as run:
+          cfg, context=ctx, refinements=[((4, 4, 12, 12), (16, 4, 24, 12))],
+          level_blocks=[(4, 4), (4, 4)]) as run:
       run.initialize(np.broadcast_to(rest, (9, *run.patch.coarse_shape)),
                      np.broadcast_to(rest, (9, *run.patch.fine_shape)))  # (9, 32, 64)
       run.step(32)
       coarse_g, fine_g = run.download()  # coarse: NaN on every covered cell;
                                          # fine: NaN on every non-ACTIVE cell
   ```
-- Legacy three-level scalar chain: add `child=(x0, y0, x1, y1)` in
-  **L1 window cells** (the patch's fine grid including its 4-cell ghost band) and
-  `child_block`, with `fine_block` and `base_block`. L2 must sit inside L1's
-  ACTIVE cells, away from L1's ghost band and transfer footprint (checked before
-  any GPU allocation). `initialize` and `download` take/return three arrays; each
-  level is NaN on its non-ACTIVE (and child-covered) cells. One coarse step = 2 L1
-  steps = 4 L2 steps. One L1 patch only. Optional stationary circles and Jets
-  live on L2 via `child_objects` and `child_jets`; `fine_objects`/`jets` retain
-  their L1 meaning and must be empty in this mode. All object and Jet geometry
-  uses the **L0 edge frame**, not the child bounds' L1 window frame. L2 has
-  dx=dt=0.25, so prescribed lattice velocity is unchanged. Each whole body and
-  its BFL stencil must be inside L2 ACTIVE cells, outside the transfer footprint.
-  A Jet may select one physical body across multiple blocks, never two bodies.
+- Three-level nested windows: pass two scalar rectangles in `refinements` —
+  L1 in L0 cells, L2 in **L1 window cells** (the parent's fine grid including its
+  4-cell ghost band) — plus all three `level_blocks` entries. L2 must sit inside
+  L1's ACTIVE cells, away from L1's ghost band and transfer footprint (checked
+  before any GPU allocation). `initialize` and `download` take/return three
+  arrays; each level is NaN on its non-ACTIVE (and child-covered) cells. One
+  coarse step = 2 L1 steps = 4 L2 steps. One L1 window only in this scalar
+  chain. Optional stationary circles and Jets live on L2 via `level_objects[2]`
+  and `level_jets[2]`; L1 bodies/Jets retain their L1 meaning and must be empty
+  in this mode. All object and Jet geometry uses the **L0 edge frame**, not the
+  child bounds' L1 window frame. L2 has dx=dt=0.25, so prescribed lattice
+  velocity is unchanged. Each whole body and its BFL stencil must be inside L2
+  ACTIVE cells, outside the transfer footprint. A Jet may select one physical
+  body across multiple blocks, never two bodies.
 
   ```python
   from CelerisLab.body.geometry.circle import CircleGeometry
@@ -188,9 +190,9 @@ Scope and limits:
   jet = Jet(Tube((6.5, 8.), 0., 3., .3), Parallel(0.), amplitude=.002)
   rest = np.asarray(compute_feq_d2q9(1.0, 0.0, 0.0), np.float32)[:, None, None]
   with CudaContext(0) as ctx, FixedPatchStepper(
-          cfg, (4, 4, 12, 12), context=ctx, fine_block=(4, 4), base_block=(4, 4),
-          child=(8, 8, 16, 16), child_block=(4, 4),
-          child_objects=[body], child_jets=[jet]) as run:
+          cfg, context=ctx, refinements=[(4, 4, 12, 12), (8, 8, 16, 16)],
+          level_blocks=[(4, 4)] * 3, level_objects=[(), (), (body,)],
+          level_jets=[(), (), (jet,)]) as run:
       shapes = (run.patch.coarse_shape, run.patch.fine_shape, run.child_patch.fine_shape)
       run.initialize(*(np.broadcast_to(rest, (9, *s)) for s in shapes))  # 16², 24², 24²
       run.step(4)
@@ -198,16 +200,16 @@ Scope and limits:
       run.step(4)
       coarse_g, l1_g, l2_g = run.download()  # solids/covered/nonleaves are NaN
   ```
-- Level lists: `refinements[l]` describes level `l+1` in allocated level-`l`
-  integer cells; use a rectangle list for a union, or a scalar rectangle for a
-  single window (a scalar child requires a scalar parent). `level_blocks`,
-  `level_objects` and `level_jets` each have one entry per level, including L0.
-  These lists are mutually exclusive with the legacy patch keywords. Three or
-  more levels require all-blocked storage and bodies/Jets only on the finest
-  level. Every adjacent pair must preserve inherited ACTIVE coverage, valid
-  donors/receivers, surviving pulls and the upstream transfer footprint.
-  Level `l` takes `2**l` steps per base interval. This describes implemented
-  topology and scheduling, not accuracy or convergence qualification.
+- Level lists (the only constructor): `refinements[l]` describes level `l+1`
+  in allocated level-`l` integer cells; use a rectangle list for a union, or a
+  scalar rectangle for a single window (a scalar child requires a scalar
+  parent). `level_blocks`, `level_objects` and `level_jets` each have one entry
+  per level, including L0. Three or more levels require all-blocked storage and
+  bodies/Jets only on the finest level. Every adjacent pair must preserve
+  inherited ACTIVE coverage, valid donors/receivers, surviving pulls and the
+  upstream transfer footprint. Level `l` takes `2**l` steps per base interval.
+  This describes implemented topology and scheduling, not accuracy or
+  convergence qualification.
 - Blocked body levels can enable `observe_raw_exchange=True` and call
   `read_raw_exchange(body_id, level=...)` after a completed base interval.
   It reduces unique cut-link observations into wall/aperture/total force,

docs/compatibility.md

@@ -0,0 +1,35 @@
+# Compatibility and API changes
+
+Changes that can make a previously accepted setup or call behave differently.
+Each entry names the commit, the affected entry points and the observable result.
+
+## 87c3a31: cut-link ownership is validated for every body
+
+`validate_cut_link_ownership` (`body/preprocess/flag_overlay.py`) now runs for
+passive bodies as well as Jet/aperture bodies. Every cut link must have an
+active fluid receiver (FLUID set, SOLID clear), and each `(receiver, direction)`
+population may have only one writer.
+
+Rejected with `ValueError`, before flags, populations or device buffers change:
+
+1. Overlapping or touching passive solids whose cut links write the same
+   population ("Boundary requires unique cut-link writers").
+2. A body whose cut-link receiver is a non-fluid cell, e.g. a channel wall row
+   or the inlet/outlet column ("Cut links must have active fluid receivers").
+
+Both were previously accepted silently for passive bodies; the interpolated
+boundary then had undefined results. The rules apply to the cut links of all
+bodies together, so a passive body overlapping a Jet body is rejected as well.
+The check runs in `Simulation.initialize`, runtime body sync (`sync_bodies`),
+checkpoint save and checkpoint load (`common/checkpoint.py`). Load validates
+before restoring anything, so a checkpoint whose saved geometry violates either
+rule no longer restores; rebuild the geometry with separated bodies.
+
+## 87c3a31: `Simulation.read_surface_exchange` accepts Jet-driven bodies
+
+`read_surface_exchange(id)` previously required `SurfaceActuation` on the body.
+It now also accepts bodies that own Jet aperture links, including Jets with zero
+or negative amplitude (the aperture test is geometric). Bodies with neither
+`SurfaceActuation` nor aperture links still raise
+`ValueError("Body has no surface actuation")`. The value is unchanged: the latest
+step's wall/aperture exchange, not total nozzle force.

docs/development/root_channel_candidate.md

@@ -1,6 +1,6 @@
 # Blocked AMR root channel increment
 
-Opt in with `FixedPatchStepper(..., coarse_boundary="channel", base_block=(by,bx), fine_block=(fy,fx))`, or the existing level-list form. Every level must be blocked. The supported boundary is the existing uniform D2Q9/SRT/FP32 double-buffer channel: local Zou-He west inlet, pressure-NEQ east outlet, stationary halfway bounce-back y walls. Uniform and parabolic inlet profiles and shifted/unshifted storage are supported; the existing configuration guards still reject MRT/LES and incompatible storage. This does not change the external DFG MRT reference.
+Opt in with `FixedPatchStepper(..., coarse_boundary="channel", refinements=[bounds], level_blocks=[(by,bx), (fy,fx)])`. Every level must be blocked. The supported boundary is the existing uniform D2Q9/SRT/FP32 double-buffer channel: local Zou-He west inlet, pressure-NEQ east outlet, stationary halfway bounce-back y walls. Uniform and parabolic inlet profiles and shifted/unshifted storage are supported; the existing configuration guards still reject MRT/LES and incompatible storage. This does not change the external DFG MRT reference.
 
 `step/block_channel.cu` is a small root-only kernel. Fine kernels, transfer kernels and recursive hop scheduling are unchanged. Global coordinates and the global profile height/velocity are separate from block-local storage strides. Corner ownership and collision phases match `OneStep`:
 

docs/fixed_three_level_multipatch.md

@@ -1,8 +1,8 @@
 # Fixed three-level rectangle unions
 
-This page describes the legacy three-level entry point of `FixedPatchStepper`,
-with rectangle unions on either adjacent pair. The `refinements` level-list
-entry point also supports arbitrary depth; three or more levels require
+This page describes three-level `FixedPatchStepper` setups built from the
+`refinements` level list, with rectangle unions on either adjacent pair. The
+same constructor supports arbitrary depth; three or more levels require
 all-blocked storage and bodies/Jets only on the finest level. This is not
 general AMR: there is no per-region timestep, moving geometry, automatic
 dynamic refinement or L1 body support in the three-level form.
@@ -12,15 +12,15 @@ No CUDA kernel changes are involved.
 
 ## Explicit coordinate frames
 
-Keep the existing scalar `child=(x0,y0,x1,y1)` for a single scalar parent patch.
-Its coordinates retain their original meaning: half-open integer indices in the
-allocated L1 window, including its ghost band.
+For a single scalar parent patch, pass a scalar rectangle as `refinements[1]`:
+`(x0,y0,x1,y1)` in half-open integer indices of the allocated L1 window,
+including its ghost band.
 
-The new keyword `child_rectangles_in_parent=[...]` is mutually exclusive with
-`child`. It accepts a scalar or a set parent and interprets every rectangle in
-the actual allocated L1 array, including ghost bands and holes. It never infers
-region-local coordinates. The physical root cell edge corresponding to an L1
-index `(x,y)` is `parent.fine_origin + (x,y)/2`:
+A rectangle list as `refinements[1]` accepts a scalar or a set parent and
+interprets every rectangle in the actual allocated L1 array, including ghost
+bands and holes. It never infers region-local coordinates. The physical root
+cell edge corresponding to an L1 index `(x,y)` is
+`parent.fine_origin + (x,y)/2`:
 
 - A scalar L1 parent has a shifted window origin.
 - A set L1 parent has origin `(0,0)` and public shape `(2*ny,2*nx)`.
@@ -47,9 +47,9 @@ tile edges.
 
 ```python
 run = FixedPatchStepper(
-    cfg, [(4, 4, 20, 28), (24, 4, 44, 28)], context=ctx,
-    base_block=(4, 4), fine_block=(4, 4), child_block=(4, 4),
-    child_rectangles_in_parent=[(16, 16, 32, 40), (56, 16, 80, 40)],
+    cfg, context=ctx,
+    refinements=[[(4, 4, 20, 28), (24, 4, 44, 28)], [(16, 16, 32, 40), (56, 16, 80, 40)]],
+    level_blocks=[(4, 4)] * 3,
 )
 ```
 
@@ -57,7 +57,7 @@ This example uses a 48x32 root. `initialize` and `download` use three Q-major
 dense arrays. Public arrays and masks still scale with the full allocated
 extent; sparse device buffers do not remove that host-memory limit.
 
-Stationary whole-circle `child_objects` and `child_jets` use root edge
+Stationary whole-circle L2 bodies and Jets (`level_objects[2]`, `level_jets[2]`) use root edge
 coordinates and a single L2 boundary object. Body IDs are physical object-list
 positions, not region IDs or caller `obj_id`. Complete boundary stencils must
 avoid the transfer footprint. Change a child Jet with

docs/supported_settings.md

@@ -0,0 +1,88 @@
+# Supported settings: 2D Jet + fixed multilevel refinement (candidate 888c8038)
+
+Every entry names the guard that enforces it (file and function or kernel guard,
+at source revision 888c80381c006e976be2af8a452288729fc17c1b) and, where one
+exists, an existing test that exercises it. A setting not listed here is not
+claimed. Observations at the end are tied to the exact case in which they were
+measured; they are not general bounds.
+
+## Lattice, storage, collision
+
+| Setting | Supported | Guard |
+|---|---|---|
+| Dimension / lattice | 2D D2Q9 only (3D refinement is future work) | `lbm/refinement.py` `FixedPatch2D.from_config`; `#error` guards in `lbm/kernels/blocks.cu` and `lbm/kernels/refinement.cu` |
+| Precision / streaming | FP32 storage, double buffer | same guards |
+| Collision on refinement levels | SRT or MRT; TRT and LES rejected | `lbm/refinement.py` `FixedPatch2D.from_config` (re-checked for fine boundaries in `prepare_fine_boundary`) |
+| Shifted (ddf_shifting) storage | only when every level is blocked; dense or mixed rejected | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__`, `FixedPatchStepper._check_storage`; `#error` guard in `lbm/kernels/refinement.cu` (dense patch kernel unshifted only) |
+| Fine relaxation | 1/tau_fine must lie inside the configured omega guards | `lbm/refinement.py` `FixedPatch2D.from_config` |
+| Uniform blocked backend `PeriodicBlockStepper` | doubly periodic, SRT only, stationary passive circles | `lbm/block_stepper.py` `PeriodicBlockStepper.__init__`, `_static_circles` |
+
+Tests: `tests/unit/test_refinement_mrt.py`, `tests/validation/multigrid/test_refinement_mrt_gpu.py`,
+`tests/unit/test_shifted_storage.py`, `tests/unit/test_shifted_block_operators.py`.
+
+## Refinement hierarchy (`FixedPatchStepper(cfg, *, context, refinements, ...)`)
+
+| Setting | Supported | Guard |
+|---|---|---|
+| Depth | at least one refinement level; one entry per level in `level_blocks`/`level_objects`/`level_jets` | `lbm/refinement_stepper.py` `_check_levels` |
+| Single L1 rectangle | two-coarse-cell margin to the domain edge | `lbm/refinement.py` `FixedPatch2D.from_config` |
+| Several L1 rectangles (union) | require `level_blocks[0]` and `level_blocks[1]`; transfers must close (dependency checks, no fixed gap rule) | `lbm/refinement_stepper.py` `_check_levels`; `lbm/refinement.py` `_check_parent_pull_closure` |
+| Nested levels | depth >= 3 must be all blocked; a blocked level requires blocked levels below; a scalar child needs a scalar parent | `lbm/refinement_stepper.py` `_check_levels`; `lbm/hierarchy.py` `build_hierarchy` |
+| Nested transfer support | child only over inherited ACTIVE cells, donors/receivers uncovered ACTIVE, no child over the parent restriction footprint | `lbm/refinement.py` `_nested_kind` |
+| Tile shape | positive `(by, bx)`; the level allocation must be divisible by the tile (no partial tiles yet) | `lbm/refinement_stepper.py` `_check_levels`, `lbm/refinement.py` `FixedPatchSet2D.fine_layout` (shape); `lbm/blocks.py` `BlockLayout2D.build` (divisibility) |
+| Coarse boundary | `periodic`, `slot_reservoir` or `channel` | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__` |
+| `slot_reservoir` | on-site velocity bottom slot, reservoir top, velocity 0, 2 <= slot width < nx, patch rows 4..ny-4 | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__`; `lbm/refinement.py` `prepare_fine_boundary` |
+| `channel` | all levels blocked; Zou-He inlet, pressure-NEQ outlet, stationary bounce-back walls; transfers away from the closure | `lbm/channel.py` `validate_channel_config`, `validate_channel_transfers` |
+| Explosion | `constant` (default); `projected_linear` only periodic, all-blocked SRT, no bodies/Jets, no LES (experimental) | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__` |
+| L2 runtime remap (`remap_child`) | periodic three-level L2 sets only, no L0/L1 bodies, identical pinned boundary | `lbm/refinement_stepper.py` `FixedPatchStepper.remap_child` |
+
+Tests: `tests/unit/test_refinement_three_level_multipatch.py`, `tests/validation/multigrid/test_refinement_multipatch.py`,
+`test_refinement_three_level.py`, `test_refinement_three_level_multipatch_gpu.py`, `test_refinement_blocked_base.py`,
+`test_refinement_blocked_fine.py`, `test_channel_root_gpu.py`, `tests/unit/test_channel_root.py`,
+`tests/unit/test_projected_pair_stepper.py`, `test_runtime_child_remap_gpu.py`, `test_pinned_body_remap_gpu.py`.
+
+## Bodies and Jets on refined levels
+
+| Setting | Supported | Guard |
+|---|---|---|
+| Bodies | stationary circles without sensor/force-region roles, entirely inside the patch union, fully resolved, disjoint, stencil away from the two-hop transfer footprint | `lbm/refinement.py` `prepare_fine_boundary` |
+| Level | not on L0; only on the finest level of a nested hierarchy | `lbm/refinement_stepper.py` `_check_levels` |
+| Jet condition | velocity only; `pressure` is reserved and raises NotImplementedError | `body/jets.py` `Jet.__post_init__` |
+| Jet geometry | each Jet selects resolved links of exactly one stationary body; no Jet/Jet or Jet/SurfaceActuation overlap; abs(amplitude) < 1 lattice | `body/jets.py` `Jet.__post_init__`, `jet_link_shapes`, `_surface_links` |
+| Jets with shifted storage | only in the all-blocked stepper with SRT and no LES | `body/jets.py` `_surface_links` |
+| SurfaceActuation | unshifted storage only | `body/objects.py` `SimObject.validate_surface_actuation` |
+| Raw boundary exchange (force) | blocked body levels only; the window ledger requires `observe_raw_exchange` | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__` |
+| Checkpoint with Jets | refused | `simulation.py` `Simulation.save_checkpoint` |
+
+Raw exchange is the cut-link momentum reaction (wall + aperture), not a validated
+nozzle thrust or stress load. Jet amplitude is a lattice velocity, not a flow rate.
+
+Tests: `tests/unit/test_jets.py`, `test_jets_review.py`, `test_shifted_jets.py`, `test_raw_exchange.py`,
+`tests/validation/multigrid/test_refinement_jets.py`, `test_shifted_jets_gpu.py`, `test_raw_exchange_gpu.py`,
+`test_exchange_window.py`, `test_fine_surface_gpu.py`.
+
+## Observations from specific tested cases (not general bounds)
+
+- SRT channel root vs uniform OneStep, strict-parity test
+  `tests/validation/multigrid/test_channel_root_gpu.py::ChannelRootGPU::test_uniform_matched_root`:
+  in one of its 18 sub-cases (2 storage modes x 3 profile/clamp/damp settings x 3 blocks; which
+  one is not recorded), 1 of the 4320 compared values differs, max abs 2.33e-10, max rel 1.77e-7
+  (about 2 ULP). The test remains failed; its cause and any general bound are unresolved.
+- MRT root parity test
+  `tests/validation/multigrid/test_refinement_mrt_gpu.py::RefinementMRTGPU::test_uniform_matched_mrt_root`
+  (12 sub-cases: 2 storage modes x 2 profiles x 3 blocks, default outlet settings) passed bitwise
+  in the Horsea focused GPU run at K1 7813334 (2026-10-21).
+- Periodic Taylor-Green Re 100, measured at 36ffc8b1 (V3 matched pair, N128 depth 2):
+  every per-sample screen passed; the constant explosion's aggregate worst velocity L2,
+  0.004418, exceeds its preregistered 0.002 budget, while projected_linear passed that budget.
+- Curved-wall pressure (pilot harness, exact-q analytic annulus, tau 0.8): the coarse swirl r1
+  case exceeded its terminal density-gain cap (1.105e-4 vs 1e-4), so the frozen 9-case matrix
+  is NOT_ESTABLISHED. No validated fix exists. This is a pilot harness, not production code.
+- DFG 2D-1 Re 20 (V4, SRT, t* 60, pressure-NEQ outlet and local Zou-He inlet): refined and uniform
+  runs were NOT_STEADY under the frozen steady rule, so no interval verdict was given.
+
+## Not supported in this release
+
+3D, dynamic/automatic AMR (tagging, time-varying hierarchies beyond `remap_child`),
+moving bodies on refined levels, pressure Jets, Jet checkpoints, LES or TRT on refined
+levels, partial (non-dividing) tiles.

docs/validation/dfg_re20.md

@@ -0,0 +1,49 @@
+# DFG 2D-1 (Re 20) against the official reference intervals
+
+`tests/validation/run_dfg_re20.py` runs the steady DFG 2D-1 case on the refined
+channel root (D40 root, D80 wide fine patch) and on the uniform D80 grid. Geometry,
+inflow, initial state, layout and observation are those of the pinned Re100 runner
+(`run_refined_dfg.py`, `re100_helpers.py`, both unchanged); only the lattice
+viscosity is set for Re 20 (root tau 0.56, fine 0.62, uniform 0.62, lattice speed 0.01).
+
+Label: "compared against the official reference bands", not a full benchmark
+certification: the outlet is pressure-NEQ (`neq_extrap`), not do-nothing, and the
+inlet is local Zou-He.
+
+## Quantities
+
+- Cd = 2 Fx / (rho U^2 D), Cl = 2 Fy / (rho U^2 D), interval means over each sample.
+- p_diff = p(0.15, 0.2) - p(0.25, 0.2), linear extrapolation to the exact surface
+  points (`re100_helpers.pressure_probes`). The helper returns dp / (rho U^2) in
+  scaled units; p_diff = dp * 0.04 for U = 0.2 m/s, rho = 1. Both are reported.
+- Official intervals (Schaefer & Turek 1996): Cd [5.57, 5.59], Cl [0.0104, 0.0110],
+  p_diff [0.1172, 0.1176].
+
+## Bounds and steady rule (frozen before any run)
+
+- End time t* = 60 (30 s physical): 240,000 refined root steps, 480,000 uniform steps.
+- One sample every t* = 0.5 (0.25 s) at the end of a synchronized root step; Cd and
+  Cl are interval means. Trailing window t* 50..60 inclusive (21 samples).
+- Steady rule (our engineering choice, not official): in the window, for each of Cd,
+  Cl, p_diff, max - min <= 10 % and |least-squares slope x window| <= 2 % of its
+  official interval width (range 2e-3 / 6e-5 / 4e-5; slope 4e-4 / 1.2e-5 / 8e-6).
+  Absolute limits, so the small Cl is never a divisor. Never relaxed after a run.
+- The verdict uses the window mean; the distance to the nearest interval edge is
+  reported for each quantity (negative outside).
+- Force sign: fluid on cylinder, x along the inflow, y up from the bottom wall; Cd > 0
+  and the official Cl > 0 (`cd_positive` is reported).
+- Not steady at t* = 60: status `NOT_STEADY`, no interval verdict. Health limits
+  (Mach <= 0.2, rho in [0.5, 1.5]) stop the run with `UNHEALTHY`.
+- Steady: `WITHIN_INTERVALS` or `OUTSIDE_INTERVALS`, with the per-quantity verdict.
+
+## Run
+
+```
+python tests/validation/run_dfg_re20.py --output <new-dir> --plan
+CUDA_VISIBLE_DEVICES=<n> python tests/validation/run_dfg_re20.py --case both \
+    --collision SRT --output <new-dir> --max-wall-seconds <cap>
+```
+
+The uniform D80 control separates method error from refinement error. `--plan` writes `plan.json` only. A GPU run writes `plan.json` and `result.json`
+(per case: lattice parameters, every sample, window ranges, status, verdict).
+Host contract tests: `tests/unit/test_dfg_re20.py`.

docs/validation/refined_dfg_window.md

@@ -56,3 +56,62 @@ Root must independently review and admit any execution. An admitted paired pilot
 the existing watchdog, device selection, `--gpu-gates-approved`, positive
 `--max-wall-seconds` and a new output directory, adding only `--exchange-window` to the
 selected collision's pilot command. No GPU or remote action was performed for this patch.
+
+## Public Jet diagnostic option
+
+`--collision MRT --exchange-window --jet-scenario zero|fixed|reverse` selects
+unshifted FP32, no LES, constant transfer on both existing public backends. Default
+`none` and the historical manifests remain unchanged. `zero` registers both Jets at
+zero amplitude and retains every existing DFG verdict, including pressure. `fixed`
+uses outward radial parabolic sectors at pi/2 and 3pi/2, half-width pi/12, peaks
+(+.001,-.001). Both use the existing D40/D80 geometry and U=.01. Sector origins are
+(80.5,81) in root edge coordinates and (160,160.5) in uniform node coordinates;
+radial extent is [0,1.005R]. `reverse` keeps those peaks through t*=300, updates both
+public setters at that completed boundary, and continues to 305 with opposite signs.
+
+Fixed diagnostic windows are 250–260, 260–270, 270–280, 280–290 and 290–300; reversal
+adds 300–305. Histories retain exact interval endpoints and root/uniform lattice
+impulses. The refined window uses total exchange; uniform uses public `read_force`
+batch averages. Driven windows report mean/min/max/RMS Cd/Cl without applying DFG
+reference or periodicity gates. Exit zero means healthy execution, not qualification;
+`matched_case_pass` remains false. No independent CV or stress-load claim is made.
+
+This opt-in records current source/helper hashes with the existing manifest format,
+checkout revision/dirty state and runner hash. Pilot/full must match that snapshot
+and the exact scenario. It does not relabel current source as historical evidence.
+Use the existing bounded plan/pilot/full commands with the three option flags above;
+full requires the matching completed t*=5 paired pilot. `step_seconds` includes force
+readback; observation/overhead and total wall time are separate, not kernel throughput.
+
+Add `--observe-cv` to record independent scalar control-volume samples from the
+already downloaded public macros. The pilot exercises this path at t*=.02..5;
+full records 501 endpoints at 290..300. Each `cv_history.jsonl` row contains actual
+root time, preceding interval start/duration/impulse, and two inventories/outward
+fluxes ordered [mass,Px,Py]. Exclude the first row's preceding impulse: that row
+supplies the initial inventory. Uniform native impulse converts by .5² to root
+units; refined ledger impulse is already in root units. Explicit native-aligned
+contours differ by .25 root cells in x and are recorded in the plan. Cropping
+retains the two-cell face-gradient stencils. No fields or per-link arrays are saved.
+`cv_observation_seconds` is the CV portion of observation/overhead time. Helper hash
+and the observation flag must match pilot/full. Time integration, temporal
+coarsening and any physical assessment remain separate offline work.
+
+## Bounded resolution and placement controls
+
+Defaults remain `--case both --uniform-diameter 80 --placement wide`. Single-case
+runs use `--case uniform` or `--case refined` and never report a matched-case pass.
+Their own reference/periodicity or driven-diagnostic results remain available.
+A single zero-Jet full run exits successfully only when its periodicity and the
+unchanged complete DFG reference verdict pass; driven completion is not qualification.
+
+The D160 option requires `--case uniform --uniform-diameter 160` and an explicit
+public Jet scenario. U and Jet speeds stay .01 and ±.001; viscosity is .016,
+grid 3521×658, center (320,320.5), radius80. CV uses root spacing/time factor .25,
+origin (.375,.75), impulse factor .0625, root viscosity .004, and x-face shift .125.
+
+`--case refined --placement compact` selects root bounds (48,48,128,128), fine
+168×168 and origin (46,46), with the same D40/fineD80 body/Jet resolution. Its CV
+contours are (52,52,112,112) and (50,50,120,120); complete face-gradient stencils
+are clear of both solids and the transfer footprint. Compact is refined-only.
+CV metadata gives each grid's spacing/origin/contours explicitly. Case, uniform
+diameter and placement must match the newly completed pilot in every receipt.

examples/fixed_patch_circle.py

@@ -90,8 +90,8 @@ def main():
     # The caller owns the CUDA context; the stepper must be closed first (the
     # nested with-statement does that).
     with CudaContext(0) as context, FixedPatchStepper(
-            cfg, BOUNDS, context=context, coarse_boundary=args.mode,
-            fine_objects=[body], fine_block=fine_block) as run:
+            cfg, context=context, coarse_boundary=args.mode, refinements=[BOUNDS],
+            level_blocks=[None, fine_block], level_objects=[(), (body,)]) as run:
         p = run.patch
         # Rest equilibrium is already postcollision G, so no initial collision is
         # needed. Non-leaf entries (covered, ghost, solid) are discarded by initialize.

examples/fixed_patch_jet.py

@@ -148,8 +148,8 @@ def main():
                   samples=[], status='running', last_synchronized_coarse_step=0)
     start, run = time.monotonic(), None
     try:
-        with CudaContext(0) as context, FixedPatchStepper(cfg, BOUNDS, context=context,
-                coarse_boundary='slot_reservoir') as run:
+        with CudaContext(0) as context, FixedPatchStepper(cfg, context=context,
+                coarse_boundary='slot_reservoir', refinements=[BOUNDS]) as run:
             p = run.patch
             masks = [np.ones(p.coarse_shape, bool), p.fine_kind == ACTIVE]
             masks[0][[0, -1]] = False

examples/fixed_periodic_patch.py

@@ -48,7 +48,7 @@ def main():
                   nx=16, ny=16, viscosity=.1, velocity=1e-3, rho=1.0,
                   lattice_model="D2Q9", collision="SRT", streaming="double_buffer",
                   store_precision="FP32", ddf_shifting=False, les_enabled=False)
-    with CudaContext(0) as context, PeriodicPatchStepper(cfg, (4, 4, 12, 12), context=context) as run:
+    with CudaContext(0) as context, PeriodicPatchStepper(cfg, context=context, refinements=[(4, 4, 12, 12)]) as run:
         p = run.patch
         coarse = shear_equilibrium(p.coarse_shape, (0, 0), 1.0, cfg.ny)
         fine = shear_equilibrium(p.fine_shape, p.fine_origin, .5, cfg.ny)

src/CelerisLab/body/jets.py

@@ -207,22 +207,30 @@ def surface_link_velocity(objects, cfg, flags, fluid_idx, dir, body_id, rx, ry,
                           jets=(), hits=None, *, _shifted_blocked=False):
     """Per-link prescribed surface velocity (n, 2) FP32 and aperture mask.
 
+    See _surface_links, which also returns the unit-amplitude jet shapes.
+    """
+    return _surface_links(objects, cfg, flags, fluid_idx, dir, body_id, rx, ry, jets, hits,
+                          _shifted_blocked=_shifted_blocked)[:2]
+
+
+def _surface_links(objects, cfg, flags, fluid_idx, dir, body_id, rx, ry,
+                   jets=(), hits=None, *, _shifted_blocked=False):
+    """(velocity, aperture, jet shape, jet id); the last two are None without jets.
+
     Host-only and storage-neutral: rows follow the given physical cut-link
-    order. Returns (None, None) when neither an object carries surface
+    order. Returns all None when neither an object carries surface
     actuation nor a jet is given. Jets need true FP64 ``hits`` (n, 2) and may
     not share a link with each other or with a legacy aperture.
     Shifted storage is privately enabled only by the all-blocked stepper.
     """
     enabled = [obj for obj in objects if obj.surface_actuation is not None]
     if not enabled and not jets:
-        return None, None
-    from ..lbm.descriptors import FLUID, SOLID
-    if np.any((flags[fluid_idx] & FLUID) == 0) or np.any(flags[fluid_idx] & SOLID):
-        raise ValueError("Surface cut links must have active fluid receivers")
-    if len(np.unique(fluid_idx.astype(np.uint64)*cfg.nq+dir)) != len(dir):
-        raise ValueError("Surface boundary requires unique cut-link writers")
+        return None, None, None, None
+    from .preprocess.flag_overlay import validate_cut_link_ownership
+    validate_cut_link_ownership(flags, fluid_idx, dir, cfg.nq)
     surface_velocity = np.zeros((len(dir), 2), np.float32)
     surface_aperture = np.zeros(len(dir), bool)
+    shape = jet_id = None
     for obj in enabled:
         obj.validate_surface_actuation(cfg)
         selected = body_id == obj.obj_id
@@ -255,4 +263,4 @@ def surface_link_velocity(objects, cfg, flags, fluid_idx, dir, body_id, rx, ry,
             raise ValueError("Jet overlaps a SurfaceActuation aperture")
         surface_velocity[driven] = jet_link_velocity(jets, shape, jet_id)[driven]
         surface_aperture |= driven
-    return surface_velocity, surface_aperture
+    return surface_velocity, surface_aperture, shape, jet_id

src/CelerisLab/body/manager.py

@@ -27,7 +27,7 @@ import pycuda.driver as cuda
 
 from .objects import SimObject
 from .registry import BodyRegistry
-from .preprocess.flag_overlay import merge_flag_masks
+from .preprocess.flag_overlay import merge_flag_masks, validate_cut_link_ownership
 from ..cuda.compiler_v2 import obs_layout
 from ..cuda.obs_layout import ObsLayout
 from ..lbm.descriptors import (
@@ -369,9 +369,9 @@ class ObjectManager:
 
         Args:
             objects: List of SimObject instances to process.
-            domain_flags: Optional uint16 flags for donor-cell classification.
+            domain_flags: Optional uint16 flags for receiver/donor classification.
             sensor_cell_counts_out: Int array to fill with per-object sensor
-                cell counts.  If provided, ``fill(0)`` is called first.
+                cell counts after the full geometry has passed validation.
             hits_out: List that receives one FP64 (n, 2) array of true hit
                 points row-aligned with the cut-link columns (for Jets).
 
@@ -394,8 +394,8 @@ class ObjectManager:
         active_fluid = ((domain_flags & FLUID) != 0) & ((domain_flags & SOLID) == 0)
         force_owners: set[int] = set()
 
-        if sensor_cell_counts_out is not None:
-            sensor_cell_counts_out.fill(0)
+        sensor_counts = (np.zeros_like(sensor_cell_counts_out)
+                         if sensor_cell_counts_out is not None else None)
 
         ez = None
         if self.cfg and self.cfg.is_d3q19:
@@ -430,9 +430,9 @@ class ObjectManager:
                 if len(cells) > 0:
                     s_cells.append(cells)
                     s_ids.append(ids)
-                    if (sensor_cell_counts_out is not None
-                            and 0 <= obj.obj_id < sensor_cell_counts_out.size):
-                        sensor_cell_counts_out[obj.obj_id] = int(len(cells))
+                    if (sensor_counts is not None
+                            and 0 <= obj.obj_id < sensor_counts.size):
+                        sensor_counts[obj.obj_id] = int(len(cells))
             if obj.is_force_region and hasattr(obj, 'get_sensor_list'):
                 cells, ids = obj.get_sensor_list(self.nx, self.ny, self.nz)
                 keep = active_fluid[cells]
@@ -461,8 +461,9 @@ class ObjectManager:
         r_fr_ids = np.concatenate(fr_ids) if fr_ids else np.zeros(0, dtype=np.int32)
 
         r_hit = np.concatenate(cl_hit) if cl_hit else np.zeros((0, 2))
-        if hits_out is not None:
-            hits_out.append(r_hit)
+        # Passive walls need the same ownership guarantees as driven walls.
+        # Overlapping solids are not a supported union-boundary operation.
+        validate_cut_link_ownership(domain_flags, r_idx, r_dir, self.nq)
         if self.jets or any(obj.surface_actuation is not None for obj in objects):
             # Reject unresolved apertures/ownership before sync or restart mutates state.
             from ..lbm.curved_links import CurvedLinkSoA
@@ -470,6 +471,11 @@ class ObjectManager:
             links.assign_host(r_idx, r_dir, r_q, r_bid, r_rx, r_ry, r_rz, r_fallback)
             links.assign_surface(objects, self.cfg, domain_flags, jets=self.jets, hits=r_hit)
 
+        if sensor_cell_counts_out is not None:
+            sensor_cell_counts_out[:] = sensor_counts
+        if hits_out is not None:
+            hits_out.append(r_hit)
+
         return (
             r_idx, r_dir, r_q, r_bid,
             r_rx, r_ry, r_rz, r_fallback,
@@ -487,8 +493,6 @@ class ObjectManager:
                 the caller has already prepared and uploaded the final flag field
                 for this geometry, such as during initialization.
         """
-        self._telemetry_field = field
-
         next_flags = self.build_flags(field.build_channel_flags()) if rebuild_flags else field.flag
         hits = []
 
@@ -499,6 +503,7 @@ class ObjectManager:
             fr_cells, fr_obj_id,
         ) = self.build_compact_lists(domain_flags=next_flags, hits_out=hits)
 
+        self._telemetry_field = field
         if rebuild_flags:
             field.flag = next_flags
             field.upload_flags()

src/CelerisLab/body/preprocess/flag_overlay.py

@@ -1,6 +1,6 @@
 # CelerisLab/body/preprocess/flag_overlay.py
 """
-Flag overlay: merge per-object flag masks onto a clean domain base.
+Flag overlay and cut-link ownership checks against the merged domain.
 """
 
 from __future__ import annotations
@@ -12,6 +12,15 @@ import numpy as np
 from ...lbm.descriptors import FLUID, SOLID, SENSOR_FLAG, FRC_REGION
 
 
+def validate_cut_link_ownership(flags, fluid_idx, directions, nq):
+    """Reject solid receivers and multiple writers to one population."""
+    if np.any((flags[fluid_idx] & FLUID) == 0) or np.any(flags[fluid_idx] & SOLID):
+        raise ValueError("Cut links must have active fluid receivers")
+    writers = fluid_idx.astype(np.uint64) * nq + directions
+    if len(np.unique(writers)) != len(directions):
+        raise ValueError("Boundary requires unique cut-link writers")
+
+
 def merge_flag_masks(
     base: np.ndarray,
     masks: List[np.ndarray],

src/CelerisLab/cuda/compiler_v2.py

@@ -54,14 +54,18 @@ def _write(path: str, content: str):
 
 
 def generate_config(cfg: LBMConfig, n_objects: int = 0,
-                    *, build_root: Optional[str] = None) -> str:
+                    *, build_root: Optional[str] = None, block_root_bc: int = 0) -> str:
     """Snapshot kernels and render headers in a fresh private directory.
 
     Returns its path. Pass it explicitly to compile_kernel for interleaved
     builds. Legacy generate_config(); compile_kernel() remains context-local.
     The caller owns cleanup_build_dir() once the loaded module is no longer
     needed for recompilation/debugging. No package files are changed.
+    block_root_bc specializes a blocked root module; fine/dense builds use 0.
     """
+    if (isinstance(block_root_bc, bool) or not isinstance(block_root_bc, int)
+            or block_root_bc not in (0, 1, 2)):
+        raise ValueError("block_root_bc must be an integer mode 0, 1 or 2")
     macros = cfg.to_macros()
     if isinstance(n_objects, bool) or not isinstance(n_objects, int) or n_objects < 0:
         raise ValueError("n_objects must be a nonnegative integer")
@@ -71,14 +75,15 @@ def generate_config(cfg: LBMConfig, n_objects: int = 0,
     try:
         shutil.copytree(_KERNEL_DIR, build_dir, dirs_exist_ok=True,
                         ignore=shutil.ignore_patterns("*.ptx", "*.cubin", "*.tmp"))
-        _render_config(macros, n_objects, os.path.join(build_dir, "config"))
+        _render_config(macros, n_objects, os.path.join(build_dir, "config"),
+                       block_root_bc=block_root_bc)
         digest = hashlib.sha256()
         for path in sorted(Path(build_dir).rglob("*")):
             if path.is_file():
                 digest.update(str(path.relative_to(build_dir)).encode())
                 digest.update(path.read_bytes())
         metadata = {"source_sha256": digest.hexdigest(), "config": asdict(cfg),
-                    "n_objects": n_objects}
+                    "n_objects": n_objects, "block_root_bc": block_root_bc}
         _write(os.path.join(build_dir, ".celeris-build"), json.dumps(metadata, sort_keys=True))
     except BaseException:
         shutil.rmtree(build_dir)
@@ -101,7 +106,7 @@ def cleanup_build_dir(build_dir: Optional[str]) -> None:
         _last_ptx_path.set(None)
 
 
-def _render_config(m: dict, n_objects: int, cdir: str):
+def _render_config(m: dict, n_objects: int, cdir: str, *, block_root_bc: int = 0):
     _write(os.path.join(cdir, "config_grid.h"), f"""{_HEADER}\
 // Layer: Global / Grid
 #ifndef CELERIS_CONFIG_GRID_H
@@ -165,6 +170,9 @@ def _render_config(m: dict, n_objects: int, cdir: str):
 #define STORE_PRECISION  {m['STORE_PRECISION']}
 #define USE_DDF_SHIFTING {m['USE_DDF_SHIFTING']}
 
+// Internal blocked root mode: 0=periodic/fine, 1=slot/reservoir, 2=channel.
+#define BLOCK_ROOT_BC {block_root_bc}
+
 #define USE_LES {m['USE_LES']}
 #define LES_CS  {m['LES_CS']}
 #define LES_CLOSED_FORM {m['LES_CLOSED_FORM']}

src/CelerisLab/lbm/block_stepper.py

@@ -204,11 +204,21 @@ class PeriodicBlockStepper:
                 raise
 
     def step(self, n=1):
-        """Advance n steps: halo refresh, then core pull/collide, per step."""
+        """Advance n steps: halo refresh, then core pull/collide, per step.
+
+        All n steps are enqueued in order on the one stream and the host waits
+        once before returning (same launches and bits as n step(1) calls).
+        Asynchronous device errors surface at that single wait, so they are not
+        attributed to the step that caused them; call step(1) repeatedly to
+        localize one. On any failure the instance is failed and step_count keeps
+        its entry value.
+        """
         if isinstance(n, bool) or not isinstance(n, Integral) or n < 0:
             raise ValueError("n must be a nonnegative integer")
         with self._lock, self.ctx.activate():
             self._require("ready")
+            if n == 0:
+                return
             kw = dict(block=self._block, stream=self._stream)
             try:
                 self._status = "advancing"
@@ -224,9 +234,9 @@ class PeriodicBlockStepper:
                     self._step(self._current, self._spare, np.uint32(self.layout.halo),
                                self._tables.get("flags", np.uintp(0)), np.uintp(0),  # no kind
                                grid=self._grid, **kw)
-                    self._stream.synchronize()
-                    self._current, self._spare = self._spare, self._current
-                    self._step_count += 1
+                    self._current, self._spare = self._spare, self._current  # host handles only
+                self._stream.synchronize()
+                self._step_count += n
                 self._status = "ready"
             except BaseException:
                 self._status = "failed"

src/CelerisLab/lbm/curved_links.py

@@ -94,12 +94,10 @@ class CurvedLinkSoA:
         """
         self.surface_velocity = self.aperture = None
         self.jets, self.jet_shape, self.jet_id = tuple(jets), None, None
-        self.surface_velocity, self.aperture = surface_link_velocity(
+        from ..body.jets import _surface_links
+        self.surface_velocity, self.aperture, self.jet_shape, self.jet_id = _surface_links(
             objects, cfg, flags, self.fluid_idx, self.dir, self.body_id, self.rx, self.ry,
             jets=self.jets, hits=hits)
-        if self.jets:
-            from ..body.jets import jet_link_shapes
-            self.jet_shape, self.jet_id = jet_link_shapes(self.jets, hits)
 
     def set_jet_amplitude(self, index, amplitude, stream=None):
         """Rewrite jet rows from cached shapes; re-upload only the velocity column.

src/CelerisLab/lbm/kernels/blocks.cu

@@ -1,5 +1,6 @@
 // CelerisLab/lbm/kernels/blocks.cu
-// One periodic level of fixed-size blocks (lbm/blocks.py layout). Separate
+// One fixed-size blocked level (lbm/blocks.py layout), specialized at compile
+// time for periodic/fine, slot/reservoir or channel stepping. Separate
 // compilation entry: NX, NY are the block window including halo, so index_f
 // and linear_index address one block; blockIdx.z selects the block, whose
 // populations start at b*NQ*TOTAL_CELLS (block-contiguous, Q-major inside).
@@ -12,6 +13,7 @@
 #include "core/params.cuh"
 #include "core/flags.cuh"
 #include "core/obs.cuh"
+#include "core/cell_kind.cuh"
 #include "operators/macro.cuh"
 #include "operators/collision_srt.cuh"
 #if COLLISION_MODEL == 2
@@ -19,188 +21,29 @@
 #endif
 #include "streaming/pull_double_buffer.cuh"
 #include "boundary/curved_boundary.cuh"
-#include "boundary/planar_velocity.cuh"
 
 #if DIM != 2 || NQ != 9 || (COLLISION_MODEL != 0 && COLLISION_MODEL != 2) || STORE_PRECISION != 0 || USE_LES || STREAMING_MODEL != 0
 #error "Blocks require FP32 D2Q9 SRT or MRT double buffers without LES"
 #endif
-
-// Fixed-patch cell kinds, matching refinement.py ACTIVE/INNER_GHOST/OUTER_GHOST.
-constexpr unsigned char KIND_ACTIVE = 0u, KIND_INNER = 1u, KIND_OUTER = 2u;
-
-// Zero-initialized for ordinary blocks/fine levels. Set only for parent walls;
-// NX/NY remain local storage strides, never physical wall coordinates.
-struct BlockWall { unsigned int nx, ny, slot_width; float amplitude; };
-extern "C" { __constant__ BlockWall d_block_wall; }
-
-// Halo refresh: block cell ids b*TOTAL_CELLS+c, sources are always core cells.
-extern "C" __global__ void BlockHaloCopy(
-    fpxx* f, const uint64_t* dst, const uint64_t* src, unsigned int count)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const uint64_t d = dst[row], s = src[row];
-    fpxx* to = f + (d / TOTAL_CELLS) * NQ * TOTAL_CELLS;
-    const fpxx* from = f + (s / TOTAL_CELLS) * NQ * TOTAL_CELLS;
-    for (unsigned int q = 0; q < NQ; ++q)
-        to[index_f(d % TOTAL_CELLS, q)] = from[index_f(s % TOTAL_CELLS, q)];
-}
-
-// Fixed-patch explosion into blocked fine storage: target = owner block cell
-// id of a patch-ghost cell; donors are parent block cell ids. Dense parent
-// storage is the one-window case with coarse_cells = full domain size.
-extern "C" __global__ void BlockExplode(
-    const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
-    unsigned int count, unsigned long coarse_cells)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const uint64_t t = target[row];
-    fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
-    const uint64_t d = donor[row];
-    const fpxx* from = coarse + (d / coarse_cells) * NQ * coarse_cells;
-    for (unsigned int q = 0; q < NQ; ++q)
-        to[index_f(t % TOTAL_CELLS, q)] = from[q * coarse_cells + d % coarse_cells];
-}
-
-// Experimental Schornbaum--Ruede (arXiv:1508.07982, Algorithm 2)
-// and Chen et al. (doi:10.1016/j.physa.2005.09.036) population interpolation.
-// Independent implementation of the mathematical stencil; no upstream source copied.
-// Stencil cells are owner cores; a disabled axis repeats its center twice.
-extern "C" __global__ void BlockExplodeProjected(
-    const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
-    unsigned int count, unsigned long coarse_cells,
-    const uint64_t* neighbors, const float* offsets)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const uint64_t t = target[row], d = donor[row];
-    fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
-    for (unsigned int q = 0; q < NQ; ++q) {
-        float v[4];
-        for (unsigned int a = 0; a < 4; ++a) {
-            const uint64_t n = neighbors[4ul * row + a];
-            v[a] = coarse[(n / coarse_cells) * NQ * coarse_cells
-                          + q * coarse_cells + n % coarse_cells];
-        }
-        float gx = .5f * (v[1] - v[0]), gy = .5f * (v[3] - v[2]);
-        const float ex = d_cx[q], ey = d_cy[q];
-        if (q != 0) {
-            const float along = (ex * gx + ey * gy) / (ex * ex + ey * ey);
-            gx -= ex * along;
-            gy -= ey * along;
-        }
-        to[index_f(t % TOTAL_CELLS, q)] =
-            coarse[(d / coarse_cells) * NQ * coarse_cells + q * coarse_cells + d % coarse_cells]
-            + offsets[2ul * row] * gx + offsets[2ul * row + 1] * gy;
-    }
-}
-
-// Targets are receiver-local SOURCE slots for the immediately following pull.
-// Keep child summation order identical to PatchCoarseStep; no collision here.
-extern "C" __global__ void BlockCoalesce(
-    const fpxx* fine, fpxx* coarse, const uint64_t* target, const uint64_t* children,
-    unsigned int count)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const uint64_t* s = children + 4ul * row;
-    coarse[target[row]] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
-}
-
-// Static cut links, receiver-local staging: cell is the receiver's local core
-// cell, so the existing BFL writes the solid-source window image that this
-// block's next pull reads (no halo copy between this kernel and BlockStep).
-// surface_velocity (nullable, 2 floats per link): prescribed fluid velocity on
-// fixed geometry, same convention as CurvedBoundaryKernel. Optional raw rows
-// (jx,jy,torque,dm) are unique physical links, never body-observation slots.
-extern "C" __global__ void BlockCurvedBoundary(
-    fpxx* f, const uint32_t* block, const uint32_t* cell, const unsigned char* dir,
-    const float* q, const unsigned char* fallback, unsigned int count,
-    const float* surface_velocity, const float* rx, const float* ry, float* raw)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    float Uw = 0.0f, Vw = 0.0f;
-    if (surface_velocity != nullptr) {
-        const float u = surface_velocity[2 * row], v = surface_velocity[2 * row + 1];
-        if (u != 0.0f || v != 0.0f) { Uw = u; Vw = v; }
-    }
-    apply_bouzidi_link(dir[row], q[row], cell[row], f + (unsigned long)block[row] * NQ * TOTAL_CELLS,
-                       Uw, Vw, rx ? rx[row] : 0.0f, ry ? ry[row] : 0.0f,
-                       fallback[row], nullptr, -1, raw ? raw + 4ul * row : nullptr,
-                       raw ? raw + 4ul * row + 3 : nullptr);
-}
-
-// Same-stream BFL -> this row-owned accumulation -> pull step. No atomics,
-// no cross-row reduction and no history allocation proportional to substeps.
-extern "C" __global__ void BlockAccumulateExchange(
-    const float* raw, double* integral, unsigned int count, double h, double dt)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const unsigned long k = 4ul * row;
-    const double impulse = h * dt;
-    integral[k] += impulse * (double)raw[k];
-    integral[k + 1] += impulse * (double)raw[k + 1];
-    integral[k + 2] += h * impulse * (double)raw[k + 2];
-    integral[k + 3] += h * h * (double)raw[k + 3];
-}
-
-// Pull/collide/store core cells only; halo cells of out are left untouched.
-// flags (block cell ids, nullable): solid cores store NaN and never collide.
-// kind (block cell ids, nullable = all ACTIVE): OUTER skips the pull and
-// stores NaN, INNER pulls without collision, ACTIVE pulls and collides.
-extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int halo,
-                                     const uint16_t* flags, const unsigned char* kind) {
-    unsigned int x, y; unsigned long k;
-    index_from_thread(x, y, k);
-    if (x < halo || y < halo || x >= NX - halo || y >= NY - halo) return;
-    const unsigned long base = (unsigned long)blockIdx.z * NQ * TOTAL_CELLS;
-    const unsigned long cell = (unsigned long)blockIdx.z * TOTAL_CELLS + k;
-    const unsigned char cell_kind = kind != nullptr ? kind[cell] : KIND_ACTIVE;
-    float f[NQ];
-    if (cell_kind == KIND_OUTER || (flags != nullptr && is_solid(flags[cell]))) {
-        for (int q = 0; q < NQ; ++q) f[q] = NAN;
-    } else {
-        unsigned long j[NQ];
-        compute_neighbors(x, y, j);  // Core neighbors stay inside the window: no wrap.
-        stream_pull_load(k, f, in + base, j);
-        if (cell_kind == KIND_ACTIVE) {
-            if (d_block_wall.nx != 0) {
-                const unsigned int bx = NX - 2 * halo, by = NY - 2 * halo;
-                const unsigned int nbx = d_block_wall.nx / bx;
-                const unsigned int gx = (blockIdx.z % nbx) * bx + x - halo;
-                const unsigned int gy = (blockIdx.z / nbx) * by + y - halo;
-                if (gy == 1)
-                    reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity(
-                        gx, d_block_wall.nx, d_block_wall.slot_width, d_block_wall.amplitude));
-                else if (gy == d_block_wall.ny - 2) reconstruct_top_reservoir(f, RHO);
-            }
-#if COLLISION_MODEL == 2
-            // Uniform MRT sequence (collide_dispatch, zero forcing); SRT path below unchanged.
-            float rho, ux, uy;
-            compute_rho_u(f, rho, ux, uy);
-            collide_mrt_no_force(f, rho, ux, uy,
-                fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega)));
-#else
-            float rho, ux, uy, feq[NQ];
-            compute_rho_u(f, rho, ux, uy);
-#if USE_DDF_SHIFTING
-            // Preserve the density perturbation before adding the unit reference.
-            const float drho = f[0]+f[1]+f[2]+f[3]+f[4]+f[5]+f[6]+f[7]+f[8];
-            const float jx = f[1]-f[2]+f[5]-f[6]+f[7]-f[8];
-            const float jy = f[3]-f[4]+f[5]-f[6]-f[7]+f[8];
-            compute_feq_from_momentum(drho, jx, jy, ux, uy, feq);
-#else
-            compute_feq(rho, ux, uy, feq);
+#ifndef BLOCK_ROOT_BC
+#error "Blocks require a generated BLOCK_ROOT_BC mode"
+#elif BLOCK_ROOT_BC < 0 || BLOCK_ROOT_BC > 2
+#error "BLOCK_ROOT_BC must be 0, 1 or 2"
 #endif
-            collide_srt_no_force(f, feq, d_params.omega);
-#endif
-        }
-    }
-    stream_pull_store(k, f, out + base);
-}
 
-// Root channel closure; fine and periodic levels keep BlockStep.
+#if BLOCK_ROOT_BC == 1
+#include "boundary/block_wall.cuh"
+#endif
+#if BLOCK_ROOT_BC != 2
+#include "operators/block_collision.cuh"
+#endif
+
+// Kernels, in definition order: transfers, boundary staging, then the one full
+// step selected for this module (BlockStep, or BlockChannelStep for a channel root).
+#include "step/block_transfer.cu"
+#include "step/block_boundary.cu"
+#if BLOCK_ROOT_BC != 2
+#include "step/block_step.cu"
+#else
 #include "step/block_channel.cu"
+#endif

src/CelerisLab/lbm/kernels/boundary/block_wall.cuh

@@ -0,0 +1,20 @@
+// Physical parent walls; NX/NY remain local storage strides.
+#ifndef CELERIS_BOUNDARY_BLOCK_WALL_CUH
+#define CELERIS_BOUNDARY_BLOCK_WALL_CUH
+
+#include "planar_velocity.cuh"
+
+struct BlockWall { unsigned int nx, ny, slot_width; float amplitude; };
+extern "C" { __constant__ BlockWall d_block_wall; }
+
+__device__ __forceinline__ void block_wall(float* f, unsigned int x, unsigned int y,
+                                         unsigned int halo) {
+    unsigned int gx, gy;
+    block_global_xy(x, y, halo, d_block_wall.nx, gx, gy);
+    if (gy == 1)
+        reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity(
+            gx, d_block_wall.nx, d_block_wall.slot_width, d_block_wall.amplitude));
+    else if (gy == d_block_wall.ny - 2) reconstruct_top_reservoir(f, RHO);
+}
+
+#endif

src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh

@@ -15,10 +15,13 @@
 // D2Q9 y-wall half-way bounce-back for pull double-buffer streaming.
 // ---------------------------------------------------------------------------
 #if LATTICE_MODEL == LATTICE_D2Q9
-__device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y,
-                                                  float* __restrict__ f,
-                                                  const fpxx* __restrict__ fi_in,
-                                                  unsigned long k)
+// Rows 1 and top_row (= physical ny - 2) next to walls at rows 0 and ny - 1.
+// y is the physical row; k addresses the node in fi_in (a block window may
+// have NX/NY strides different from the physical extent).
+__device__ inline void apply_wall_bb_d2q9_rows(unsigned int y, unsigned int top_row,
+                                               float* __restrict__ f,
+                                               const fpxx* __restrict__ fi_in,
+                                               unsigned long k)
 {
     if (y == 1) {
         // Directions sourced from y=0 wall in pull step: +y, +x+y, -x+y.
@@ -26,14 +29,22 @@ __device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y,
         f[5] = load_ddf(fi_in, index_f(k, 6u));
         f[8] = load_ddf(fi_in, index_f(k, 7u));
     }
-    else if (y == (unsigned int)(NY - 2)) {
-        // Directions sourced from y=NY-1 wall in pull step: -y, -x-y, +x-y.
+    else if (y == top_row) {
+        // Directions sourced from the y = ny - 1 wall in pull step: -y, -x-y, +x-y.
         f[4] = load_ddf(fi_in, index_f(k, 3u));
         f[6] = load_ddf(fi_in, index_f(k, 5u));
         f[7] = load_ddf(fi_in, index_f(k, 8u));
     }
 }
 
+__device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y,
+                                                  float* __restrict__ f,
+                                                  const fpxx* __restrict__ fi_in,
+                                                  unsigned long k)
+{
+    apply_wall_bb_d2q9_rows(y, (unsigned int)(NY - 2), f, fi_in, k);
+}
+
 // ---------------------------------------------------------------------------
 // D2Q9 y-wall free-slip (specular reflection) for pull double-buffer streaming.
 //

src/CelerisLab/lbm/kernels/config/config_method.h

@@ -8,6 +8,9 @@
 #define STORE_PRECISION  0
 #define USE_DDF_SHIFTING 0
 
+// Internal blocked root mode: 0=periodic/fine, 1=slot/reservoir, 2=channel.
+#define BLOCK_ROOT_BC 0
+
 #define USE_LES 0
 #define LES_CS  0.160000f
 #define LES_CLOSED_FORM 1

src/CelerisLab/lbm/kernels/core/cell_kind.cuh

@@ -0,0 +1,16 @@
+// CelerisLab – core/cell_kind.cuh
+// Fixed-patch cell kinds stored per block cell, matching lbm/refinement.py
+// ACTIVE/INNER_GHOST/OUTER_GHOST: OUTER is invalid (stores NaN, no pull),
+// INNER streams without collision, ACTIVE pulls and collides.
+#ifndef CELERIS_CORE_CELL_KIND_CUH
+#define CELERIS_CORE_CELL_KIND_CUH
+
+constexpr unsigned char KIND_ACTIVE = 0u, KIND_INNER = 1u, KIND_OUTER = 2u;
+
+// NaN populations mark a cell without valid state (OUTER ghost or solid core).
+__device__ __forceinline__ void fill_invalid(float* f)
+{
+    for (int q = 0; q < NQ; ++q) f[q] = NAN;
+}
+
+#endif

src/CelerisLab/lbm/kernels/core/layout.cuh

@@ -53,6 +53,18 @@ __device__ __forceinline__ void index_from_thread(unsigned int& x,
     k = (unsigned long)y * NX + x;
 }
 
+// Global cell of window cell (x, y) in block blockIdx.z of a row-major,
+// full-domain block layout domain_nx cells wide. NX/NY are the local strides.
+__device__ __forceinline__ void block_global_xy(unsigned int x, unsigned int y,
+                                                unsigned int halo, unsigned int domain_nx,
+                                                unsigned int& gx, unsigned int& gy)
+{
+    const unsigned int bx = NX - 2 * halo, by = NY - 2 * halo;
+    const unsigned int nbx = domain_nx / bx;
+    gx = (blockIdx.z % nbx) * bx + x - halo;
+    gy = (blockIdx.z / nbx) * by + y - halo;
+}
+
 // ---------------------------------------------------------------------------
 // Neighbor list (periodic BCs via modular arithmetic)
 // ---------------------------------------------------------------------------

src/CelerisLab/lbm/kernels/operators/block_collision.cuh

@@ -0,0 +1,31 @@
+// Collision used by periodic/fine and slot/reservoir BlockStep.
+#ifndef CELERIS_OPERATORS_BLOCK_COLLISION_CUH
+#define CELERIS_OPERATORS_BLOCK_COLLISION_CUH
+
+#include "../core/params.cuh"
+#include "macro.cuh"
+#include "collision_srt.cuh"
+#if COLLISION_MODEL == 2
+#include "collision_mrt.cuh"
+#endif
+
+__device__ __forceinline__ void block_collide(float* f) {
+#if COLLISION_MODEL == 2
+    collide_mrt_unforced(f);  // SRT path below unchanged.
+#else
+    float rho, ux, uy, feq[NQ];
+    compute_rho_u(f, rho, ux, uy);
+#if USE_DDF_SHIFTING
+    // Preserve the density perturbation before adding the unit reference.
+    const float drho = f[0]+f[1]+f[2]+f[3]+f[4]+f[5]+f[6]+f[7]+f[8];
+    const float jx = f[1]-f[2]+f[5]-f[6]+f[7]-f[8];
+    const float jy = f[3]-f[4]+f[5]-f[6]-f[7]+f[8];
+    compute_feq_from_momentum(drho, jx, jy, ux, uy, feq);
+#else
+    compute_feq(rho, ux, uy, feq);
+#endif
+    collide_srt_no_force(f, feq, d_params.omega);
+#endif
+}
+
+#endif

src/CelerisLab/lbm/kernels/operators/bounce_back_swap.cuh

@@ -0,0 +1,14 @@
+// CelerisLab – operators/bounce_back_swap.cuh
+// Full-way bounce-back at a wall node: swap every opposite pair (paired
+// order q, q+1). Needs only NQ; shared by the uniform and blocked steps.
+#ifndef CELERIS_OPERATORS_BOUNCE_BACK_SWAP_CUH
+#define CELERIS_OPERATORS_BOUNCE_BACK_SWAP_CUH
+
+__device__ __forceinline__ void bounce_back_swap(float* f) {
+    #pragma unroll
+    for (int i = 1; i < NQ; i += 2) {
+        float t = f[i]; f[i] = f[i+1]; f[i+1] = t;
+    }
+}
+
+#endif

src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh

@@ -114,6 +114,16 @@ __device__ __forceinline__ void collide_mrt_no_force(float* __restrict__ g,
     collide_mrt(g, rho, ux, uy, Fin, omega);
 }
 
+// Unforced MRT with the clamped device omega, as in collide_dispatch with zero
+// forcing. Shared by the blocked, channel-root and dense-patch collisions.
+__device__ __forceinline__ void collide_mrt_unforced(float* __restrict__ g)
+{
+    float rho, ux, uy;
+    compute_rho_u(g, rho, ux, uy);
+    collide_mrt_no_force(g, rho, ux, uy,
+        fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega)));
+}
+
 #elif LATTICE_MODEL == LATTICE_D3Q19
 
 __device__ __forceinline__ void collide_mrt(float* __restrict__ g,

src/CelerisLab/lbm/kernels/operators/helpers.cuh

@@ -5,12 +5,7 @@
 #ifndef CELERIS_OPERATORS_HELPERS_CUH
 #define CELERIS_OPERATORS_HELPERS_CUH
 
-__device__ __forceinline__ void bounce_back_swap(float* f) {
-    #pragma unroll
-    for (int i = 1; i < NQ; i += 2) {
-        float t = f[i]; f[i] = f[i+1]; f[i+1] = t;
-    }
-}
+#include "bounce_back_swap.cuh"
 
 __device__ __forceinline__ void collide_dispatch(
     float* f, float rho_n,

src/CelerisLab/lbm/kernels/refinement.cu

@@ -9,6 +9,7 @@
 #include "core/params.cuh"
 #include "core/flags.cuh"
 #include "core/obs.cuh"
+#include "core/cell_kind.cuh"
 #include "operators/macro.cuh"
 #include "operators/collision_srt.cuh"
 #if COLLISION_MODEL == 2
@@ -35,96 +36,5 @@ __device__ unsigned int incoming_counts[NQ * TOTAL_CELLS];
 }
 #endif
 
-__device__ __forceinline__ void patch_collide(float* f, unsigned long k) {
-#if COLLISION_MODEL == 2
-    float rho, ux, uy;
-    compute_rho_u(f, rho, ux, uy);
-    collide_mrt_no_force(f, rho, ux, uy,
-        fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega)));
-#else
-    float rho, ux, uy, feq[NQ];
-    compute_rho_u(f, rho, ux, uy);
-    compute_feq(rho, ux, uy, feq);
-    collide_srt_no_force(f, feq, d_params.omega);
-#endif
-#ifdef CELERIS_REFINEMENT_AUDIT
-    ++collision_counts[k];
-#endif
-}
-
-extern "C" __global__ void PatchExplode(
-    const float* coarse, float* fine, const uint32_t* targets,
-    const uint32_t* donors, unsigned int count, unsigned long coarse_cells)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    for (unsigned int q = 0; q < NQ; ++q)
-        fine[index_f(targets[row], q)] = coarse[q * coarse_cells + donors[row]];
-#ifdef CELERIS_REFINEMENT_AUDIT
-    ++explosion_counts[targets[row]];
-#endif
-}
-
-extern "C" __global__ void PatchFineStep(const float* in, float* out, const uint16_t* flags) {
-    unsigned int x, y; unsigned long k;
-    index_from_thread(x, y, k);
-    if (x >= NX || y >= NY) return;
-    float f[NQ];
-    // Outer halo is invalid after EACH hop, including when reusing the buffer
-    // that held explosion data. Inner halo only streams; it never collides.
-    if (x < 2 || y < 2 || x >= NX - 2 || y >= NY - 2
-            || (flags != nullptr && is_solid(flags[k]))) {
-        for (int q = 0; q < NQ; ++q) f[q] = NAN;
-    } else {
-        unsigned long j[NQ];
-        compute_neighbors(x, y, j);  // These interior targets never wrap.
-        stream_pull_load(k, f, in, j);
-        if (x >= 4 && y >= 4 && x < NX - 4 && y < NY - 4)
-            patch_collide(f, k);
-    }
-    stream_pull_store(k, f, out);
-}
-
-extern "C" __global__ void PatchCoarseStep(
-    const float* in, float* out, const float* fine,
-    const uint64_t* targets, const uint64_t* children, unsigned int count,
-    unsigned int x0, unsigned int y0, unsigned int x1, unsigned int y1,
-    unsigned int wall_aware)
-{
-    unsigned int x, y; unsigned long k;
-    index_from_thread(x, y, k);
-    if (x >= NX || y >= NY) return;
-    // Both buffers were poisoned at initialization; covered cells stay invalid.
-    if (x >= x0 && x < x1 && y >= y0 && y < y1) return;
-    if (wall_aware && (y == 0 || y == NY - 1)) return;
-    unsigned long j[NQ];
-    compute_neighbors(x, y, j);
-    float f[NQ];
-    stream_pull_load(k, f, in, j);
-    if (x >= x0 - 1 && x <= x1 && y >= y0 - 1 && y <= y1) {
-        for (unsigned int q = 1; q < NQ; ++q) {
-            uint64_t slot = index_f(k, q);
-            unsigned int lo = 0, hi = count;
-            while (lo < hi) {
-                unsigned int mid = lo + (hi - lo) / 2;
-                if (targets[mid] < slot) lo = mid + 1;
-                else hi = mid;
-            }
-            if (lo < count && targets[lo] == slot) {
-                const uint64_t* s = children + 4ul * lo;
-                f[q] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
-#ifdef CELERIS_REFINEMENT_AUDIT
-                ++incoming_counts[slot];
-#endif
-            }
-        }
-    }
-    // Includes the eight grazing-corner links with UNCOVERED coarse donors.
-    // All replacement arrivals must be in place before computing rho/u.
-    if (wall_aware) {
-        if (y == 1) reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity(x));
-        else if (y == NY - 2) reconstruct_top_reservoir(f, RHO);
-    }
-    patch_collide(f, k);
-    stream_pull_store(k, f, out);
-}
+// Kernels: patch collision helper, PatchExplode, PatchFineStep, PatchCoarseStep.
+#include "step/patch_step.cu"

src/CelerisLab/lbm/kernels/step/block_boundary.cu

@@ -0,0 +1,42 @@
+// CelerisLab/lbm/kernels/step/block_boundary.cu
+// Receiver-local curved-boundary staging and its raw-exchange accumulation.
+// Part of the blocks.cu compilation entry; relies on its includes and guards.
+
+// Static cut links, receiver-local staging: cell is the receiver's local core
+// cell, so the existing BFL writes the solid-source window image that this
+// block's next pull reads (no halo copy between this kernel and BlockStep).
+// surface_velocity (nullable, 2 floats per link): prescribed fluid velocity on
+// fixed geometry, same convention as CurvedBoundaryKernel. Optional raw rows
+// (jx,jy,torque,dm) are unique physical links, never body-observation slots.
+extern "C" __global__ void BlockCurvedBoundary(
+    fpxx* f, const uint32_t* block, const uint32_t* cell, const unsigned char* dir,
+    const float* q, const unsigned char* fallback, unsigned int count,
+    const float* surface_velocity, const float* rx, const float* ry, float* raw)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    float Uw = 0.0f, Vw = 0.0f;
+    if (surface_velocity != nullptr) {
+        const float u = surface_velocity[2 * row], v = surface_velocity[2 * row + 1];
+        if (u != 0.0f || v != 0.0f) { Uw = u; Vw = v; }
+    }
+    apply_bouzidi_link(dir[row], q[row], cell[row], f + (unsigned long)block[row] * NQ * TOTAL_CELLS,
+                       Uw, Vw, rx ? rx[row] : 0.0f, ry ? ry[row] : 0.0f,
+                       fallback[row], nullptr, -1, raw ? raw + 4ul * row : nullptr,
+                       raw ? raw + 4ul * row + 3 : nullptr);
+}
+
+// Same-stream BFL -> this row-owned accumulation -> pull step. No atomics,
+// no cross-row reduction and no history allocation proportional to substeps.
+extern "C" __global__ void BlockAccumulateExchange(
+    const float* raw, double* integral, unsigned int count, double h, double dt)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const unsigned long k = 4ul * row;
+    const double impulse = h * dt;
+    integral[k] += impulse * (double)raw[k];
+    integral[k + 1] += impulse * (double)raw[k + 1];
+    integral[k + 2] += h * impulse * (double)raw[k + 2];
+    integral[k + 3] += h * h * (double)raw[k + 3];
+}

src/CelerisLab/lbm/kernels/step/block_channel.cu

@@ -1,5 +1,12 @@
+#if BLOCK_ROOT_BC != 2
+#error "BlockChannelStep requires BLOCK_ROOT_BC == 2"
+#endif
+
 // Root-only blocked counterpart of OneStep's D2Q9/SRT channel closure.
 // NX/NY are LOCAL strides; all physical boundary decisions use global coordinates.
+#include "../core/cell_kind.cuh"
+#include "../operators/bounce_back_swap.cuh"
+#include "../boundary/bounce_back.cuh"
 #include "../boundary/inlet/common.cuh"
 #include "../boundary/inlet/zou_he_local.cuh"
 #include "../boundary/outlet/pressure_neq.cuh"
@@ -15,15 +22,13 @@ extern "C" __global__ void BlockChannelStep(
     unsigned int x, y; unsigned long k;
     index_from_thread(x, y, k);
     if (x < halo || y < halo || x >= NX - halo || y >= NY - halo) return;
-    const unsigned int bx = NX - 2 * halo, by = NY - 2 * halo;
-    const unsigned int nbx = d_block_channel.nx / bx;
-    const unsigned int gx = (blockIdx.z % nbx) * bx + x - halo;
-    const unsigned int gy = (blockIdx.z / nbx) * by + y - halo;
+    unsigned int gx, gy;
+    block_global_xy(x, y, halo, d_block_channel.nx, gx, gy);
     const unsigned long base = (unsigned long)blockIdx.z * NQ * TOTAL_CELLS;
     const unsigned long cell = (unsigned long)blockIdx.z * TOTAL_CELLS + k;
     float f[NQ];
     if (kind != nullptr && kind[cell] == KIND_OUTER) {
-        for (int q = 0; q < NQ; ++q) f[q] = NAN;
+        fill_invalid(f);
     } else {
         unsigned long j[NQ];
         compute_neighbors(x, y, j);
@@ -31,35 +36,21 @@ extern "C" __global__ void BlockChannelStep(
         const bool wall = gy == 0 || gy == d_block_channel.ny - 1;
         const bool outlet = gx == d_block_channel.nx - 1;
         if (wall) {
-            // Wall owns all four corners. Ghosts pull and swap, never collide.
-            for (int q = 1; q < NQ; q += 2) {
-                const float t = f[q]; f[q] = f[q + 1]; f[q + 1] = t;
-            }
+            bounce_back_swap(f);  // Wall owns all four corners; ghosts never collide.
         } else if (gx == 0) {
             apply_zou_he_left_velocity_inlet_d2q9(
                 f, inlet_target_u((float)gy, d_block_channel.ny, d_block_channel.velocity), 0.0f);
         } else if (outlet) {
             // Immutable input donor, receiver-local (also valid across a block seam).
             float donor[NQ];
-            for (int q = 0; q < NQ; ++q)
-                donor[q] = load_ddf(in + base, index_f(k - 1, q));
+            load_node(in + base, k - 1, donor);
             apply_pressure_outlet_d2q9(f, donor);
-        } else if (gy == 1) {
-            f[3] = load_ddf(in + base, index_f(k, 4));
-            f[5] = load_ddf(in + base, index_f(k, 6));
-            f[8] = load_ddf(in + base, index_f(k, 7));
-        } else if (gy == d_block_channel.ny - 2) {
-            f[4] = load_ddf(in + base, index_f(k, 3));
-            f[6] = load_ddf(in + base, index_f(k, 5));
-            f[7] = load_ddf(in + base, index_f(k, 8));
+        } else {
+            apply_wall_bb_d2q9_rows(gy, d_block_channel.ny - 2, f, in + base, k);
         }
         if (!wall && !outlet) {
 #if COLLISION_MODEL == 2
-            // Match uniform MRT arithmetic (collide_dispatch, zero forcing).
-            float rho, ux, uy;
-            compute_rho_u(f, rho, ux, uy);
-            collide_mrt_no_force(f, rho, ux, uy,
-                fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega)));
+            collide_mrt_unforced(f);
 #else
             // Match uniform SRT arithmetic, including its shifted equilibrium path.
             float rho, ux, uy, feq[NQ], forcing[NQ] = {};

src/CelerisLab/lbm/kernels/step/block_step.cu

@@ -0,0 +1,35 @@
+// CelerisLab/lbm/kernels/step/block_step.cu
+// Generic blocked pull/collide/store for periodic/fine and slot/reservoir levels.
+// Part of the blocks.cu compilation entry; relies on its includes and guards.
+#if BLOCK_ROOT_BC == 2
+#error "BlockStep is not built for the channel root (BLOCK_ROOT_BC == 2)"
+#endif
+
+// Pull/collide/store core cells only; halo cells of out are left untouched.
+// flags (block cell ids, nullable): solid cores store NaN and never collide.
+// kind (block cell ids, nullable = all ACTIVE): OUTER skips the pull and
+// stores NaN, INNER pulls without collision, ACTIVE pulls and collides.
+extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int halo,
+                                     const uint16_t* flags, const unsigned char* kind) {
+    unsigned int x, y; unsigned long k;
+    index_from_thread(x, y, k);
+    if (x < halo || y < halo || x >= NX - halo || y >= NY - halo) return;
+    const unsigned long base = (unsigned long)blockIdx.z * NQ * TOTAL_CELLS;
+    const unsigned long cell = (unsigned long)blockIdx.z * TOTAL_CELLS + k;
+    const unsigned char cell_kind = kind != nullptr ? kind[cell] : KIND_ACTIVE;
+    float f[NQ];
+    if (cell_kind == KIND_OUTER || (flags != nullptr && is_solid(flags[cell]))) {
+        fill_invalid(f);
+    } else {
+        unsigned long j[NQ];
+        compute_neighbors(x, y, j);  // Core neighbors stay inside the window: no wrap.
+        stream_pull_load(k, f, in + base, j);
+        if (cell_kind == KIND_ACTIVE) {
+#if BLOCK_ROOT_BC == 1
+            block_wall(f, x, y, halo);
+#endif
+            block_collide(f);
+        }
+    }
+    stream_pull_store(k, f, out + base);
+}

src/CelerisLab/lbm/kernels/step/block_transfer.cu

@@ -0,0 +1,78 @@
+// CelerisLab/lbm/kernels/step/block_transfer.cu
+// Same-level halo refresh and 2:1 explosion/coalescence between blocked levels.
+// Part of the blocks.cu compilation entry; relies on its includes and guards.
+
+// Halo refresh: block cell ids b*TOTAL_CELLS+c, sources are always core cells.
+extern "C" __global__ void BlockHaloCopy(
+    fpxx* f, const uint64_t* dst, const uint64_t* src, unsigned int count)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const uint64_t d = dst[row], s = src[row];
+    fpxx* to = f + (d / TOTAL_CELLS) * NQ * TOTAL_CELLS;
+    const fpxx* from = f + (s / TOTAL_CELLS) * NQ * TOTAL_CELLS;
+    for (unsigned int q = 0; q < NQ; ++q)
+        to[index_f(d % TOTAL_CELLS, q)] = from[index_f(s % TOTAL_CELLS, q)];
+}
+
+// Fixed-patch explosion into blocked fine storage: target = owner block cell
+// id of a patch-ghost cell; donors are parent block cell ids. Dense parent
+// storage is the one-window case with coarse_cells = full domain size.
+extern "C" __global__ void BlockExplode(
+    const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
+    unsigned int count, unsigned long coarse_cells)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const uint64_t t = target[row];
+    fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
+    const uint64_t d = donor[row];
+    const fpxx* from = coarse + (d / coarse_cells) * NQ * coarse_cells;
+    for (unsigned int q = 0; q < NQ; ++q)
+        to[index_f(t % TOTAL_CELLS, q)] = from[q * coarse_cells + d % coarse_cells];
+}
+
+// Experimental Schornbaum--Ruede (arXiv:1508.07982, Algorithm 2)
+// and Chen et al. (doi:10.1016/j.physa.2005.09.036) population interpolation.
+// Independent implementation of the mathematical stencil; no upstream source copied.
+// Stencil cells are owner cores; a disabled axis repeats its center twice.
+extern "C" __global__ void BlockExplodeProjected(
+    const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
+    unsigned int count, unsigned long coarse_cells,
+    const uint64_t* neighbors, const float* offsets)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const uint64_t t = target[row], d = donor[row];
+    fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
+    for (unsigned int q = 0; q < NQ; ++q) {
+        float v[4];
+        for (unsigned int a = 0; a < 4; ++a) {
+            const uint64_t n = neighbors[4ul * row + a];
+            v[a] = coarse[(n / coarse_cells) * NQ * coarse_cells
+                          + q * coarse_cells + n % coarse_cells];
+        }
+        float gx = .5f * (v[1] - v[0]), gy = .5f * (v[3] - v[2]);
+        const float ex = d_cx[q], ey = d_cy[q];
+        if (q != 0) {
+            const float along = (ex * gx + ey * gy) / (ex * ex + ey * ey);
+            gx -= ex * along;
+            gy -= ey * along;
+        }
+        to[index_f(t % TOTAL_CELLS, q)] =
+            coarse[(d / coarse_cells) * NQ * coarse_cells + q * coarse_cells + d % coarse_cells]
+            + offsets[2ul * row] * gx + offsets[2ul * row + 1] * gy;
+    }
+}
+
+// Targets are receiver-local SOURCE slots for the immediately following pull.
+// Keep child summation order identical to PatchCoarseStep; no collision here.
+extern "C" __global__ void BlockCoalesce(
+    const fpxx* fine, fpxx* coarse, const uint64_t* target, const uint64_t* children,
+    unsigned int count)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const uint64_t* s = children + 4ul * row;
+    coarse[target[row]] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
+}

src/CelerisLab/lbm/kernels/step/patch_step.cu

@@ -0,0 +1,95 @@
+// CelerisLab/lbm/kernels/step/patch_step.cu
+// Dense fixed-patch collision, explosion and the fine/coarse steps.
+// Part of the refinement.cu compilation entry; relies on its includes, guards
+// and the optional CELERIS_REFINEMENT_AUDIT counters declared there.
+
+__device__ __forceinline__ void patch_collide(float* f, unsigned long k) {
+#if COLLISION_MODEL == 2
+    collide_mrt_unforced(f);
+#else
+    float rho, ux, uy, feq[NQ];
+    compute_rho_u(f, rho, ux, uy);
+    compute_feq(rho, ux, uy, feq);
+    collide_srt_no_force(f, feq, d_params.omega);
+#endif
+#ifdef CELERIS_REFINEMENT_AUDIT
+    ++collision_counts[k];
+#endif
+}
+
+extern "C" __global__ void PatchExplode(
+    const float* coarse, float* fine, const uint32_t* targets,
+    const uint32_t* donors, unsigned int count, unsigned long coarse_cells)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    for (unsigned int q = 0; q < NQ; ++q)
+        fine[index_f(targets[row], q)] = coarse[q * coarse_cells + donors[row]];
+#ifdef CELERIS_REFINEMENT_AUDIT
+    ++explosion_counts[targets[row]];
+#endif
+}
+
+extern "C" __global__ void PatchFineStep(const float* in, float* out, const uint16_t* flags) {
+    unsigned int x, y; unsigned long k;
+    index_from_thread(x, y, k);
+    if (x >= NX || y >= NY) return;
+    float f[NQ];
+    // Outer halo is invalid after EACH hop, including when reusing the buffer
+    // that held explosion data. Inner halo only streams; it never collides.
+    if (x < 2 || y < 2 || x >= NX - 2 || y >= NY - 2
+            || (flags != nullptr && is_solid(flags[k]))) {
+        fill_invalid(f);
+    } else {
+        unsigned long j[NQ];
+        compute_neighbors(x, y, j);  // These interior targets never wrap.
+        stream_pull_load(k, f, in, j);
+        if (x >= 4 && y >= 4 && x < NX - 4 && y < NY - 4)
+            patch_collide(f, k);
+    }
+    stream_pull_store(k, f, out);
+}
+
+extern "C" __global__ void PatchCoarseStep(
+    const float* in, float* out, const float* fine,
+    const uint64_t* targets, const uint64_t* children, unsigned int count,
+    unsigned int x0, unsigned int y0, unsigned int x1, unsigned int y1,
+    unsigned int wall_aware)
+{
+    unsigned int x, y; unsigned long k;
+    index_from_thread(x, y, k);
+    if (x >= NX || y >= NY) return;
+    // Both buffers were poisoned at initialization; covered cells stay invalid.
+    if (x >= x0 && x < x1 && y >= y0 && y < y1) return;
+    if (wall_aware && (y == 0 || y == NY - 1)) return;
+    unsigned long j[NQ];
+    compute_neighbors(x, y, j);
+    float f[NQ];
+    stream_pull_load(k, f, in, j);
+    if (x >= x0 - 1 && x <= x1 && y >= y0 - 1 && y <= y1) {
+        for (unsigned int q = 1; q < NQ; ++q) {
+            uint64_t slot = index_f(k, q);
+            unsigned int lo = 0, hi = count;
+            while (lo < hi) {
+                unsigned int mid = lo + (hi - lo) / 2;
+                if (targets[mid] < slot) lo = mid + 1;
+                else hi = mid;
+            }
+            if (lo < count && targets[lo] == slot) {
+                const uint64_t* s = children + 4ul * lo;
+                f[q] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
+#ifdef CELERIS_REFINEMENT_AUDIT
+                ++incoming_counts[slot];
+#endif
+            }
+        }
+    }
+    // Includes the eight grazing-corner links with UNCOVERED coarse donors.
+    // All replacement arrivals must be in place before computing rho/u.
+    if (wall_aware) {
+        if (y == 1) reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity(x));
+        else if (y == NY - 2) reconstruct_top_reservoir(f, RHO);
+    }
+    patch_collide(f, k);
+    stream_pull_store(k, f, out);
+}

src/CelerisLab/lbm/kernels/streaming/pull_double_buffer.cuh

@@ -21,6 +21,14 @@
 //   i.e. the population traveling in direction i at node n
 //   arrived from the node at n - c_i  = neighbor in direction opp(i)
 // ---------------------------------------------------------------------------
+// load_node: all NQ populations of node n as stored (no streaming).
+__device__ __forceinline__ void load_node(const fpxx* __restrict__ fi, unsigned long n,
+                                          float* __restrict__ f)
+{
+    for (int q = 0; q < NQ; ++q)
+        f[q] = load_ddf(fi, index_f(n, q));
+}
+
 __device__ __forceinline__ void stream_pull_load(unsigned long n,
                                                  float* __restrict__ f,
                                                  const fpxx* __restrict__ fi,

src/CelerisLab/lbm/refinement_stepper.py

@@ -65,52 +65,8 @@ def _level_name(level):
     return ("base", "fine", "child")[level] if level < 3 else f"level {level}"
 
 
-def _legacy_levels(bounds, *, fine_block, base_block, child, child_block, child_rectangles_in_parent,
-                   fine_objects, jets, child_objects, child_jets):
-    """Normalise the original two/three-level keywords to per-level lists (original errors)."""
-    if child is not None and child_rectangles_in_parent is not None:
-        raise ValueError("child and child_rectangles_in_parent are mutually exclusive")
-    nested = child is not None or child_rectangles_in_parent is not None
-    bounds = tuple(bounds)
-    if not is_rectangle(bounds) and (fine_block is None or base_block is None):
-        raise ValueError("several patches require fine_block and base_block")
-    fine_objects, jets = tuple(fine_objects), tuple(jets)
-    child_objects, child_jets = tuple(child_objects), tuple(child_jets)
-    if not nested and (child_objects or child_jets):
-        raise ValueError("child_objects/child_jets require a nested child")
-    if base_block is not None and fine_block is None:
-        raise ValueError("base_block requires fine_block")
-    refinements, blocks = [bounds], [base_block, fine_block]
-    objects, level_jets = [(), fine_objects], [(), jets]
-    if nested:
-        if child is not None and not is_rectangle(bounds):
-            raise ValueError("scalar child requires one parent patch; use child_rectangles_in_parent")
-        if fine_block is None or base_block is None or child_block is None:
-            raise ValueError("a nested child needs fine_block, base_block and child_block")
-        if fine_objects or jets:
-            raise ValueError("a nested child supports bodies/Jets only on L2; use child_objects/child_jets")
-        refinements.append(tuple(child) if child is not None else tuple(child_rectangles_in_parent))
-        blocks.append(child_block)
-        objects.append(child_objects)
-        level_jets.append(child_jets)
-    return refinements, blocks, objects, level_jets
-
-
-def _initial_levels(level_g, coarse_g, fine_g, child_g):
-    """One tuple of per-level arrays from initialize()'s positional or legacy named form."""
-    if coarse_g is None and fine_g is None:
-        if child_g is not None and len(level_g) != 2:
-            raise TypeError("child_g needs exactly the coarse and fine arrays before it")
-        return tuple(level_g) + (() if child_g is None else (child_g,))
-    if level_g:
-        raise TypeError("pass level arrays positionally or as coarse_g/fine_g/child_g, not both")
-    if coarse_g is None or fine_g is None:
-        raise TypeError("coarse_g and fine_g are both required in the named form")
-    return (coarse_g, fine_g) + (() if child_g is None else (child_g,))
-
-
 def _check_levels(refinements, blocks, objects, jets):
-    """Candidate limits shared by the list and the legacy entry points."""
+    """Candidate limits of the level lists, checked before any device work."""
     depth = len(refinements) + 1
     if depth < 2:
         raise ValueError("at least one refinement level is required")
@@ -119,13 +75,13 @@ def _check_levels(refinements, blocks, objects, jets):
             raise ValueError(f"{name} needs one entry per level ({depth})")
     for level, block in enumerate(blocks):
         if block is not None and len(block) != 2:
-            raise ValueError(f"{_level_name(level)}_block must be (by, bx)")
+            raise ValueError(f"level_blocks[{level}] must be (by, bx)")
     if objects[0] or jets[0]:
         raise ValueError("base-level bodies/Jets are not supported")
     if any(blocks[level - 1] is not None and blocks[level] is None for level in range(1, depth)):
         raise ValueError("a blocked level requires blocked refined levels below it")
     if not is_rectangle(refinements[0]) and (blocks[0] is None or blocks[1] is None):
-        raise ValueError("several patches require fine_block and base_block")
+        raise ValueError("several L1 patches require level_blocks[0] and level_blocks[1]")
     if depth > 2:
         if any(block is None for block in blocks):
             raise ValueError("three or more levels must all be blocked")
@@ -134,7 +90,14 @@ def _check_levels(refinements, blocks, objects, jets):
 
 
 class FixedPatchStepper:
-    """D2Q9/SRT/FP32, one unforced internal 2:1 patch in an x-periodic box.
+    """D2Q9 SRT/MRT FP32 fixed 2:1 refinement hierarchy, built from level lists.
+
+    ``refinements[l]`` gives level l+1 in level-l integer cells: one rectangle
+    (x0, y0, x1, y1) (FixedPatch2D, a fine window with its own origin) or a
+    sequence of rectangles (FixedPatchSet2D, a dependency-validated union on one
+    level-wide grid, public fine arrays (2*ny, 2*nx), NaN outside ACTIVE cells);
+    see lbm.hierarchy. ``level_blocks``, ``level_objects`` and ``level_jets``
+    hold one entry per level, L0 included (L0 bodies/Jets are not supported).
 
     ``coarse_boundary`` is ``periodic`` (default) or ``slot_reservoir``:
     a bottom on-site slot and top fixed-density reservoir from ``cfg``.
@@ -145,68 +108,55 @@ class FixedPatchStepper:
     actual channel closure footprint. Channel remapping is not supported.
     Periodic mode does not apply the wall/slot boundary fields in ``cfg``.
     Outer walls remain coarse; fine halo donors must be ordinary fluid.
-    ``fine_objects`` accepts fixed circles in coarse physical coordinates;
-    velocities are in coarse lattice units (dx=dt=1). Independent local copies
-    are projected at fine cell centers. Their complete stencil must be active
-    and disjoint from the actual two-hop transfer footprint.
-    These use the edge/cell-center frame (coarse centers I+.5), whereas ordinary Simulation objects use the integer-node frame.
 
-    ``initialize(coarse_g, fine_g)`` accepts physical Q-major (9,ny,nx) arrays
-    already in postcollision phase. It does not collide or build channel flags.
-    Every ``step()`` publishes only after both levels complete on one stream.
-    On a device/launch failure the instance becomes unusable: close and recreate
-    it, rather than pretending the partially advanced fine buffer rolled back.
+    Bodies are fixed circles in root (L0) physical coordinates; velocities are
+    in root lattice units (dx=dt=1). Independent local copies are projected at
+    fine cell centers; their complete stencil must be active and disjoint from
+    the actual two-hop transfer footprint. These use the edge/cell-center frame
+    (coarse centers I+.5), whereas ordinary Simulation objects use the
+    integer-node frame. Jets (body.jets.Jet) use the same frame and drive cut
+    links selected at their true hits; ``set_jet_amplitude(index, amplitude,
+    level=l)`` re-uploads the per-link velocity column only.
+
+    ``initialize(*level_g)`` accepts one physical Q-major (9, ny, nx) array per
+    level, already in postcollision phase. It does not collide or build channel
+    flags. Every ``step()`` publishes only after all levels complete on one
+    stream. On a device/launch failure the instance becomes unusable: close and
+    recreate it, rather than pretending a partially advanced buffer rolled back.
+
+    A ``level_blocks`` entry (by, bx) stores that level in blocks.cu storage:
+    the ENTIRE level shape, including the four-cell patch ghost band, is
+    partitioned by a non-periodic BlockLayout2D with one-cell same-level halos.
+    Patch kinds stay per cell (OUTER invalid, INNER stream only, ACTIVE
+    collide); halos are gathered before each hop. Public arrays keep the dense
+    shapes; only the allocation differs. A blocked level requires blocked
+    refined levels below it; several L1 rectangles and depth >= 3 require all
+    levels blocked. Blocked bodies are stationary circles, optionally
+    surface-actuated: prepare_fine_boundary decides legality on the dense grid,
+    then cut links are staged receiver-local per hop: halo -> BFL -> step, same
+    buffer. A blocked parent gathers halos before receiver-local coalescence
+    staging and pull/collision; covered cells are poisoned on every step.
 
-    ``fine_block=(by, bx)`` (opt-in) stores the fine level in
-    blocks.cu storage: the ENTIRE fine_shape, including the four-cell patch
-    ghost band, is partitioned by a non-periodic BlockLayout2D with one-cell
-    same-level halos. Patch kinds stay per cell (OUTER invalid, INNER stream
-    only, ACTIVE collide); halos are gathered before each fine hop. Public
-    arrays keep the dense (9, ny, nx) fine shape; only the allocation differs.
-    Blocked fine_objects are stationary circles, optionally surface-actuated:
-    prepare_fine_boundary decides legality on the dense fine grid, then cut
-    links are staged receiver-local per hop: halo -> BFL -> step, same buffer.
-    ``base_block=(by, bx)`` also blocks the parent; it requires ``fine_block``.
-    Parent halos precede receiver-local coalescence staging and pull/collision;
-    covered cells are explicitly poisoned on every step. Public shapes stay dense.
     ``observe_raw_exchange=True`` records blocked physical cut-link exchange
     without changing populations. read_raw_exchange reports the last complete
     base interval or last substep in base-lattice units per unit depth. This
     is boundary momentum reaction, not a validated nozzle/stress load. Dense
-    body layouts are not supported by this new diagnostic in this candidate.
+    body levels are not supported by this diagnostic.
     ``observe_window_exchange=True`` (requires observe_raw_exchange) adds a second
     per-link FP64 ledger accumulated by the same kernel between
     start_exchange_window() and read_exchange_window(); disabled by default.
-    ``jets`` (body.jets.Jet) use the same edge frame as ``fine_objects`` and
-    drive fine cut links selected at their true hits; ``set_jet_amplitude``
-    re-uploads the per-link velocity column only.
 
-    Several patches: pass ``bounds`` as a sequence of (x0, y0, x1, y1)
-    rectangles (FixedPatchSet2D, dependency-validated union). This requires
-    fine_block and base_block; fine_objects/jets are projected once onto the
-    level-wide fine grid (each circle inside the union, body ids = object
-    positions, jet and writer checks level-wide), and
-    uses one level-wide fine grid (2*ny, 2*nx): public fine arrays have that
-    shape, NaN outside the patches' ACTIVE cells. The step sequence is the
-    blocked one above, with unique union transfers in one launch each.
-
-    Three-level separated unions: ``child_rectangles_in_parent=[...]`` is
-    mutually exclusive with scalar ``child`` and accepts either parent form.
-    Coordinates index the allocated L1 array, including ghost bands and holes:
-    root edge = patch.fine_origin + L1 index / 2. A parent set has origin zero;
-    a scalar parent retains its shifted window origin. The child set's public
-    extent is twice the allocated L1 shape, even though storage is sparse.
-    Each pair requires gap >= 4 in its own parent-cell units. All regions share
-    one buffer pair per level and the fixed 1:2:4 schedule; L2-only bodies/Jets
-    below apply unchanged, with object IDs assigned over the whole level.
-
-    Three-level single chain: ``child`` bounds use L1 window cells; all three
-    levels must be blocked. ``child_objects`` and ``child_jets`` use the L0
-    edge frame, with L2 dx=dt=.25 and unchanged lattice velocity. L1 bodies
-    and Jets remain unsupported in a chain. Every whole circle and its BFL
-    stencil must fit strictly away from the L2 transfer footprint. L2 runs
-    halo -> BFL -> step four times per base step. Runtime child Jet changes
-    require ``set_jet_amplitude(index, amplitude, level=2)``.
+    Nested levels: a child set's coordinates index the allocated parent array,
+    including ghost bands and holes: root edge = parent fine_origin + index / 2
+    (zero origin for a parent set; a scalar parent keeps its shifted window
+    origin). A scalar child requires a scalar parent. Rectangle unions impose no
+    fixed gap rule: adjacent, overlapping and concave unions are accepted when
+    their actual transfer dependencies close (FixedPatchSet2D /
+    nest_patch_set). All regions share one buffer pair per level and the fixed
+    schedule of 2**l hops per base step at level l. Bodies/Jets are supported only on the
+    finest level of a nested hierarchy (L2 dx=dt=.25 in root units, unchanged
+    lattice velocity); every whole circle and its BFL stencil must fit strictly
+    away from that level's transfer footprint.
 
     ``explosion="projected_linear"`` selects experimental transverse linear
     reconstruction on each adjacent pair. Default "constant" is unchanged.
@@ -214,43 +164,21 @@ class FixedPatchStepper:
     bodies/Jets or LES; this backend supplies no forcing. Geometry is generic,
     but union/multilevel GPU and matched-physics qualification remain pending.
 
-    Level lists (any depth): ``refinements[l]`` gives level l+1 in level-l
-    integer cells (scalar window or rectangle list, see lbm.hierarchy);
-    ``level_blocks``/``level_objects``/``level_jets`` hold one entry per level,
-    L0 included. They are mutually exclusive with the legacy keywords above,
-    which are normalised to the same lists. One recursive hop advances every
-    level: a child explodes from its parent's current hop source, takes two
-    hops, then the parent gathers halos, stages restriction, applies BFL and
-    steps. Depth >= 3 must be fully blocked with bodies/Jets on the finest level.
+    One recursive hop advances every level: a child explodes from its parent's
+    current hop source, takes two hops, then the parent gathers halos, stages
+    restriction, applies BFL and steps.
     """
-    def __init__(self, cfg, bounds=None, *, context, coarse_boundary="periodic", fine_objects=(),
-                 fine_block=None, base_block=None, jets=(), child=None, child_block=None,
-                 child_objects=(), child_jets=(), child_rectangles_in_parent=None,
-                 refinements=None, level_blocks=None, level_objects=None, level_jets=None,
+    def __init__(self, cfg, *, context, refinements, coarse_boundary="periodic",
+                 level_blocks=None, level_objects=None, level_jets=None,
                  observe_raw_exchange=False, observe_window_exchange=False, explosion="constant"):
         if HALO != 4:
             raise RuntimeError("refinement.cu requires a four-cell fine halo")
         if coarse_boundary not in ("periodic", "slot_reservoir", "channel"):
             raise ValueError("coarse_boundary must be periodic, slot_reservoir or channel")
-        if refinements is None:
-            if bounds is None:
-                raise ValueError("bounds or refinements is required")
-            if level_blocks is not None or level_objects is not None or level_jets is not None:
-                raise ValueError("level_blocks/level_objects/level_jets require refinements")
-            refinements, level_blocks, level_objects, level_jets = _legacy_levels(
-                bounds, fine_block=fine_block, base_block=base_block, child=child,
-                child_block=child_block, child_rectangles_in_parent=child_rectangles_in_parent,
-                fine_objects=fine_objects, jets=jets, child_objects=child_objects,
-                child_jets=child_jets)
-        else:
-            if (bounds is not None or child is not None or child_rectangles_in_parent is not None
-                    or fine_block is not None or base_block is not None or child_block is not None
-                    or tuple(fine_objects) or tuple(jets) or tuple(child_objects) or tuple(child_jets)):
-                raise ValueError("refinements and the legacy patch keywords are mutually exclusive")
-            depth = len(refinements) + 1
-            level_blocks = [None] * depth if level_blocks is None else list(level_blocks)
-            level_objects = [()] * depth if level_objects is None else list(level_objects)
-            level_jets = [()] * depth if level_jets is None else list(level_jets)
+        depth = len(refinements) + 1
+        level_blocks = [None] * depth if level_blocks is None else list(level_blocks)
+        level_objects = [()] * depth if level_objects is None else list(level_objects)
+        level_jets = [()] * depth if level_jets is None else list(level_jets)
         level_objects = [tuple(o) for o in level_objects]
         level_jets = [tuple(j) for j in level_jets]
         level_blocks = [None if b is None else tuple(b) for b in level_blocks]
@@ -434,7 +362,10 @@ class FixedPatchStepper:
             raise RuntimeError("level omega/tau mismatch")
         host = self._boundary_hosts[index]
         n_objects = len(host.objects) if name == "PatchFineStep" and host is not None else 0
-        build = compiler.generate_config(config, n_objects=n_objects)
+        block_root_bc = 0
+        if index == 0 and layout is not None:
+            block_root_bc = 2 if self._channel else 1 if self._wall_aware else 0
+        build = compiler.generate_config(config, n_objects=n_objects, block_root_bc=block_root_bc)
         try:
             ptx = compiler.compile_kernel(self.ctx.sm_arch, source=source, build_dir=build)
             module = compiler.load_module(ptx)
@@ -540,13 +471,13 @@ class FixedPatchStepper:
         if self._ddf_shifting and any(layout is None for layout in self._layouts):
             raise ValueError("shifted populations require all-blocked levels")
 
-    def initialize(self, *level_g, coarse_g=None, fine_g=None, child_g=None):
+    def initialize(self, *level_g):
         """Upload physical postcollision G, encoding before any FP32 rounding.
 
         Shifted input should be physical FP64 to retain small perturbations.
         Public initialization resets the step count; non-leaves are discarded.
         """
-        given = _initial_levels(level_g, coarse_g, fine_g, child_g)
+        given = tuple(level_g)
         with self._lock, self.ctx.activate():
             self._require("new", "ready")
             self._check_storage()
@@ -670,10 +601,9 @@ class FixedPatchStepper:
                 old, new, source, base_count=count, level_steps=(count, 2*count, 4*count),
                 phase="postcollision_g", **options)
             candidate = type(self)(
-                cfg, pairs[0].bounds, context=self.ctx, base_block=base_block,
-                fine_block=fine_block, child_block=child_block,
-                child_rectangles_in_parent=new.child.bounds,
-                child_objects=child_objects, child_jets=jets,
+                cfg, context=self.ctx, refinements=[pairs[0].bounds, new.child.bounds],
+                level_blocks=[base_block, fine_block, child_block],
+                level_objects=[(), (), child_objects], level_jets=[(), (), jets],
                 observe_raw_exchange=self._observe_raw_exchange,
                 observe_window_exchange=self._observe_window, explosion=self._explosion)
             try:
@@ -1019,8 +949,8 @@ class FixedPatchStepper:
             if exc[0] is None:
                 raise
 
-    # Read-only compatibility views over the level lists. Internal code must not use them;
-    # REMOVE AFTER the existing tests/docs migrate to hierarchy/_layouts/_level_* lists.
+    # Read-only compatibility views over the level lists, kept on purpose for the existing
+    # tests and docs. Internal code must not use them; it reads hierarchy/_layouts/_level_*.
     patch = property(lambda self: self.hierarchy.pairs[0])
     child_patch = property(lambda self: self.hierarchy.pairs[1] if self.hierarchy.depth > 2 else None)
     base_layout = property(lambda self: self._layouts[0])

src/CelerisLab/lbm/refinement_tables.py

@@ -8,7 +8,7 @@ from dataclasses import dataclass, replace
 import numpy as np
 
 from .blocks import blocked_link_tables
-from ..body.jets import jet_link_shapes, surface_link_velocity
+from ..body.jets import _surface_links
 from .descriptors import D2Q9_EX, D2Q9_EY
 from .refinement import ACTIVE, OUTER_GHOST
 
@@ -142,11 +142,13 @@ def _prepare_blocked_boundary(boundary, layout, cfg, jets, *, _shifted_blocked=F
     tables.update(body_id=np.asarray(body_id, np.int32), rx=np.asarray(rx, np.float32),
                   ry=np.asarray(ry, np.float32), aperture=np.zeros(len(dirs), bool))
     actuated = bool(jets) or any(o.surface_actuation is not None for o in boundary.objects)
+    jet_cache = None
     if actuated:
         cfg = replace(cfg)
         cfg.validate()
-        tables["surface_velocity"], tables["aperture"] = surface_link_velocity(
+        tables["surface_velocity"], tables["aperture"], shape, jet_id = _surface_links(
             boundary.objects, cfg, boundary.flags, fluid_idx, dirs, body_id, rx, ry,
             jets=jets, hits=boundary.curved_hits, _shifted_blocked=_shifted_blocked)
-    return _BlockedBoundary(boundary, tables, jets, actuated,
-                            jet_link_shapes(jets, boundary.curved_hits) if jets else None)
+        if jets:
+            jet_cache = (shape, jet_id)
+    return _BlockedBoundary(boundary, tables, jets, actuated, jet_cache)

src/CelerisLab/simulation.py

@@ -336,13 +336,19 @@ class Simulation:
     def read_surface_exchange(self, id: int):
         """Latest step's wall/aperture (Fx,Fy,Tz) exchange, not total nozzle force.
 
+        Supports SurfaceActuation and Jet-driven bodies, including zero-amplitude
+        apertures. Passive bodies without an aperture are rejected.
         Unlike read_body this is not accumulated/time-averaged. It resets after
         geometry sync or checkpoint load; advance before interpreting it.
         """
-        if self.bodies.get(id).surface_actuation is None:
+        obj = self.bodies.get(id)
+        curved = self.field.curved
+        has_aperture = (curved.aperture is not None
+                        and np.any(curved.aperture & (curved.body_id == id)))
+        if obj.surface_actuation is None and not has_aperture:
             raise ValueError("Body has no surface actuation")
         self.stream.synchronize()
-        return self.field.curved.surface_exchange(id)
+        return curved.surface_exchange(id)
 
     @context_bound
     def read_body(self, id: int, *,
@@ -451,18 +457,21 @@ class Simulation:
     @context_bound
     def initialize(self):
         """Initialize flow field and sync objects to GPU."""
-        # Recompile if objects were added after construction.
         self._assert_allocation_contract()
+        # Validate the complete geometry before flags, populations or modules
+        # change. Compact-list construction publishes no outputs on failure.
+        next_flags = self.bodies.build_flags(self.field.build_channel_flags())
+        if self.bodies.jets and self.bodies.count == 0:
+            raise ValueError("Jets need bodies with cut links")
+        self.bodies.build_compact_lists_for(self.bodies.objects, next_flags)
+
+        # Recompile if objects were added after construction.
         if self._read_generated_n_objects() != self.bodies.count:
             self.recompile()
 
-        self.field.flag = self.field.build_channel_flags()
-        if self.bodies.count > 0:
-            self.field.flag = self.bodies.build_flags(self.field.flag)
+        self.field.flag = next_flags
         self.field.upload_flags()
 
-        if self.bodies.jets and self.bodies.count == 0:
-            raise ValueError("Jets need bodies with cut links")
         self.stepper.initialize()
         if self.bodies.count > 0:
             self.bodies.sync_to_gpu(self.field, rebuild_flags=False)

tests/integration/test_checkpoint_bodies_gpu.py

@@ -0,0 +1,66 @@
+"""Checkpoint save -> load round trip for a scene with passive bodies. Requires GPU.
+
+A fresh Simulation with the same scene, restored from the checkpoint must hold the same populations,
+flags, body count and step count, and must stay bitwise identical to the original
+after both continue for the same number of steps. Checkpoints with Jets are
+refused (documented limitation) and are not covered here.
+"""
+
+import os
+import tempfile
+import unittest
+
+import numpy as np
+
+from ._cases import make_sim
+
+# Three separated cylinders (pinball-like); cut-link receivers stay in bulk fluid.
+CIRCLES = (((40, 32), 6), ((56, 22), 6), ((56, 42), 6))
+
+
+def scene():
+    sim = make_sim()
+    sim.initialize()
+    for center, radius in CIRCLES:
+        sim.add_body("circle", center=center, radius=radius)
+    sim.sync_bodies()
+    return sim
+
+
+class CheckpointBodiesGPU(unittest.TestCase):
+    def test_round_trip_is_bitwise_and_continues_identically(self):
+        sim = scene()
+        sim.run(60)
+        with tempfile.NamedTemporaryFile(suffix=".h5", delete=False) as f:
+            path = f.name
+        restored = None
+        try:
+            saved = sim.save_checkpoint(path)
+            # The loader requires the same body count and geometry: rebuild the scene.
+            restored = scene()
+            restored.load_checkpoint(saved)
+
+            self.assertEqual(restored.bodies.count, len(CIRCLES))
+            self.assertEqual(restored.stepper.step_count, sim.stepper.step_count)
+            np.testing.assert_array_equal(restored.get_flags(), sim.get_flags())
+            a, b = sim.get_ddf(), restored.get_ddf()
+            self.assertEqual(a.dtype, b.dtype)
+            self.assertEqual(a.tobytes(), b.tobytes())
+
+            sim.run(40)
+            restored.run(40)
+            self.assertEqual(restored.stepper.step_count, sim.stepper.step_count)
+            a, b = sim.get_ddf(), restored.get_ddf()
+            self.assertTrue(np.isfinite(a).all())
+            self.assertEqual(a.tobytes(), b.tobytes())
+            for i in range(len(CIRCLES)):
+                self.assertTrue(np.all(np.isfinite(restored.read_force(i))))
+        finally:
+            if restored is not None:
+                restored.close()
+            os.unlink(path)
+            sim.close()
+
+
+if __name__ == "__main__":
+    unittest.main()

tests/unit/cpp/channel_root.cpp

@@ -16,6 +16,7 @@
 #define NZ 1
 #define NT 1
 #define COLLISION_MODEL 0
+#define BLOCK_ROOT_BC 2
 #define USE_LES 0
 #define RHO 1.f
 #define U0 .025f
@@ -40,7 +41,7 @@ inline float __fdiv_rn(float a,float b){return a/b;}
 #include "boundary/inlet_outlet.cuh"
 #include "step/one_step_double.cu"
 #else
-constexpr unsigned char KIND_OUTER=2;
+#include "core/cell_kind.cuh"
 #include "step/block_channel.cu"
 #endif
 int main(){

tests/unit/historical_tree.py

@@ -0,0 +1,46 @@
+"""Check historical source manifests against their own recorded commit, not today's checkout."""
+
+import io
+from pathlib import Path
+import subprocess
+import tarfile
+
+ROOT = Path(__file__).resolve().parents[2]
+
+
+def _require(commit):
+    """Missing recorded history is an explicit failure, never a skip or a repin."""
+    found = subprocess.run(
+        ["git", "-C", str(ROOT), "cat-file", "-e", commit + "^{commit}"],
+        capture_output=True,
+    )
+    if found.returncode:
+        raise AssertionError(
+            f"Required history missing: commit {commit} is not in this clone; "
+            "fetch it instead of repinning or substituting"
+        )
+
+
+def tracked(commit, prefix="src"):
+    """Tracked paths under prefix at the recorded commit."""
+    _require(commit)
+    return set(
+        subprocess.check_output(
+            ["git", "-C", str(ROOT), "ls-tree", "-r", "--name-only", commit, prefix],
+            text=True,
+        ).splitlines()
+    )
+
+
+def historical_tree(commit, dest):
+    """Extract the tracked tree of the recorded commit into dest."""
+    _require(commit)
+    archive = subprocess.check_output(
+        ["git", "-C", str(ROOT), "archive", "--format=tar", commit]
+    )
+    dest = Path(dest)
+    dest.mkdir(parents=True, exist_ok=True)
+    extra = {"filter": "data"} if hasattr(tarfile, "data_filter") else {}
+    with tarfile.open(fileobj=io.BytesIO(archive)) as tar:
+        tar.extractall(dest, **extra)
+    return dest

tests/unit/test_block_step_sync.py

@@ -0,0 +1,104 @@
+"""Host contract of PeriodicBlockStepper.step(n): one stream wait per call (mocked driver)."""
+import ast
+from contextlib import nullcontext
+import importlib.util
+from pathlib import Path
+import types
+from unittest.mock import Mock
+import unittest
+
+import numpy as np
+
+from CelerisLab.config import LBMConfig
+from CelerisLab.lbm.blocks import BlockLayout2D
+
+
+def host_module():
+    """block_stepper with pycuda/compiler replaced by mocks; no GPU context."""
+    # Locate the source without importing it: importing block_stepper loads pycuda.driver,
+    # which CPU-only import-boundary tests (test_build_contracts) must never see.
+    path = Path(importlib.util.find_spec("CelerisLab.lbm.block_stepper").origin)
+    tree = ast.parse(path.read_text())
+    tree.body = [node for node in tree.body if not (
+        isinstance(node, ast.Import) and any(a.name == "pycuda.driver" for a in node.names)
+        or isinstance(node, ast.ImportFrom) and node.module == "cuda")]
+    module = types.ModuleType("CelerisLab.lbm._test_host_block_stepper")
+    module.__package__ = "CelerisLab.lbm"
+    module.__dict__.update(cuda=Mock(), compiler=Mock())
+    exec(compile(tree, str(path), "exec"), module.__dict__)
+    return module
+
+
+def backend():
+    host, log = host_module(), []
+    host.cuda.mem_get_info.return_value = (1 << 40, 1 << 40)
+    host.cuda.mem_alloc.side_effect = lambda size: Mock(size=size)  # distinct buffers
+    module = Mock()
+    module.get_global.return_value = (Mock(), 20)
+    module.get_function.side_effect = lambda name: Mock(
+        side_effect=lambda *a, **k: log.append((name, a[:2])))
+    host.compiler.load_module.return_value = module
+    cfg = LBMConfig(nx=16, ny=16, viscosity=.1)
+    run = host.PeriodicBlockStepper(cfg, BlockLayout2D.build(16, 16, 8, 8),
+                                    context=Mock(activate=nullcontext, sm_arch="sm_70"))
+    run._stream.synchronize.side_effect = lambda: log.append(("sync", ()))
+    run.initialize(np.full((9, 16, 16), 1 / 9))
+    log.clear()
+    return run, log
+
+
+class BlockStepSync(unittest.TestCase):
+    def test_step_n_waits_once_and_matches_n_single_steps(self):
+        many, log_many = backend()
+        many.step(3)
+        single, log_single = backend()
+        for _ in range(3):
+            single.step(1)
+        names = [name for name, _ in log_many]
+        self.assertEqual(names.count("sync"), 1)
+        self.assertEqual(names[-1], "sync")
+        self.assertEqual(names.count("BlockStep"), 3)
+        # Same launch order and buffer roles as three step(1) calls.
+        strip = lambda log: [(n, tuple(id(x) for x in a)) for n, a in log if n != "sync"]
+        roles = lambda run, log: [(n, tuple({id(run._current): "c", id(run._spare): "s"}.get(x, "-")
+                                            for x in ids)) for n, ids in strip(log)]
+        self.assertEqual([n for n, _ in strip(log_many)], [n for n, _ in strip(log_single)])
+        self.assertEqual(roles(many, log_many), roles(single, log_single))
+        self.assertIsNot(many._current, many._spare)
+        self.assertEqual((many.step_count, single.step_count), (3, 3))
+
+    def test_step_zero_launches_and_waits_for_nothing(self):
+        run, log = backend()
+        run.step(0)
+        self.assertEqual((log, run.step_count), ([], 0))
+
+    def test_failure_keeps_entry_step_count(self):
+        run, log = backend()
+        run.step(1)
+        run._stream.synchronize.side_effect = RuntimeError("device fault")
+        with self.assertRaises(RuntimeError):
+            run.step(2)
+        self.assertEqual((run.step_count, run._status), (1, "failed"))
+
+    def test_launch_failure_mid_batch_fails_closed_without_waiting(self):
+        run, log = backend()
+        calls = []
+        step = run._step
+
+        def failing(*args, **kwargs):
+            calls.append(1)
+            if len(calls) == 2:
+                raise RuntimeError("launch failure")
+            return step(*args, **kwargs)
+
+        run._step = failing
+        with self.assertRaises(RuntimeError):
+            run.step(3)
+        self.assertNotIn("sync", [name for name, _ in log])
+        self.assertEqual((run.step_count, run._status), (0, "failed"))
+        with self.assertRaises(RuntimeError):
+            run.download()
+
+
+if __name__ == "__main__":
+    unittest.main()

tests/unit/test_build_contracts.py

@@ -83,6 +83,32 @@ assert abs(sum(compute_feq_d2q9(1.,0.,0.)) - 1.) < 1e-6
 
 
 class BuildContracts(unittest.TestCase):
+    def test_block_root_modes_are_rendered_and_recorded(self):
+        with tempfile.TemporaryDirectory() as root:
+            for mode in (0, 1, 2):
+                with self.subTest(mode=mode):
+                    build = compiler.generate_config(LBMConfig(), build_root=root, block_root_bc=mode)
+                    try:
+                        method = Path(build, "config/config_method.h").read_text()
+                        self.assertIn(f"#define BLOCK_ROOT_BC {mode}\n", method)
+                        metadata = json.loads(Path(build, ".celeris-build").read_text())
+                        self.assertEqual(metadata["block_root_bc"], mode)
+                        self.assertNotIn("block_root_bc", metadata["config"])
+                    finally:
+                        compiler.cleanup_build_dir(build)
+            build = compiler.generate_config(LBMConfig(), build_root=root)
+            try:
+                self.assertEqual(json.loads(Path(build, ".celeris-build").read_text())["block_root_bc"], 0)
+            finally:
+                compiler.cleanup_build_dir(build)
+
+    def test_invalid_block_root_mode_fails_before_snapshot(self):
+        for mode in (-1, 3, True, False, 1.0, "1", None):
+            with self.subTest(mode=mode), patch.object(compiler.tempfile, "mkdtemp") as allocate:
+                with self.assertRaisesRegex(ValueError, "block_root_bc"):
+                    compiler.generate_config(LBMConfig(), block_root_bc=mode)
+                allocate.assert_not_called()
+
     @staticmethod
     def fake_compile(command, check):
         source = Path(command[-3])

tests/unit/test_channel_root.py

@@ -50,7 +50,7 @@ def test_masks_and_minimal_footprint():
 def test_actual_closure_read_conflicts():
     pair, layout, targets = transfer((6, 6, 18, 14))
     validate_channel_transfers(pair, layout, targets)
-    for target in (layout.slot(layout.owner[5, -1]-1, 2),
+    for target in (layout.slot(int(layout.owner[5, -1]) - 1, 2),  # integer under NumPy 1 and 2
                    layout.slot(layout.owner[1, 4], 4),
                    layout.slot(layout.owner[-2, 4], 3)):
         with pytest.raises(ValueError, match="immutable"):
@@ -80,8 +80,8 @@ def test_evolved_ghost_snapshot_and_physical_download():
     host.cuda.memcpy_dtoh.side_effect = lambda a, ptr: np.copyto(a, storage[id(ptr)])
     context = Mock(activate=nullcontext, sm_arch="sm_70")
     cfg = LBMConfig(nx=24, ny=20, viscosity=.1, velocity=.025)
-    backend = host.FixedPatchStepper(cfg, (6, 6, 18, 14), context=context,
-                                    coarse_boundary="channel", base_block=(4, 6), fine_block=(4, 4))
+    backend = host.FixedPatchStepper(cfg, context=context, coarse_boundary="channel",
+                                    refinements=[(6, 6, 18, 14)], level_blocks=[(4, 6), (4, 4)])
     try:
         weights = np.array([4/9]+[1/9]*4+[1/36]*4)[:, None, None]
         arrays = [np.broadcast_to(weights, level.shape).copy() for level in backend._levels]
@@ -97,5 +97,6 @@ def test_evolved_ghost_snapshot_and_physical_download():
         with pytest.raises(ValueError, match="finite"):
             backend.initialize(*arrays)
         assert backend._levels[0].kernel == host.compiler.load_module.return_value.get_function("BlockChannelStep")
+        assert [call.kwargs["block_root_bc"] for call in host.compiler.generate_config.call_args_list] == [2, 0]
     finally:
         backend.close()

tests/unit/test_dfg_re20.py

@@ -0,0 +1,79 @@
+"""CPU-only contract tests for run_dfg_re20. Not CUDA or DFG acceptance evidence."""
+
+import sys
+from pathlib import Path
+import unittest
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "validation"))
+import run_dfg_re20 as v  # noqa: E402
+import run_refined_dfg as rr  # noqa: E402
+
+h, _ = rr.helpers(rr.DEFAULT_SOURCE)
+
+
+def rows(values, tstar_end=v.TSTAR_MAX):
+    n = int(round(tstar_end / v.SAMPLE_TSTAR))
+    return [dict(tstar=(i + 1) * v.SAMPLE_TSTAR, **values(i)) for i in range(n)]
+
+
+class Re20Contract(unittest.TestCase):
+    def test_lattice_parameters_and_bounds(self):
+        root, uniform = v.lattice_parameters(40), v.lattice_parameters(80)
+        self.assertAlmostEqual(root["tau"], 0.56, places=12)
+        self.assertAlmostEqual(uniform["tau"], 0.62, places=12)
+        self.assertEqual((root["steps"], root["interval"]), (240000, 2000))
+        self.assertEqual((uniform["steps"], uniform["interval"]), (480000, 4000))
+        self.assertAlmostEqual(v.plan()["refined_fine_tau"], 0.62, places=12)
+
+    def test_config_only_changes_viscosity(self):
+        base = v._re100_config(h, 40, "SRT", "none")
+        cfg, body, p = v.config(h, 40, "SRT")
+        self.assertAlmostEqual(cfg["physics"]["viscosity"], 0.01 * 40 / 20, places=15)
+        self.assertEqual((p["re"], p["viscosity"]), (20.0, cfg["physics"]["viscosity"]))
+        base[0]["physics"]["viscosity"] = cfg["physics"]["viscosity"]
+        self.assertEqual(base[0], cfg)
+        self.assertEqual(base[1], body)
+        with self.assertRaisesRegex(ValueError, "no Jets"):
+            v.config(h, 40, "MRT", "zero")
+
+    def test_units(self):
+        s = v.sample([2.0, -0.5], 4, 40, 3.0)
+        den = rr.SPEED**2 * 40 / 2
+        self.assertAlmostEqual(s["cd"], 2.0 / 4 / den)
+        self.assertAlmostEqual(s["cl"], -0.5 / 4 / den)
+        self.assertAlmostEqual(s["p_diff"], 3.0 * 0.04)
+        self.assertEqual(s["dp_scaled"], 3.0)
+
+    def test_steady_and_verdict(self):
+        inside = dict(cd=5.58, cl=0.0106, p_diff=0.1174)
+        out = v.summarize(rows(lambda i: inside))
+        self.assertEqual((out["status"], out["steady"], out["window_samples"]),
+                         ("WITHIN_INTERVALS", True, 21))
+        self.assertTrue(out["cd_positive"])
+        self.assertAlmostEqual(out["edge_distance"]["cl"], 0.0002)
+        self.assertTrue(all(out["verdict"].values()))
+        low = v.summarize(rows(lambda i: dict(inside, cl=0.0100)))
+        self.assertEqual(low["status"], "OUTSIDE_INTERVALS")
+        self.assertFalse(low["verdict"]["cl"])
+
+    def test_slow_drift_gives_no_verdict(self):
+        # Range 2e-5 passes (limit 6e-5); slope * window 2e-5 exceeds 1.2e-5.
+        out = v.summarize(rows(lambda i: dict(cd=5.58, cl=0.0106 + 1e-6 * i, p_diff=0.1174)))
+        self.assertEqual((out["status"], out["verdict"]), ("NOT_STEADY", None))
+        self.assertLessEqual(out["ranges"]["cl"], out["range_limits"]["cl"])
+        self.assertGreater(out["drifts"]["cl"], out["slope_limits"]["cl"])
+
+    def test_noise_range(self):
+        quiet = v.summarize(rows(lambda i: dict(cd=5.58, cl=0.0106 + 2e-5 * (i % 2), p_diff=0.1174)))
+        self.assertTrue(quiet["steady"])  # within the range limit, no trend
+        noisy = v.summarize(rows(lambda i: dict(cd=5.58, cl=0.0106 + 1e-4 * (i % 2), p_diff=0.1174)))
+        self.assertEqual(noisy["status"], "NOT_STEADY")
+        self.assertGreater(noisy["ranges"]["cl"], noisy["range_limits"]["cl"])
+
+    def test_incomplete_run_gives_no_verdict(self):
+        out = v.summarize(rows(lambda i: dict(cd=5.58, cl=0.0106, p_diff=0.1174), tstar_end=30))
+        self.assertEqual((out["status"], out["verdict"]), ("INCOMPLETE", None))
+
+
+if __name__ == "__main__":
+    unittest.main()

tests/unit/test_jet_circle_cv.py

@@ -8,6 +8,7 @@ import numpy as np
 ROOT = Path(__file__).resolve().parents[2]
 REPO = ROOT
 sys.path.insert(0, str(ROOT / "tests/validation"))
+sys.path.insert(0, str(ROOT / "tests/unit"))
 sys.path.insert(0, str(REPO / "src"))
 from run_jet_circle_cv import (
     CX,
@@ -16,8 +17,10 @@ from run_jet_circle_cv import (
     case,
     state,
     cv,
+    cv_from_macros,
     aperture_probe,
     measured_aperture_flux,
+    PIN,
     verify,
 )
 
@@ -28,6 +31,39 @@ def eq(rho, u, v):
 
 
 class Host(unittest.TestCase):
+    def test_public_macros_pressure_advection_and_cropped_frames(self):
+        for h, origin, shape, rect in (
+            (.5, (0., 0.), (48, 64), (5., 4., 25., 20.)),
+            (.5, (.25, .5), (48, 64), (5.25, 4., 25.25, 20.)),
+            (.25, (.375, .75), (96, 128), (5.125, 4., 25.125, 20.)),
+        ):
+            y, x = np.indices(shape)
+            x = origin[0] + h * (x + .5)
+            k, u, v = 1e-5, .02, -.01
+            rho = 1 + 3 * k * x  # p=(rho-1)/3=k*x
+            ux, uy, fluid = np.full(shape, u), np.full(shape, v), np.ones(shape, bool)
+            result = cv_from_macros(rho, ux, uy, fluid, origin, rect, .004, h)
+            expected = 320 * np.array([3*k*u, k + 3*k*u*u, 3*k*u*v])
+            np.testing.assert_allclose(result[1], expected, rtol=1e-11, atol=1e-13)
+            np.testing.assert_allclose(result, cv(eq(rho, u, v), fluid, origin, rect, .004, h), atol=1e-12)
+            x0, y0, x1, y1 = np.rint((np.array(rect)-np.tile(origin, 2))/h).astype(int)
+            selection = np.s_[y0-2:y1+2, x0-2:x1+2]
+            cropped_origin = np.asarray(origin) + h*np.array([x0-2, y0-2])
+            cropped = cv_from_macros(rho[selection], ux[selection], uy[selection],
+                                    fluid[selection], cropped_origin, rect, .004, h)
+            np.testing.assert_array_equal(cropped, result)
+
+    def test_public_macros_full_normal_viscous_stress(self):
+        y, x = np.indices((48, 64))
+        x = (x + .5) * .5
+        rho, ux, uy, fluid = np.ones_like(x), 1e-4*x*x, np.zeros_like(x), np.ones(x.shape, bool)
+        args = (rho, ux, uy, fluid, (0., 0.), (5., 4., 25., 20.))
+        _, viscous = cv_from_macros(*args, .004)
+        _, inviscid = cv_from_macros(*args, 0.)
+        np.testing.assert_allclose(viscous-inviscid, [0, -4*.004*1e-4*320, 0], atol=1e-13)
+        with self.assertRaisesRegex(ValueError, "matching 2D"):
+            cv_from_macros(rho, ux, uy[:-1], fluid, (0, 0), (5, 4, 25, 20), .004)
+
     def test_descriptor_and_uniform_closed_flux(self):
         from CelerisLab.lbm.descriptors import D2Q9_EX, D2Q9_EY
 
@@ -132,7 +168,13 @@ class Host(unittest.TestCase):
         self.assertEqual(a["steps"] * 4, b["steps"])
 
     def test_source_pin(self):
-        verify(REPO)
+        from tempfile import TemporaryDirectory
+        from historical_tree import historical_tree
+
+        with self.assertRaisesRegex(RuntimeError, "Source hash mismatch"):
+            verify(REPO)
+        with TemporaryDirectory() as directory:
+            verify(historical_tree(PIN, directory))
 
 
 if __name__ == "__main__":

tests/unit/test_jets.py

@@ -188,6 +188,55 @@ def jet_case(device, jets):
     return field, bodies
 
 
+def test_public_surface_exchange_for_jet_preserves_aperture_at_zero_and_reverse(device):
+    from CelerisLab.simulation import Simulation
+    jet = Jet(Sector(CENTER, 0., .6, 3.), Radial(CENTER, inward=False), "quartic", 2e-3)
+    field, bodies = jet_case(device, [jet])
+    sim = SimpleNamespace(bodies=bodies, field=field, _initialized=True,
+                          stream=SimpleNamespace(synchronize=device.synchronize))
+    cl = field.curved
+    assert bodies.get(0).surface_actuation is None
+    aperture, velocity = cl.aperture, cl.surface_velocity.copy()
+    assert aperture.any() and (~aperture).any()
+    # Synthetic per-link exchanges test readback/splitting, not kernel physics.
+    values = np.arange(cl.count * 3, dtype=np.float32).reshape(-1, 3)
+    device.htod(cl.link_exchange_gpu, values)
+    for amplitude in (2e-3, 0., -1e-3):
+        Simulation.set_jet_amplitude(sim, 0, amplitude)
+        result = Simulation.read_surface_exchange(sim, 0)
+        assert cl.aperture is aperture
+        np.testing.assert_array_equal(aperture, cl.jet_id >= 0)
+        np.testing.assert_allclose(cl.surface_velocity, velocity * (amplitude / 2e-3))
+        for name, mask in (("wall", ~aperture), ("aperture", aperture)):
+            np.testing.assert_array_equal(result[name], values[mask].sum(axis=0, dtype=np.float64))
+
+
+def test_public_surface_exchange_keeps_legacy_and_passive_guards(device):
+    from CelerisLab.simulation import Simulation
+    field = make_field(device)
+    bodies = manager(field, [circle(SurfaceActuation(0., .6, .002), radius=2.)])
+    bodies.sync_to_gpu(field)
+    sim = SimpleNamespace(bodies=bodies, field=field,
+                          stream=SimpleNamespace(synchronize=device.synchronize))
+    for value in Simulation.read_surface_exchange(sim, 0).values():
+        np.testing.assert_array_equal(value, 0.)
+    field.curved.free_gpu_columns()
+    with pytest.raises(RuntimeError, match="diagnostic is allocated"):
+        Simulation.read_surface_exchange(sim, 0)
+
+    # An exchange buffer elsewhere in the scene must not admit a passive body.
+    for jets in ([], [Jet(Sector(CENTER, 0., .6, 3.), Radial(CENTER))]):
+        field, bodies = jet_case(device, jets)
+        passive = bodies.add(circle(center=(2., 3.), radius=.6))
+        bodies.sync_to_gpu(field)
+        sim.bodies, sim.field = bodies, field
+        with pytest.raises(ValueError, match="Body has no surface actuation"):
+            Simulation.read_surface_exchange(sim, passive)
+        if not jets:
+            with pytest.raises(ValueError, match="Body has no surface actuation"):
+                Simulation.read_surface_exchange(sim, 0)
+
+
 def test_manager_passes_true_hits_and_reuploads_only_velocity(device):
     # CPU memory double only: checks host/device bookkeeping, not GPU numerics.
     jet = Jet(Sector(CENTER, 0., .6, 3.), Radial(CENTER, inward=False), "quartic", 2e-3)

tests/unit/test_les_flow_runner.py

@@ -15,6 +15,10 @@ sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "validation"))
 from run_uniform_les_flow import source_guard, decode_macros, HERE, HEADER
 
 ROOT = HERE.parents[1]
+sys.path.insert(0, str(ROOT / "tests/unit"))
+from historical_tree import historical_tree, tracked
+
+A135 = "a135252ae78b47b57d7e77d19001700b3bedac8b"
 
 
 class GuardTests(unittest.TestCase):
@@ -42,19 +46,23 @@ class GuardTests(unittest.TestCase):
             decode_macros(invalid, w, cx, cy)
 
     def test_current_corrected_pins(self):
-        self.assertEqual(
-            source_guard(ROOT, "corrected", "a135252")["base_commit"],
-            "a135252ae78b47b57d7e77d19001700b3bedac8b",
-        )
+        with TemporaryDirectory() as directory:
+            root = historical_tree(A135, directory)
+            with patch("run_uniform_les_flow.subprocess.check_output", return_value="historical-tree\n"):
+                identity = source_guard(root, "corrected", "a135252")
+            # base_commit comes from the pins file; the mock only feeds checkout_commit.
+            self.assertEqual(identity["base_commit"], A135)
+            self.assertEqual(identity["checkout_commit"], "historical-tree")
 
     def test_corrected_and_tamper_guard(self):
         pins = json.loads((HERE / "data/les_flow/source_pins_a135252.json").read_text())
         with TemporaryDirectory() as directory:
-            root = Path(directory)
+            origin = historical_tree(A135, Path(directory) / "origin")
+            root = Path(directory) / "fixture"
             for name in pins["sources"]:
                 dest = root / name
                 dest.parent.mkdir(parents=True, exist_ok=True)
-                shutil.copyfile(ROOT / name, dest)
+                shutil.copyfile(origin / name, dest)
             with patch(
                 "run_uniform_les_flow.subprocess.check_output",
                 return_value="isolated-fixture\n",
@@ -87,23 +95,26 @@ class GuardTests(unittest.TestCase):
 
     def test_current_manifest_pins_actual_corrected_tree(self):
         pins = json.loads((HERE / "data/les_flow/source_pins_a135252.json").read_text())
-        tracked = subprocess.check_output(["git", "-C", str(ROOT), "ls-files", "src"], text=True)
-        self.assertEqual(set(pins["sources"]), set(tracked.splitlines()))
-        for name, expected in pins["sources"].items():
-            self.assertEqual(hashlib.sha256((ROOT / name).read_bytes()).hexdigest(), expected)
-        self.assertEqual(pins["sources"][HEADER], pins["candidate_header_sha256"])
-        self.assertNotEqual(pins["sources"][HEADER], pins["baseline_header_sha256"])
-        with self.assertRaisesRegex(RuntimeError, "Source pin mismatch"):
-            source_guard(ROOT, "baseline", "a135252")
+        self.assertEqual(pins["base_commit"], A135)
+        self.assertEqual(set(pins["sources"]), tracked(A135))
+        with TemporaryDirectory() as directory:
+            root = historical_tree(A135, directory)
+            for name, expected in pins["sources"].items():
+                self.assertEqual(hashlib.sha256((root / name).read_bytes()).hexdigest(), expected)
+            self.assertEqual(pins["sources"][HEADER], pins["candidate_header_sha256"])
+            self.assertNotEqual(pins["sources"][HEADER], pins["baseline_header_sha256"])
+            with self.assertRaisesRegex(RuntimeError, "Source pin mismatch"):
+                source_guard(root, "baseline", "a135252")
 
     def test_named_baseline_accepts_only_reversed_header(self):
         pins = json.loads((HERE / "data/les_flow/source_pins_a135252.json").read_text())
         with TemporaryDirectory() as directory:
-            root = Path(directory)
+            origin = historical_tree(A135, Path(directory) / "origin")
+            root = Path(directory) / "fixture"
             for name in pins["sources"]:
                 dest = root / name
                 dest.parent.mkdir(parents=True, exist_ok=True)
-                shutil.copyfile(ROOT / name, dest)
+                shutil.copyfile(origin / name, dest)
             subprocess.run(["git", "apply", "--reverse", str(HERE / "les/data/smagorinsky_normalization.patch")], cwd=root, check=True)
             with patch("run_uniform_les_flow.subprocess.check_output", return_value="baseline-fixture\n"):
                 source_guard(root, "baseline", "a135252")

tests/unit/test_projected_explosion.py

@@ -162,7 +162,7 @@ class ProjectedKernelTests(unittest.TestCase):
         targets = child.owner.ravel()[p.ghost_fine]
         donors = parent.owner.ravel()[p.ghost_coarse]
         source = (Path(__file__).resolve().parents[2] /
-                  "src/CelerisLab/lbm/kernels/blocks.cu").read_text()
+                  "src/CelerisLab/lbm/kernels/step/block_transfer.cu").read_text()
         kernel = source.split('extern "C" __global__ void BlockExplodeProjected(',1)[1]
         kernel = 'void BlockExplodeProjected(' + kernel.split('// Targets are receiver-local SOURCE slots',1)[0]
         def array(name, values, kind):

tests/unit/test_raw_exchange.py

@@ -157,8 +157,9 @@ class RawExchange(unittest.TestCase):
         host = host_module()
         ctx = Mock()
         with self.assertRaisesRegex(ValueError, 'blocked body'):
-            host.FixedPatchStepper(LBMConfig(nx=32, ny=32, viscosity=.1), (4, 4, 28, 28),
-                                   context=ctx, fine_objects=[body(radius=2.15)], observe_raw_exchange=True)
+            host.FixedPatchStepper(LBMConfig(nx=32, ny=32, viscosity=.1), context=ctx,
+                                   refinements=[(4, 4, 28, 28)],
+                                   level_objects=[(), (body(radius=2.15),)], observe_raw_exchange=True)
         ctx.activate.assert_not_called()
 
     def test_all_depth_weights(self):

tests/unit/test_re100.py

@@ -5,6 +5,7 @@ import pytest
 
 ROOT = Path(__file__).resolve().parents[2]
 sys.path.insert(0, str(ROOT / "tests/validation"))
+sys.path.insert(0, str(ROOT / "tests/unit"))
 import re100_helpers as h
 import periodic as p
 
@@ -62,13 +63,16 @@ def test_periodic_synthetic():
     assert not p.assess(rows()[:100])["periodic"]
 
 
-def test_production_unmodified():
+def test_production_unmodified(tmp_path):
+    from historical_tree import historical_tree
+
     m = json.loads(
         (ROOT / "tests/validation/data/dfg_re100/source_manifest.json").read_text()
     )
     assert m["commit"] == h.PIN
+    origin = historical_tree(m["commit"], tmp_path)
     for name, digest in m["production_sha256"].items():
-        assert hashlib.sha256((ROOT / name).read_bytes()).hexdigest() == digest, name
+        assert hashlib.sha256((origin / name).read_bytes()).hexdigest() == digest, name
 
 
 def test_reference():

tests/unit/test_refined_dfg.py

@@ -7,7 +7,10 @@ import numpy as np
 import unittest
 
 sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "validation"))
+sys.path.insert(0, str(Path(__file__).resolve().parent))
 import run_refined_dfg as r
+from historical_tree import historical_tree
+from tempfile import TemporaryDirectory
 
 h, periodic = r.helpers(r.DEFAULT_SOURCE)
 
@@ -17,11 +20,24 @@ def test_source_pin():
     assert historical["commit"].startswith("cd3a657")
     mrt_history = json.loads((r.HERE / "data/refined_dfg/source_manifest_mrt.json").read_text())
     assert mrt_history["commit"].startswith("3ea131e")
-    assert r.verify_source(r.DEFAULT_SOURCE, "MRT", "window")["commit"].startswith("a233c90")
+    with unittest.TestCase().assertRaisesRegex(RuntimeError, "Source mismatch"):
+        r.verify_source(r.DEFAULT_SOURCE, "MRT", "window")
     with unittest.TestCase().assertRaisesRegex(RuntimeError, "Source mismatch"):
         r.verify_source(r.DEFAULT_SOURCE, "MRT")
     with unittest.TestCase().assertRaisesRegex(RuntimeError, "Source mismatch"):
         r.verify_source(r.DEFAULT_SOURCE, "SRT")
+    with TemporaryDirectory() as directory:
+        for collision, mode, prefix in (
+            ("SRT", "last_base_step", "cd3a657"),
+            ("MRT", "last_base_step", "3ea131e"),
+            ("MRT", "window", "a233c90"),
+        ):
+            name = {"SRT": "source_manifest.json", "MRT": "source_manifest_mrt.json"}[collision]
+            name = "source_manifest_window.json" if mode == "window" else name
+            commit = json.loads((r.HERE / "data/refined_dfg" / name).read_text())["commit"]
+            assert commit.startswith(prefix)
+            root = historical_tree(commit, Path(directory) / commit)
+            assert r.verify_source(root, collision, mode)["commit"] == commit
 
 
 def test_matched_units_and_planes():
@@ -141,28 +157,62 @@ def test_axis_specific_sampling_phase():
     np.testing.assert_allclose(fine_y - ((k + 43) - 0.5) * 0.00125, 0, atol=1e-15)
 
 
+def frozen_manifest(observation_mode):
+    """Frozen MRT manifest JSON as recorded, without hashing the working tree."""
+    name = "source_manifest_window.json" if observation_mode == "window" else "source_manifest_mrt.json"
+    return json.loads((r.HERE / "data/refined_dfg" / name).read_text())
+
+
+def selected_manifest(collision, observation_mode):
+    """verify_source's manifest selection with its working-tree hash comparison neutralised.
+
+    Source identity is owned by test_source_pin; this exercises only which frozen manifest
+    verify_source selects and returns for (collision, observation_mode).
+    """
+    from unittest.mock import patch
+
+    class Digest(str):
+        def __ne__(self, other):
+            return False
+
+    class Hash:
+        def __init__(self, data=b""):
+            pass
+
+        def hexdigest(self):
+            return Digest()
+
+    with patch.object(r.hashlib, "sha256", Hash):
+        return r.verify_source(r.DEFAULT_SOURCE, collision, observation_mode)
+
+
 def test_full_requires_completed_pilot():
     with unittest.TestCase().assertRaisesRegex(ValueError, "pilot-results"):
-        r.verify_pilot(None, r.verify_source(r.DEFAULT_SOURCE, "MRT", "window"))
+        r.verify_pilot(None, frozen_manifest("window"))
 
 
-def pilot_fixture(path, observation_mode="window"):
+def pilot_fixture(path, observation_mode="window", jet_scenario="none", case_selection="both", uniform_diameter=80, placement="wide"):
     """Synthetic contract fixture only; never a physical result or run receipt."""
-    manifest = (r.verify_source(r.DEFAULT_SOURCE, "MRT", "window")
-                if observation_mode == "window" else
-                json.loads((r.HERE / "data/refined_dfg/source_manifest_mrt.json").read_text()))
+    # Receipt-logic fixture: the frozen manifest verify_source returns for a matching tree.
+    # Source identity itself is checked only by test_source_pin.
+    manifest = (r.verify_source(r.DEFAULT_SOURCE, "MRT", "window", jet_scenario)
+                if observation_mode == "window" and jet_scenario != "none" else
+                frozen_manifest(observation_mode))
+    selection = dict(case_selection=case_selection, uniform_diameter=uniform_diameter, placement=placement)
     (path / "plan.json").write_text(
         json.dumps(
             dict(
                 source_manifest=manifest,
+                **selection,
                 collision="MRT",
+                jet_scenario=jet_scenario,
                 observation_mode=observation_mode,
-                bounds_root_edge=list(r.BOUNDS),
+                bounds_root_edge=list(r.PLACEMENTS[placement]),
                 runner_sha256=r.runner_hash(),
             )
         )
     )
-    for case, diameter in (("uniform", 80), ("refined", 40)):
+    for case, diameter in r.selected_cases(case_selection, uniform_diameter):
         dest = path / case
         dest.mkdir()
         interval = round(0.02 * diameter / r.SPEED)
@@ -189,7 +239,9 @@ def pilot_fixture(path, observation_mode="window"):
             json.dumps(
                 dict(
                     case=case,
+                    **selection,
                     collision="MRT",
+                    jet_scenario=jet_scenario,
                     observation_mode=observation_mode,
                     mode="pilot",
                     status="pilot_complete",
@@ -206,7 +258,9 @@ def pilot_fixture(path, observation_mode="window"):
             json.dumps(
                 dict(
                     case=case,
+                    **selection,
                     collision="MRT",
+                    jet_scenario=jet_scenario,
                     observation_mode=observation_mode,
                     mode="pilot",
                     step=interval,
@@ -341,7 +395,7 @@ def test_record_actual_checkout_separately():
     actual = checkout_provenance(r.DEFAULT_SOURCE)
     assert len(actual["checkout_revision"]) == 40
     assert isinstance(actual["checkout_dirty"], bool)
-    assert r.verify_source(r.DEFAULT_SOURCE, "MRT", "window")["production_origin_commit"].startswith("d930")
+    assert selected_manifest("MRT", "window")["production_origin_commit"].startswith("d930")
 
 
 def test_window_plan_preserves_physics():
@@ -351,7 +405,7 @@ def test_window_plan_preserves_physics():
         assert baseline.pop("observation_mode") == "last_base_step"
         assert window.pop("observation_mode") == "window"
         assert baseline == window
-        assert r.verify_source(r.DEFAULT_SOURCE, collision, "window")["observation_mode"] == "window"
+        assert selected_manifest(collision, "window")["observation_mode"] == "window"
     with unittest.TestCase().assertRaises(ValueError):
         r.verify_source(r.DEFAULT_SOURCE, "MRT", "unknown")
 
@@ -451,6 +505,169 @@ def test_window_contract_rejection():
     assert not backend.events
 
 
+def test_public_jet_configuration_and_geometry():
+    for diameter, center, radius in ((40, (80.5, 81), 20), (80, (160, 160.5), 40)):
+        cfg, _, _ = r.config(h, diameter, "MRT", "fixed")
+        assert not cfg["method"]["ddf_shifting"] and not cfg["method"]["les"]["enabled"]
+        jets = r.public_jets(center, radius, "fixed")
+        assert [j.amplitude for j in jets] == [.001, -.001]
+        assert all(j.field.inward is False and j.profile == "parabolic" for j in jets)
+        assert all(j.region.r_max == radius * 1.005 for j in jets)
+        assert all(j.amplitude == 0 for j in r.public_jets(center, radius, "zero"))
+    cfg, _, _, host = r.host_geometry(r.DEFAULT_SOURCE, h, "MRT", "fixed")
+    assert not cfg.ddf_shifting and len(host.boundary_cells) > 0
+    assert r.public_jets((0, 0), 1, "none") == ()
+    with unittest.TestCase().assertRaises(ValueError):
+        r.config(h, 40, "SRT", "fixed")
+
+
+def test_public_jet_reversal_boundary_and_windows():
+    from types import SimpleNamespace
+
+    calls = []
+    handle = SimpleNamespace(set_jet_amplitude=lambda i, a: calls.append((i, a)))
+    for diameter in (40, 80):
+        boundary = round(250 * diameter / r.SPEED)
+        assert r.jet_interval(handle, "fixed", boundary - 3, 80, diameter) == (3, None)
+        switch = round(300 * diameter / r.SPEED)
+        assert r.jet_interval(handle, "reverse", switch - 3, 80, diameter) == (3, None)
+        assert not calls
+        n, event = r.jet_interval(handle, "reverse", switch, 80, diameter)
+        assert n == 80 and event["step"] == switch and event["tstar"] == 300
+        assert calls == [(0, -.001), (1, .001)]
+        calls.clear()
+        assert r.jet_interval(handle, "reverse", switch + 80, 80, diameter) == (80, None)
+        assert r.jet_interval(handle, "fixed", switch, 80, diameter) == (80, None)
+        assert not calls
+    rows = [dict(tstar=t, tstar_force=t - .5, cd=2.0, cl=-.1) for t in range(251, 306)]
+    windows = r.jet_diagnostics(rows, "reverse")
+    assert len(windows) == 6 and all(w["complete"] for w in windows)
+    assert all(w["cd"]["mean"] == 2 and np.isclose(w["cl"]["mean"], -.1) for w in windows)
+    assert windows[-1]["covered_tstar"] == 5
+    with unittest.TestCase().assertRaisesRegex(ValueError, "crosses"):
+        r.jet_diagnostics([dict(tstar=250.5, tstar_force=250, cd=1, cl=0)], "fixed")
+    assert r.plan(h, "MRT", "window", "reverse")["modes"]["full_tstar"] == 305
+
+
+def test_public_jet_source_and_pilot_separation():
+    from tempfile import TemporaryDirectory
+
+    with TemporaryDirectory() as tmp:
+        path = Path(tmp)
+        manifest = pilot_fixture(path, jet_scenario="fixed")
+        assert manifest["source_capture"].startswith("current checkout")
+        assert len(manifest["checkout_revision"]) == 40
+        assert r.verify_pilot(path, manifest, "MRT", "window", "fixed") == 3606.0
+        for scenario in ("none", "zero", "reverse"):
+            with unittest.TestCase().assertRaisesRegex(ValueError, "Jet scenario"):
+                r.verify_pilot(path, manifest, "MRT", "window", scenario)
+        for relative in ("uniform/summary.json", "refined/ready.json"):
+            receipt = path / relative
+            record = json.loads(receipt.read_text())
+            receipt.write_text(json.dumps(dict(record, jet_scenario="zero")))
+            with unittest.TestCase().assertRaises(ValueError):
+                r.verify_pilot(path, manifest, "MRT", "window", "fixed")
+            receipt.write_text(json.dumps(record))
+    with unittest.TestCase().assertRaises(ValueError):
+        r.verify_source(r.DEFAULT_SOURCE, "MRT", "last_base_step", "fixed")
+
+
+def test_cv_snapshot_common_units_and_window_endpoints():
+    records = []
+    for refined, shape, origin, diameter, impulse in (
+        (True, (240, 568), (38., 22.), 40, [4., -8.]),
+        (False, (330, 1761), (.25, .5), 80, [16., -32.]),
+        (False, (658, 3521), (.375, .75), 160, [64., -128.]),
+    ):
+        rho, ux, uy, fluid = np.ones(shape), np.full(shape, .01), np.zeros(shape), np.ones(shape, bool)
+        begin, n = round(290*diameter/r.SPEED), round(.02*diameter/r.SPEED)
+        record = r.cv_snapshot(rho, ux, uy, fluid, origin, refined, begin, n, diameter, impulse)
+        records.append(record)
+        np.testing.assert_allclose(record["outward_flux"], 0, atol=1e-12)
+        assert r.cv_selected(begin, diameter, "full") and not r.cv_selected(begin-n, diameter, "full")
+        end = round(300*diameter/r.SPEED)
+        assert r.cv_selected(end, diameter, "full") and not r.cv_selected(end+n, diameter, "full")
+        assert r.cv_selected(n, diameter, "pilot")
+        assert sum(r.cv_selected(step, diameter, "full") for step in range(begin-n, end+n+1, n)) == 501
+    for key in ("time_root", "interval_start_root", "interval_dt_root", "interval_impulse_root", "inventory"):
+        for record in records[1:]:
+            np.testing.assert_allclose(records[0][key], record[key], rtol=0, atol=1e-10)
+    assert records[0]["interval_dt_root"] == 80 and np.isclose(records[0]["tstar"], 290.02)
+    np.testing.assert_array_equal(records[0]["interval_impulse_root"], [4, -8])
+
+
+def test_cv_pilot_mode_and_helper_provenance():
+    from tempfile import TemporaryDirectory
+
+    manifest = r.verify_source(r.DEFAULT_SOURCE, "MRT", "window", "fixed", True)
+    assert "tests/validation/run_jet_circle_cv.py" in manifest["sha256"]
+    with TemporaryDirectory() as tmp:
+        path = Path(tmp)
+        old = pilot_fixture(path, jet_scenario="fixed")
+        with unittest.TestCase().assertRaisesRegex(ValueError, "CV observation mode"):
+            r.verify_pilot(path, old, "MRT", "window", "fixed", True)
+        for relative in ("plan.json", "uniform/summary.json", "refined/summary.json",
+                         "uniform/ready.json", "refined/ready.json"):
+            file = path / relative
+            record = json.loads(file.read_text())
+            record.update(observe_cv=True, cv_samples=250)
+            if relative == "plan.json":
+                record["source_manifest"] = manifest
+            file.write_text(json.dumps(record))
+        assert r.verify_pilot(path, manifest, "MRT", "window", "fixed", True) == 3606
+
+
+def test_compact_geometry_and_cv_stencils():
+    from CelerisLab.lbm.refinement import ACTIVE
+    from CelerisLab.lbm.descriptors import SOLID
+
+    _, pair, _, host = r.host_geometry(r.DEFAULT_SOURCE, h, "MRT", "fixed", "compact")
+    assert pair.fine_shape == (168, 168) and tuple(pair.fine_origin) == (46, 46)
+    fluid = (pair.fine_kind == ACTIVE) & ((host.flags.reshape(pair.fine_shape) & SOLID) == 0)
+    rho, ux, uy = np.ones(pair.fine_shape), np.full(pair.fine_shape, .01), np.zeros(pair.fine_shape)
+    record = r.cv_snapshot(rho, ux, uy, fluid, pair.fine_origin, True, 1160000, 80, 40, [1, 2], "compact")
+    np.testing.assert_allclose(record["outward_flux"], 0, atol=1e-12)
+    stencil = np.zeros(pair.fine_shape, bool)
+    for rect in r.cv_contours(True, 40, "compact"):
+        x0, y0, x1, y1 = np.rint((rect - np.tile(pair.fine_origin, 2))/.5).astype(int)
+        stencil[y0-2:y1+2, x0-2:x0+2] = stencil[y0-2:y1+2, x1-2:x1+2] = True
+        stencil[y0-2:y0+2, x0-2:x1+2] = stencil[y1-2:y1+2, x0-2:x1+2] = True
+    assert fluid[stencil].all() and not stencil.ravel()[host.transfer_fine_cells].any()
+
+
+def test_d160_config_and_single_case_receipts():
+    from tempfile import TemporaryDirectory
+    from unittest.mock import patch
+
+    cfg, _, p = r.config(h, 160, "MRT", "fixed")
+    assert (p["nx"], p["ny"], p["center"], p["radius"]) == (3521, 658, [320, 320.5], 80)
+    assert p["viscosity"] == .016 and p["mean_speed"] == .01 and not cfg["method"]["ddf_shifting"]
+    assert r.uniform_origin(160) == (.375, .75)
+    np.testing.assert_array_equal(r.cv_contours(False, 160), r.CV_CONTOURS + [.125, 0, .125, 0])
+    for scenario, expected_exit in (("fixed", 0), ("zero", 2)):
+        with TemporaryDirectory() as tmp:
+            path = Path(tmp)
+            manifest = pilot_fixture(path, jet_scenario=scenario, case_selection="uniform", uniform_diameter=160)
+            assert r.verify_pilot(path, manifest, "MRT", "window", scenario, False, "uniform", 160) == 1803
+            for selection, diameter, placement in (("both", 160, "wide"), ("uniform", 80, "wide"), ("uniform", 160, "compact")):
+                with unittest.TestCase().assertRaises(ValueError):
+                    r.verify_pilot(path, manifest, "MRT", "window", scenario, False, selection, diameter, placement)
+            # Synthetic dispatch/result check: no CUDA imports or execution.
+            output = path / "full"
+            argv = ["runner", "--collision", "MRT", "--exchange-window", "--jet-scenario", scenario,
+                    "--case", "uniform", "--uniform-diameter", "160", "--mode", "full",
+                    "--pilot-results", str(path), "--output", str(output), "--max-wall-seconds", "10", "--gpu-gates-approved"]
+            result = dict(case="uniform", healthy=True, periodic=scenario == "zero", jet_diagnostic_windows=[],
+                          status="driven_diagnostics_complete" if scenario == "fixed" else "periodic",
+                          reference_comparison=dict(reference_agreement=False))
+            with patch.object(sys, "argv", argv), patch.dict(r.os.environ, CUDA_VISIBLE_DEVICES="mock-no-gpu"), \
+                    patch.object(r, "run_case", return_value=result) as run:
+                assert r.main() == expected_exit and run.call_count == 1
+            comparison = json.loads((output / "comparison.json").read_text())
+            assert comparison["matched_case_pass"] is False and "refined_vs_uniform" not in comparison
+            assert comparison["status"].startswith("single_case_")
+
+
 if __name__ == "__main__":
     suite = unittest.TestSuite(
         unittest.FunctionTestCase(fn)

tests/unit/test_refinement_mrt.py

@@ -22,13 +22,16 @@ class RefinementMRTHost(unittest.TestCase):
 
     def test_srt_statements_and_guards_preserved(self):
         """MRT is an additive compile-time branch; SRT statements and other guards remain."""
-        blocks = (KERNELS / "blocks.cu").read_text()
+        blocks = (KERNELS / "blocks.cu").read_text() + (KERNELS / "step/block_step.cu").read_text()
         channel = (KERNELS / "step/block_channel.cu").read_text()
-        refinement = (KERNELS / "refinement.cu").read_text()
+        block_collision = (KERNELS / "operators/block_collision.cuh").read_text()
+        self.assertIn('#include "operators/block_collision.cuh"', blocks)
+        self.assertEqual(blocks.count("block_collide(f);"), 1)
+        refinement = (KERNELS / "refinement.cu").read_text() + (KERNELS / "step/patch_step.cu").read_text()
         for text in ("compute_feq_from_momentum(drho, jx, jy, ux, uy, feq);",
                      "compute_feq(rho, ux, uy, feq);",
                      "collide_srt_no_force(f, feq, d_params.omega);"):
-            self.assertIn(text, blocks)
+            self.assertIn(text, block_collision)
         self.assertIn("collide_srt(f, feq, forcing, omega);", channel)
         self.assertIn("collide_srt_no_force(f, feq, d_params.omega);", refinement)
         for source in (blocks, refinement):
@@ -37,8 +40,14 @@ class RefinementMRTHost(unittest.TestCase):
             self.assertIn("STORE_PRECISION != 0", source)
             self.assertIn('#if COLLISION_MODEL == 2\n#include "operators/collision_mrt.cuh"\n#endif', source)
         self.assertIn("USE_DDF_SHIFTING ||", refinement)  # dense/mixed shifted stays rejected
-        for source in (blocks, channel, refinement):
-            self.assertEqual(source.count("collide_mrt_no_force("), 1)
+        # One shared unforced MRT body; each blocked/patch path calls it exactly once.
+        mrt = (KERNELS / "operators/collision_mrt.cuh").read_text()
+        self.assertEqual(mrt.count("void collide_mrt_unforced("), 1)
+        for source in (blocks + "\n" + block_collision, channel, refinement):
+            self.assertEqual(source.count("collide_mrt_unforced(f);"), 1)
+            self.assertNotIn("collide_mrt_no_force(", source)
+            self.assertNotIn("compute_rho_u(f, rho, ux, uy);\n            collide_mrt", source)
+            self.assertNotRegex(source, r"compute_rho_u\(f, rho, ux, uy\);\s*collide_mrt")
 
 
 if __name__ == "__main__": unittest.main()

tests/unit/test_refinement_three_level_multipatch.py

@@ -173,8 +173,10 @@ def bodies_and_jets(speed=.002):
     return bodies,jets
 
 
-KW=dict(base_block=(4,4),fine_block=(4,4),child_block=(4,4),
-        child_rectangles_in_parent=CHILDREN)
+def levels(parents=PARENTS,children=CHILDREN,blocks=((4,4),)*3,objects=((),(),()),jets=((),(),())):
+    """Level-list constructor keywords for the L0 set, L1 set and L2 set used here."""
+    return dict(refinements=[parents,children],level_blocks=list(blocks),
+                level_objects=list(objects),level_jets=list(jets))
 
 
 def mock_backend(*, children=CHILDREN, bodies=False, fail_compile=False, fail_upload=False):
@@ -201,9 +203,8 @@ def mock_backend(*, children=CHILDREN, bodies=False, fail_compile=False, fail_up
         module.cuda.memcpy_htod.side_effect=upload
     objects,jets=bodies_and_jets() if bodies else ((),())
     try:
-        run=module.FixedPatchStepper(CFG,PARENTS,context=context,
-            child_objects=objects,child_jets=jets,
-            **dict(KW,child_rectangles_in_parent=children))
+        run=module.FixedPatchStepper(CFG,context=context,
+            **levels(children=children,objects=((),(),objects),jets=((),(),jets)))
     except RuntimeError:
         for allocation in allocations:
             allocation.free.assert_called_once_with()
@@ -214,28 +215,29 @@ def mock_backend(*, children=CHILDREN, bodies=False, fail_compile=False, fail_up
 class ConstructorAndSchedule(unittest.TestCase):
     def test_invalid_inputs_fail_before_context_or_allocations(self):
         context=Mock()
-        cases=[dict(child=(8,8,16,16)),dict(child_rectangles_in_parent=[]),
-               dict(child_rectangles_in_parent=[(16,16,32)]),
-               dict(child_rectangles_in_parent=[(16.5,16,32,40)]),
-               dict(child_rectangles_in_parent=[(16,16,32,24),(16,27,32,55)]),
-               dict(child_rectangles_in_parent=[(40,16,48,40)]),
-               dict(child_rectangles_in_parent=[(8,16,32,40)]),
-               dict(child_block=None),dict(fine_block=None),dict(base_block=None),
-               dict(child_block=(5,4)),dict(fine_block=(5,4)),
-               dict(child_block=(4,)),dict(child_rectangles_in_parent=(16,16,32,40))]
+        b=(4,4)
+        cases=[dict(children=(8,8,16,16)),dict(children=[]),
+               dict(children=[(16,16,32)]),
+               dict(children=[(16.5,16,32,40)]),
+               dict(children=[(16,16,32,24),(16,27,32,55)]),
+               dict(children=[(40,16,48,40)]),
+               dict(children=[(8,16,32,40)]),
+               dict(blocks=(b,b,None)),dict(blocks=(b,None,b)),dict(blocks=(None,b,b)),
+               dict(blocks=(b,b,(5,4))),dict(blocks=(b,(5,4),b)),
+               dict(blocks=(b,b,(4,))),dict(children=(16,16,32,40))]
         for changes in cases:
             with self.subTest(changes=changes),self.assertRaises((ValueError,TypeError)):
-                host.FixedPatchStepper(CFG,PARENTS,context=context,**dict(KW,**changes))
+                host.FixedPatchStepper(CFG,context=context,**levels(**changes))
         with self.assertRaisesRegex(ValueError,'margin'):
-            host.FixedPatchStepper(CFG,[(1,4,20,28),(23,4,44,28)],context=context,**KW)
+            host.FixedPatchStepper(CFG,context=context,**levels(parents=[(1,4,20,28),(23,4,44,28)]))
         objects,jets=bodies_and_jets()
-        with self.assertRaisesRegex(ValueError,'only on L2'):
-            host.FixedPatchStepper(CFG,PARENTS,context=context,fine_objects=objects,**KW)
+        with self.assertRaisesRegex(ValueError,'only on the finest level'):
+            host.FixedPatchStepper(CFG,context=context,**levels(objects=((),objects,())))
         from CelerisLab.body.jets import Jet,Tube,Parallel
         spanning=Jet(Tube((10.5,14.),0.,25.,.3),Parallel(0.),amplitude=0.)
         with self.assertRaisesRegex(ValueError,'one jet'):
-            host.FixedPatchStepper(CFG,PARENTS,context=context,child_objects=objects,
-                                  child_jets=[spanning],**KW)
+            host.FixedPatchStepper(CFG,context=context,
+                                  **levels(objects=((),(),objects),jets=((),(),[spanning])))
         context.activate.assert_not_called()
         host.cuda.mem_alloc.assert_not_called()
 
@@ -313,21 +315,23 @@ class ConstructorAndSchedule(unittest.TestCase):
 
 
 
-class InitialLevels(unittest.TestCase):
-    """initialize() entry adapter: positional per-level arrays or the legacy names, never both."""
+class RemovedLegacyAPI(unittest.TestCase):
+    """The pre-level-list constructor/initialize forms are gone and fail loudly."""
 
-    def test_forms(self):
-        a, b, c = object(), object(), object()
-        f = host._initial_levels
-        self.assertEqual(f((a, b), None, None, None), (a, b))
-        self.assertEqual(f((a, b, c), None, None, None), (a, b, c))
-        self.assertEqual(f((a, b), None, None, c), (a, b, c))      # old initialize(c, f, child_g=g)
-        self.assertEqual(f((), a, b, None), (a, b))
-        self.assertEqual(f((), a, b, c), (a, b, c))
-        for args in (((a,), None, None, c), ((a, b, c), None, None, c), ((a,), a, b, None),
-                     ((), a, None, None), ((), None, b, c)):
-            with self.subTest(args=args), self.assertRaises(TypeError):
-                f(*args)
+    def test_removed_constructor_forms_raise_type_error(self):
+        import inspect
+        context = Mock()
+        with self.assertRaises(TypeError):  # positional bounds
+            host.FixedPatchStepper(CFG, PARENTS, context=context)
+        removed = dict(bounds=PARENTS, fine_block=(4, 4), base_block=(4, 4), fine_objects=(),
+                       jets=(), child=(8, 8, 16, 16), child_block=(4, 4), child_objects=(),
+                       child_jets=(), child_rectangles_in_parent=CHILDREN)
+        for name, value in removed.items():
+            with self.subTest(keyword=name), self.assertRaises(TypeError):
+                host.FixedPatchStepper(CFG, context=context, **levels(), **{name: value})
+        context.activate.assert_not_called()
+        self.assertEqual(list(inspect.signature(host.FixedPatchStepper.initialize).parameters),
+                         ["self", "level_g"])  # no coarse_g/fine_g/child_g
 
 
 if __name__=='__main__':

tests/unit/test_shifted_jets.py

@@ -54,14 +54,13 @@ class ShiftedJets(unittest.TestCase):
                 (Jet(Tube((10.,14.),0.,26.,.3),Parallel(0.)),))]
         for objects,jets in cases:
             with self.subTest(objects=len(objects),jets=len(jets)),self.assertRaises(ValueError):
-                module.FixedPatchStepper(CFG,backend.PARENTS,context=None,
-                    base_block=(4,4),fine_block=(4,4),child_block=(4,4),
-                    child_rectangles_in_parent=A,child_objects=objects,child_jets=jets)
+                module.FixedPatchStepper(CFG,context=None,refinements=[backend.PARENTS,A],
+                    level_blocks=[(4,4)]*3,level_objects=[(),(),objects],level_jets=[(),(),jets])
         for blocks in ((None,None,None),((4,4),None,(4,4))):
             with self.assertRaises(ValueError):
-                module.FixedPatchStepper(CFG,backend.PARENTS,context=None,
-                    base_block=blocks[0],fine_block=blocks[1],child_block=blocks[2],
-                    child_rectangles_in_parent=A,child_objects=(body(),),child_jets=(jet(),))
+                module.FixedPatchStepper(CFG,context=None,refinements=[backend.PARENTS,A],
+                    level_blocks=list(blocks),level_objects=[(),(),(body(),)],
+                    level_jets=[(),(),(jet(),)])
         module.cuda.mem_alloc.assert_not_called()
         module.compiler.compile_kernel.assert_not_called()
 

tests/unit/test_shifted_storage.py

@@ -78,13 +78,15 @@ class ShiftedStorage(unittest.TestCase):
 
     def test_shifted_density_and_tag_validation(self):
         old,new = topology(A),topology(B)
-        a = tuple(np.where(m[None], np.float32(-.001), np.nan).repeat(9,axis=0) for m in old.leaf_masks)
+        a = tuple(np.where(m[None], np.float32(-.001), np.float32(np.nan)).repeat(9,axis=0)
+                  for m in old.leaf_masks)  # FP32 under NumPy 1 and 2 promotion rules
         remap(old,new,a)
         with self.assertRaisesRegex(ValueError,'representation'): remap(old,new,a,False)
         with self.assertRaisesRegex(ValueError,'FP32'): remap(old,new,tuple(x.astype(np.float64) for x in a))
         a[0][:,old.leaf_masks[0]] = -1
         with self.assertRaisesRegex(ValueError,'positive density'): remap(old,new,a)
-        rest = tuple(np.where(m[None],np.float32(0),np.nan).repeat(9,axis=0) for m in old.leaf_masks)
+        rest = tuple(np.where(m[None],np.float32(0),np.float32(np.nan)).repeat(9,axis=0)
+                     for m in old.leaf_masks)
         remap(old,new,rest)
         with self.assertRaisesRegex(ValueError,'positive density'):
             remap_synchronized(old,new,rest,base_count=7,level_steps=(7,14,28),phase='postcollision_g')
@@ -129,9 +131,10 @@ class ShiftedStorage(unittest.TestCase):
     def test_dense_or_mixed_shifted_rejected_before_compilation(self):
         module=backend.host_module()
         cfg=replace(backend.CFG,ddf_shifting=True)
-        for kwargs in ({},dict(base_block=(4,4)),dict(fine_block=(4,4))):
-            with self.assertRaisesRegex(ValueError,'all-blocked|base_block requires fine_block'):
-                module.FixedPatchStepper(cfg,(4,4,20,28),context=None,**kwargs)
+        for blocks in ([None,None],[(4,4),None],[None,(4,4)]):
+            with self.assertRaisesRegex(ValueError,'all-blocked|blocked refined levels below'):
+                module.FixedPatchStepper(cfg,context=None,refinements=[(4,4,20,28)],
+                                         level_blocks=blocks)
         module.compiler.compile_kernel.assert_not_called()
 
     def test_uniform_block_codec_and_raw_snapshot(self):

tests/unit/test_state_regressions.py

@@ -238,6 +238,81 @@ def test_overlapping_force_regions_rejected(device):
         bodies.build_compact_lists(bodies.build_flags(field.flag))
 
 
+@pytest.mark.parametrize("second_x,error", [(14., "unique cut-link writers"),
+                                           (17., "active fluid receivers")])
+def test_passive_overlap_rejected_before_publishing_lists(second_x, error):
+    cfg = LBMConfig(nx=32, ny=24)
+    cfg.validate()
+    bodies = ObjectManager(cfg.nx, cfg.ny, cfg.nz, cfg.nq, cfg)
+    for x in (14., second_x):
+        bodies.add(obj(x=x, y=12., r=4.3))
+    bodies.add(obj(x=4., y=4., r=2., is_sensor=True))
+    counts = np.full(bodies.count, 91, np.int32)
+    hits = []
+    flags = bodies.build_flags(np.full(cfg.nx * cfg.ny, FLUID, np.uint16))
+    for domain_flags in (flags, None):
+        with pytest.raises(ValueError, match=error):
+            bodies.build_compact_lists_for(bodies.objects, domain_flags,
+                                          sensor_cell_counts_out=counts, hits_out=hits)
+        np.testing.assert_array_equal(counts, 91)
+        assert hits == []
+
+
+@pytest.mark.parametrize("centers", [((10., 12.), (22., 12.)),
+                                    ((10.5, 12.), (19.5, 12.))])
+def test_separated_passive_bodies_keep_valid_writers(centers):
+    # The second case has a subcell physical gap with a resolved fluid column.
+    cfg = LBMConfig(nx=32, ny=24)
+    cfg.validate()
+    bodies = ObjectManager(cfg.nx, cfg.ny, cfg.nz, cfg.nq, cfg)
+    for x, y in centers:
+        bodies.add(obj(x=x, y=y, r=4.3))
+    flags = bodies.build_flags(np.full(cfg.nx * cfg.ny, FLUID, np.uint16))
+    indices, directions, _, ids, *_ = bodies.build_compact_lists(flags)
+    assert set(ids) == {0, 1}
+    assert np.all(flags[indices] & FLUID) and not np.any(flags[indices] & SOLID)
+    assert len(np.unique(indices.astype(np.uint64) * cfg.nq + directions)) == len(indices)
+
+
+@pytest.mark.parametrize("entry", ["sync", "initialize", "runtime_edit"])
+@pytest.mark.parametrize("second_x", [5., 6.])
+def test_passive_overlap_failure_keeps_live_state(device, entry, second_x):
+    from CelerisLab.simulation import Simulation
+    field = make_field(device)
+    bodies = manager(field, [obj(), obj(x=9., y=3., r=1., is_sensor=True)])
+    bodies.sync_to_gpu(field)
+    overlapping = obj(x=second_x)
+    if entry == "runtime_edit":
+        bodies.stage_add(overlapping)
+    else:
+        bodies.add(overlapping)
+    objects = bodies.objects
+    flags, curved, counts = field.flag.copy(), field.curved.fluid_idx.copy(), bodies.sensor_cell_counts.copy()
+    action = bodies.action.copy()
+    buffers = [(block, bytes(block.data), block.freed) for block in device.blocks]
+    writes = device.writes
+    sim = SimpleNamespace(bodies=bodies, field=field, lbm_cfg=field.cfg,
+                          _initialized=entry == "runtime_edit",
+                          _assert_allocation_contract=lambda: None,
+                          stream=SimpleNamespace(synchronize=device.synchronize))
+    with pytest.raises(ValueError, match="active fluid receivers|unique cut-link writers"):
+        if entry == "sync":
+            bodies.sync_to_gpu(field)
+        elif entry == "initialize":
+            Simulation.initialize(sim)
+        else:
+            Simulation.sync_bodies(sim)
+    assert bodies.objects == objects
+    assert bodies.has_pending_edit() == (entry == "runtime_edit")
+    assert sim._initialized == (entry == "runtime_edit")
+    assert device.writes == writes and len(device.blocks) == len(buffers)
+    assert all(bytes(block.data) == data and block.freed == freed for block, data, freed in buffers)
+    np.testing.assert_array_equal(field.flag, flags)
+    np.testing.assert_array_equal(field.curved.fluid_idx, curved)
+    np.testing.assert_array_equal(bodies.sensor_cell_counts, counts)
+    np.testing.assert_array_equal(bodies.action, action)
+
+
 def test_action_upload_and_scratch_are_pinned(device):
     field = make_field(device, streaming="esopull")
     bodies = manager(field, [obj(is_force_region=True)])

tests/unit/test_validation_source_provenance.py

@@ -12,13 +12,15 @@ from source_provenance import checkout_provenance
 import re100_helpers
 import run_jet_circle_cv
 
+sys.path.insert(0, str(ROOT / "tests/unit"))
+from historical_tree import historical_tree, tracked
+
 ORIGIN = "47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5"
 
 
-def test_manifests_cover_current_source_and_preserve_origin():
-    names = subprocess.check_output(
-        ["git", "-C", str(ROOT), "ls-files", "src"], text=True
-    ).splitlines()
+def test_manifests_cover_current_source_and_preserve_origin(tmp_path):
+    assert re100_helpers.PIN == run_jet_circle_cv.PIN == ORIGIN
+    manifests = {}
     for group, hashes_key, revision_key in (
         ("dfg_re100", "production_sha256", "commit"),
         ("jet_circle_cv", "files", "revision"),
@@ -29,18 +31,22 @@ def test_manifests_cover_current_source_and_preserve_origin():
             ).read_text()
         )
         assert manifest[revision_key] == ORIGIN
-        assert set(manifest[hashes_key]) == set(names)
         assert manifest["source_origin_commits"][-1] == ORIGIN
         assert (
             manifest["historical_checkpoints"]["production"]
             == "d79e3f6f0a424f5bcc3a1a6d610313e49daf97a2"
         )
         assert "not inherited from d79" in manifest["experiment"]
-        for name, expected in manifest[hashes_key].items():
+        manifests[group] = manifest[hashes_key]
+    # History-dependent checks last, so a missing commit masks nothing above.
+    names = tracked(ORIGIN)
+    origin = historical_tree(ORIGIN, tmp_path)
+    for group, hashes in manifests.items():
+        assert set(hashes) == set(names), group
+        for name, expected in hashes.items():
             assert (
-                hashlib.sha256((ROOT / name).read_bytes()).hexdigest() == expected
+                hashlib.sha256((origin / name).read_bytes()).hexdigest() == expected
             ), name
-    assert re100_helpers.PIN == run_jet_circle_cv.PIN == ORIGIN
 
 
 def test_runtime_checkout_identity_is_separate_from_source_origin(tmp_path):

tests/validation/multigrid/test_block_stepper.py

@@ -192,6 +192,7 @@ class CpuContracts(unittest.TestCase):
         backend._lock = RLock()
         backend._status = "ready"
         backend._step_count = 0
+        backend._ddf_shifting = False  # unshifted storage, as for the default config
         backend._stream = Mock()
         backend._module = Mock()
         backend._copy = backend._step = Mock()

tests/validation/multigrid/test_channel_root_gpu.py

@@ -27,7 +27,7 @@ def state(ny, nx):
 
 def compile_module(cfg, context, source):
     cfg.validate()
-    build = compiler.generate_config(cfg)
+    build = compiler.generate_config(cfg, block_root_bc=2 if source == "blocks.cu" else 0)
     try:
         module = compiler.load_module(compiler.compile_kernel(context.sm_arch, source=source, build_dir=build))
     finally:
@@ -106,9 +106,10 @@ class ChannelRootGPU(unittest.TestCase):
                     cfg = LBMConfig(nx=48, ny=32, viscosity=.1, velocity=.015, ddf_shifting=shifted)
                     reference = None
                     for block in ((32,48), (8,8), (4,1)):
-                        with FixedPatchStepper(cfg, (8,6,40,26), context=context,
-                                coarse_boundary="channel", base_block=block, fine_block=(8,8),
-                                fine_objects=(body,) if bodies else (), jets=(jet,) if bodies else ()) as backend:
+                        with FixedPatchStepper(cfg, context=context, coarse_boundary="channel",
+                                refinements=[(8,6,40,26)], level_blocks=[block, (8,8)],
+                                level_objects=[(), (body,) if bodies else ()],
+                                level_jets=[(), (jet,) if bodies else ()]) as backend:
                             arrays = [state(*level.shape[1:]) for level in backend._levels]
                             backend.initialize(*arrays)
                             backend.step(4)

tests/validation/multigrid/test_exchange_window.py

@@ -25,8 +25,8 @@ EPS = np.finfo(np.float64).eps
 
 
 def build(m, window=True, raw=True):
-    return FixedPatchStepper(config("periodic"), BOUNDS, context=m["context"], base_block=(3, 4),
-                             fine_block=(16, 20), fine_objects=[circle()],
+    return FixedPatchStepper(config("periodic"), context=m["context"], refinements=[BOUNDS],
+                             level_blocks=[(3, 4), (16, 20)], level_objects=[(), (circle(),)],
                              observe_raw_exchange=raw, observe_window_exchange=window)
 
 
@@ -109,10 +109,11 @@ class ExchangeWindowGPU(unittest.TestCase):
         from test_raw_exchange_gpu import BOUNDS as RB, CFG, circle as rc, jet, state
         from CelerisLab.cuda.context import CudaContext
         with CudaContext(0) as ctx:
-            kw = dict(context=ctx, fine_block=(16, 20), fine_objects=[rc()], jets=[jet()],
+            kw = dict(context=ctx, refinements=[RB], level_blocks=[None, (16, 20)],
+                      level_objects=[(), (rc(),)], level_jets=[(), (jet(),)],
                       observe_raw_exchange=True)
-            with FixedPatchStepper(CFG, RB, observe_window_exchange=True, **kw) as windowed, \
-                 FixedPatchStepper(CFG, RB, **kw) as plain:
+            with FixedPatchStepper(CFG, observe_window_exchange=True, **kw) as windowed, \
+                 FixedPatchStepper(CFG, **kw) as plain:
                 arrays = [state(level.shape) for level in windowed._levels]
                 windowed.initialize(*arrays)
                 plain.initialize(*arrays)
@@ -200,11 +201,11 @@ class ExchangeWindowShiftedMRTGPU(unittest.TestCase):
         from CelerisLab.cuda.context import CudaContext
         cfg = LBMConfig(nx=32, ny=24, viscosity=.1, collision="MRT", ddf_shifting=True)
         with CudaContext(0) as ctx:
-            kw = dict(context=ctx, base_block=(12, 8), fine_block=(16, 20), fine_objects=[rc()],
-                      observe_raw_exchange=True)
-            with FixedPatchStepper(cfg, RB, observe_window_exchange=True, **kw) as windowed, \
-                 FixedPatchStepper(cfg, RB, observe_window_exchange=True, **kw) as split, \
-                 FixedPatchStepper(cfg, RB, **kw) as plain:
+            kw = dict(context=ctx, refinements=[RB], level_blocks=[(12, 8), (16, 20)],
+                      level_objects=[(), (rc(),)], observe_raw_exchange=True)
+            with FixedPatchStepper(cfg, observe_window_exchange=True, **kw) as windowed, \
+                 FixedPatchStepper(cfg, observe_window_exchange=True, **kw) as split, \
+                 FixedPatchStepper(cfg, **kw) as plain:
                 arrays = [state(level.shape) for level in windowed._levels]
                 for run in (windowed, split, plain):
                     run.initialize(*arrays)

tests/validation/multigrid/test_fine_surface_gpu.py

@@ -140,7 +140,8 @@ def run_patch(speed, *, audit=True):
         build = (patch.object(compiler, "compile_kernel", audit_compiler(compiler.compile_kernel))
                  if audit else nullcontext())
         with build:
-            backend = implementation.FixedPatchStepper(cfg, BOUNDS, context=context, fine_objects=[source])
+            backend = implementation.FixedPatchStepper(cfg, context=context, refinements=[BOUNDS],
+                                             level_objects=[(), (source,)])
         with backend, context.activate():
             c, f = backend._levels
             cl = backend._fine_curved
@@ -281,7 +282,8 @@ def test_fine_surface_four_hops(speed):
 @pytest.mark.parametrize("missing", ["surface_velocity_gpu", "link_exchange_gpu"])
 def test_missing_surface_device_buffer_rejects_before_launch(missing):
     backend = implementation.FixedPatchStepper.__new__(implementation.FixedPatchStepper)
-    backend.fine_boundary = SimpleNamespace(objects=[object_spec(.002)])
+    backend._boundary_hosts = [None, SimpleNamespace(objects=[object_spec(.002)])]
+    backend._level_jets = [(), ()]
     backend._fine_curved = SimpleNamespace(fluid_idx_gpu=object(),
                                           surface_velocity_gpu=object(), link_exchange_gpu=object())
     setattr(backend._fine_curved, missing, None)
@@ -325,6 +327,7 @@ def test_cross_interface_donor_rejects_before_cuda():
     context.activate.side_effect = AssertionError("must reject before CUDA activation")
     with patch.object(cuda, "mem_alloc") as allocate:
         with pytest.raises(ValueError, match="donor crosses"):
-            implementation.FixedPatchStepper(cfg, BOUNDS, context=context, fine_objects=[source])
+            implementation.FixedPatchStepper(cfg, context=context, refinements=[BOUNDS],
+                                             level_objects=[(), (source,)])
         context.activate.assert_not_called()
         allocate.assert_not_called()

tests/validation/multigrid/test_multiple_pinned_remap_gpu.py

@@ -48,9 +48,8 @@ def make(context):
         for (x, y), amplitude in zip(centers, (.003, -.002))
     )
     run = FixedPatchStepper(
-        CFG, PARENTS, context=context, base_block=(4, 4), fine_block=(4, 4),
-        child_block=(4, 4), child_rectangles_in_parent=START,
-        child_objects=objects, child_jets=jets,
+        CFG, context=context, refinements=[PARENTS, START], level_blocks=[(4, 4)] * 3,
+        level_objects=[(), (), objects], level_jets=[(), (), jets],
     )
     return run, objects
 

tests/validation/multigrid/test_pinned_body_remap_gpu.py

@@ -25,9 +25,8 @@ B = ((16,16,32,40), (60,16,76,40))
 def make(ctx):
     body = SimObject(71, geometry=CircleGeometry(12.,14.,.8),center=(12.,14.),radius=.8)
     jet = Jet(Tube((10.5,14.),0.,3.,.3),Parallel(0.),'parabolic',.003)
-    run = FixedPatchStepper(CFG,PARENTS,context=ctx,base_block=(4,4),fine_block=(4,4),
-                           child_block=(4,4),child_rectangles_in_parent=A,
-                           child_objects=(body,),child_jets=(jet,))
+    run = FixedPatchStepper(CFG,context=ctx,refinements=[PARENTS,A],level_blocks=[(4,4)]*3,
+                           level_objects=[(),(),(body,)],level_jets=[(),(),(jet,)])
     return run,(body,)
 
 

tests/validation/multigrid/test_raw_exchange_gpu.py

@@ -90,14 +90,18 @@ def census(run, index, buffer):
     return totals(dense, leaf), external, independent
 
 
+
+def fine(block):
+    """Two-level circle+jet construction on BOUNDS; block None keeps the dense fine level."""
+    return dict(refinements=[BOUNDS], level_blocks=[None, block],
+                level_objects=[(), (circle(),)], level_jets=[(), (jet(),)])
+
 class RawExchangeGPU(unittest.TestCase):
     def test_row_ledger_interval_jet_switch_and_no_population_effect(self):
         with CudaContext(0) as ctx:
             for block in ((32, 40), (16, 20), (8, 10)):
-                with FixedPatchStepper(CFG, BOUNDS, context=ctx, fine_block=block,
-                        fine_objects=[circle()], jets=[jet()], observe_raw_exchange=True) as measured, \
-                     FixedPatchStepper(CFG, BOUNDS, context=ctx, fine_block=block,
-                        fine_objects=[circle()], jets=[jet()]) as plain:
+                with FixedPatchStepper(CFG, context=ctx, **fine(block), observe_raw_exchange=True) as measured, \
+                     FixedPatchStepper(CFG, context=ctx, **fine(block)) as plain:
                     arrays = [state(l.shape) for l in measured._levels]
                     measured.initialize(*arrays)
                     plain.initialize(*arrays)
@@ -146,9 +150,8 @@ class RawExchangeGPU(unittest.TestCase):
         """Same fine resolution/initial data with dense storage: a numerical
         reference for cut-link rows, not a uniform-whole-domain convergence study."""
         with CudaContext(0) as ctx:
-            with FixedPatchStepper(CFG, BOUNDS, context=ctx, fine_objects=[circle()], jets=[jet()]) as dense, \
-                 FixedPatchStepper(CFG, BOUNDS, context=ctx, fine_block=(16, 20),
-                    fine_objects=[circle()], jets=[jet()], observe_raw_exchange=True) as blocked:
+            with FixedPatchStepper(CFG, context=ctx, **fine(None)) as dense, \
+                 FixedPatchStepper(CFG, context=ctx, **fine((16, 20)), observe_raw_exchange=True) as blocked:
                 arrays = [state(l.shape) for l in blocked._levels]
                 dense.initialize(*arrays)
                 blocked.initialize(*arrays)
@@ -229,8 +232,7 @@ class RawExchangeGPU(unittest.TestCase):
         layout = BlockLayout2D.build(48, 64, 8, 8)
         with CudaContext(0) as ctx:
             with PeriodicBlockStepper(fine_cfg, layout, context=ctx, objects=[obj]) as uniform, \
-                 FixedPatchStepper(CFG, BOUNDS, context=ctx, fine_block=(16, 20),
-                    fine_objects=[circle()], jets=[jet()], observe_raw_exchange=True) as amr:
+                 FixedPatchStepper(CFG, context=ctx, **fine((16, 20)), observe_raw_exchange=True) as amr:
                 amr.initialize(*(state(l.shape) for l in amr._levels))
                 uniform.initialize(state((9, 48, 64)))
                 b = amr._level_boundaries[1]

tests/validation/multigrid/test_refinement_blocked_base.py

@@ -83,11 +83,10 @@ def mocked_backend(mode="periodic", fail_alloc=None):
 class BlockedBaseCPU(unittest.TestCase):
     def test_reject_before_cuda(self):
         context = Mock(activate=Mock(side_effect=AssertionError("CUDA reached")))
-        for kwargs in ({"base_block": (12, 8)},
-                       {"base_block": (7, 8), "fine_block": FINE_BLOCK},
-                       {"base_block": (12,), "fine_block": FINE_BLOCK}):
-            with self.subTest(kwargs=kwargs), self.assertRaises(ValueError):
-                FixedPatchStepper(config("periodic"), BOUNDS, context=context, **kwargs)
+        for blocks in ([(12, 8), None], [(7, 8), FINE_BLOCK], [(12,), FINE_BLOCK]):
+            with self.subTest(blocks=blocks), self.assertRaises(ValueError):
+                FixedPatchStepper(config("periodic"), context=context, refinements=[BOUNDS],
+                                  level_blocks=blocks)
         context.activate.assert_not_called()
 
     def test_all_arrival_rows_donors_and_kind_masks(self):
@@ -193,9 +192,10 @@ class BlockedBaseCPU(unittest.TestCase):
 
     def test_launch_order_stream_wall_constants_and_buffer_swaps(self):
         with mocked_backend("slot_reservoir") as m:
-            backend = FixedPatchStepper(config("slot_reservoir"), BOUNDS, context=m["context"],
-                                        coarse_boundary="slot_reservoir", fine_block=FINE_BLOCK,
-                                        base_block=(3, 4), fine_objects=[circle(actuation(.002))])
+            backend = FixedPatchStepper(config("slot_reservoir"), context=m["context"],
+                                        coarse_boundary="slot_reservoir", refinements=[BOUNDS],
+                                        level_blocks=[(3, 4), FINE_BLOCK],
+                                        level_objects=[(), (circle(actuation(.002)),)])
             try:
                 backend.initialize(state(NY, NX, 1), state(*backend.patch.fine_shape, 2))
                 coarse, fine = backend._levels
@@ -223,6 +223,8 @@ class BlockedBaseCPU(unittest.TestCase):
                     self.assertIs(collide[4], coarse.kind)
                 walls = [a for ptr, a in m["uploads"] if isinstance(ptr, str) and ptr == "d_block_wall"]
                 self.assertEqual(len(walls), 1)
+                self.assertEqual([call.kwargs["block_root_bc"] for call in
+                                  compiler.generate_config.call_args_list], [1, 0])
                 width, height, slot, speed = struct.unpack("=IIIf", walls[0])
                 self.assertEqual((width, height, slot), (NX, NY, 8))
                 self.assertAlmostEqual(speed, .001)
@@ -236,16 +238,16 @@ class BlockedBaseCPU(unittest.TestCase):
         for fail_at in (1, 6, 10):
             with self.subTest(fail_at=fail_at), mocked_backend(fail_alloc=fail_at) as m:
                 with self.assertRaisesRegex(MemoryError, "injected allocation"):
-                    FixedPatchStepper(config("periodic"), BOUNDS, context=m["context"],
-                                      fine_block=FINE_BLOCK, base_block=(3, 4))
+                    FixedPatchStepper(config("periodic"), context=m["context"], refinements=[BOUNDS],
+                                      level_blocks=[(3, 4), FINE_BLOCK])
                 self.assertEqual(len(m["allocations"]), fail_at)
                 for allocation in m["allocations"]:
                     allocation.free.assert_called_once_with()
 
     def test_failed_parent_stage_cannot_publish_or_resume(self):
         with mocked_backend() as m:
-            backend = FixedPatchStepper(config("periodic"), BOUNDS, context=m["context"],
-                                        fine_block=FINE_BLOCK, base_block=(3, 4))
+            backend = FixedPatchStepper(config("periodic"), context=m["context"], refinements=[BOUNDS],
+                                        level_blocks=[(3, 4), FINE_BLOCK])
             backend.initialize(state(NY, NX, 1), state(*backend.patch.fine_shape, 2))
             m["functions"]["BlockCoalesce"].side_effect = RuntimeError("injected parent stage failure")
             current = backend._levels[0].current
@@ -289,9 +291,10 @@ class BlockedBaseGPU(unittest.TestCase):
                     for base_block in (None, *layouts):
                         label = f"{mode} width={width} speed={speed} base={base_block}"
                         snapshots = []
-                        with FixedPatchStepper(cfg, BOUNDS, context=context, coarse_boundary=mode,
-                                               fine_block=FINE_BLOCK, base_block=base_block,
-                                               fine_objects=() if speed is None else [circle(actuation(speed))]) as backend:
+                        with FixedPatchStepper(cfg, context=context, coarse_boundary=mode,
+                                               refinements=[BOUNDS], level_blocks=[base_block, FINE_BLOCK],
+                                               level_objects=[(), () if speed is None
+                                                              else (circle(actuation(speed)),)]) as backend:
                             backend.initialize(*arrays)
                             for step in (0, 1, 2, 8):
                                 if step == 1 and base_block is not None:

tests/validation/multigrid/test_refinement_blocked_fine.py

@@ -105,9 +105,10 @@ class BlockedFineCPU(unittest.TestCase):
         return Mock(activate=Mock(side_effect=AssertionError("CUDA reached")), sm_arch="sm_70")
 
     def test_rejections_before_cuda(self):
-        for kw in ({"fine_block": (7, 7)}, {"fine_block": (16,)}):
+        for kw in ({"level_blocks": [None, (7, 7)]}, {"level_blocks": [None, (16,)]}):
             with self.subTest(kw=kw), self.assertRaises(ValueError):
-                FixedPatchStepper(config("periodic"), BOUNDS, context=self.forbidden_context(), **kw)
+                FixedPatchStepper(config("periodic"), context=self.forbidden_context(),
+                                  refinements=[BOUNDS], **kw)
 
     def test_single_block_layout_allocates_no_empty_halo_tables(self):
         # A 1x1 non-periodic fine layout has zero halo pairs; CUDA rejects mem_alloc(0).
@@ -125,8 +126,8 @@ class BlockedFineCPU(unittest.TestCase):
                 patch.object(compiler, "compile_kernel"), patch.object(compiler, "cleanup_build_dir"), \
                 patch.object(compiler, "load_module", return_value=module), \
                 patch.object(cuda, "mem_alloc", side_effect=alloc), patch.object(cuda, "memcpy_htod"):
-            backend = FixedPatchStepper(config("periodic"), BOUNDS, context=context,
-                                        fine_block=FINE_BLOCKS[0])
+            backend = FixedPatchStepper(config("periodic"), context=context, refinements=[BOUNDS],
+                                        level_blocks=[None, FINE_BLOCKS[0]])
             self.assertEqual(len(backend.fine_layout.halo_dst), 0)
             self.assertIsNone(backend._levels[1].halo_dst)
             backend._gather_halo(backend._levels[1], Mock())  # no launch without pairs
@@ -177,8 +178,9 @@ class BlockedFineCPU(unittest.TestCase):
                 patch.object(cuda, "mem_alloc", side_effect=lambda size: Mock()), \
                 patch.object(cuda, "memcpy_htod"), \
                 patch.object(curved_links, "CurvedLinkSoA", side_effect=AssertionError("dense SoA built")):
-            backend = FixedPatchStepper(config("periodic"), BOUNDS, context=context,
-                                        fine_objects=objects, fine_block=FINE_BLOCKS[1])
+            backend = FixedPatchStepper(config("periodic"), context=context, refinements=[BOUNDS],
+                                        level_blocks=[None, FINE_BLOCKS[1]],
+                                        level_objects=[(), tuple(objects)])
             backend.initialize(state(NY, NX, 1), state(*backend.patch.fine_shape, 2))
         return backend, calls
 
@@ -289,8 +291,9 @@ class BlockedFineGPU(unittest.TestCase):
             results = []
             with CudaContext(0) as context:
                 for fine_block in (None, *FINE_BLOCKS):
-                    with FixedPatchStepper(cfg, BOUNDS, context=context, coarse_boundary=mode,
-                                           fine_block=fine_block) as backend:
+                    with FixedPatchStepper(cfg, context=context, coarse_boundary=mode,
+                                           refinements=[BOUNDS],
+                                           level_blocks=[None, fine_block]) as backend:
                         backend.initialize(g_coarse, g_fine)
                         backend.step(STEPS)
                         self.assertEqual(backend.step_count, STEPS)
@@ -318,8 +321,9 @@ class BlockedFineGPU(unittest.TestCase):
             results = []
             with CudaContext(0) as context:
                 for fine_block in (None, *FINE_BLOCKS):
-                    with FixedPatchStepper(cfg, BOUNDS, context=context, coarse_boundary=mode,
-                                           fine_objects=[circle()], fine_block=fine_block) as backend:
+                    with FixedPatchStepper(cfg, context=context, coarse_boundary=mode,
+                                           refinements=[BOUNDS], level_blocks=[None, fine_block],
+                                           level_objects=[(), (circle(),)]) as backend:
                         backend.initialize(g_coarse, g_fine)
                         backend.step(STEPS)
                         results.append((backend.download(), raw(backend)))
@@ -346,9 +350,9 @@ class BlockedFineGPU(unittest.TestCase):
                 results = []
                 with CudaContext(0) as context:
                     for fine_block in (None, *FINE_BLOCKS):
-                        with FixedPatchStepper(cfg, BOUNDS, context=context, coarse_boundary=mode,
-                                               fine_objects=[circle(actuation(speed))],
-                                               fine_block=fine_block) as backend:
+                        with FixedPatchStepper(cfg, context=context, coarse_boundary=mode,
+                                               refinements=[BOUNDS], level_blocks=[None, fine_block],
+                                               level_objects=[(), (circle(actuation(speed)),)]) as backend:
                             backend.initialize(g_coarse, g_fine)
                             backend.step(STEPS)
                             results.append((backend.download(), raw(backend)))

tests/validation/multigrid/test_refinement_followup.py

@@ -94,7 +94,8 @@ class FaultGuardsCPU(unittest.TestCase):
                              ("slot_reservoir", (12, 4, 20, 21))):
             with self.subTest(mode=mode, bounds=bounds):
                 with self.assertRaises(ValueError):
-                    implementation.FixedPatchStepper(cfg, bounds, context=context, coarse_boundary=mode)
+                    implementation.FixedPatchStepper(cfg, context=context, coarse_boundary=mode,
+                                                     refinements=[bounds])
                 context.activate.assert_not_called()
         # Both exact channel limits and the original periodic two-cell margin
         # must reach context activation, with no allocation or device required.
@@ -102,11 +103,14 @@ class FaultGuardsCPU(unittest.TestCase):
                              ("periodic", (2, 2, 6, 5))):
             with self.subTest(mode=mode, bounds=bounds):
                 with self.assertRaisesRegex(RuntimeError, "validated before CUDA"):
-                    implementation.FixedPatchStepper(cfg, bounds, context=context, coarse_boundary=mode)
+                    implementation.FixedPatchStepper(cfg, context=context, coarse_boundary=mode,
+                                                     refinements=[bounds])
 
     def fake_backend(self):
         backend = PeriodicPatchStepper.__new__(PeriodicPatchStepper)
-        backend.fine_layout = None  # dense fine storage
+        backend._layouts = [None, None]  # dense storage on both levels
+        backend._ddf_shifting = False  # unshifted on every level, as _check_storage expects
+        backend._level_cfgs = [LBMConfig(nx=16, ny=16, viscosity=.1)] * 2  # ddf_shifting off
         backend.ctx = Mock(activate=nullcontext)
         backend._lock = RLock()
         backend._status = "ready"
@@ -154,12 +158,23 @@ class FaultGuardsCPU(unittest.TestCase):
         context = Mock(activate=nullcontext, sm_arch="sm_70")
         with patch.object(implementation, "HALO", 2):
             with self.assertRaisesRegex(RuntimeError, "four-cell"):
-                PeriodicPatchStepper(cfg, (4, 4, 12, 12), context=context)
+                PeriodicPatchStepper(cfg, context=context, refinements=[(4, 4, 12, 12)])
         p = implementation.FixedPatch2D.from_config(cfg, (4, 4, 12, 12))
-        with patch.object(cuda, "Stream"), patch.object(implementation.FixedPatch2D, "from_config",
-                                                       return_value=replace(p, tau_coarse=.9)):
-            with self.assertRaisesRegex(RuntimeError, "omega/tau"):
-                PeriodicPatchStepper(cfg, p.bounds, context=context)
+        from CelerisLab.lbm import hierarchy
+        bad = patch.object(implementation.FixedPatch2D, "from_config",
+                           return_value=replace(p, tau_coarse=.9))
+        # A pair tau inconsistent with its level config is rejected twice over, both before
+        # any device allocation or upload: first by the hierarchy chain check, and, if that
+        # is bypassed, by the per-level omega/tau check in _build_level.
+        with patch.object(cuda, "Stream"), bad, patch.object(cuda, "mem_alloc") as alloc, \
+                patch.object(cuda, "memcpy_htod") as upload:
+            with self.assertRaisesRegex(RuntimeError, "config chain is inconsistent"):
+                PeriodicPatchStepper(cfg, context=context, refinements=[p.bounds])
+            with patch.object(hierarchy, "_check_chain"), \
+                    self.assertRaisesRegex(RuntimeError, "omega/tau"):
+                PeriodicPatchStepper(cfg, context=context, refinements=[p.bounds])
+            alloc.assert_not_called()
+            upload.assert_not_called()
         module = Mock()
         module.get_global.return_value = (Mock(), 16)
         with patch.object(cuda, "Stream"), patch.object(compiler, "generate_config"), \
@@ -167,7 +182,7 @@ class FaultGuardsCPU(unittest.TestCase):
                 patch.object(compiler, "load_module", return_value=module), \
                 patch.object(cuda, "memcpy_htod") as upload:
             with self.assertRaisesRegex(RuntimeError, "layout mismatch"):
-                PeriodicPatchStepper(cfg, p.bounds, context=context)
+                PeriodicPatchStepper(cfg, context=context, refinements=[p.bounds])
             upload.assert_not_called()
 
 
@@ -241,7 +256,7 @@ class NormalBuildGPU(unittest.TestCase):
         max_error = max_bc_error = 0.0
         with CudaContext(0) as context:
             with patch.object(compiler, "generate_config", side_effect=record_config):
-                backend = implementation.FixedPatchStepper(cfg, bounds, context=context,
+                backend = implementation.FixedPatchStepper(cfg, context=context, refinements=[bounds],
                                                             coarse_boundary="slot_reservoir")
             with backend:
                 self.assertEqual(len(generated), 2)
@@ -311,7 +326,8 @@ class NormalBuildGPU(unittest.TestCase):
                             for a in (initial.coarse_f, initial.fine_f)))
         max_pop = max_drift = max_moment = 0.0
         with CudaContext(0) as context, PeriodicPatchStepper(
-                LBMConfig(nx=16, ny=16, viscosity=.1), (4, 4, 12, 12), context=context) as backend:
+                LBMConfig(nx=16, ny=16, viscosity=.1), context=context,
+                refinements=[(4, 4, 12, 12)]) as backend:
             backend.initialize(*postcollision(state, spec, maps))
             before = totals(backend.download(), maps)
             for step in range(1, 9):
@@ -342,7 +358,7 @@ class NormalBuildGPU(unittest.TestCase):
         fine_mask[4:-4, 4:-4] = True
         initial[1][:, ~fine_mask] = np.nan
         expected = scalar_interval(*[a.astype(np.float64) for a in initial], bounds)
-        with CudaContext(0) as context, PeriodicPatchStepper(cfg, bounds, context=context) as backend:
+        with CudaContext(0) as context, PeriodicPatchStepper(cfg, context=context, refinements=[bounds]) as backend:
             self.assertEqual(backend.patch.fine_shape, (14, 16))
             p = backend.patch
             allocated = 2 * 9 * (17*13+16*14) * 4 + sum(a.nbytes for a in

tests/validation/multigrid/test_refinement_gpu.py

@@ -84,13 +84,13 @@ class CoupledGPU(unittest.TestCase):
 
         cases = (("dense_unshifted", False, {}),
                  ("blocked_shifted", True,
-                  dict(base_block=(4, 4), fine_block=(6, 6))))
+                  dict(level_blocks=[(4, 4), (6, 6)])))
         with CudaContext(0) as context:
             for label, shifted, storage in cases:
                 with self.subTest(storage=label):
                     cfg = LBMConfig(nx=16, ny=16, viscosity=.0192,
                                     ddf_shifting=shifted)
-                    with FixedPatchStepper(cfg, (4, 4, 12, 12), context=context,
+                    with FixedPatchStepper(cfg, context=context, refinements=[(4, 4, 12, 12)],
                                            **storage) as backend:
                         self.assertAlmostEqual(backend.patch.tau_coarse, spec.tau_coarse)
                         self.assertAlmostEqual(backend.patch.tau_fine, spec.tau_fine)
@@ -148,7 +148,7 @@ class CoupledGPU(unittest.TestCase):
         maps = ref.build_maps(spec)
         with CudaContext(0) as context:
             with PeriodicPatchStepper(LBMConfig(nx=16, ny=16, viscosity=.1),
-                                      (4, 4, 12, 12), context=context) as backend:
+                                      context=context, refinements=[(4, 4, 12, 12)]) as backend:
                 backend.initialize(*postcollision(ref.make_state(spec), spec, maps))
                 coarse, fine = backend._levels
 
@@ -187,7 +187,7 @@ class CoupledGPU(unittest.TestCase):
 
         with CudaContext(0) as context:
             with patch.object(compiler, "compile_kernel", audit_compile):
-                backend = PeriodicPatchStepper(cfg, (4, 4, 12, 12), context=context)
+                backend = PeriodicPatchStepper(cfg, context=context, refinements=[(4, 4, 12, 12)])
             with backend:
                 # Production allocations: four DDF arrays and four compact tables.
                 device_bytes = 2 * 9 * (16**2 + 24**2) * 4
@@ -270,7 +270,7 @@ class CoupledGPU(unittest.TestCase):
                 backend.download()
 
             # Compile and execute the actual non-audited entry as well: one step.
-            with PeriodicPatchStepper(cfg, (4, 4, 12, 12), context=context) as plain:
+            with PeriodicPatchStepper(cfg, context=context, refinements=[(4, 4, 12, 12)]) as plain:
                 rest = postcollision(ref.make_state(spec), spec, maps)
                 plain.initialize(*rest)
                 plain.step()

tests/validation/multigrid/test_refinement_jets.py

@@ -41,8 +41,9 @@ def mocked(objects, jets, fine_block):
             patch.object(compiler, "load_module", return_value=module), \
             patch.object(cuda, "mem_alloc", side_effect=lambda size: Mock()), \
             patch.object(cuda, "memcpy_htod", htod):
-        backend = FixedPatchStepper(config("periodic"), BOUNDS, context=context,
-                                    fine_objects=objects, fine_block=fine_block, jets=jets)
+        backend = FixedPatchStepper(config("periodic"), context=context, refinements=[BOUNDS],
+                                    level_blocks=[None, fine_block], level_objects=[(), tuple(objects)],
+                                    level_jets=[(), tuple(jets)])
     return backend, htod
 
 
@@ -149,8 +150,9 @@ class JetStepperCPU(unittest.TestCase):
     def test_jet_overlapping_legacy_aperture_rejected_before_cuda(self):
         with patch.object(cuda, "Stream", side_effect=AssertionError("CUDA touched")):
             with self.assertRaisesRegex(ValueError, "SurfaceActuation"):
-                FixedPatchStepper(config("periodic"), BOUNDS, context=Mock(),
-                                  fine_objects=[circle(actuation(SPEED))], jets=[jet()])
+                FixedPatchStepper(config("periodic"), context=Mock(), refinements=[BOUNDS],
+                                  level_objects=[(), (circle(actuation(SPEED)),)],
+                                  level_jets=[(), (jet(),)])
 
 
 def run(mode, fine_block, *, objects, jets, switch=None):
@@ -158,8 +160,9 @@ def run(mode, fine_block, *, objects, jets, switch=None):
     cfg = config(mode)
     patch_ = FixedPatch2D.from_config(cfg, BOUNDS)
     with CudaContext(0) as context, FixedPatchStepper(
-            cfg, BOUNDS, context=context, coarse_boundary=mode, fine_objects=objects,
-            fine_block=fine_block, jets=jets) as backend:
+            cfg, context=context, coarse_boundary=mode, refinements=[BOUNDS],
+            level_blocks=[None, fine_block], level_objects=[(), tuple(objects)],
+            level_jets=[(), tuple(jets)]) as backend:
         backend.initialize(state(NY, NX, 1), state(*patch_.fine_shape, 2))
         if switch is None:
             backend.step(STEPS)

tests/validation/multigrid/test_refinement_mrt_gpu.py

@@ -42,8 +42,9 @@ class RefinementMRTGPU(unittest.TestCase):
                                     collision="MRT", ddf_shifting=shifted)
                     reference = None
                     for block in ((32, 48), (8, 8), (4, 1)):
-                        with FixedPatchStepper(cfg, (8, 6, 40, 26), context=context, coarse_boundary=boundary,
-                                               base_block=block, fine_block=(8, 8), fine_objects=(body,)) as backend:
+                        with FixedPatchStepper(cfg, context=context, coarse_boundary=boundary,
+                                               refinements=[(8, 6, 40, 26)], level_blocks=[block, (8, 8)],
+                                               level_objects=[(), (body,)]) as backend:
                             backend.initialize(*[state(*level.shape[1:]) for level in backend._levels])
                             backend.step(8)
                             raw = backend._download_raw()
@@ -63,8 +64,8 @@ class RefinementMRTGPU(unittest.TestCase):
             results = {}
             layouts = ((None, None), *((base, FINE_BLOCK) for base in (None, *BASE_BLOCKS)))
             for base_block, fine_block in layouts:
-                with FixedPatchStepper(cfg, BOUNDS, context=context, fine_block=fine_block,
-                                       base_block=base_block) as backend:
+                with FixedPatchStepper(cfg, context=context, refinements=[BOUNDS],
+                                       level_blocks=[base_block, fine_block]) as backend:
                     backend.initialize(*[fixture.state(*level.shape[1:], seed) for seed, level in enumerate(backend._levels)])
                     backend.step(4)
                     results[base_block, fine_block] = raw(backend)
@@ -75,9 +76,9 @@ class RefinementMRTGPU(unittest.TestCase):
     def test_shifted_all_blocked_three_level_mrt_smoke(self):
         """Per-level shear rate omega_l = 1/(3*2**l*nu + 1/2); finite ACTIVE leaves."""
         cfg = LBMConfig(nx=16, ny=16, viscosity=.1, collision="MRT", ddf_shifting=True)
-        kw = dict(fine_block=(4, 4), base_block=(4, 4), child=(8, 8, 16, 16), child_block=(4, 4))
+        kw = dict(refinements=[(4, 4, 12, 12), (8, 8, 16, 16)], level_blocks=[(4, 4)] * 3)
         with CudaContext(0) as context:
-            with FixedPatchStepper(cfg, (4, 4, 12, 12), context=context, **kw) as backend:
+            with FixedPatchStepper(cfg, context=context, **kw) as backend:
                 self.assertEqual(len(backend._level_cfgs), 3)
                 for level, level_cfg in enumerate(backend._level_cfgs):
                     checked = replace(level_cfg)

tests/validation/multigrid/test_refinement_multipatch.py

@@ -26,6 +26,12 @@ FACE = ((4, 4, 8, 12), (12, 4, 16, 12))
 S1 = ((4, 4, 12, 12), (20, 4, 28, 12))
 
 
+def levels(bounds, blocks=(None, None), objects=(), jets=(), **rest):
+    """Two-level constructor keywords: one L1 rectangle or set, per-level blocks, L1 bodies/Jets."""
+    return dict(refinements=[bounds], level_blocks=list(blocks),
+                level_objects=[(), tuple(objects)], level_jets=[(), tuple(jets)], **rest)
+
+
 def level_fine(patch_set, windows):
     """Level-wide fine array with each patch window filled from `windows`."""
     out = np.full((9, *patch_set.fine_shape), np.nan, np.float32)
@@ -45,21 +51,21 @@ class MultipatchCPU(unittest.TestCase):
     def test_rejections_before_cuda(self):
         cfg = config("periodic")
         context = object()  # never activated: every case must fail first
-        for kw in (dict(), dict(fine_block=(8, 4)), dict(base_block=(4, 4)),
-                   dict(fine_block=(8, 4), base_block=(4, 4), fine_objects=[circle()]),
-                   dict(fine_block=(5, 4), base_block=(4, 4))):
+        for kw in (dict(), dict(blocks=(None, (8, 4))), dict(blocks=((4, 4), None)),
+                   dict(blocks=((4, 4), (8, 4)), objects=[circle()]),
+                   dict(blocks=((4, 4), (5, 4)))):
             with self.subTest(kw=kw), self.assertRaises(ValueError):
-                FixedPatchStepper(cfg, ((4, 4, 8, 12), (14, 4, 18, 12)), context=context, **kw)
+                FixedPatchStepper(cfg, context=context, **levels(((4, 4, 8, 12), (14, 4, 18, 12)), **kw))
         with self.assertRaisesRegex(ValueError, "margin"):
-            FixedPatchStepper(cfg, ((1, 4, 8, 12), (10, 4, 14, 12)), context=context,
-                              fine_block=(8, 4), base_block=(4, 4))
+            FixedPatchStepper(cfg, context=context,
+                              **levels(((1, 4, 8, 12), (10, 4, 14, 12)), blocks=((4, 4), (8, 4))))
         from CelerisLab.body.jets import Jet, Parallel, Tube
         cfg32 = LBMConfig(nx=32, ny=16, viscosity=.1)
         span = Jet(Tube((5., 8.), 0., 22., .6), Parallel(0.), amplitude=0.)
         with self.assertRaisesRegex(ValueError, "one jet"):  # same obj_id cannot bypass it
-            FixedPatchStepper(cfg32, ((2, 2, 14, 14), (18, 2, 30, 14)), context=context,
-                              fine_block=(4, 4), base_block=(4, 4),
-                              fine_objects=[body(8., 8.), body(24., 8.)], jets=[span])
+            FixedPatchStepper(cfg32, context=context,
+                              **levels(((2, 2, 14, 14), (18, 2, 30, 14)), blocks=((4, 4), (4, 4)),
+                                       objects=[body(8., 8.), body(24., 8.)], jets=[span]))
 
 
 class MultipatchGPU(unittest.TestCase):
@@ -67,7 +73,7 @@ class MultipatchGPU(unittest.TestCase):
 
     def run_steps(self, context, cfg, bounds, coarse, fine, **kw):
         out = []
-        with FixedPatchStepper(cfg, bounds, context=context, **kw) as backend:
+        with FixedPatchStepper(cfg, context=context, **levels(bounds, **kw)) as backend:
             backend.initialize(coarse, fine)
             for step in self.STEPS:
                 backend.step(step - backend.step_count)
@@ -82,7 +88,7 @@ class MultipatchGPU(unittest.TestCase):
                 p = FixedPatch2D.from_config(cfg, BOUNDS)
                 s = FixedPatchSet2D.from_config(cfg, (BOUNDS,))
                 coarse, fine = state(NY, NX, 1), state(*p.fine_shape, 2)
-                kw = dict(coarse_boundary=mode, fine_block=(8, 4), base_block=(3, 4))
+                kw = dict(coarse_boundary=mode, blocks=((3, 4), (8, 4)))
                 old = self.run_steps(context, cfg, BOUNDS, coarse, fine, **kw)
                 new = self.run_steps(context, cfg, (BOUNDS,), coarse, level_fine(s, [fine]), **kw)
                 for step, a, b in zip(self.STEPS, old, new):
@@ -99,7 +105,7 @@ class MultipatchGPU(unittest.TestCase):
         one = FixedPatchSet2D.from_config(cfg16, (S1[0],))
         two = FixedPatchSet2D.from_config(cfg32, S1)
         coarse, fine = state(16, 16, 3), state(*one.patches[0].fine_shape, 4)
-        kw = dict(fine_block=(4, 4), base_block=(4, 4))
+        kw = dict(blocks=((4, 4), (4, 4)))
         with CudaContext(0) as context:
             ref = self.run_steps(context, cfg16, (S1[0],), coarse, level_fine(one, [fine]), **kw)
             tiled = np.concatenate([coarse, coarse], axis=2)
@@ -118,7 +124,7 @@ class MultipatchGPU(unittest.TestCase):
         s = FixedPatchSet2D.from_config(cfg, FACE)
         coarse = state(16, 24, 5)
         fine = level_fine(s, [state(*p.fine_shape, 6 + i) for i, p in enumerate(s.patches)])
-        kw = dict(fine_block=(4, 4), base_block=(4, 4))
+        kw = dict(blocks=((4, 4), (4, 4)))
         leaf = s.fine_kind == ACTIVE
         support = s.fine_layout((4, 4)).owned
         outer = support & (s.fine_kind == OUTER_GHOST)
@@ -164,7 +170,7 @@ class MultipatchBodiesGPU(unittest.TestCase):
 
     def run_steps(self, context, cfg, bounds, coarse, fine, switch=None, **kw):
         out = []
-        with FixedPatchStepper(cfg, bounds, context=context, **kw) as backend:
+        with FixedPatchStepper(cfg, context=context, **levels(bounds, **kw)) as backend:
             backend.initialize(coarse, fine)
             for step in self.STEPS:
                 if switch is not None and step == 8:
@@ -182,8 +188,7 @@ class MultipatchBodiesGPU(unittest.TestCase):
                 p = FixedPatch2D.from_config(cfg, BOUNDS)
                 s = FixedPatchSet2D.from_config(cfg, (BOUNDS,))
                 coarse, fine = state(NY, NX, 1), state(*p.fine_shape, 2)
-                kw = dict(coarse_boundary=mode, fine_block=(8, 4), base_block=(3, 4),
-                          fine_objects=[circle()])
+                kw = dict(coarse_boundary=mode, blocks=((3, 4), (8, 4)), objects=[circle()])
                 old = self.run_steps(context, cfg, BOUNDS, coarse, fine, **kw)
                 new = self.run_steps(context, cfg, (BOUNDS,), coarse, level_fine(s, [fine]), **kw)
                 for step, a, b in zip(self.STEPS, old, new):
@@ -202,17 +207,17 @@ class MultipatchBodiesGPU(unittest.TestCase):
         two = FixedPatchSet2D.from_config(cfg32, self.PATCHES)
         coarse, fine = state(16, 16, 3), state(*one.patches[0].fine_shape, 4)
         jet = lambda x: Jet(Tube((x, 8.), 0., 6., .6), Parallel(0.), amplitude=.002)
-        kw = dict(fine_block=(4, 4), base_block=(4, 4))
+        kw = dict(blocks=((4, 4), (4, 4)))
         with CudaContext(0) as context:
             for jets in ((), (jet(5.),)):
                 ref = self.run_steps(context, cfg16, self.PATCHES[:1], coarse, level_fine(one, [fine]),
                                      switch=[(0, -.001)] if jets else None,
-                                     fine_objects=[body(8., 8.)], jets=jets, **kw)
+                                     objects=[body(8., 8.)], jets=jets, **kw)
                 both = jets and (jet(5.), jet(21.))
                 got = self.run_steps(context, cfg32, self.PATCHES, np.concatenate([coarse, coarse], 2),
                                      level_fine(two, [fine, fine]),
                                      switch=[(0, -.001), (1, -.001)] if jets else None,
-                                     fine_objects=[body(8., 8., 5), body(24., 8., 5)],
+                                     objects=[body(8., 8., 5), body(24., 8., 5)],
                                      jets=both or (), **kw)
                 for step, a, b in zip(self.STEPS, ref, got):
                     label = f"jets={len(jets)} step {step}"

tests/validation/multigrid/test_refinement_three_level.py

@@ -27,7 +27,16 @@ from test_refinement_blocked_base import raw
 
 CFG = LBMConfig(nx=16, ny=16, viscosity=.1)
 L1, L2 = (4, 4, 12, 12), (8, 8, 16, 16)
-KW = dict(fine_block=(4, 4), base_block=(4, 4), child=L2, child_block=(4, 4))
+B = (4, 4)
+
+
+def levels(l1=L1, l2=L2, blocks=(B, B, B), objects=((), (), ()), jets=((), (), ())):
+    """Level-list keywords for the scalar L1 patch with one scalar L2 child."""
+    return dict(refinements=[l1, l2], level_blocks=list(blocks),
+                level_objects=list(objects), level_jets=list(jets))
+
+
+KW = levels()
 
 
 def child_body(spec=None):
@@ -52,26 +61,29 @@ def mocked_child():
             patch.object(compiler, "load_module", return_value=module), \
             patch.object(cuda, "mem_alloc", side_effect=lambda size: Mock()), \
             patch.object(cuda, "memcpy_htod"):
-        return FixedPatchStepper(CFG, L1, context=context, child_objects=[child_body()],
-                                 child_jets=[child_jet()], **KW)
+        return FixedPatchStepper(CFG, context=context,
+                                 **levels(objects=((), (), (child_body(),)), jets=((), (), (child_jet(),))))
 
 
 class ThreeLevelCPU(unittest.TestCase):
     def test_rejections_before_cuda(self):
         context = object()  # never activated: every case must fail first
-        bad = (dict(KW, child_block=None), dict(KW, fine_block=None),
-               dict(KW, fine_objects=[circle()]), dict(KW, child=(4, 8, 16, 16)),
-               dict(KW, child_block=(5, 4)))
+        bad = (levels(blocks=(B, B, None)), levels(blocks=(B, None, B)),
+               levels(objects=((), (circle(),), ())), levels(l2=(4, 8, 16, 16)),
+               levels(blocks=(B, B, (5, 4))))
         for kw in bad:
             with self.subTest(kw=kw), self.assertRaises(ValueError):
-                FixedPatchStepper(CFG, L1, context=context, **kw)
-        for kw in (dict(child_objects=[child_body()]), dict(child_jets=[child_jet()]),
-                   dict(KW, child_jets=[child_jet()])):
+                FixedPatchStepper(CFG, context=context, **kw)
+        # L2 bodies/Jets without an L2 level (one entry too many), and an L2 Jet without a body.
+        two_level = dict(refinements=[L1], level_blocks=[None, None])
+        for kw in (dict(two_level, level_objects=[(), (), (child_body(),)]),
+                   dict(two_level, level_jets=[(), (), (child_jet(),)]),
+                   levels(jets=((), (), (child_jet(),)))):
             with self.assertRaises(ValueError):
-                FixedPatchStepper(CFG, L1, context=context, **kw)
-        with self.assertRaises(ValueError):  # several L1 patches + child: not yet
-            FixedPatchStepper(LBMConfig(nx=32, ny=16, viscosity=.1), ((4, 4, 12, 12), (20, 4, 28, 12)),
-                              context=context, **KW)
+                FixedPatchStepper(CFG, context=context, **kw)
+        with self.assertRaises(ValueError):  # several L1 patches + scalar child: not yet
+            FixedPatchStepper(LBMConfig(nx=32, ny=16, viscosity=.1), context=context,
+                              **levels(l1=((4, 4, 12, 12), (20, 4, 28, 12))))
 
 
     def test_four_halo_bfl_steps_and_fresh_middle_feeds(self):
@@ -80,11 +92,12 @@ class ThreeLevelCPU(unittest.TestCase):
         f, g = run._levels[1:]
         b = g.boundary
         events, feeds = [], []
-        advance, gather = run._advance_child, run._gather_halo
-        def record_feed(buffer, **kw):
-            self.assertIs(buffer, (f.current, f.spare)[len(feeds)])
-            feeds.append(buffer)
-            advance(buffer, **kw)
+        advance, gather = run._advance, run._gather_halo
+        def record_feed(index, buffer, **kw):  # index 2: L1 hop feeds L2
+            if index == 2:
+                self.assertIs(buffer, (f.current, f.spare)[len(feeds)])
+                feeds.append(buffer)
+            advance(index, buffer, **kw)
         def record_halo(level, buffer, **kw):
             if level is g:
                 events.append(("halo", buffer))
@@ -95,7 +108,7 @@ class ThreeLevelCPU(unittest.TestCase):
             self.assertIs(kind, g.kind)
             events.append(("step", src))
         g.kernel.side_effect = record_step
-        with patch.object(run, "_advance_child", side_effect=record_feed), \
+        with patch.object(run, "_advance", side_effect=record_feed), \
                 patch.object(run, "_gather_halo", side_effect=record_halo):
             run.step()
         self.assertEqual(feeds, [f.current, f.spare])
@@ -105,8 +118,8 @@ class ThreeLevelCPU(unittest.TestCase):
         self.assertGreater(len(np.unique(b.tables["block"])), 1)
         # Negative control: reusing L1's old source at the midpoint must fail.
         feeds.clear()
-        with patch.object(run, "_advance_child", side_effect=lambda buffer, **kw:
-                          record_feed(f.current, **kw)):
+        with patch.object(run, "_advance", side_effect=lambda index, buffer, **kw:
+                          record_feed(index, f.current if index == 2 else buffer, **kw)):
             with self.assertRaises(AssertionError):
                 run.step()
         run.close()
@@ -155,7 +168,7 @@ class ThreeLevelGPU(unittest.TestCase):
     def test_equilibrium_is_kept_on_every_level(self):
         from CelerisLab.cuda.context import CudaContext
         feq = np.asarray(compute_feq_d2q9(1.0, .03, -.02), np.float32)[:, None, None]
-        with CudaContext(0) as ctx, FixedPatchStepper(CFG, L1, context=ctx, **KW) as run:
+        with CudaContext(0) as ctx, FixedPatchStepper(CFG, context=ctx, **KW) as run:
             shapes = ((16, 16), run.patch.fine_shape, run.child_patch.fine_shape)
             run.initialize(*(np.broadcast_to(feq, (9, *s)).copy() for s in shapes))
             run.step(8)
@@ -168,15 +181,16 @@ class ThreeLevelGPU(unittest.TestCase):
 
     def test_nonequilibrium_masks_and_fresh_middle_feeds(self):
         from CelerisLab.cuda.context import CudaContext
-        with CudaContext(0) as ctx, FixedPatchStepper(CFG, L1, context=ctx, **KW) as run:
+        with CudaContext(0) as ctx, FixedPatchStepper(CFG, context=ctx, **KW) as run:
             p1, p2 = run.patch, run.child_patch
             run.initialize(state(16, 16, 1), state(*p1.fine_shape, 2), state(*p2.fine_shape, 3))
             f = run._levels[1]
-            feeds, original = [], run._advance_child
-            def record(buffer, **kw):
-                feeds.append(int(buffer))
-                original(buffer, **kw)
-            with patch.object(run, "_advance_child", side_effect=record):
+            feeds, original = [], run._advance
+            def record(index, buffer, **kw):  # index 2: L1 hop feeds L2
+                if index == 2:
+                    feeds.append(int(buffer))
+                original(index, buffer, **kw)
+            with patch.object(run, "_advance", side_effect=record):
                 cur, spare = int(f.current), int(f.spare)
                 run.step(1)
             self.assertEqual(feeds, [cur, spare])  # current G, then the new G of hop 1
@@ -203,8 +217,9 @@ class ThreeLevelBodyGPU(unittest.TestCase):
             CFG, velocity=0., wall_normal_velocity=.001, wall_slot_width=4,
             y_wall_bc="on_site_velocity", top_wall_bc="reservoir")
         with CudaContext(0) as ctx, FixedPatchStepper(
-                cfg, L1, context=ctx, coarse_boundary=mode, child_objects=[child_body(spec)],
-                child_jets=jets, **dict(KW, child_block=block)) as run:
+                cfg, context=ctx, coarse_boundary=mode,
+                **levels(blocks=(B, B, block), objects=((), (), (child_body(spec),)),
+                         jets=((), (), tuple(jets)))) as run:
             run.initialize(*(state(*level.shape[1:], i + 1) for i, level in enumerate(run._levels)))
             snapshots = []
             for stop in (0, 1, 2, 4, 8):

tests/validation/multigrid/test_refinement_three_level_multipatch_gpu.py

@@ -78,10 +78,10 @@ class ThreeLevelMultipatchGPU(unittest.TestCase):
     def run_case(self, ctx, cfg=CFG, parents=PARENTS, children=CHILDREN,
                  mode='periodic', blocks=((4, 4), (4, 4)), objects=(), drive=(), switch=False,
                  scalar_child=None):
-        kwargs = dict(child_rectangles_in_parent=children) if scalar_child is None else dict(child=scalar_child)
-        with FixedPatchStepper(config(cfg, mode), parents, context=ctx, coarse_boundary=mode,
-                base_block=(4, 4), fine_block=blocks[0], child_block=blocks[1],
-                child_objects=objects, child_jets=drive, **kwargs) as run:
+        child = children if scalar_child is None else scalar_child
+        with FixedPatchStepper(config(cfg, mode), context=ctx, coarse_boundary=mode,
+                refinements=[parents, child], level_blocks=[(4, 4), blocks[0], blocks[1]],
+                level_objects=[(), (), objects], level_jets=[(), (), drive]) as run:
             frame = frames(run)
             run.initialize(*(physical_state(level.shape[1:], *f)
                              for level, f in zip(run._levels, frame)))

tests/validation/multigrid/test_runtime_child_remap_gpu.py

@@ -26,9 +26,8 @@ FP64_ATOL = 3e-6
 
 
 def make(ctx):
-    return FixedPatchStepper(CFG, PARENTS, context=ctx, base_block=(4, 4),
-                             fine_block=(4, 4), child_block=(4, 4),
-                             child_rectangles_in_parent=A)
+    return FixedPatchStepper(CFG, context=ctx, refinements=[PARENTS, A],
+                             level_blocks=[(4, 4)] * 3)
 
 
 def initial(run, uniform):

tests/validation/multigrid/test_shifted_jets_gpu.py

@@ -136,8 +136,8 @@ class ShiftedJetsGPU(unittest.TestCase):
                   Parallel(0.),amplitude=.013)
         for rho in (.99,1.01):
             with self.subTest(rho=rho),CudaContext(0) as ctx:
-                with FixedPatchStepper(cfg,(8,6,24,18),context=ctx,base_block=(4,4),fine_block=(4,4),
-                        fine_objects=(obj,),jets=(whole,),observe_raw_exchange=True) as run:
+                with FixedPatchStepper(cfg,context=ctx,refinements=[(8,6,24,18)],level_blocks=[(4,4)]*2,
+                        level_objects=[(),(obj,)],level_jets=[(),(whole,)],observe_raw_exchange=True) as run:
                     run._initialize_raw(initial(run,rho,u=.013,v=0.))
                     expected=run._download_raw();ledger=Ledger(self,run,1)
                     self.assertTrue(run._level_boundaries[1].tables['aperture'].all())
@@ -153,7 +153,8 @@ class ShiftedJetsGPU(unittest.TestCase):
         obj=circle((16.,12.),2.15);j=drive((16.,12.),2.15)
         for rho in (.99,1.01):
             with self.subTest(rho=rho),CudaContext(0) as ctx,ExitStack() as owned:
-                settings=dict(bounds=(8,6,24,18),context=ctx,base_block=(4,4),fine_block=(4,4),fine_objects=(obj,),jets=(j,))
+                settings=dict(context=ctx,refinements=[(8,6,24,18)],level_blocks=[(4,4)]*2,
+                              level_objects=[(),(obj,)],level_jets=[(),(j,)])
                 measured=owned.enter_context(FixedPatchStepper(cfg,observe_raw_exchange=True,**settings))
                 plain=owned.enter_context(FixedPatchStepper(cfg,**settings))
                 measured._initialize_raw(initial(measured,rho));plain._initialize_raw(initial(plain,rho))
@@ -186,8 +187,9 @@ class ShiftedJetsGPU(unittest.TestCase):
         parents=((4,4,20,28),(24,4,44,28));a=((16,16,32,40),(56,16,72,40));b=(a[0],(60,16,76,40))
         obj=circle((12.,14.),.8);j=drive((12.,14.),.8)
         with CudaContext(0) as ctx,ExitStack() as owned:
-            old=owned.enter_context(FixedPatchStepper(cfg,parents,context=ctx,base_block=(4,4),fine_block=(4,4),
-                child_block=(4,4),child_rectangles_in_parent=a,child_objects=(obj,),child_jets=(j,),observe_raw_exchange=True))
+            old=owned.enter_context(FixedPatchStepper(cfg,context=ctx,refinements=[parents,a],
+                level_blocks=[(4,4)]*3,level_objects=[(),(),(obj,)],level_jets=[(),(),(j,)],
+                observe_raw_exchange=True))
             old._initialize_raw(initial(old,.99));ledger=Ledger(self,old,2)
             for speed in (.002,-.001):old.set_jet_amplitude(0,speed,level=2);ledger.step()
             before=old._download_raw();count=old.step_count

tests/validation/multigrid/test_shifted_remap_gpu.py

@@ -22,9 +22,8 @@ B=((20,16,36,40),)
 
 
 def make(ctx):
-    return FixedPatchStepper(CFG,PARENTS,context=ctx,base_block=(4,4),
-                             fine_block=(4,4),child_block=(4,4),
-                             child_rectangles_in_parent=A)
+    return FixedPatchStepper(CFG,context=ctx,refinements=[PARENTS,A],
+                             level_blocks=[(4,4)]*3)
 
 
 def initial(run):

tests/validation/run_dfg_re20.py

@@ -0,0 +1,222 @@
+"""Steady DFG 2D-1 (Re 20): refined and uniform runs against the official reference intervals.
+
+Not a full benchmark certification. Declared deviations from the official setup:
+pressure-NEQ outlet (``neq_extrap``) instead of do-nothing, local Zou-He inlet.
+Geometry, inflow profile, initial state, refined layout and observation reuse the
+pinned Re100 runner (``run_refined_dfg``) and ``re100_helpers`` unchanged; only the
+lattice viscosity is set for Re 20 (nu = U_lattice * D / 20 at the root level).
+
+Units. The helpers use scaled DFG units with mean inflow speed 1, so
+Cd = 2 Fx / (rho U^2 D) and Cl = 2 Fy / (rho U^2 D) are scale-free, and
+``pressure_probes`` returns dp / (rho U^2). The official p_diff = p(0.15, 0.2) -
+p(0.25, 0.2) uses U = 0.2 m/s, rho = 1, so p_diff = dp * 0.04. Both are reported;
+only p_diff is compared with the official interval. Forces use the existing sign
+convention: F is the force of the fluid on the cylinder, i.e. the refined raw
+exchange impulse (window ledger, ``total.impulse``) or the uniform ``read_force``,
+with x along the inflow and y up from the bottom wall. Cd must be positive, and the
+official Cl is positive (centre 0.20 below the centreline 0.205), so a flipped sign
+shows up as Cl outside its interval.
+
+Sampling. One sample every SAMPLE_TSTAR (t* = U t / D; 0.25 s), taken at the end of
+a completed root step with all levels synchronized; the step index is recorded. Cd
+and Cl are means over the sample interval (window exchange ledger refined, per-batch
+forces uniform), which filters acoustic and step-parity oscillation; p_diff is the
+end-of-interval value at the exact surface points.
+
+Steady rule (our engineering choice, frozen before any run; not official). Over the
+trailing WINDOW_TSTAR (t* 50..60 inclusive, 21 samples), for each of Cd, Cl and
+p_diff: range (max - min) <= 10 % and |least-squares slope * window| <= 2 % of the
+width of its official interval. Absolute limits, so the small Cl is never a divisor;
+the slope term bounds slow monotone drift that the range alone would pass. If the
+rule fails at TSTAR_MAX the run reports "not steady" and gives no interval verdict;
+limits are never relaxed after a run. The verdict uses the window mean, and the
+distance of each mean to its nearest interval edge is reported (negative outside).
+"""
+
+import argparse
+import json
+import math
+import os
+from pathlib import Path
+import sys
+import time
+from tempfile import TemporaryDirectory
+
+import numpy as np
+
+HERE = Path(__file__).resolve().parent
+sys.path.insert(0, str(HERE))
+import run_refined_dfg as rr  # noqa: E402
+
+RE = 20.0
+U_PHYS = 0.2                   # m/s mean inflow; rho = 1
+TSTAR_MAX = 60.0               # 30 s physical
+SAMPLE_TSTAR = 0.5             # 0.25 s
+WINDOW_TSTAR = 10.0            # trailing 5 s, 21 samples
+# Schaefer & Turek (1996), 2D-1 lower/upper bounds.
+BANDS = dict(cd=(5.57, 5.59), cl=(0.0104, 0.0110), p_diff=(0.1172, 0.1176))
+STEADY_RANGE = {k: (hi - lo) / 10 for k, (lo, hi) in BANDS.items()}
+STEADY_SLOPE = {k: (hi - lo) / 50 for k, (lo, hi) in BANDS.items()}
+DEVIATIONS = ("pressure-NEQ outlet (neq_extrap), not do-nothing", "local Zou-He inlet")
+
+_re100_config = rr.config
+
+
+def config(h, diameter, collision="SRT", jet_scenario="none"):
+    """The pinned Re100 configuration with the lattice viscosity set for Re 20."""
+    if jet_scenario != "none":
+        raise ValueError("DFG 2D-1 runs have no Jets")
+    cfg, body, p = _re100_config(h, diameter, collision, jet_scenario)
+    nu = p["mean_speed"] * p["diameter"] / RE
+    cfg["physics"]["viscosity"] = nu
+    return cfg, body, dict(p, viscosity=nu, re=RE)
+
+
+def lattice_parameters(diameter, speed=rr.SPEED):
+    nu = speed * diameter / RE
+    return dict(viscosity=nu, tau=0.5 + 3 * nu, steps=math.ceil(TSTAR_MAX * diameter / speed),
+                interval=round(SAMPLE_TSTAR * diameter / speed))
+
+
+def sample(impulse, n, diameter, dp):
+    den = rr.SPEED**2 * diameter / 2
+    return dict(cd=float(impulse[0] / n / den), cl=float(impulse[1] / n / den),
+                dp_scaled=float(dp), p_diff=float(dp) * U_PHYS**2)
+
+
+def summarize(rows):
+    """Steady rule over the trailing window, then the interval verdict (or none)."""
+    if not rows or rows[-1]["tstar"] < TSTAR_MAX - 1e-9:
+        return dict(status="INCOMPLETE", steady=False, verdict=None)
+    window = [r for r in rows if r["tstar"] >= TSTAR_MAX - WINDOW_TSTAR - 1e-9]
+    t = np.array([r["tstar"] for r in window])
+    ranges = {k: max(r[k] for r in window) - min(r[k] for r in window) for k in BANDS}
+    drifts = {k: abs(np.polyfit(t, [r[k] for r in window], 1)[0]) * WINDOW_TSTAR for k in BANDS}
+    steady = all(ranges[k] <= STEADY_RANGE[k] and drifts[k] <= STEADY_SLOPE[k] for k in BANDS)
+    final = {k: float(np.mean([r[k] for r in window])) for k in BANDS}
+    edge = {k: min(final[k] - BANDS[k][0], BANDS[k][1] - final[k]) for k in BANDS}
+    result = dict(window_samples=len(window), ranges=ranges, drifts=drifts,
+                  range_limits=STEADY_RANGE, slope_limits=STEADY_SLOPE, steady=steady,
+                  window_mean=final, edge_distance=edge, bands=BANDS,
+                  deviations=DEVIATIONS, cd_positive=final["cd"] > 0)
+    if not steady:
+        return dict(result, status="NOT_STEADY", verdict=None)
+    inside = {k: BANDS[k][0] <= final[k] <= BANDS[k][1] for k in BANDS}
+    return dict(result, status="WITHIN_INTERVALS" if all(inside.values()) else "OUTSIDE_INTERVALS",
+                verdict=inside)
+
+
+def healthy(row):
+    return row["max_mach"] <= 0.2 and row["rho_min"] >= 0.5 and row["rho_max"] <= 1.5
+
+
+def run_case(a, h, case):
+    from contextlib import ExitStack
+    from CelerisLab.cuda.context import CudaContext
+    from CelerisLab.lbm.descriptors import FLUID
+    from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+    from CelerisLab.simulation import Simulation
+
+    refined = case == "refined"
+    cfgdict, bodydict, p = config(h, 40 if refined else 80, a.collision)
+    lp = lattice_parameters(p["diameter"])
+    rows = []
+    with ExitStack() as stack:
+        if refined:
+            cfg, pair, body, host = rr.host_geometry(a.source, h, a.collision)
+            context = stack.enter_context(CudaContext(0))
+            backend = stack.enter_context(FixedPatchStepper(
+                cfg, context=context, coarse_boundary="channel", refinements=[rr.PLACEMENTS["wide"]],
+                level_blocks=[rr.BASE_BLOCK, rr.FINE_BLOCK], level_objects=[(), (body,)],
+                observe_raw_exchange=True, observe_window_exchange=True))
+            from CelerisLab.lbm.channel import channel_masks
+            from CelerisLab.lbm.refinement import ACTIVE
+            from CelerisLab.lbm.descriptors import SOLID
+            masks = [channel_masks(pair.coarse_shape)[0] & ~pair.covered,
+                     (pair.fine_kind == ACTIVE)
+                     & ((host.flags.reshape(pair.fine_shape) & SOLID) == 0)]
+            backend.initialize(rr.initial_state(pair.coarse_shape, (0.0, 0.0), 1.0, root=True),
+                               rr.initial_state(pair.fine_shape, pair.fine_origin, 0.5))
+            probe_p = dict(p, center=[(80.5 - pair.fine_origin[0]) / 0.5 - 0.5,
+                                      (81.0 - pair.fine_origin[1]) / 0.5 - 0.5], radius=40.0)
+        else:
+            tmp = Path(stack.enter_context(TemporaryDirectory()))
+            (tmp / "lbm.json").write_text(json.dumps(cfgdict))
+            (tmp / "body.json").write_text(json.dumps(bodydict))
+            sim = stack.enter_context(Simulation(str(tmp / "lbm.json"), str(tmp / "body.json"),
+                                                 device_id=0))
+            sim.initialize()
+            probe_p = p
+            fluid = (sim.field.flag.reshape(p["ny"], p["nx"]) & FLUID) != 0
+        begin = 0
+        while begin < lp["steps"]:
+            rr.check_deadline(a.deadline)
+            n = min(lp["interval"], lp["steps"] - begin)
+            if refined:
+                impulse = rr.advance_refined_interval(backend, n, a.deadline, "window")
+                fields = [rr.macros(g) for g in backend.download()]
+                rho, ux, uy = fields[1]
+                fluid = masks[1]
+                vals = [tuple(v[m] for v in f) for m, f in zip(masks, fields)]
+            else:
+                impulse, done = np.zeros(2), 0
+                while done < n:
+                    batch = min(100, n - done)
+                    sim.run(batch)
+                    impulse += sim.read_force(0).astype(np.float64) * batch
+                    done += batch
+                m = sim.get_macroscopic()
+                rho, ux, uy = [np.asarray(m[k], np.float64).reshape(p["ny"], p["nx"])
+                               for k in ("rho", "ux", "uy")]
+                vals = [(rho[fluid], ux[fluid], uy[fluid])]
+            begin += n
+            row = dict(step=begin, tstar=begin * rr.SPEED / p["diameter"],
+                       **sample(impulse, n, p["diameter"],
+                                h.pressure_probes(rho, fluid, probe_p)["dp"]),
+                       rho_min=float(min(v[0].min() for v in vals)),
+                       rho_max=float(max(v[0].max() for v in vals)),
+                       max_mach=float(max(np.sqrt(3 * (v[1]**2 + v[2]**2)).max() for v in vals)))
+            rr.ensure_finite(row)
+            rows.append(row)
+            if not healthy(row):
+                break
+    summary = summarize(rows) if healthy(rows[-1]) else dict(status="UNHEALTHY", steady=False,
+                                                             verdict=None)
+    return dict(case=case, collision=a.collision, lattice=lp, rows=rows, **summary)
+
+
+def plan(collision="SRT"):
+    return dict(re=RE, collision=collision, tstar_max=TSTAR_MAX, sample_tstar=SAMPLE_TSTAR,
+                window_tstar=WINDOW_TSTAR, bands=BANDS, range_limits=STEADY_RANGE, slope_limits=STEADY_SLOPE,
+                deviations=DEVIATIONS, pressure_scale=U_PHYS**2,
+                refined_root=lattice_parameters(40), refined_fine_tau=0.5 + 3 * 2 * rr.SPEED * 40 / RE,
+                uniform=lattice_parameters(80),
+                label="compared against the official reference bands (non-canonical neq_extrap outlet, zou_he_local inlet); not a full benchmark certification")
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__.splitlines()[0])
+    parser.add_argument("--case", choices=["refined", "uniform", "both"], default="refined")
+    parser.add_argument("--collision", choices=["SRT", "MRT"], default="SRT")
+    parser.add_argument("--source", type=Path, default=rr.DEFAULT_SOURCE)
+    parser.add_argument("--output", type=Path, required=True)
+    parser.add_argument("--plan", action="store_true", help="Write the frozen plan; no CUDA")
+    parser.add_argument("--max-wall-seconds", type=float)
+    a = parser.parse_args()
+    a.source = a.source.resolve()
+    a.output.mkdir(parents=True, exist_ok=False)
+    (a.output / "plan.json").write_text(json.dumps(plan(a.collision), indent=2, allow_nan=False))
+    if a.plan:
+        return
+    if not a.max_wall_seconds or a.max_wall_seconds <= 0 or not os.environ.get("CUDA_VISIBLE_DEVICES"):
+        raise SystemExit("GPU run needs --max-wall-seconds > 0 and explicit CUDA_VISIBLE_DEVICES")
+    h, _ = rr.helpers(a.source)
+    rr.config = config  # host_geometry builds its root configuration through this name
+    a.deadline = time.monotonic() + a.max_wall_seconds
+    cases = ["refined", "uniform"] if a.case == "both" else [a.case]
+    results = [run_case(a, h, case) for case in cases]
+    (a.output / "result.json").write_text(json.dumps(results, indent=2, allow_nan=False))
+
+
+if __name__ == "__main__":
+    main()

tests/validation/run_jet_circle_cv.py

@@ -74,7 +74,25 @@ def cv(g, fluid, origin, rectangle, nu, h=0.5):
     Cauchy stress from centred velocity gradients, NOT postcollision non-equilibrium.
     Rectangle faces lie on cell edges; fluid mask excludes the staircase cylinder.
     """
-    rho, mom, vel = moments(g)
+    return _cv_quantities(*moments(g), fluid, origin, rectangle, nu, h)
+
+
+def cv_from_macros(rho, ux, uy, fluid, origin, rectangle, nu, h=0.5):
+    """Same CV calculation from public macros, in one common length/time frame.
+
+    ``origin`` is the lower cell edge, ``h`` the cell width, and ``nu`` uses
+    that frame. Returns [mass, Px, Py] inventory and outward total flux.
+    Crop fields/mask together and translate origin before calling if desired.
+    No boundary-exchange observation enters this independent estimator.
+    """
+    rho = np.asarray(rho, np.float64)
+    if rho.ndim != 2 or any(np.shape(v) != rho.shape for v in (ux, uy, fluid)):
+        raise ValueError("CV macros and fluid mask must have matching 2D shapes")
+    vel = np.asarray([ux, uy], np.float64)
+    return _cv_quantities(rho, rho[None] * vel, vel, fluid, origin, rectangle, nu, h)
+
+
+def _cv_quantities(rho, mom, vel, fluid, origin, rectangle, nu, h):
     x0, y0, x1, y1 = np.rint((np.array(rectangle) - np.tile(origin, 2)) / h).astype(int)
     np.testing.assert_allclose(
         np.array([x0, y0, x1, y1]) * h + np.tile(origin, 2), rectangle, atol=1e-12
@@ -261,12 +279,11 @@ def run_case(c, ctx, output, provenance, continuation=False):
         PeriodicBlockStepper(fine_cfg, layout, context=ctx, objects=[fobj]) as uniform,
         FixedPatchStepper(
             cfg,
-            c["bounds"],
             context=ctx,
-            base_block=(8, 8),
-            fine_block=(8, 8),
-            fine_objects=[obj],
-            jets=jets,
+            refinements=[c["bounds"]],
+            level_blocks=[(8, 8), (8, 8)],
+            level_objects=[(), (obj,)],
+            level_jets=[(), tuple(jets)],
             observe_raw_exchange=True,
         ) as amr,
     ):

tests/validation/run_refined_dfg.py

@@ -17,8 +17,109 @@ HERE = Path(__file__).resolve().parent
 DEFAULT_SOURCE = HERE.parents[1]
 BOUNDS = (40, 24, 320, 140)
 BASE_BLOCK = (2, 881)  # Exact 166 x 881 domain; no geometric padding.
-FINE_BLOCK = (8, 8)  # 240 x 568 including four fine halo cells.
+FINE_BLOCK = (8, 8)  # Wide 240 x 568 or compact 168 x 168, including fine halos.
 SPEED = 0.01
+JET_PEAK = 0.001
+JET_WINDOWS = [(t, t + 10) for t in range(250, 300, 10)]
+CV_CONTOURS = np.array([(48., 44., 160., 120.), (44., 36., 180., 128.)])
+PLACEMENTS = {"wide": BOUNDS, "compact": (48, 48, 128, 128)}
+COMPACT_CV_CONTOURS = np.array([(52., 52., 112., 112.), (50., 50., 120., 120.)])
+
+
+def selected_cases(selection="both", uniform_diameter=80):
+    return [(case, diameter) for case, diameter in (("uniform", uniform_diameter), ("refined", 40))
+            if selection in ("both", case)]
+
+
+def uniform_origin(diameter):
+    h = 40 / diameter
+    return (.5 - .5*h, 1 - h)
+
+
+def cv_contours(refined, diameter, placement="wide"):
+    contours = CV_CONTOURS if placement == "wide" else COMPACT_CV_CONTOURS
+    return contours if refined else contours + [20/diameter, 0., 20/diameter, 0.]
+
+
+def cv_selected(step, diameter, mode):
+    return mode == "pilot" or round(290 * diameter / SPEED) <= step <= round(300 * diameter / SPEED)
+
+
+def cv_snapshot(rho, ux, uy, fluid, origin, refined, begin, n, diameter, impulse, placement="wide"):
+    """Cropped public macros only; all inventories, fluxes and impulses use root units."""
+    from run_jet_circle_cv import cv_from_macros
+
+    rectangles = cv_contours(refined, diameter, placement)
+    h_cell = .5 if refined else 40 / diameter
+    origin = np.asarray(origin, dtype=float)
+    x0, y0, x1, y1 = np.rint((rectangles[-1] - np.tile(origin, 2)) / h_cell).astype(int)
+    selection = np.s_[y0-2:y1+2, x0-2:x1+2]
+    cropped_origin = origin + h_cell * np.array([x0-2, y0-2])
+    fields = [v[selection] for v in (rho, ux, uy, fluid)]
+    values = [cv_from_macros(*fields, cropped_origin, rect, .004, h=h_cell)
+              for rect in rectangles]
+    h_native = 1.0 if refined else 40 / diameter
+    return dict(step=begin+n, tstar=(begin+n)*SPEED/diameter,
+                time_root=(begin+n)*h_native, interval_start_root=begin*h_native,
+                interval_dt_root=n*h_native,
+                interval_impulse_root=(np.asarray(impulse)*h_native**2).tolist(),
+                inventory=[inventory.tolist() for inventory, _ in values],
+                outward_flux=[flux.tolist() for _, flux in values])
+
+
+def public_jets(center, radius, scenario):
+    """Same physical apertures; caller supplies its public geometry frame."""
+    if scenario == "none":
+        return ()
+    from CelerisLab.body.jets import Jet, Sector, Radial
+
+    peak = 0.0 if scenario == "zero" else JET_PEAK
+    return tuple(
+        Jet(Sector(center, angle, math.pi / 12, radius * 1.005),
+            Radial(center, inward=False), "parabolic", amplitude)
+        for angle, amplitude in ((math.pi / 2, peak), (3 * math.pi / 2, -peak))
+    )
+
+
+def jet_interval(handle, scenario, begin, n, diameter):
+    """Apply a reversal only at a completed interval boundary, before advancing."""
+    if scenario != "none":
+        boundaries = [round(t * diameter / SPEED) for t in (250, 260, 270, 280, 290, 300, 305)]
+        n = min([n] + [step - begin for step in boundaries if step > begin])
+    if scenario != "reverse":
+        return n, None
+    switch = round(300 * diameter / SPEED)
+    if begin < switch:
+        return min(n, switch - begin), None
+    if begin == switch:
+        handle.set_jet_amplitude(0, -JET_PEAK)
+        handle.set_jet_amplitude(1, JET_PEAK)
+        return n, dict(step=begin, tstar=300.0, amplitudes=[-JET_PEAK, JET_PEAK])
+    return n, None
+
+
+def jet_diagnostics(rows, scenario):
+    """Raw interval-weighted diagnostics only; no driven-flow qualification gate."""
+    windows = JET_WINDOWS + ([(300, 305)] if scenario == "reverse" else [])
+    result = []
+    for start, end in windows:
+        selected = [r for r in rows if start <= r["tstar_force"] < end]
+        if not selected:
+            continue
+        if any(2 * r["tstar_force"] - r["tstar"] < start - 1e-9
+               or r["tstar"] > end + 1e-9 for r in selected):
+            raise ValueError("Jet diagnostic interval crosses a window boundary")
+        weights = np.array([2 * (r["tstar"] - r["tstar_force"]) for r in selected])
+        record = dict(start_tstar=start, end_tstar=end, samples=len(selected),
+                      covered_tstar=float(weights.sum()),
+                      complete=bool(np.isclose(weights.sum(), end - start, rtol=0, atol=1e-9)))
+        for key in ("cd", "cl"):
+            values = np.array([r[key] for r in selected])
+            record[key] = dict(mean=float(np.average(values, weights=weights)),
+                               min=float(values.min()), max=float(values.max()),
+                               rms=float(np.sqrt(np.average(values**2, weights=weights))))
+        result.append(record)
+    return result
 
 
 def helpers(source):
@@ -39,26 +140,34 @@ def helpers(source):
     return h, periodic
 
 
-def config(h, diameter, collision="SRT"):
+def config(h, diameter, collision="SRT", jet_scenario="none"):
     cfg, body, p = h.make_config(diameter, SPEED, collision, "neq_extrap", "zou_he_local")
     cfg["method"]["outlet"]["srt_neq_damp"] = 0.5
+    if jet_scenario != "none":
+        if collision != "MRT":
+            raise ValueError("Public Jet scenarios require MRT")
+        cfg["method"]["ddf_shifting"] = False
     return cfg, body, p
 
 
-def plan(h, collision="SRT", observation_mode="last_base_step"):
-    _, _, c = config(h, 40, collision)
-    _, _, f = config(h, 80, collision)
-    x0, y0, x1, y1 = BOUNDS
+def plan(h, collision="SRT", observation_mode="last_base_step", jet_scenario="none", uniform_diameter=80, placement="wide"):
+    _, _, c = config(h, 40, collision, jet_scenario)
+    _, _, f = config(h, uniform_diameter, collision, jet_scenario)
+    bounds = PLACEMENTS[placement]
+    x0, y0, x1, y1 = bounds
     fine_active = 4 * (x1 - x0) * (y1 - y0)
     coarse_uncovered = c["nx"] * c["ny"] - fine_active // 4
     return dict(
         status="prepared_not_run",
         collision=collision,
+        jet_scenario=jet_scenario,
+        ddf_shifting=jet_scenario == "none",
         observation_mode=observation_mode,
         root_gpu_gates_required=True,
         root=c,
         uniform=f,
-        bounds_root_edge=list(BOUNDS),
+        bounds_root_edge=list(bounds),
+        fine_origin_root_edge=[x0 - 2, y0 - 2],
         physical_patch_m=[
             (x0 - 0.5) * c["dx"],
             (y0 - 1) * c["dx"],
@@ -78,19 +187,31 @@ def plan(h, collision="SRT", observation_mode="last_base_step"):
         physical_viscosity_m2_s=0.001,
         finest_dx_m=f["dx"],
         finest_dt_s=f["dt"],
+        refined_finest_dx_m=c["dx"]/2,
+        refined_finest_dt_s=c["dt"]/2,
         approximate_leaf_cell_ratio=(coarse_uncovered + fine_active)
         / (f["nx"] * f["ny"]),
         approximate_collision_work_ratio=(coarse_uncovered + 2 * fine_active)
-        / (2 * f["nx"] * f["ny"]),
+        / ((uniform_diameter/40) * f["nx"] * f["ny"]),
         cost_caveat="Geometry-only estimates include ghosts/solids; exclude block halos, transfers, observation and launch overhead. Not measured speedup.",
-        sampling_offset="For origin (38,22), fine physical x=(j+75.5)*dx80 versus uniform i*dx80: half-fine-cell x offset. Fine physical y=(k+42.5)*dx80 equals uniform (l-.5)*dx80 at l=k+43: y samples align. Body, wall and source planes coincide.",
+        sampling_offset="Uniform nodes map to root edges X=.5+i*(40/D), Y=1+(j-.5)*(40/D). Refined centers are fine_origin_root_edge+(index+.5)*.5. Body, wall and source planes coincide; compare observables, not pointwise arrays.",
         outlet=(
             "SRT pressure NEQ, all nine fneq damped by .5, no backflow clamp; noncanonical DFG outlet"
             if collision == "SRT" else
             "MRT pressure NEQ, unknowns q2/6/8 extrapolated; no backflow clamp; noncanonical DFG outlet"
         ),
-        modes={"pilot_tstar": 5.0, "full_tstar": 300.0},
+        modes={"pilot_tstar": 5.0, "full_tstar": 305.0 if jet_scenario == "reverse" else 300.0},
         observation_dense_tstar=0.02,
+        **(dict(jets=dict(
+            uniform_origin_node=f["center"], refined_origin_root_edge=[80.5, 81.0],
+            angles_rad=[math.pi / 2, 3 * math.pi / 2], half_width_rad=math.pi / 12,
+            radial_extent_over_radius=[0.0, 1.005],
+            field="Radial(outward)", profile="parabolic", condition="velocity",
+            initial_amplitudes=[0.0, 0.0] if jet_scenario == "zero" else [JET_PEAK, -JET_PEAK],
+            reversal_tstar=300.0 if jet_scenario == "reverse" else None,
+            diagnostic_windows_tstar=JET_WINDOWS + ([(300, 305)] if jet_scenario == "reverse" else []),
+            transfer="constant", driven_qualification="not assessed",
+        )) if jet_scenario != "none" else {}),
     )
 
 
@@ -126,7 +247,7 @@ def initial_state(shape, origin, spacing, height=164.0, speed=SPEED, root=False)
     return w * (1 + 3 * cu + 4.5 * cu * cu - 1.5 * velocity * velocity)
 
 
-def host_geometry(source, h, collision="SRT"):
+def host_geometry(source, h, collision="SRT", jet_scenario="none", placement="wide"):
     """Production donor/transfer and complete cut-link checks, without CUDA."""
     from CelerisLab.config import load_lbm_config
     from CelerisLab.lbm.refinement import FixedPatch2D, prepare_fine_boundary
@@ -139,12 +260,12 @@ def host_geometry(source, h, collision="SRT"):
     from CelerisLab.body.geometry.circle import CircleGeometry
     from CelerisLab.body.objects import SimObject
 
-    cfgdict, _, _ = config(h, 40, collision)
+    cfgdict, _, _ = config(h, 40, collision, jet_scenario)
     with TemporaryDirectory() as tmp:
         p = Path(tmp) / "cfg.json"
         p.write_text(json.dumps(cfgdict))
         cfg = load_lbm_config(str(p))
-    pair = FixedPatch2D.from_config(cfg, BOUNDS)
+    pair = FixedPatch2D.from_config(cfg, PLACEMENTS[placement])
     layout = BlockLayout2D.build(cfg.ny, cfg.nx, *BASE_BLOCK)
     validate_channel_config(cfg, [BASE_BLOCK, FINE_BLOCK])
     targets = blocked_base_tables(pair, layout, np.argsort(pair.coalesce_coarse))[1]
@@ -153,18 +274,41 @@ def host_geometry(source, h, collision="SRT"):
     body = SimObject(
         0, geometry=CircleGeometry(80.5, 81.0, 20.0), center=(80.5, 81.0), radius=20.0
     )
-    host = prepare_fine_boundary(cfg, pair, (body,), _shifted_blocked=True)
+    host = prepare_fine_boundary(
+        cfg, pair, (body,), jets=public_jets((80.5, 81.0), 20.0, jet_scenario),
+        _shifted_blocked=jet_scenario == "none",
+    )
     return cfg, pair, body, host
 
 
-def verify_source(source, collision="SRT", observation_mode="last_base_step"):
+def verify_source(source, collision="SRT", observation_mode="last_base_step", jet_scenario="none", observe_cv=False):
     if collision not in ("SRT", "MRT"):
         raise ValueError("Matched collision must be SRT or MRT")
     if observation_mode not in ("last_base_step", "window"):
         raise ValueError("Unknown refined observation mode")
+    if observe_cv and jet_scenario == "none":
+        raise ValueError("--observe-cv requires an explicit public Jet scenario")
     name = ("source_manifest_window.json" if observation_mode == "window" else
             "source_manifest.json" if collision == "SRT" else "source_manifest_mrt.json")
     expected = json.loads((HERE / "data/refined_dfg" / name).read_text())
+    if jet_scenario != "none":
+        if collision != "MRT" or observation_mode != "window":
+            raise ValueError("Public Jet scenarios require MRT and --exchange-window")
+        import subprocess
+        sys.path.insert(0, str(HERE))
+        from source_provenance import checkout_provenance
+
+        paths = set(expected["sha256"])
+        if observe_cv:
+            paths.add("tests/validation/run_jet_circle_cv.py")
+        paths.update(subprocess.check_output(
+            ["git", "-C", str(source), "ls-files", "src"], text=True).splitlines())
+        return dict(
+            source_capture="current checkout, not historical DFG qualification",
+            **checkout_provenance(source), observation_mode=observation_mode,
+            sha256={name: hashlib.sha256((source / name).read_bytes()).hexdigest()
+                    for name in sorted(paths)},
+        )
     for name, digest in expected["sha256"].items():
         if hashlib.sha256((source / name).read_bytes()).hexdigest() != digest:
             raise RuntimeError("Source mismatch: " + name)
@@ -219,29 +363,32 @@ def run_case(a, h, periodic, case):
     from contextlib import ExitStack
 
     refined = case == "refined"
-    cfgdict, bodydict, p = config(h, 40 if refined else 80, a.collision)
-    cap = 5.0 if a.mode == "pilot" else 300.0
+    cfgdict, bodydict, p = config(h, 40 if refined else a.uniform_diameter, a.collision, a.jet_scenario)
+    cap = 5.0 if a.mode == "pilot" else (305.0 if a.jet_scenario == "reverse" else 300.0)
     max_steps = math.ceil(cap * p["diameter"] / SPEED)
     rows = []
     healthy = True
     step_seconds = 0.0
+    cv_seconds, cv_samples = 0.0, 0
+    amplitude_changes = []
     started = time.monotonic()
     check_deadline(a.deadline)
     dest = a.output / case
     dest.mkdir()
     with ExitStack() as stack:
+        cv_out = stack.enter_context((dest / "cv_history.jsonl").open("w")) if a.observe_cv else None
         if refined:
-            cfg, pair, body, host = host_geometry(a.source, h, a.collision)
+            cfg, pair, body, host = host_geometry(a.source, h, a.collision, a.jet_scenario, a.placement)
             context = stack.enter_context(CudaContext(0))
             backend = stack.enter_context(
                 FixedPatchStepper(
                     cfg,
-                    BOUNDS,
                     context=context,
                     coarse_boundary="channel",
-                    base_block=BASE_BLOCK,
-                    fine_block=FINE_BLOCK,
-                    fine_objects=(body,),
+                    refinements=[PLACEMENTS[a.placement]],
+                    level_blocks=[BASE_BLOCK, FINE_BLOCK],
+                    level_objects=[(), (body,)],
+                    level_jets=[(), public_jets((80.5, 81.0), 20.0, a.jet_scenario)],
                     observe_raw_exchange=True,
                     **({"observe_window_exchange": True} if a.observation_mode == "window" else {}),
                 )
@@ -274,6 +421,9 @@ def run_case(a, h, periodic, case):
             sim = stack.enter_context(
                 Simulation(str(tmp / "lbm.json"), str(tmp / "body.json"), device_id=0)
             )
+            for index, jet in enumerate(public_jets(p["center"], p["radius"], a.jet_scenario)):
+                if sim.add_jet(jet) != index:
+                    raise RuntimeError("Unexpected public Jet registration index")
             sim.initialize()
             probe_p = p
             fluid = (sim.field.flag.reshape(p["ny"], p["nx"]) & FLUID) != 0
@@ -290,6 +440,10 @@ def run_case(a, h, periodic, case):
                     else 0.5
                 )
                 n = min(max(1, round(width * p["diameter"] / SPEED)), max_steps - begin)
+                n, change = jet_interval(backend if refined else sim, a.jet_scenario,
+                                         begin, n, p["diameter"])
+                if change is not None:
+                    amplitude_changes.append(change)
                 impulse = np.zeros(2)
                 tick = time.monotonic()
                 if refined:
@@ -359,6 +513,19 @@ def run_case(a, h, periodic, case):
                     / SPEED,
                     **h.pressure_probes(rho, fluid, probe_p),
                 )
+                if a.jet_scenario != "none":
+                    row.update(step_start=begin, tstar_start=begin * SPEED / p["diameter"],
+                               interval_steps=n, impulse_x=float(impulse[0]),
+                               impulse_y=float(impulse[1]))
+                if a.observe_cv and cv_selected(begin+n, p["diameter"], a.mode):
+                    tick = time.monotonic()
+                    record = cv_snapshot(rho, ux, uy, fluid,
+                                         pair.fine_origin if refined else uniform_origin(p["diameter"]),
+                                         refined, begin, n, p["diameter"], impulse, a.placement)
+                    cv_out.write(json.dumps(record, allow_nan=False) + "\n")
+                    cv_out.flush()
+                    cv_samples += 1
+                    cv_seconds += time.monotonic() - tick
                 check_deadline(a.deadline)
                 ensure_finite(row)
                 rows.append(row)
@@ -402,14 +569,24 @@ def run_case(a, h, periodic, case):
         steps=begin,
         startup_seconds=startup,
         step_seconds=step_seconds,
+        observation_and_overhead_seconds=time.monotonic() - started - startup - step_seconds,
         wall_seconds=time.monotonic() - started,
         force="raw stationary-body link momentum reaction, per-unit-depth; average over entire observation interval; normalized in each base lattice",
         mass_units="rho_reference * m^2 per unit depth",
         flux_units="rho_reference * m^2/s per unit depth",
         pressure="postcollision end-sample; physical pressure difference / (rho_reference * 1 m/s squared)",
         suite_validated=False,
-        **periodic.assess(rows),
+        **(periodic.assess(rows) if a.jet_scenario in ("none", "zero") else
+           dict(periodic=False, checks={}, cycles=[])),
     )
+    if a.jet_scenario != "none":
+        result.update(jet_diagnostic_windows=jet_diagnostics(rows, a.jet_scenario),
+                      amplitude_changes=amplitude_changes,
+                      pressure_classification=("DFG reference verdict retained, including pressure"
+                          if a.jet_scenario == "zero" else "diagnostic only; no official DFG reference gate"))
+    if a.observe_cv:
+        result.update(cv_observation_seconds=cv_seconds, cv_samples=cv_samples,
+                      cv_qualification="not assessed; scalar observations only")
     if not healthy:
         result["periodic"] = False
     reference = json.loads(
@@ -425,7 +602,8 @@ def run_case(a, h, periodic, case):
         else (
             "pilot_complete"
             if a.mode == "pilot"
-            else ("periodic" if result["periodic"] else "incomplete_periodicity")
+            else ("driven_diagnostics_complete" if a.jet_scenario in ("fixed", "reverse") else
+                  "periodic" if result["periodic"] else "incomplete_periodicity")
         )
     )
     (dest / "summary.json").write_text(json.dumps(result, indent=2, allow_nan=False))
@@ -464,27 +642,39 @@ def check_deadline(deadline):
         )
 
 
-def verify_pilot(path, manifest, collision="SRT", observation_mode="last_base_step"):
+def verify_pilot(path, manifest, collision="SRT", observation_mode="last_base_step", jet_scenario="none", observe_cv=False,
+                 case_selection="both", uniform_diameter=80, placement="wide"):
     if path is None:
         raise ValueError(
             "Full comparison requires --pilot-results from a completed matched pilot"
         )
     previous = read_finite_json(path / "plan.json")
+    options = dict(case_selection=(case_selection, "both"), uniform_diameter=(uniform_diameter, 80), placement=(placement, "wide"))
+    if any(previous.get(k, default) != expected for k, (expected, default) in options.items()):
+        raise ValueError("Pilot case, resolution or placement differs")
+    if previous.get("observe_cv", False) != observe_cv:
+        raise ValueError("Pilot CV observation mode differs")
+    if previous.get("jet_scenario", "none") != jet_scenario:
+        raise ValueError("Pilot Jet scenario differs")
     if previous.get("observation_mode") != observation_mode:
         raise ValueError("Pilot observation mode differs")
     if previous.get("collision") != collision:
         raise ValueError("Pilot collision differs")
     if previous["source_manifest"] != manifest or previous["bounds_root_edge"] != list(
-        BOUNDS
+        PLACEMENTS[placement]
     ):
         raise ValueError("Pilot source or frozen geometry differs")
     if previous.get("runner_sha256") != runner_hash():
         raise ValueError("Pilot runner differs")
     calibrated_seconds = 0.0
-    for case, diameter in (("uniform", 80), ("refined", 40)):
+    for case, diameter in selected_cases(case_selection, uniform_diameter):
         summary = read_finite_json(path / case / "summary.json")
         if (
-            summary.get("observation_mode") != observation_mode
+            any(summary.get(k, default) != expected for k, (expected, default) in options.items())
+            or summary.get("observe_cv", False) != observe_cv
+            or (observe_cv and summary.get("cv_samples") != 250)
+            or summary.get("jet_scenario", "none") != jet_scenario
+            or summary.get("observation_mode") != observation_mode
             or summary.get("collision") != collision
             or summary.get("case") != case
             or summary.get("mode") != "pilot"
@@ -537,7 +727,10 @@ def verify_pilot(path, manifest, collision="SRT", observation_mode="last_base_st
                 raise ValueError("Unhealthy pilot observations")
         ready = read_finite_json(path / case / "ready.json")
         if (
-            ready.get("observation_mode") != observation_mode
+            any(ready.get(k, default) != expected for k, (expected, default) in options.items())
+            or ready.get("observe_cv", False) != observe_cv
+            or ready.get("jet_scenario", "none") != jet_scenario
+            or ready.get("observation_mode") != observation_mode
             or ready.get("collision") != collision
             or ready.get("case") != case
             or ready.get("mode") != "pilot"
@@ -555,9 +748,9 @@ def verify_pilot(path, manifest, collision="SRT", observation_mode="last_base_st
             or not 0 < stepping <= wall
         ):
             raise ValueError("Invalid pilot calibration timings")
-        # Full physical duration is 60 times pilot; dense pilot sampling makes
+        # Full duration is 60 (or 61 with reversal) times pilot; dense sampling makes
         # this conservative for the sparser early full observations. Not a guarantee.
-        calibrated_seconds += startup + 60 * (wall - startup)
+        calibrated_seconds += startup + (61 if jet_scenario == "reverse" else 60) * (wall - startup)
     return 3 * calibrated_seconds
 
 
@@ -566,6 +759,13 @@ def main():
     parser.add_argument("--exchange-window", action="store_true",
                         help="Opt in to refined window-ledger observations on the pinned candidate")
     parser.add_argument("--collision", choices=["SRT", "MRT"], default="SRT")
+    parser.add_argument("--case", choices=["both", "uniform", "refined"], default="both")
+    parser.add_argument("--uniform-diameter", type=int, choices=[80, 160], default=80)
+    parser.add_argument("--placement", choices=list(PLACEMENTS), default="wide")
+    parser.add_argument("--jet-scenario", choices=["none", "zero", "fixed", "reverse"],
+                        default="none", help="Opt in to public unshifted MRT Jets; reverse switches at t*=300 and ends at 305")
+    parser.add_argument("--observe-cv", action="store_true",
+                        help="Record independent cropped-macro CV samples during pilot or full t*=290..300")
     parser.add_argument("--source", type=Path, default=DEFAULT_SOURCE)
     parser.add_argument("--output", type=Path, required=True)
     parser.add_argument("--mode", choices=["plan", "pilot", "full"], default="plan")
@@ -577,15 +777,26 @@ def main():
         help="Required explicit campaign wall-time ceiling for GPU execution",
     )
     a = parser.parse_args()
+    if a.uniform_diameter == 160 and a.case != "uniform":
+        raise ValueError("D160 requires --case uniform; no matched-resolution claim")
+    if a.placement == "compact" and a.case != "refined":
+        raise ValueError("Compact placement requires --case refined")
+    if (a.uniform_diameter != 80 or a.placement != "wide") and a.jet_scenario == "none":
+        raise ValueError("Resolution/placement variants require an explicit public Jet scenario")
     a.source = a.source.resolve()
     a.observation_mode = "window" if a.exchange_window else "last_base_step"
-    manifest = verify_source(a.source, a.collision, a.observation_mode)
+    manifest = verify_source(a.source, a.collision, a.observation_mode, a.jet_scenario, a.observe_cv)
     h, periodic = helpers(a.source)
     from source_provenance import checkout_provenance
 
     a.provenance = checkout_provenance(a.source)
+    a.provenance.update(jet_scenario=a.jet_scenario, ddf_shifting=a.jet_scenario == "none")
+    a.provenance.update(case_selection=a.case, uniform_diameter=a.uniform_diameter, placement=a.placement)
+    if a.observe_cv:
+        a.provenance["observe_cv"] = True
     calibration = (
-        verify_pilot(a.pilot_results, manifest, a.collision, a.observation_mode)
+        verify_pilot(a.pilot_results, manifest, a.collision, a.observation_mode, a.jet_scenario, a.observe_cv,
+                     a.case, a.uniform_diameter, a.placement)
         if a.mode == "full" else None
     )
     if a.mode != "plan" and (
@@ -602,17 +813,31 @@ def main():
     if a.output.exists() and any(a.output.iterdir()):
         raise ValueError("Output must be empty/new")
     a.output.mkdir(parents=True, exist_ok=True)
-    specification = plan(h, a.collision, a.observation_mode)
+    specification = plan(h, a.collision, a.observation_mode, a.jet_scenario, a.uniform_diameter, a.placement)
     specification.update(
         source_manifest=manifest,
         **a.provenance,
-        gates=periodic.GATES,
-        reference_gates=periodic.REFERENCE_GATES,
+        gates=periodic.GATES if a.jet_scenario in ("none", "zero") else None,
+        reference_gates=periodic.REFERENCE_GATES if a.jet_scenario in ("none", "zero") else None,
         runner_sha256=runner_hash(),
         wall_time_budget_seconds=budget,
         full_pilot_calibration_seconds=calibration,
     )
-    host_geometry(a.source, h, a.collision)
+    if a.observe_cv:
+        specification["cv_observation"] = dict(
+            full_tstar=[290, 300], pilot_tstar=[.02, 5], spacing_tstar=.02,
+            refined_contours_root=cv_contours(True, 40, a.placement).tolist(),
+            uniform_contours_root=cv_contours(False, a.uniform_diameter, a.placement).tolist(),
+            uniform_grid_origin_root=uniform_origin(a.uniform_diameter),
+            refined_grid_origin_root=specification["fine_origin_root_edge"],
+            cell_width_root={"uniform": 40/a.uniform_diameter, "refined": .5},
+            viscosity_root=.004, components=["mass", "Px", "Py"],
+            units="root lattice per unit depth; outward flux per root time",
+            interval="Impulse belongs to the interval ending at each snapshot; first retained endpoint (pilot .02, full 290) supplies initial inventory only: exclude its preceding impulse",
+            qualification="not assessed; trapezoidal integration and time-coarsening are offline",
+        )
+    if a.case != "uniform":
+        host_geometry(a.source, h, a.collision, a.jet_scenario, a.placement)
     (a.output / "plan.json").write_text(json.dumps(specification, indent=2))
     if a.mode == "plan":
         return 0
@@ -621,8 +846,9 @@ def main():
             "GPU execution blocked: root GPU gates approval and explicit device selection required"
         )
     results = []
+    cases = selected_cases(a.case, a.uniform_diameter)
     a.deadline = time.monotonic() + budget
-    for case in ("uniform", "refined"):
+    for case, _ in cases:
         try:
             result = run_case(a, h, periodic, case)
         except TimeoutError as exc:
@@ -650,7 +876,12 @@ def main():
         status="pilot_only" if a.mode == "pilot" else "incomplete",
         matched_case_pass=False,
     )
-    if len(results) == 2 and all(r["healthy"] and r["periodic"] for r in results):
+    if a.jet_scenario in ("fixed", "reverse"):
+        comparison.update(status=("pilot_only" if a.mode == "pilot" else "driven_diagnostics_complete")
+                          if len(results) == len(cases) and all(r["healthy"] for r in results) else "invalid_state",
+                          qualification="not assessed; execution completion is not physical acceptance",
+                          jet_diagnostic_windows={r["case"]: r["jet_diagnostic_windows"] for r in results})
+    if a.case == "both" and len(results) == 2 and all(r["healthy"] and r["periodic"] for r in results):
         comparison.update(
             status="periodic_comparison",
             refined_vs_uniform=periodic.compare(
@@ -660,14 +891,20 @@ def main():
         comparison["matched_case_pass"] = comparison["refined_vs_uniform"][
             "reference_agreement"
         ]
+    if a.case != "both":
+        comparison.update(status="single_case_" + results[0]["status"],
+                          own_reference_comparison=results[0].get("reference_comparison"))
     (a.output / "comparison.json").write_text(
         json.dumps(comparison, indent=2, allow_nan=False)
     )
     return (
         0
-        if len(results) == 2
+        if len(results) == len(cases)
         and all(r["healthy"] for r in results)
-        and (a.mode == "pilot" or comparison["matched_case_pass"])
+        and (a.mode == "pilot" or comparison["matched_case_pass"]
+             or a.jet_scenario in ("fixed", "reverse")
+             or (a.case != "both" and results[0]["periodic"]
+                 and results[0].get("reference_comparison", {}).get("reference_agreement", False)))
         else 2
     )
 

tests/validation/taylor_green/test_projected_pairs.py

@@ -249,19 +249,41 @@ class ProjectedPairsTests(unittest.TestCase):
         self.assertFalse(assess(incomplete)["comparisons"][3]["passed"])
 
     def test_pinned_build_sources_match_host_generated_inputs(self):
-        from CelerisLab.cuda import compiler_v2
+        """Check the pinned fingerprints using NUMERICAL_REVISION's renderer and kernels.
+        The imported LBMConfig is unchanged from that revision; this is a host
+        build-input check, not a reconstruction of the historical CUDA toolchain."""
+        import importlib.util
+        import tarfile
 
-        for size in (128, 256):
-            cfg = replace(run_case.config(size, 100), nx=10, ny=10)
-            cfg.validate()
-            root = compiler_v2.generate_config(cfg)
-            try:
-                actual = json.loads((Path(root) / ".celeris-build").read_text())
-                expected = build_metadata(size)
-                self.assertEqual(actual, {key: expected[key] for key in
-                                          ("source_sha256", "config", "n_objects")})
-            finally:
-                compiler_v2.cleanup_build_dir(root)
+        kernels = "src/CelerisLab/lbm/kernels"
+        renderer = "src/CelerisLab/cuda/compiler_v2.py"
+        if subprocess.run(["git", "-C", str(REPO), "cat-file", "-e",
+                           f"{NUMERICAL_REVISION}^{{commit}}"], capture_output=True).returncode:
+            self.fail(f"{NUMERICAL_REVISION} is not in this clone; the pinned build "
+                      "sources cannot be checked")
+        archive = subprocess.run(["git", "-C", str(REPO), "archive", "--format=tar",
+                                  NUMERICAL_REVISION, kernels, renderer],
+                                 capture_output=True, check=True).stdout
+        with tempfile.TemporaryDirectory() as snapshot:
+            with tarfile.open(fileobj=io.BytesIO(archive)) as tar:
+                tar.extractall(snapshot)
+            spec = importlib.util.spec_from_file_location(
+                "CelerisLab.cuda._compiler_a135252", Path(snapshot) / renderer)
+            compiler = importlib.util.module_from_spec(spec)
+            spec.loader.exec_module(compiler)
+            with patch.object(compiler, "_KERNEL_DIR", str(Path(snapshot) / kernels)):
+                for size in (128, 256):
+                    with self.subTest(size=size):
+                        cfg = replace(run_case.config(size, 100), nx=10, ny=10)
+                        cfg.validate()
+                        root = compiler.generate_config(cfg)
+                        try:
+                            actual = json.loads((Path(root) / ".celeris-build").read_text())
+                            expected = build_metadata(size)
+                            self.assertEqual(actual, {key: expected[key] for key in
+                                                      ("source_sha256", "config", "n_objects")})
+                        finally:
+                            compiler.cleanup_build_dir(root)
 
     def test_source_preflight_requires_clean_exact_revision_and_tree(self):
         import oracle

tests/validation/time_refined_dfg.py

@@ -173,8 +173,9 @@ def time_refined(a, h, sched):
         cfg, pair, body, host = r.host_geometry(a.source, h, "MRT")
         context = stack.enter_context(CudaContext(0))
         backend = stack.enter_context(FixedPatchStepper(
-            cfg, r.BOUNDS, context=context, coarse_boundary="channel", base_block=r.BASE_BLOCK,
-            fine_block=r.FINE_BLOCK, fine_objects=(body,), observe_raw_exchange=True))
+            cfg, context=context, coarse_boundary="channel", refinements=[r.BOUNDS],
+            level_blocks=[r.BASE_BLOCK, r.FINE_BLOCK], level_objects=[(), (body,)],
+            observe_raw_exchange=True))
         masks = [channel_masks(pair.coarse_shape)[0] & ~pair.covered,  # as run_refined_dfg
                  (pair.fine_kind == ACTIVE) & ((host.flags.reshape(pair.fine_shape) & SOLID) == 0)]
         fresh = (r.initial_state(pair.coarse_shape, (0.0, 0.0), 1.0, root=True),

tests/validation/time_window_dfg.py

@@ -130,8 +130,8 @@ def main():
         cfg, pair, body, host = r.host_geometry(a.source, h, "MRT")
         context = stack.enter_context(CudaContext(0))
         backend = stack.enter_context(FixedPatchStepper(
-            cfg, r.BOUNDS, context=context, coarse_boundary="channel", base_block=r.BASE_BLOCK,
-            fine_block=r.FINE_BLOCK, fine_objects=(body,), observe_raw_exchange=True,
+            cfg, context=context, coarse_boundary="channel", refinements=[r.BOUNDS],
+            level_blocks=[r.BASE_BLOCK, r.FINE_BLOCK], level_objects=[(), (body,)], observe_raw_exchange=True,
             observe_window_exchange=True))
         masks = [channel_masks(pair.coarse_shape)[0] & ~pair.covered,
                  (pair.fine_kind == ACTIVE) & ((host.flags.reshape(pair.fine_shape) & SOLID) == 0)]