docs: record six-case physical validation evidence for 750e8a4

Preserve native gate outcomes and source/output hashes. No simulation or production source changes.

README.md

@@ -112,9 +112,8 @@ 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, context=ctx, refinements=[(8, 6, 24, 18)],
-                                              level_blocks=[None, (16, 20)],
-                                              level_objects=[(), (body,)]) as run:
+with CudaContext(0) as ctx, FixedPatchStepper(cfg, (8, 6, 24, 18), context=ctx,
+                                              fine_objects=[body], fine_block=(16, 20)) 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)),
@@ -125,8 +124,8 @@ with CudaContext(0) as ctx, FixedPatchStepper(cfg, context=ctx, refinements=[(8,
 
 Scope and limits:
 - The examples below use D2Q9, SRT, FP32, double buffer, unshifted, no LES or
-  forcing. Fixed 2:1 hierarchies use level lists of arbitrary depth; two- and
-  three-level examples below use that same constructor.
+  forcing. Fixed 2:1 hierarchies can use level lists of arbitrary depth; the
+  legacy two/three-level forms are shown below.
 - `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)`),
@@ -134,51 +133,50 @@ 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.
-- `level_blocks=[None, (by, bx)]` stores the fine level in fixed-size blocks; `by`/`bx` must
+- `fine_block=(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
-  `refinements[0]`, with both levels blocked (`level_blocks`). Transfers are built once
+  `bounds`, with both `fine_block` and `base_block`. 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. L1 bodies/Jets (`level_objects[1]`/`level_jets[1]`) work
+  with tiles. Order of rectangles does not matter. `fine_objects`/`jets` 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, add further `refinements` entries as shown below.
+  For nesting, use the level lists or legacy child forms 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, context=ctx, refinements=[((4, 4, 12, 12), (16, 4, 24, 12))],
-          level_blocks=[(4, 4), (4, 4)]) as run:
+          cfg, ((4, 4, 12, 12), (16, 4, 24, 12)), context=ctx,
+          fine_block=(4, 4), base_block=(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
   ```
-- 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.
+- 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.
 
   ```python
   from CelerisLab.body.geometry.circle import CircleGeometry
@@ -190,9 +188,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, context=ctx, refinements=[(4, 4, 12, 12), (8, 8, 16, 16)],
-          level_blocks=[(4, 4)] * 3, level_objects=[(), (), (body,)],
-          level_jets=[(), (), (jet,)]) as run:
+          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:
       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)
@@ -200,16 +198,16 @@ Scope and limits:
       run.step(4)
       coarse_g, l1_g, l2_g = run.download()  # solids/covered/nonleaves are NaN
   ```
-- 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.
+- 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.
 - 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/development/root_channel_candidate.md

@@ -1,6 +1,6 @@
 # Blocked AMR root channel increment
 
-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.
+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.
 
 `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 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
+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
 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
 
-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.
+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.
 
-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`:
+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 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, context=ctx,
-    refinements=[[(4, 4, 20, 28), (24, 4, 44, 28)], [(16, 16, 32, 40), (56, 16, 80, 40)]],
-    level_blocks=[(4, 4)] * 3,
+    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)],
 )
 ```
 
@@ -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 L2 bodies and Jets (`level_objects[2]`, `level_jets[2]`) use root edge
+Stationary whole-circle `child_objects` and `child_jets` 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

@@ -1,87 +0,0 @@
-# Supported settings: 2D Jet + fixed multilevel refinement (candidate c95a7ba)
-
-Every entry names the guard that enforces it (file:line at c95a7ba) 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:65`, `lbm/kernels/blocks.cu:25`, `lbm/kernels/refinement.cu:26` |
-| Precision / streaming | FP32 storage, double buffer | same guards |
-| Collision on refinement levels | SRT or MRT; TRT and LES rejected | `lbm/refinement.py:65` (re-checked for fine boundaries at `lbm/refinement.py:457`) |
-| Shifted (ddf_shifting) storage | only when every level is blocked; dense or mixed rejected | `lbm/refinement_stepper.py:198`, `:469`; `lbm/kernels/refinement.cu:26` (dense patch kernel unshifted only) |
-| Fine relaxation | 1/tau_fine must lie inside the configured omega guards | `lbm/refinement.py:80` |
-| Uniform blocked backend `PeriodicBlockStepper` | doubly periodic, SRT only, stationary passive circles | `lbm/block_stepper.py:84`, `:91`, `:45`-`:67` |
-
-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:72`, `:75` |
-| Single L1 rectangle | two-coarse-cell margin to the domain edge | `lbm/refinement.py:73` |
-| 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:84`; `lbm/refinement.py:283`-`:295` |
-| 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:82`, `:87`; `lbm/hierarchy.py:54` |
-| Nested transfer support | child only over inherited ACTIVE cells, donors/receivers uncovered ACTIVE, no child over the parent restriction footprint | `lbm/refinement.py:391`-`:407` |
-| Tile shape | positive `(by, bx)`; the level allocation must be divisible by the tile (no partial tiles yet) | `lbm/refinement_stepper.py:78`, `lbm/refinement.py:246` (shape); `lbm/blocks.py:134` (divisibility) |
-| Coarse boundary | `periodic`, `slot_reservoir` or `channel` | `lbm/refinement_stepper.py:177` |
-| `slot_reservoir` | on-site velocity bottom slot, reservoir top, velocity 0, 2 <= slot width < nx, patch rows 4..ny-4 | `lbm/refinement_stepper.py:205`, `:212`; `lbm/refinement.py:546` |
-| `channel` | all levels blocked; Zou-He inlet, pressure-NEQ outlet, stationary bounce-back walls; transfers away from the closure | `lbm/channel.py:22`, `:26`, `:43`, `:54` |
-| Explosion | `constant` (default); `projected_linear` only periodic, all-blocked SRT, no bodies/Jets, no LES (experimental) | `lbm/refinement_stepper.py:188`, `:193` |
-| L2 runtime remap (`remap_child`) | periodic three-level L2 sets only, no L0/L1 bodies, identical pinned boundary | `lbm/refinement_stepper.py:557`, `:567`-`:581` |
-
-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:474`-`:542` |
-| Level | not on L0; only on the finest level of a nested hierarchy | `lbm/refinement_stepper.py:80`, `:89` |
-| Jet condition | velocity only; `pressure` is reserved and raises NotImplementedError | `body/jets.py:148`, `:150` |
-| 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:155`, `:175`, `:177`, `:255`, `:261`, `:263` |
-| Jets with shifted storage | only in the all-blocked stepper with SRT and no LES | `body/jets.py:249` |
-| SurfaceActuation | unshifted storage only | `body/objects.py:155` |
-| Raw boundary exchange (force) | blocked body levels only; the window ledger requires `observe_raw_exchange` | `lbm/refinement_stepper.py:264`, `:269` |
-| Checkpoint with Jets | refused | `simulation.py:687` |
-
-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/evidence/physical-750e8a4-20261020/GROK-PHYSICAL-HANDOFF.md

@@ -0,0 +1,76 @@
+待 Grok 独立复核：750e8a4 本轮六例物理验证
+==========================================
+
+固定候选 `750e8a483ec47dcf472e528edc1cf2936f0e7e83`，工作树 `/home/frank14f/CelerisLab-release-prep-20261020` 干净。六例全部完成、退出码 0、无超时；调度器和六个子进程均已消失，最终 `nvidia-smi` 无计算进程。206 个源码/测试/文档哈希与启动 manifest 一致，各例运行前后 source map 相同，原生结果与外部摘要一致；V2 原生源码捕获及 runner 哈希也一致。V4 采用外部 manifest 绑定，没有冒称 runner 内部具备同样 provenance。没有新实验、阈值变更、源码变更、main 合并或公开发布。
+
+六例状态及进程墙钟秒数：V2 zero refined `periodic` 511.409，zero uniform `periodic` 325.254；fixed refined `driven_diagnostics_complete` 530.173，fixed uniform 同状态 333.016；V4 refined `NOT_STEADY` 315.714，uniform `NOT_STEADY` 114.183。墙钟包含初始化和观测开销，仅用于执行记录，不能作正式性能比较。
+
+**V2 zero：MRT、unshifted，均达到 t*=300。** 每例 3480 条 history；refined 1.2M steps，uniform 2.4M steps。各提取 14 个 t*>250 的周期，末五周期全部原生周期门槛通过，各自与冻结参考比较通过。
+
+| 周期指标 | refined | uniform | 原门槛 |
+|---|---:|---:|---:|
+| 末五周期范围 | 281.663875–298.232508 | 281.590556–298.153468 | 起点>=250，5 个连续周期 |
+| period min / max | 3.311073 / 3.316868 | 3.311314 / 3.314196 | 每周期严格在 (2,5) |
+| period ptp/mean | 0.00174873 | 0.000869972 | <=0.005 |
+| 末周期距最后力采样 | 1.757492 | 1.836532 | [0,5] |
+| 连续周期最大间隙 | 0 | 0 | <1e-8 |
+| 最小 Cl peak-to-peak | 2.003137 | 1.998667 | >=0.2 |
+| Cd_min / Cd_max 跨周期 range | 8.14323e-5 / 1.39714e-4 | 2.60406e-4 / 8.66420e-5 | 每项<=0.005 |
+| Cl_min / Cl_max 跨周期 range | 1.96279e-4 / 1.15290e-4 | 1.59490e-4 / 1.82477e-4 | 每项<=0.005 |
+| dp_min / dp_max 跨周期 range | 5.50774e-4 / 5.87194e-4 | 8.12699e-4 / 7.99311e-4 | 每项<=0.005 |
+| 周期质量均值 ptp/abs(mean) | 3.61945e-8 | 8.51505e-9 | <=1e-4 |
+| 最大周期均值 flux_relative | 6.35676e-6 | 5.22752e-5 | <=0.005 |
+
+末周期实际值：refined period=3.312817588、St=0.301857852、Cd=[3.160863689,3.234990541]、Cl=[-1.019161308,0.984224152]、dp=[2.404928344,2.499337133]；uniform period=3.311313790、St=0.301994937、Cd=[3.161468623,3.231458726]、Cl=[-1.016677997,0.982020469]、dp=[2.420541455,2.513318511]。
+
+| 末周期参考误差 | refined | uniform | 原门槛 |
+|---|---:|---:|---:|
+| Cd_min / Cd_max 相对误差 | 0.00107520 / 0.00235425 | 0.000884026 / 0.00125993 | 每项<=0.02 |
+| Cl_min / Cl_max 绝对误差 | 0.00212432 / 0.00234646 | 0.00460763 / 0.00455014 | 每项<=0.03 |
+| dp_min / dp_max 相对误差 | 0.00638037 / 0.00306172 | 0.0000703422 / 0.00251519 | 每项<=0.02 |
+| St 相对误差 | 0.0000558054 | 0.000509970 | <=0.02 |
+| Cd / Cl / dp waveform RMS | 0.00472589 / 0.00216188 / 0.0117979 | 0.00317231 / 0.00337428 / 0.00382206 | 描述指标，不是 gate |
+
+`matched_case_pass=false` 原样保留：四次 V2 调用均选择单 case，runner 只有 `--case both` 的两结果原生比较分支才赋值这一字段。它不等于 zero 的各自周期或参考比较失败，也不能由执行者自行改为 true；本轮没有运行该 paired-case 原生比较。fixed 没有官方 DFG 参考 verdict，完成执行不等于物理验收。
+
+**Fixed Jet 健康与 CV。** 两例达到 t*=300，五个 250–300 诊断窗均完整。原健康规则为 physical leaf 数据及 impulse 有限、rho_min>=0.5、rho_max<=1.5、Mach<=0.2。所有 3480 条 history 均有限，观测数值如下。
+
+| 指标 | refined | uniform |
+|---|---:|---:|
+| 全序列 rho min / max | 0.994218661 / 1.004163150 | 0.993391573 / 1.004678726 |
+| 全序列 max Mach | 0.0422925943 | 0.0421810197 |
+| mass t*=0.5 / t*=300 | 0.8931541481 / 0.8931214909 | 0.8936680602 / 0.8936524063 |
+| 上述端点质量相对变化 | -3.65640e-5 | -1.75164e-5 |
+| t*=290–300 mass ptp/mean | 8.40337e-5 | 6.31000e-5 |
+| 290–300 Cd / Cl 均值 | 3.184020005 / -0.580990812 | 3.182612828 / -0.582235721 |
+
+CV 各 501 个端点，间距 t*=0.02，积分区间 290–300（root time 40000）。排除首端点之前的 impulse，只累加后 500 个区间；root reaction impulse refined=(254.721600438,-46.479264945)，uniform=(254.609026243,-46.578857642)。两条矩形轮廓 refined 为 (48,44,160,120)/(44,36,180,128)，uniform 的 x 坐标均加 0.25。以原式 `J_CV=-ΔP-∫outward_momentum_flux dt`，残差为 `J_CV-J_raw`，系数除以 `duration*0.5*U_root²*D_root=80`。这里是 wall+aperture 原始流体对物体反作用，不是已修正的 nozzle thrust。
+
+| 动量残差系数 (x,y) | refined | uniform |
+|---|---|---|
+| 内轮廓、stride1 | (2.24958e-4,-1.00611e-4) | (1.78190e-4,-7.24811e-5) |
+| 外轮廓、stride1 | (7.88038e-5,-4.35879e-5) | (-8.30841e-5,-1.99246e-5) |
+| 外减内 | (-1.46155e-4,5.70231e-5) | (-2.61274e-4,5.25565e-5) |
+| stride2 相对 stride1 最大分量差 | 1.18312e-5 | 2.19402e-5 |
+| stride4 相对 stride1 最大分量差 | 4.44884e-5 | 6.26044e-5 |
+
+CV 未定义验收门槛，以上仅为离线积分诊断。质量记录缺少 raw `mass_into_fluid` Jet 源账本，无法扣除质量源判断闭合。`ΔM+∫outward_mass_flux dt` 仅是表观质量源：refined 内/外=-0.0172316912/-0.646617293 root mass；uniform=3.519017515/3.575029022。相对初始 CV 库存分别 (-2.37505e-6,-5.74501e-5)/(4.84891e-4,3.17573e-4)。轮廓间质量差异保留供审查，不能据此宣称质量闭合通过。
+
+**V4：SRT，无 Jet，t*=60。** uniform 480k/refined 240k steps，各 120 个样本。末窗 t*=50–60 共 21 点；drift 是原 runner 的 `abs(OLS slope per t*)*10`。
+
+| case / 指标 | 末窗 range | 原 range 上限 | 末窗 drift | 原 drift 上限 |
+|---|---:|---:|---:|---:|
+| refined Cd | 9.86614e-4 | 0.002 | 4.68834e-6 | 0.0004 |
+| refined Cl | 6.95374e-5 | 6e-5 | 3.28188e-6 | 1.2e-5 |
+| refined p_diff | 1.45066e-4 | 4e-5 | 9.28999e-6 | 8e-6 |
+| uniform Cd | 1.19061e-3 | 0.002 | 1.60374e-4 | 0.0004 |
+| uniform Cl | 7.93289e-5 | 6e-5 | 5.92289e-6 | 1.2e-5 |
+| uniform p_diff | 1.27157e-4 | 4e-5 | 2.80735e-5 | 8e-6 |
+
+两例均有 Cl range、p_diff range/drift 超限，保持 `steady=false,status=NOT_STEADY,verdict=null`。末窗均值 refined=(Cd 5.588876919,Cl 0.010732806,p_diff 0.117811546)，uniform=(5.586439289,0.010680915,0.117402607)。原区间 Cd=[5.57,5.59]、Cl=[0.0104,0.011]、p_diff=[0.1172,0.1176]；refined 压差均值在区间外，uniform 均值在区间内，但均不能越过稳态 gate。
+
+120 点全序列（t*=0.5–60）存在初始瞬态和反复起伏：refined Cd/Cl/p_diff 的增量符号变化次数为 82/80/82，uniform 为 79/78/81；末窗仍为 13/13/14 和 12/13/11。末窗有符号 OLS slope per t* refined=(-4.68834e-7,-3.28188e-7,-9.28999e-7)，uniform=(-1.60374e-5,+5.92289e-7,-2.80735e-6)。这只是当前采样分辨率下的非单调时序描述，不能确定振荡频率、物理/数值根因，也不是发散或瓶颈结论。完整 JSON 保留全序列及逐 10 t* 窗统计。边界偏差仍为 pressure-NEQ/neq_extrap outlet 与 local Zou-He inlet，不能称为原始 do-nothing DFG 边界。
+
+**复核边界与证据。** 本轮不是发布通过：短套件仍为 767 passed/24 failed/1 skipped；V1 历史 pin 阻塞，V3 四例原 TG dry-run 均被 frozen-source guard 拒绝、尚无 CUDA 计算或 TG 结果；未接纳待修订 `56004f44`，没有修改 dtype/NumPy/数值核心。下一步由 Claude 提交修订和解释、Grok 独立复核，再由集成任务接受明确通过审查的 SHA。
+
+原始证据根 `/home/frank14f/CelerisLab-release-evidence-20261020/physical-750e8a4/followup-queue`。`final-physical-digest.json` 约 99 KB，含全部六例原生小摘要、门槛、CV 抽稀结果、源/输出哈希、进程检查和限制；后处理脚本 `final_physical_digest.py` 不导入或运行模拟器。原始 history/log 只保留在 Horsea，没有上传巨量日志。

docs/validation/evidence/physical-750e8a4-20261020/final_physical_digest.py

@@ -0,0 +1,89 @@
+"""Read-only postprocessing of the six existing, authorized simulation outputs."""
+import csv
+import hashlib
+import json
+import math
+from pathlib import Path
+import subprocess
+import datetime
+import numpy as np
+
+Q = Path('/home/frank14f/CelerisLab-release-evidence-20261020/physical-750e8a4/followup-queue')
+R = Path('/home/frank14f/CelerisLab-release-prep-20261020')
+SHA = '750e8a483ec47dcf472e528edc1cf2936f0e7e83'
+def read(p):
+    return json.loads(p.read_text())
+def sha(p):
+    return hashlib.sha256(p.read_bytes()).hexdigest()
+def cmd(*a):
+    return subprocess.check_output(a, text=True).strip()
+def span(a):
+    return float(np.ptp(a))
+def finite(x):
+    if isinstance(x, dict): return all(finite(v) for v in x.values())
+    if isinstance(x, list): return all(finite(v) for v in x)
+    return not isinstance(x, (int,float)) or math.isfinite(x)
+def stats(rows, key):
+    t=np.array([v['tstar'] for v in rows]); v=np.array([r[key] for r in rows])
+    d=np.diff(v); nz=np.sign(d[d!=0])
+    return dict(samples=len(v),first=float(v[0]),last=float(v[-1]),min=float(min(v)),max=float(max(v)),mean=float(np.mean(v)),range=span(v),ols_slope_per_tstar=float(np.polyfit(t-t[0],v,1)[0]),increment_sign_changes=int(np.count_nonzero(nz[1:]!=nz[:-1])),total_variation=float(np.abs(d).sum()))
+
+state=read(Q/'state.json'); manifest=read(Q/'manifest.json'); execution=read(Q/'execution.json')
+out=dict(schema=1,utc=datetime.datetime.now(datetime.timezone.utc).isoformat(),sha=SHA,evidence_root=str(Q),postprocessor_sha256=sha(Path(__file__)),scope='Offline read-only digest; no new simulation, gate edit, acceptance promotion, or performance claim',verification={},cases={})
+v=out['verification']; v['all_six_completed']=len(state['jobs'])==6 and all(x['status']=='COMPLETED' for x in state['jobs'].values())
+v['all_exit_zero_no_timeout']=all(x['exit_code']==0 and not x['timeout'] for x in state['jobs'].values())
+v['head']=cmd('git','-C',str(R),'rev-parse','HEAD');v['git_status_porcelain']=cmd('git','-C',str(R),'status','--porcelain');v['clean_fixed_head']=v['head']==SHA and not v['git_status_porcelain']
+v['tree']=cmd('git','-C',str(R),'rev-parse','HEAD^{tree}');v['src_tree']=cmd('git','-C',str(R),'rev-parse','HEAD:src')
+v['tracked_source_map_matches_manifest']=all(sha(R/p)==h for p,h in manifest['files_sha256'].items());v['tracked_source_map_count']=len(manifest['files_sha256'])
+v['execution_final_source_matches_manifest']=all(execution['source_after'][k]==manifest[k] for k in ['sha','tree','src_tree','git_status','git_diff_sha256','files_sha256'])
+pids=[state['scheduler_pid']]+[x['pid'] for x in state['jobs'].values()]
+v['tracked_processes']={}
+for pid in pids:
+    p=subprocess.run(['ps','-p',str(pid),'-o','pid=,ppid=,stat=,comm='],capture_output=True,text=True)
+    v['tracked_processes'][str(pid)]=p.stdout.strip() or 'absent'
+v['no_tracked_process_running']=all(s=='absent' or ' Z' in s for s in v['tracked_processes'].values())
+v['gpu_compute_processes']=cmd('nvidia-smi','--query-compute-apps=gpu_uuid,pid,process_name,used_gpu_memory','--format=csv')
+out['hashes']={p.name:sha(p) for p in [Q/'manifest.json',Q/'execution.json',Q/'state.json',Q/'queue-plan.json',Q/'receipt-preflight.json',Q/'authorized_followup_queue.py',Q/'v4_external_timing.py']}
+for label,sj in state['jobs'].items():
+    p=Q/label; r=read(p/'result.json'); c=out['cases'][label]=dict(status=sj['native_status'],exit_code=sj['exit_code'],timeout=sj['timeout'],process_wall_seconds=sj['wall_seconds'],pid=sj['pid'],gpu_uuid=sj['gpu_uuid'],source_unchanged=r['source_unchanged'])
+    c['source_matches_manifest']=all(r['source'][k]==manifest[k] and r['source_after'][k]==manifest[k] for k in ['sha','tree','src_tree','git_status','git_diff_sha256','files_sha256'])
+    c['external_manifest_hash_matches']=r['external_manifest_sha256']==sha(Q/'manifest.json')
+    c['hashes']={str(z.relative_to(p)):sha(z) for z in [p/'command.json',p/'result.json',*sorted((p/'output').rglob('*.json')),*sorted((p/'output').rglob('*.csv')),*sorted((p/'output').rglob('*.jsonl'))]}
+    if label.startswith('v4'):
+        n=read(p/'output/result.json')[0]; rows=n['rows']
+        captured=r['native_results'][0]
+        c['native_receipt_matches_external_result']=all(n[k]==z for k,z in captured.items() if k not in ['last_row','rows_count']) and captured['last_row']==rows[-1] and captured['rows_count']==len(rows)
+        c['external_receipt_history_handling']='V4 external receipt stores row count and last row; full native rows are separately hashed'
+        c['native_all_finite']=finite(n)
+        c['native']={k:z for k,z in n.items() if k not in ['rows','lattice']}
+        c['full_series']={k:stats(rows,k) for k in ['cd','cl','p_diff']}
+        c['ten_tstar_windows']={str(a):{k:stats([z for z in rows if a<=z['tstar']<=a+10],k) for k in ['cd','cl','p_diff']} for a in range(0,60,10)}
+        c['last_window']=dict(tstar=[50,60],metrics={k:stats(rows[-21:],k) for k in ['cd','cl','p_diff']})
+        c['health']=dict(rho_min=min(z['rho_min'] for z in rows),rho_max=max(z['rho_max'] for z in rows),max_mach=max(z['max_mach'] for z in rows))
+        continue
+    case=label.rsplit('-',1)[1]; base=p/'output'/case;n=read(base/'summary.json');plan=read(p/'output/plan.json')
+    c['native_receipt_matches_external_result']={k:z for k,z in n.items() if k!='history'}==r['native_summary']
+    c['external_receipt_history_handling']='V2 external receipt omits the history path; complete history.csv is separately hashed'
+    c['native_all_finite']=finite(n)
+    c['native_runner_hash_matches']=plan['runner_sha256']==sha(R/'tests/validation/run_refined_dfg.py')
+    sm=plan['source_manifest'];c['native_source_capture_matches']=sm['checkout_revision']==SHA and not sm['checkout_dirty'] and all(sha(R/f)==h==manifest['files_sha256'][f] for f,h in sm['sha256'].items())
+    c['native_source_capture_file_count']=len(sm['sha256'])
+    c['native']={k:z for k,z in n.items() if k not in ['history','cycles','jet_diagnostic_windows']}; c['comparison']=read(p/'output/comparison.json');c['gates']=plan['gates'];c['reference_gates']=plan['reference_gates'];c['diagnostic_windows']=n['jet_diagnostic_windows']
+    rows=[{k:float(z) for k,z in row.items()} for row in csv.DictReader((base/'history.csv').open())];late=[x for x in rows if x['tstar']>=290]
+    c['health']=dict(rows=len(rows),tstar=[rows[0]['tstar'],rows[-1]['tstar']],all_finite=finite(rows),rho_min=min(z['rho_min'] for z in rows),rho_max=max(z['rho_max'] for z in rows),max_mach=max(z['max_mach'] for z in rows),mass_first=rows[0]['mass'],mass_last=rows[-1]['mass'],mass_change_fraction=(rows[-1]['mass']-rows[0]['mass'])/rows[0]['mass'],mass_late_range_fraction=span([z['mass'] for z in late])/np.mean([z['mass'] for z in late]))
+    if n['cycles']:
+        tail=n['cycles'][-5:]; periods=[z['period'] for z in tail]; masses=[z['mass_time_mean'] for z in tail]
+        c['periodic_actual']=dict(cycles_total=len(n['cycles']),tail_cycles=len(tail),tail_start=tail[0]['start'],tail_end=tail[-1]['end'],period_min=min(periods),period_max=max(periods),period_mean=float(np.mean(periods)),period_relative_range=span(periods)/np.mean(periods),last_cycle_age=rows[-1]['tstar_force']-tail[-1]['end'],consecutive_max_gap=max(abs(a['end']-b['start']) for a,b in zip(tail,tail[1:])),minimum_lift_amplitude=min(z['cl_amplitude'] for z in tail),extrema_ranges={k+'_'+e:span([z[k+'_'+e] for z in tail]) for k in ['cd','cl','dp'] for e in ['min','max']},mass_relative_range=span(masses)/abs(np.mean(masses)),max_cycle_flux_relative=max(z['flux_relative'] for z in tail),last_cycle={k:z for k,z in tail[-1].items() if not k.endswith('_waveform')})
+    cv=[json.loads(x) for x in (base/'cv_history.jsonl').read_text().splitlines()];t=np.array([x['time_root'] for x in cv]);inv=np.array([x['inventory'] for x in cv]);flux=np.array([x['outward_flux'] for x in cv]);j=np.array([x['interval_impulse_root'] for x in cv[1:]]).sum(axis=0);duration=t[-1]-t[0];norm=duration*.5*.01**2*40
+    cvout=c['cv']=dict(samples=len(cv),all_finite=finite(cv),tstar=[cv[0]['tstar'],cv[-1]['tstar']],duration_root=duration,coefficient_normalization=norm,raw_body_reaction_impulse_root=j.tolist(),first_preceding_impulse_excluded=True,contours_root=plan['cv_observation'][case+'_contours_root'],mass_qualification='Apparent mass source only: raw Jet mass_into_fluid ledger was not recorded, so mass-source-subtracted closure cannot be assessed',acceptance='offline diagnostics only; no CV pass threshold was defined',strides={})
+    for stride in [1,2,4]:
+        idx=np.arange(0,len(t),stride);assert idx[-1]==len(t)-1
+        integral=np.trapz(flux[idx],t[idx],axis=0); delta=inv[-1]-inv[0];jcv=-delta[:,1:]-integral[:,1:];res=jcv-j
+        cvout['strides'][str(stride)]=dict(samples=len(idx),momentum_delta_root=delta[:,1:].tolist(),outward_momentum_integral_root=integral[:,1:].tolist(),cv_reaction_impulse_root=jcv.tolist(),residual_impulse_root=res.tolist(),residual_coefficient=(res/norm).tolist(),apparent_mass_source_root=(delta[:,0]+integral[:,0]).tolist(),apparent_mass_source_fraction_initial_inventory=((delta[:,0]+integral[:,0])/inv[0,:,0]).tolist())
+    a=np.array(cvout['strides']['1']['residual_coefficient']);cvout['contour_residual_coefficient_difference']=(a[1]-a[0]).tolist();cvout['coarsening_coefficient_difference']={s:(np.array(cvout['strides'][s]['residual_coefficient'])-a).tolist() for s in ['2','4']}
+v['all_six_source_unchanged']=all(c['source_unchanged'] and c['source_matches_manifest'] and c['external_manifest_hash_matches'] for c in out['cases'].values())
+v['all_six_native_receipts_match']=all(c['native_receipt_matches_external_result'] and c['native_all_finite'] for c in out['cases'].values())
+v['all_v2_native_captures_match']=all(c['native_runner_hash_matches'] and c['native_source_capture_matches'] for k,c in out['cases'].items() if k.startswith('v2'))
+out['limitations']=['V1 historical source pin remains blocked','V3 original TG source guard blocked all four dry-runs before CUDA; no TG result or acceptance','Short suite remains 767 passed, 24 failed, 1 skipped; not a release pass','Matched_case_pass remains false because each invocation selected one case; native paired-case comparison branch did not execute','Fixed Jet CV is diagnostic; mass source ledger absent; no physical acceptance','V4 both NOT_STEADY with original ranges/drifts/bands; boundary deviations retained','No formal performance claim: timings include observation/profiling and startup is separately recorded']
+dest=Q/'final-physical-digest.json';dest.write_text(json.dumps(out,indent=2,allow_nan=False)+'\n')
+print(json.dumps(dict(output=str(dest),bytes=dest.stat().st_size,sha256=sha(dest),verification=v),indent=2))

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, context=context, coarse_boundary=args.mode, refinements=[BOUNDS],
-            level_blocks=[None, fine_block], level_objects=[(), (body,)]) as run:
+            cfg, BOUNDS, context=context, coarse_boundary=args.mode,
+            fine_objects=[body], fine_block=fine_block) 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, context=context,
-                coarse_boundary='slot_reservoir', refinements=[BOUNDS]) as run:
+        with CudaContext(0) as context, FixedPatchStepper(cfg, BOUNDS, context=context,
+                coarse_boundary='slot_reservoir') 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, context=context, refinements=[(4, 4, 12, 12)]) as run:
+    with CudaContext(0) as context, PeriodicPatchStepper(cfg, (4, 4, 12, 12), context=context) 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,30 +207,19 @@ 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 all None when neither an object carries surface
+    order. Returns (None, 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, None, None
+        return 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
@@ -263,4 +252,4 @@ def _surface_links(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, shape, jet_id
+    return surface_velocity, surface_aperture

src/CelerisLab/cuda/compiler_v2.py

@@ -54,18 +54,14 @@ def _write(path: str, content: str):
 
 
 def generate_config(cfg: LBMConfig, n_objects: int = 0,
-                    *, build_root: Optional[str] = None, block_root_bc: int = 0) -> str:
+                    *, build_root: Optional[str] = None) -> 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")
@@ -75,15 +71,14 @@ 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"),
-                       block_root_bc=block_root_bc)
+        _render_config(macros, n_objects, os.path.join(build_dir, "config"))
         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, "block_root_bc": block_root_bc}
+                    "n_objects": n_objects}
         _write(os.path.join(build_dir, ".celeris-build"), json.dumps(metadata, sort_keys=True))
     except BaseException:
         shutil.rmtree(build_dir)
@@ -106,7 +101,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, *, block_root_bc: int = 0):
+def _render_config(m: dict, n_objects: int, cdir: str):
     _write(os.path.join(cdir, "config_grid.h"), f"""{_HEADER}\
 // Layer: Global / Grid
 #ifndef CELERIS_CONFIG_GRID_H
@@ -170,9 +165,6 @@ def _render_config(m: dict, n_objects: int, cdir: str, *, block_root_bc: int = 0
 #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,21 +204,11 @@ class PeriodicBlockStepper:
                 raise
 
     def step(self, n=1):
-        """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.
-        """
+        """Advance n steps: halo refresh, then core pull/collide, per step."""
         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"
@@ -234,9 +224,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._current, self._spare = self._spare, self._current  # host handles only
-                self._stream.synchronize()
-                self._step_count += n
+                    self._stream.synchronize()
+                    self._current, self._spare = self._spare, self._current
+                    self._step_count += 1
                 self._status = "ready"
             except BaseException:
                 self._status = "failed"

src/CelerisLab/lbm/curved_links.py

@@ -94,10 +94,12 @@ class CurvedLinkSoA:
         """
         self.surface_velocity = self.aperture = None
         self.jets, self.jet_shape, self.jet_id = tuple(jets), None, None
-        from ..body.jets import _surface_links
-        self.surface_velocity, self.aperture, self.jet_shape, self.jet_id = _surface_links(
+        self.surface_velocity, self.aperture = surface_link_velocity(
             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,6 +1,5 @@
 // CelerisLab/lbm/kernels/blocks.cu
-// One fixed-size blocked level (lbm/blocks.py layout), specialized at compile
-// time for periodic/fine, slot/reservoir or channel stepping. Separate
+// One periodic level of fixed-size blocks (lbm/blocks.py layout). 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).
@@ -20,18 +19,7 @@
 #endif
 #include "streaming/pull_double_buffer.cuh"
 #include "boundary/curved_boundary.cuh"
-#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
-
-#if BLOCK_ROOT_BC == 1
-#include "boundary/block_wall.cuh"
-#endif
-#if BLOCK_ROOT_BC != 2
-#include "operators/block_collision.cuh"
-#endif
+#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"
@@ -40,6 +28,11 @@
 // 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)
@@ -154,7 +147,6 @@ extern "C" __global__ void BlockAccumulateExchange(
     integral[k + 3] += h * h * (double)raw[k + 3];
 }
 
-#if BLOCK_ROOT_BC != 2
 // 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
@@ -175,18 +167,40 @@ extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int hal
         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);
+            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);
+#endif
+            collide_srt_no_force(f, feq, d_params.omega);
 #endif
-            block_collide(f);
         }
     }
     stream_pull_store(k, f, out + base);
 }
 
-#endif
-
-// Root channel is a replacement full step, selected once by the host.
-#if BLOCK_ROOT_BC == 2
+// Root channel closure; fine and periodic levels keep BlockStep.
 #include "step/block_channel.cu"
-#endif

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

@@ -1,20 +0,0 @@
-// 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/config/config_method.h

@@ -8,9 +8,6 @@
 #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/layout.cuh

@@ -53,18 +53,6 @@ __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

@@ -1,24 +0,0 @@
-// Collision used by periodic/fine and slot/reservoir BlockStep.
-#ifndef CELERIS_OPERATORS_BLOCK_COLLISION_CUH
-#define CELERIS_OPERATORS_BLOCK_COLLISION_CUH
-
-__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/collision_mrt.cuh

@@ -114,16 +114,6 @@ __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/refinement.cu

@@ -37,7 +37,10 @@ __device__ unsigned int incoming_counts[NQ * TOTAL_CELLS];
 
 __device__ __forceinline__ void patch_collide(float* f, unsigned long k) {
 #if COLLISION_MODEL == 2
-    collide_mrt_unforced(f);
+    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);

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

@@ -1,7 +1,3 @@
-#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 "../boundary/inlet/common.cuh"
@@ -19,8 +15,10 @@ 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;
-    unsigned int gx, gy;
-    block_global_xy(x, y, halo, d_block_channel.nx, gx, gy);
+    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;
     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];
@@ -57,7 +55,11 @@ extern "C" __global__ void BlockChannelStep(
         }
         if (!wall && !outlet) {
 #if COLLISION_MODEL == 2
-            collide_mrt_unforced(f);
+            // 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)));
 #else
             // Match uniform SRT arithmetic, including its shifted equilibrium path.
             float rho, ux, uy, feq[NQ], forcing[NQ] = {};

src/CelerisLab/lbm/refinement_stepper.py

@@ -65,8 +65,52 @@ 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 of the level lists, checked before any device work."""
+    """Candidate limits shared by the list and the legacy entry points."""
     depth = len(refinements) + 1
     if depth < 2:
         raise ValueError("at least one refinement level is required")
@@ -75,13 +119,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_blocks[{level}] must be (by, bx)")
+            raise ValueError(f"{_level_name(level)}_block 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 L1 patches require level_blocks[0] and level_blocks[1]")
+        raise ValueError("several patches require fine_block and base_block")
     if depth > 2:
         if any(block is None for block in blocks):
             raise ValueError("three or more levels must all be blocked")
@@ -90,14 +134,7 @@ def _check_levels(refinements, blocks, objects, jets):
 
 
 class FixedPatchStepper:
-    """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).
+    """D2Q9/SRT/FP32, one unforced internal 2:1 patch in an x-periodic box.
 
     ``coarse_boundary`` is ``periodic`` (default) or ``slot_reservoir``:
     a bottom on-site slot and top fixed-density reservoir from ``cfg``.
@@ -108,55 +145,68 @@ 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.
 
-    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.
+    ``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.
 
+    ``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 levels are not supported by this diagnostic.
+    body layouts are not supported by this new diagnostic in this candidate.
     ``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.
 
-    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.
+    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)``.
 
     ``explosion="projected_linear"`` selects experimental transverse linear
     reconstruction on each adjacent pair. Default "constant" is unchanged.
@@ -164,21 +214,43 @@ class FixedPatchStepper:
     bodies/Jets or LES; this backend supplies no forcing. Geometry is generic,
     but union/multilevel GPU and matched-physics qualification remain pending.
 
-    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.
+    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.
     """
-    def __init__(self, cfg, *, context, refinements, coarse_boundary="periodic",
-                 level_blocks=None, level_objects=None, level_jets=None,
+    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,
                  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")
-        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)
+        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)
         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]
@@ -362,10 +434,7 @@ 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
-        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)
+        build = compiler.generate_config(config, n_objects=n_objects)
         try:
             ptx = compiler.compile_kernel(self.ctx.sm_arch, source=source, build_dir=build)
             module = compiler.load_module(ptx)
@@ -471,13 +540,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):
+    def initialize(self, *level_g, coarse_g=None, fine_g=None, child_g=None):
         """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 = tuple(level_g)
+        given = _initial_levels(level_g, coarse_g, fine_g, child_g)
         with self._lock, self.ctx.activate():
             self._require("new", "ready")
             self._check_storage()
@@ -601,9 +670,10 @@ class FixedPatchStepper:
                 old, new, source, base_count=count, level_steps=(count, 2*count, 4*count),
                 phase="postcollision_g", **options)
             candidate = type(self)(
-                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],
+                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,
                 observe_raw_exchange=self._observe_raw_exchange,
                 observe_window_exchange=self._observe_window, explosion=self._explosion)
             try:
@@ -752,41 +822,27 @@ class FixedPatchStepper:
                             grid=((pairs + self._block[0] - 1) // self._block[0], 1), **kw)
 
     def step(self, n=1):
-        """Advance G→G by n base intervals: level l takes 2**l hops per interval.
-
-        All n intervals are enqueued in order on the one stream, and the host waits
-        once, before returning; step(0) launches and waits for nothing. On return
-        the device work is complete, step_count and an open window cover all n
-        intervals, and the raw-exchange ledger holds the final interval only.
-
-        Asynchronous device errors surface at that single wait, so they are not
-        attributed to the interval that caused them; call step(1) repeatedly to
-        localize one (same launches, same bits). On any failure the instance is
-        failed, the exchange ledger and window are invalid, and step_count keeps
-        its value from entry to the failing call.
-        """
+        """Advance G→G by n base intervals: level l takes 2**l hops per interval."""
         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
             root = self._levels[0]
             kw = dict(block=self._block, stream=self._stream)
             try:
                 self._status = "advancing"
-                self._exchange_valid = False
                 for _ in range(n):
-                    if self._observe_raw_exchange:  # stream-ordered: holds the last interval only
+                    if self._observe_raw_exchange:
+                        self._exchange_valid = False
                         for b in self._level_boundaries:
                             if b is not None:
                                 cuda.memset_d8_async(b.interval_integral, 0,
                                                      len(b.tables["cell"]) * 32, self._stream)
                     self._hop(0, root.current, root.spare, **kw)
-                    root.current, root.spare = root.spare, root.current  # host handles only
-                self._stream.synchronize()
-                self._step_count += n
-                self._exchange_valid = self._observe_raw_exchange
+                    self._stream.synchronize()
+                    root.current, root.spare = root.spare, root.current
+                    self._step_count += 1
+                    self._exchange_valid = self._observe_raw_exchange
                 self._status = "ready"
             except BaseException:
                 self._window_valid = False

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 _surface_links
+from ..body.jets import jet_link_shapes, surface_link_velocity
 from .descriptors import D2Q9_EX, D2Q9_EY
 from .refinement import ACTIVE, OUTER_GHOST
 
@@ -142,13 +142,11 @@ 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"], shape, jet_id = _surface_links(
+        tables["surface_velocity"], tables["aperture"] = surface_link_velocity(
             boundary.objects, cfg, boundary.flags, fluid_idx, dirs, body_id, rx, ry,
             jets=jets, hits=boundary.curved_hits, _shifted_blocked=_shifted_blocked)
-        if jets:
-            jet_cache = (shape, jet_id)
-    return _BlockedBoundary(boundary, tables, jets, actuated, jet_cache)
+    return _BlockedBoundary(boundary, tables, jets, actuated,
+                            jet_link_shapes(jets, boundary.curved_hits) if jets else None)

tests/integration/test_checkpoint_bodies_gpu.py

@@ -1,66 +0,0 @@
-"""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,7 +16,6 @@
 #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

tests/unit/historical_tree.py

@@ -1,46 +0,0 @@
-"""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

@@ -1,104 +0,0 @@
-"""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,32 +83,6 @@ 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(int(layout.owner[5, -1]) - 1, 2),  # integer under NumPy 1 and 2
+    for target in (layout.slot(layout.owner[5, -1]-1, 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, context=context, coarse_boundary="channel",
-                                    refinements=[(6, 6, 18, 14)], level_blocks=[(4, 6), (4, 4)])
+    backend = host.FixedPatchStepper(cfg, (6, 6, 18, 14), context=context,
+                                    coarse_boundary="channel", base_block=(4, 6), fine_block=(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,6 +97,5 @@ 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_jet_circle_cv.py

@@ -8,7 +8,6 @@ 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,
@@ -20,7 +19,6 @@ from run_jet_circle_cv import (
     cv_from_macros,
     aperture_probe,
     measured_aperture_flux,
-    PIN,
     verify,
 )
 
@@ -168,13 +166,7 @@ class Host(unittest.TestCase):
         self.assertEqual(a["steps"] * 4, b["steps"])
 
     def test_source_pin(self):
-        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))
+        verify(REPO)
 
 
 if __name__ == "__main__":

tests/unit/test_les_flow_runner.py

@@ -15,10 +15,6 @@ 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):
@@ -46,23 +42,19 @@ class GuardTests(unittest.TestCase):
             decode_macros(invalid, w, cx, cy)
 
     def test_current_corrected_pins(self):
-        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")
+        self.assertEqual(
+            source_guard(ROOT, "corrected", "a135252")["base_commit"],
+            "a135252ae78b47b57d7e77d19001700b3bedac8b",
+        )
 
     def test_corrected_and_tamper_guard(self):
         pins = json.loads((HERE / "data/les_flow/source_pins_a135252.json").read_text())
         with TemporaryDirectory() as directory:
-            origin = historical_tree(A135, Path(directory) / "origin")
-            root = Path(directory) / "fixture"
+            root = Path(directory)
             for name in pins["sources"]:
                 dest = root / name
                 dest.parent.mkdir(parents=True, exist_ok=True)
-                shutil.copyfile(origin / name, dest)
+                shutil.copyfile(ROOT / name, dest)
             with patch(
                 "run_uniform_les_flow.subprocess.check_output",
                 return_value="isolated-fixture\n",
@@ -95,26 +87,23 @@ 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())
-        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")
+        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")
 
     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:
-            origin = historical_tree(A135, Path(directory) / "origin")
-            root = Path(directory) / "fixture"
+            root = Path(directory)
             for name in pins["sources"]:
                 dest = root / name
                 dest.parent.mkdir(parents=True, exist_ok=True)
-                shutil.copyfile(origin / name, dest)
+                shutil.copyfile(ROOT / 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_raw_exchange.py

@@ -157,9 +157,8 @@ class RawExchange(unittest.TestCase):
         host = host_module()
         ctx = Mock()
         with self.assertRaisesRegex(ValueError, 'blocked body'):
-            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)
+            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)
         ctx.activate.assert_not_called()
 
     def test_all_depth_weights(self):

tests/unit/test_re100.py

@@ -5,7 +5,6 @@ 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
 
@@ -63,16 +62,13 @@ def test_periodic_synthetic():
     assert not p.assess(rows()[:100])["periodic"]
 
 
-def test_production_unmodified(tmp_path):
-    from historical_tree import historical_tree
-
+def test_production_unmodified():
     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((origin / name).read_bytes()).hexdigest() == digest, name
+        assert hashlib.sha256((ROOT / name).read_bytes()).hexdigest() == digest, name
 
 
 def test_reference():

tests/unit/test_refined_dfg.py

@@ -7,10 +7,7 @@ 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)
 
@@ -20,24 +17,11 @@ 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")
-    with unittest.TestCase().assertRaisesRegex(RuntimeError, "Source mismatch"):
-        r.verify_source(r.DEFAULT_SOURCE, "MRT", "window")
+    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")
     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():
@@ -157,47 +141,16 @@ 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, frozen_manifest("window"))
+        r.verify_pilot(None, r.verify_source(r.DEFAULT_SOURCE, "MRT", "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."""
-    # 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))
+                if observation_mode == "window" else
+                json.loads((r.HERE / "data/refined_dfg/source_manifest_mrt.json").read_text()))
     selection = dict(case_selection=case_selection, uniform_diameter=uniform_diameter, placement=placement)
     (path / "plan.json").write_text(
         json.dumps(
@@ -395,7 +348,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 selected_manifest("MRT", "window")["production_origin_commit"].startswith("d930")
+    assert r.verify_source(r.DEFAULT_SOURCE, "MRT", "window")["production_origin_commit"].startswith("d930")
 
 
 def test_window_plan_preserves_physics():
@@ -405,7 +358,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 selected_manifest(collision, "window")["observation_mode"] == "window"
+        assert r.verify_source(r.DEFAULT_SOURCE, collision, "window")["observation_mode"] == "window"
     with unittest.TestCase().assertRaises(ValueError):
         r.verify_source(r.DEFAULT_SOURCE, "MRT", "unknown")
 

tests/unit/test_refinement_mrt.py

@@ -24,14 +24,11 @@ class RefinementMRTHost(unittest.TestCase):
         """MRT is an additive compile-time branch; SRT statements and other guards remain."""
         blocks = (KERNELS / "blocks.cu").read_text()
         channel = (KERNELS / "step/block_channel.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()
         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, block_collision)
+            self.assertIn(text, blocks)
         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):
@@ -40,14 +37,8 @@ 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
-        # 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")
+        for source in (blocks, channel, refinement):
+            self.assertEqual(source.count("collide_mrt_no_force("), 1)
 
 
 if __name__ == "__main__": unittest.main()

tests/unit/test_refinement_three_level_multipatch.py

@@ -173,10 +173,8 @@ def bodies_and_jets(speed=.002):
     return bodies,jets
 
 
-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))
+KW=dict(base_block=(4,4),fine_block=(4,4),child_block=(4,4),
+        child_rectangles_in_parent=CHILDREN)
 
 
 def mock_backend(*, children=CHILDREN, bodies=False, fail_compile=False, fail_upload=False):
@@ -203,8 +201,9 @@ 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,context=context,
-            **levels(children=children,objects=((),(),objects),jets=((),(),jets)))
+        run=module.FixedPatchStepper(CFG,PARENTS,context=context,
+            child_objects=objects,child_jets=jets,
+            **dict(KW,child_rectangles_in_parent=children))
     except RuntimeError:
         for allocation in allocations:
             allocation.free.assert_called_once_with()
@@ -215,29 +214,28 @@ 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()
-        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))]
+        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))]
         for changes in cases:
             with self.subTest(changes=changes),self.assertRaises((ValueError,TypeError)):
-                host.FixedPatchStepper(CFG,context=context,**levels(**changes))
+                host.FixedPatchStepper(CFG,PARENTS,context=context,**dict(KW,**changes))
         with self.assertRaisesRegex(ValueError,'margin'):
-            host.FixedPatchStepper(CFG,context=context,**levels(parents=[(1,4,20,28),(23,4,44,28)]))
+            host.FixedPatchStepper(CFG,[(1,4,20,28),(23,4,44,28)],context=context,**KW)
         objects,jets=bodies_and_jets()
-        with self.assertRaisesRegex(ValueError,'only on the finest level'):
-            host.FixedPatchStepper(CFG,context=context,**levels(objects=((),objects,())))
+        with self.assertRaisesRegex(ValueError,'only on L2'):
+            host.FixedPatchStepper(CFG,PARENTS,context=context,fine_objects=objects,**KW)
         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,context=context,
-                                  **levels(objects=((),(),objects),jets=((),(),[spanning])))
+            host.FixedPatchStepper(CFG,PARENTS,context=context,child_objects=objects,
+                                  child_jets=[spanning],**KW)
         context.activate.assert_not_called()
         host.cuda.mem_alloc.assert_not_called()
 
@@ -315,23 +313,21 @@ class ConstructorAndSchedule(unittest.TestCase):
 
 
 
-class RemovedLegacyAPI(unittest.TestCase):
-    """The pre-level-list constructor/initialize forms are gone and fail loudly."""
+class InitialLevels(unittest.TestCase):
+    """initialize() entry adapter: positional per-level arrays or the legacy names, never both."""
 
-    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
+    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)
 
 
 if __name__=='__main__':

tests/unit/test_shifted_jets.py

@@ -54,13 +54,14 @@ 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,context=None,refinements=[backend.PARENTS,A],
-                    level_blocks=[(4,4)]*3,level_objects=[(),(),objects],level_jets=[(),(),jets])
+                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)
         for blocks in ((None,None,None),((4,4),None,(4,4))):
             with self.assertRaises(ValueError):
-                module.FixedPatchStepper(CFG,context=None,refinements=[backend.PARENTS,A],
-                    level_blocks=list(blocks),level_objects=[(),(),(body(),)],
-                    level_jets=[(),(),(jet(),)])
+                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.cuda.mem_alloc.assert_not_called()
         module.compiler.compile_kernel.assert_not_called()
 

tests/unit/test_shifted_storage.py

@@ -78,15 +78,13 @@ 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.float32(np.nan)).repeat(9,axis=0)
-                  for m in old.leaf_masks)  # FP32 under NumPy 1 and 2 promotion rules
+        a = tuple(np.where(m[None], np.float32(-.001), np.nan).repeat(9,axis=0) for m in old.leaf_masks)
         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.float32(np.nan)).repeat(9,axis=0)
-                     for m in old.leaf_masks)
+        rest = tuple(np.where(m[None],np.float32(0),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')
@@ -131,10 +129,9 @@ 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 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)
+        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)
         module.compiler.compile_kernel.assert_not_called()
 
     def test_uniform_block_codec_and_raw_snapshot(self):

tests/unit/test_step_sync.py

@@ -1,170 +0,0 @@
-"""Host contract of FixedPatchStepper.step(n): one stream wait per call (mocked driver).
-
-All n base intervals are enqueued in order on the backend stream and the host
-waits once before step returns. Launch order and buffer arguments must equal n
-step(1) calls; completion metadata, raw-exchange and window ledgers, and failure
-invalidation are published only after that wait.
-"""
-from contextlib import nullcontext
-from unittest.mock import Mock
-import unittest
-
-import numpy as np
-
-from CelerisLab.config import LBMConfig
-from tests.unit.test_raw_exchange import body
-from tests.unit.test_refinement_three_level_multipatch import host_module
-
-BODY_LEVEL, HOPS = 3, 8  # body on L3: 2**3 hops per base interval
-
-
-def backend(window=False):
-    """Four-level host backend; every launch, memset and stream wait goes to one log."""
-    host = host_module()
-    log, allocations = [], []
-
-    def allocate(size):
-        allocations.append(Mock(size=size))
-        return allocations[-1]
-
-    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, k)))
-    host.compiler.load_module.return_value = module
-    host.cuda.mem_alloc.side_effect = allocate
-    host.cuda.memcpy_dtoh.side_effect = lambda array, pointer: array.fill(0)
-    host.cuda.memset_d8_async.side_effect = lambda *a: log.append(("memset", a, {}))
-    run = host.FixedPatchStepper(
-        LBMConfig(nx=48, ny=48, viscosity=.02), context=Mock(activate=nullcontext, sm_arch="sm_70"),
-        refinements=[[(6, 6, 26, 26)], [(20, 20, 44, 44)], [(56, 56, 72, 72)]],
-        level_blocks=[(4, 4)] * 4, level_objects=[(), (), (), (body(),)],
-        observe_raw_exchange=True, observe_window_exchange=window)
-    run._stream.synchronize.side_effect = lambda: log.append(("sync", (), {}))
-    run.initialize(*[np.full(level.shape, 1 / 9, np.float32) for level in run._levels])
-    log.clear()
-    return run, log, allocations
-
-
-def signature(log, allocations):
-    """Event names with device-pointer arguments replaced by allocation index."""
-    index = {id(a): i for i, a in enumerate(allocations)}
-    return [(name, tuple(index.get(id(a), "-") for a in args)) for name, args, _ in log]
-
-
-def names(log):
-    return [name for name, _, _ in log]
-
-
-class StepSync(unittest.TestCase):
-    def test_step1_waits_once_after_all_launches(self):
-        run, log, _ = backend()
-        run.step(1)
-        self.assertEqual(names(log).count("sync"), 1)
-        self.assertEqual(names(log)[-1], "sync")
-        self.assertEqual(names(log)[0], "memset")
-        self.assertEqual(names(log).count("BlockStep"), 1 + 2 + 4 + 8)
-        self.assertEqual(run.step_count, 1)
-        run.close()
-
-    def test_step_n_equals_n_step1_launches_with_one_wait(self):
-        single, single_log, single_alloc = backend()
-        for _ in range(3):
-            single.step(1)
-        batched, batched_log, batched_alloc = backend()
-        batched.step(3)
-        expected = [e for e in signature(single_log, single_alloc) if e[0] != "sync"]
-        got = signature(batched_log, batched_alloc)
-        self.assertEqual(got[-1][0], "sync")
-        self.assertEqual(names(batched_log).count("sync"), 1)
-        self.assertEqual(got[:-1], expected)  # same order, same ping-pong buffers
-        self.assertEqual(names(batched_log).count("memset"), 3)
-        self.assertEqual(batched.step_count, single.step_count)
-        index = [{id(a): i for i, a in enumerate(alloc)} for alloc in (single_alloc, batched_alloc)]
-        self.assertEqual(index[1][id(batched._levels[0].current)],
-                         index[0][id(single._levels[0].current)])
-        single.close()
-        batched.close()
-
-    def test_step0_launches_and_waits_for_nothing(self):
-        run, log, _ = backend()
-        run.step(2)
-        log.clear()
-        run.step(0)
-        self.assertEqual(log, [])
-        self.assertEqual(run.step_count, 2)
-        value = run.read_raw_exchange(0, level=BODY_LEVEL)  # step(0) keeps the ledger valid
-        self.assertEqual((value["time_start"], value["time_end"]), (1, 2))
-        run.close()
-
-    def test_completion_metadata_covers_all_intervals(self):
-        run, _, _ = backend()
-        with self.assertRaisesRegex(RuntimeError, "completed"):
-            run.read_raw_exchange(0, level=BODY_LEVEL)
-        run.step(3)
-        value = run.read_raw_exchange(0, level=BODY_LEVEL)
-        self.assertEqual((value["time_start"], value["time_end"], value["sample_count"]),
-                         (2, 3, HOPS))  # step(n) retains only the final base interval
-        run.step(2)
-        self.assertEqual(run.read_raw_exchange(0, level=BODY_LEVEL)["time_end"], 5)
-        run.close()
-
-    def test_window_accounting_across_batched_calls(self):
-        windows = []
-        for calls in ((1, 3), (4,), (1, 1, 1, 1)):
-            run, log, _ = backend(window=True)
-            run.step(1)
-            run.start_exchange_window()
-            log.clear()
-            for n in calls:
-                run.step(n)
-            self.assertEqual(names(log).count("BlockAccumulateExchange"), 2 * 4 * HOPS)
-            self.assertEqual(names(log).count("sync"), len(calls))
-            out = run.read_exchange_window(0, level=BODY_LEVEL)
-            windows.append((out["time_start"], out["time_end"], out["sample_count"]))
-            run.close()
-        self.assertEqual(windows, [(1., 5, 4 * HOPS)] * 3)
-
-    def test_failed_wait_publishes_nothing(self):
-        run, _, _ = backend(window=True)
-        run.step(2)
-        run.start_exchange_window()
-        run._stream.synchronize.side_effect = RuntimeError("device failure")
-        with self.assertRaisesRegex(RuntimeError, "device failure"):
-            run.step(3)
-        self.assertEqual(run.step_count, 2)  # value at entry to the failing call
-        self.assertFalse(run._exchange_valid)
-        self.assertFalse(run._window_valid)
-        for read in (lambda: run.read_raw_exchange(0, level=BODY_LEVEL),
-                     lambda: run.read_exchange_window(0, level=BODY_LEVEL), run.download,
-                     lambda: run.step(1)):
-            with self.assertRaisesRegex(RuntimeError, "failed"):
-                read()
-        run._stream.synchronize.side_effect = None
-        run.close()
-
-    def test_launch_error_mid_batch_skips_wait_and_fails(self):
-        run, log, _ = backend()
-        run.step(1)
-        log.clear()
-        hop, calls = run._hop, []
-
-        def failing(index, *args, **kw):
-            if index == 0:
-                calls.append(1)
-                if len(calls) == 2:
-                    raise RuntimeError("launch failure")
-            return hop(index, *args, **kw)
-        run._hop = failing
-        with self.assertRaisesRegex(RuntimeError, "launch failure"):
-            run.step(3)
-        self.assertNotIn("sync", names(log))
-        self.assertEqual(run.step_count, 1)
-        self.assertFalse(run._exchange_valid)
-        with self.assertRaisesRegex(RuntimeError, "failed"):
-            run.read_raw_exchange(0, level=BODY_LEVEL)
-        run.close()
-
-
-if __name__ == "__main__":
-    unittest.main()

tests/unit/test_validation_source_provenance.py

@@ -12,15 +12,13 @@ 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(tmp_path):
-    assert re100_helpers.PIN == run_jet_circle_cv.PIN == ORIGIN
-    manifests = {}
+def test_manifests_cover_current_source_and_preserve_origin():
+    names = subprocess.check_output(
+        ["git", "-C", str(ROOT), "ls-files", "src"], text=True
+    ).splitlines()
     for group, hashes_key, revision_key in (
         ("dfg_re100", "production_sha256", "commit"),
         ("jet_circle_cv", "files", "revision"),
@@ -31,22 +29,18 @@ def test_manifests_cover_current_source_and_preserve_origin(tmp_path):
             ).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"]
-        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():
+        for name, expected in manifest[hashes_key].items():
             assert (
-                hashlib.sha256((origin / name).read_bytes()).hexdigest() == expected
+                hashlib.sha256((ROOT / 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,7 +192,6 @@ 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, block_root_bc=2 if source == "blocks.cu" else 0)
+    build = compiler.generate_config(cfg)
     try:
         module = compiler.load_module(compiler.compile_kernel(context.sm_arch, source=source, build_dir=build))
     finally:
@@ -106,10 +106,9 @@ 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, 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:
+                        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:
                             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"), context=m["context"], refinements=[BOUNDS],
-                             level_blocks=[(3, 4), (16, 20)], level_objects=[(), (circle(),)],
+    return FixedPatchStepper(config("periodic"), BOUNDS, context=m["context"], base_block=(3, 4),
+                             fine_block=(16, 20), fine_objects=[circle()],
                              observe_raw_exchange=raw, observe_window_exchange=window)
 
 
@@ -109,11 +109,10 @@ 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, refinements=[RB], level_blocks=[None, (16, 20)],
-                      level_objects=[(), (rc(),)], level_jets=[(), (jet(),)],
+            kw = dict(context=ctx, fine_block=(16, 20), fine_objects=[rc()], jets=[jet()],
                       observe_raw_exchange=True)
-            with FixedPatchStepper(CFG, observe_window_exchange=True, **kw) as windowed, \
-                 FixedPatchStepper(CFG, **kw) as plain:
+            with FixedPatchStepper(CFG, RB, observe_window_exchange=True, **kw) as windowed, \
+                 FixedPatchStepper(CFG, RB, **kw) as plain:
                 arrays = [state(level.shape) for level in windowed._levels]
                 windowed.initialize(*arrays)
                 plain.initialize(*arrays)
@@ -201,11 +200,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, 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:
+            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:
                 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,8 +140,7 @@ 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, context=context, refinements=[BOUNDS],
-                                             level_objects=[(), (source,)])
+            backend = implementation.FixedPatchStepper(cfg, BOUNDS, context=context, fine_objects=[source])
         with backend, context.activate():
             c, f = backend._levels
             cl = backend._fine_curved
@@ -282,8 +281,7 @@ 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._boundary_hosts = [None, SimpleNamespace(objects=[object_spec(.002)])]
-    backend._level_jets = [(), ()]
+    backend.fine_boundary = SimpleNamespace(objects=[object_spec(.002)])
     backend._fine_curved = SimpleNamespace(fluid_idx_gpu=object(),
                                           surface_velocity_gpu=object(), link_exchange_gpu=object())
     setattr(backend._fine_curved, missing, None)
@@ -327,7 +325,6 @@ 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, context=context, refinements=[BOUNDS],
-                                             level_objects=[(), (source,)])
+            implementation.FixedPatchStepper(cfg, BOUNDS, context=context, fine_objects=[source])
         context.activate.assert_not_called()
         allocate.assert_not_called()

tests/validation/multigrid/test_multiple_pinned_remap_gpu.py

@@ -48,8 +48,9 @@ def make(context):
         for (x, y), amplitude in zip(centers, (.003, -.002))
     )
     run = FixedPatchStepper(
-        CFG, context=context, refinements=[PARENTS, START], level_blocks=[(4, 4)] * 3,
-        level_objects=[(), (), objects], level_jets=[(), (), jets],
+        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,
     )
     return run, objects
 

tests/validation/multigrid/test_pinned_body_remap_gpu.py

@@ -25,8 +25,9 @@ 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,context=ctx,refinements=[PARENTS,A],level_blocks=[(4,4)]*3,
-                           level_objects=[(),(),(body,)],level_jets=[(),(),(jet,)])
+    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,))
     return run,(body,)
 
 

tests/validation/multigrid/test_raw_exchange_gpu.py

@@ -90,18 +90,14 @@ 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, context=ctx, **fine(block), observe_raw_exchange=True) as measured, \
-                     FixedPatchStepper(CFG, context=ctx, **fine(block)) as plain:
+                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:
                     arrays = [state(l.shape) for l in measured._levels]
                     measured.initialize(*arrays)
                     plain.initialize(*arrays)
@@ -150,8 +146,9 @@ 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, context=ctx, **fine(None)) as dense, \
-                 FixedPatchStepper(CFG, context=ctx, **fine((16, 20)), observe_raw_exchange=True) as blocked:
+            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:
                 arrays = [state(l.shape) for l in blocked._levels]
                 dense.initialize(*arrays)
                 blocked.initialize(*arrays)
@@ -232,7 +229,8 @@ 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, context=ctx, **fine((16, 20)), observe_raw_exchange=True) as amr:
+                 FixedPatchStepper(CFG, BOUNDS, context=ctx, fine_block=(16, 20),
+                    fine_objects=[circle()], jets=[jet()], 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,10 +83,11 @@ 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 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)
+        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)
         context.activate.assert_not_called()
 
     def test_all_arrival_rows_donors_and_kind_masks(self):
@@ -192,10 +193,9 @@ 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"), context=m["context"],
-                                        coarse_boundary="slot_reservoir", refinements=[BOUNDS],
-                                        level_blocks=[(3, 4), FINE_BLOCK],
-                                        level_objects=[(), (circle(actuation(.002)),)])
+            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))])
             try:
                 backend.initialize(state(NY, NX, 1), state(*backend.patch.fine_shape, 2))
                 coarse, fine = backend._levels
@@ -223,8 +223,6 @@ 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)
@@ -238,16 +236,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"), context=m["context"], refinements=[BOUNDS],
-                                      level_blocks=[(3, 4), FINE_BLOCK])
+                    FixedPatchStepper(config("periodic"), BOUNDS, context=m["context"],
+                                      fine_block=FINE_BLOCK, base_block=(3, 4))
                 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"), context=m["context"], refinements=[BOUNDS],
-                                        level_blocks=[(3, 4), FINE_BLOCK])
+            backend = FixedPatchStepper(config("periodic"), BOUNDS, context=m["context"],
+                                        fine_block=FINE_BLOCK, base_block=(3, 4))
             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
@@ -291,10 +289,9 @@ class BlockedBaseGPU(unittest.TestCase):
                     for base_block in (None, *layouts):
                         label = f"{mode} width={width} speed={speed} base={base_block}"
                         snapshots = []
-                        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:
+                        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:
                             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,10 +105,9 @@ 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 ({"level_blocks": [None, (7, 7)]}, {"level_blocks": [None, (16,)]}):
+        for kw in ({"fine_block": (7, 7)}, {"fine_block": (16,)}):
             with self.subTest(kw=kw), self.assertRaises(ValueError):
-                FixedPatchStepper(config("periodic"), context=self.forbidden_context(),
-                                  refinements=[BOUNDS], **kw)
+                FixedPatchStepper(config("periodic"), BOUNDS, context=self.forbidden_context(), **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).
@@ -126,8 +125,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"), context=context, refinements=[BOUNDS],
-                                        level_blocks=[None, FINE_BLOCKS[0]])
+            backend = FixedPatchStepper(config("periodic"), BOUNDS, context=context,
+                                        fine_block=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
@@ -178,9 +177,8 @@ 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"), context=context, refinements=[BOUNDS],
-                                        level_blocks=[None, FINE_BLOCKS[1]],
-                                        level_objects=[(), tuple(objects)])
+            backend = FixedPatchStepper(config("periodic"), BOUNDS, context=context,
+                                        fine_objects=objects, fine_block=FINE_BLOCKS[1])
             backend.initialize(state(NY, NX, 1), state(*backend.patch.fine_shape, 2))
         return backend, calls
 
@@ -291,9 +289,8 @@ class BlockedFineGPU(unittest.TestCase):
             results = []
             with CudaContext(0) as context:
                 for fine_block in (None, *FINE_BLOCKS):
-                    with FixedPatchStepper(cfg, context=context, coarse_boundary=mode,
-                                           refinements=[BOUNDS],
-                                           level_blocks=[None, fine_block]) as backend:
+                    with FixedPatchStepper(cfg, BOUNDS, context=context, coarse_boundary=mode,
+                                           fine_block=fine_block) as backend:
                         backend.initialize(g_coarse, g_fine)
                         backend.step(STEPS)
                         self.assertEqual(backend.step_count, STEPS)
@@ -321,9 +318,8 @@ class BlockedFineGPU(unittest.TestCase):
             results = []
             with CudaContext(0) as context:
                 for fine_block in (None, *FINE_BLOCKS):
-                    with FixedPatchStepper(cfg, context=context, coarse_boundary=mode,
-                                           refinements=[BOUNDS], level_blocks=[None, fine_block],
-                                           level_objects=[(), (circle(),)]) as backend:
+                    with FixedPatchStepper(cfg, BOUNDS, context=context, coarse_boundary=mode,
+                                           fine_objects=[circle()], fine_block=fine_block) as backend:
                         backend.initialize(g_coarse, g_fine)
                         backend.step(STEPS)
                         results.append((backend.download(), raw(backend)))
@@ -350,9 +346,9 @@ class BlockedFineGPU(unittest.TestCase):
                 results = []
                 with CudaContext(0) as context:
                     for fine_block in (None, *FINE_BLOCKS):
-                        with FixedPatchStepper(cfg, context=context, coarse_boundary=mode,
-                                               refinements=[BOUNDS], level_blocks=[None, fine_block],
-                                               level_objects=[(), (circle(actuation(speed)),)]) as backend:
+                        with FixedPatchStepper(cfg, BOUNDS, context=context, coarse_boundary=mode,
+                                               fine_objects=[circle(actuation(speed))],
+                                               fine_block=fine_block) 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,8 +94,7 @@ class FaultGuardsCPU(unittest.TestCase):
                              ("slot_reservoir", (12, 4, 20, 21))):
             with self.subTest(mode=mode, bounds=bounds):
                 with self.assertRaises(ValueError):
-                    implementation.FixedPatchStepper(cfg, context=context, coarse_boundary=mode,
-                                                     refinements=[bounds])
+                    implementation.FixedPatchStepper(cfg, bounds, context=context, coarse_boundary=mode)
                 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.
@@ -103,14 +102,11 @@ 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, context=context, coarse_boundary=mode,
-                                                     refinements=[bounds])
+                    implementation.FixedPatchStepper(cfg, bounds, context=context, coarse_boundary=mode)
 
     def fake_backend(self):
         backend = PeriodicPatchStepper.__new__(PeriodicPatchStepper)
-        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.fine_layout = None  # dense fine storage
         backend.ctx = Mock(activate=nullcontext)
         backend._lock = RLock()
         backend._status = "ready"
@@ -158,23 +154,12 @@ 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, context=context, refinements=[(4, 4, 12, 12)])
+                PeriodicPatchStepper(cfg, (4, 4, 12, 12), context=context)
         p = implementation.FixedPatch2D.from_config(cfg, (4, 4, 12, 12))
-        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()
+        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)
         module = Mock()
         module.get_global.return_value = (Mock(), 16)
         with patch.object(cuda, "Stream"), patch.object(compiler, "generate_config"), \
@@ -182,7 +167,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, context=context, refinements=[p.bounds])
+                PeriodicPatchStepper(cfg, p.bounds, context=context)
             upload.assert_not_called()
 
 
@@ -256,7 +241,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, context=context, refinements=[bounds],
+                backend = implementation.FixedPatchStepper(cfg, bounds, context=context,
                                                             coarse_boundary="slot_reservoir")
             with backend:
                 self.assertEqual(len(generated), 2)
@@ -326,8 +311,7 @@ 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), context=context,
-                refinements=[(4, 4, 12, 12)]) as backend:
+                LBMConfig(nx=16, ny=16, viscosity=.1), (4, 4, 12, 12), context=context) as backend:
             backend.initialize(*postcollision(state, spec, maps))
             before = totals(backend.download(), maps)
             for step in range(1, 9):
@@ -358,7 +342,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, context=context, refinements=[bounds]) as backend:
+        with CudaContext(0) as context, PeriodicPatchStepper(cfg, bounds, context=context) 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(level_blocks=[(4, 4), (6, 6)])))
+                  dict(base_block=(4, 4), fine_block=(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, context=context, refinements=[(4, 4, 12, 12)],
+                    with FixedPatchStepper(cfg, (4, 4, 12, 12), context=context,
                                            **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),
-                                      context=context, refinements=[(4, 4, 12, 12)]) as backend:
+                                      (4, 4, 12, 12), context=context) 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, context=context, refinements=[(4, 4, 12, 12)])
+                backend = PeriodicPatchStepper(cfg, (4, 4, 12, 12), context=context)
             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, context=context, refinements=[(4, 4, 12, 12)]) as plain:
+            with PeriodicPatchStepper(cfg, (4, 4, 12, 12), context=context) as plain:
                 rest = postcollision(ref.make_state(spec), spec, maps)
                 plain.initialize(*rest)
                 plain.step()

tests/validation/multigrid/test_refinement_jets.py

@@ -41,9 +41,8 @@ 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"), context=context, refinements=[BOUNDS],
-                                    level_blocks=[None, fine_block], level_objects=[(), tuple(objects)],
-                                    level_jets=[(), tuple(jets)])
+        backend = FixedPatchStepper(config("periodic"), BOUNDS, context=context,
+                                    fine_objects=objects, fine_block=fine_block, jets=jets)
     return backend, htod
 
 
@@ -150,9 +149,8 @@ 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"), context=Mock(), refinements=[BOUNDS],
-                                  level_objects=[(), (circle(actuation(SPEED)),)],
-                                  level_jets=[(), (jet(),)])
+                FixedPatchStepper(config("periodic"), BOUNDS, context=Mock(),
+                                  fine_objects=[circle(actuation(SPEED))], jets=[jet()])
 
 
 def run(mode, fine_block, *, objects, jets, switch=None):
@@ -160,9 +158,8 @@ 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, context=context, coarse_boundary=mode, refinements=[BOUNDS],
-            level_blocks=[None, fine_block], level_objects=[(), tuple(objects)],
-            level_jets=[(), tuple(jets)]) as backend:
+            cfg, BOUNDS, context=context, coarse_boundary=mode, fine_objects=objects,
+            fine_block=fine_block, jets=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,9 +42,8 @@ class RefinementMRTGPU(unittest.TestCase):
                                     collision="MRT", ddf_shifting=shifted)
                     reference = None
                     for block in ((32, 48), (8, 8), (4, 1)):
-                        with FixedPatchStepper(cfg, context=context, coarse_boundary=boundary,
-                                               refinements=[(8, 6, 40, 26)], level_blocks=[block, (8, 8)],
-                                               level_objects=[(), (body,)]) as backend:
+                        with FixedPatchStepper(cfg, (8, 6, 40, 26), context=context, coarse_boundary=boundary,
+                                               base_block=block, fine_block=(8, 8), fine_objects=(body,)) as backend:
                             backend.initialize(*[state(*level.shape[1:]) for level in backend._levels])
                             backend.step(8)
                             raw = backend._download_raw()
@@ -64,8 +63,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, context=context, refinements=[BOUNDS],
-                                       level_blocks=[base_block, fine_block]) as backend:
+                with FixedPatchStepper(cfg, BOUNDS, context=context, fine_block=fine_block,
+                                       base_block=base_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)
@@ -76,9 +75,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(refinements=[(4, 4, 12, 12), (8, 8, 16, 16)], level_blocks=[(4, 4)] * 3)
+        kw = dict(fine_block=(4, 4), base_block=(4, 4), child=(8, 8, 16, 16), child_block=(4, 4))
         with CudaContext(0) as context:
-            with FixedPatchStepper(cfg, context=context, **kw) as backend:
+            with FixedPatchStepper(cfg, (4, 4, 12, 12), 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,12 +26,6 @@ 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)
@@ -51,21 +45,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(blocks=(None, (8, 4))), dict(blocks=((4, 4), None)),
-                   dict(blocks=((4, 4), (8, 4)), objects=[circle()]),
-                   dict(blocks=((4, 4), (5, 4)))):
+        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))):
             with self.subTest(kw=kw), self.assertRaises(ValueError):
-                FixedPatchStepper(cfg, context=context, **levels(((4, 4, 8, 12), (14, 4, 18, 12)), **kw))
+                FixedPatchStepper(cfg, ((4, 4, 8, 12), (14, 4, 18, 12)), context=context, **kw)
         with self.assertRaisesRegex(ValueError, "margin"):
-            FixedPatchStepper(cfg, context=context,
-                              **levels(((1, 4, 8, 12), (10, 4, 14, 12)), blocks=((4, 4), (8, 4))))
+            FixedPatchStepper(cfg, ((1, 4, 8, 12), (10, 4, 14, 12)), context=context,
+                              fine_block=(8, 4), base_block=(4, 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, 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]))
+            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])
 
 
 class MultipatchGPU(unittest.TestCase):
@@ -73,7 +67,7 @@ class MultipatchGPU(unittest.TestCase):
 
     def run_steps(self, context, cfg, bounds, coarse, fine, **kw):
         out = []
-        with FixedPatchStepper(cfg, context=context, **levels(bounds, **kw)) as backend:
+        with FixedPatchStepper(cfg, bounds, context=context, **kw) as backend:
             backend.initialize(coarse, fine)
             for step in self.STEPS:
                 backend.step(step - backend.step_count)
@@ -88,7 +82,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, blocks=((3, 4), (8, 4)))
+                kw = dict(coarse_boundary=mode, fine_block=(8, 4), base_block=(3, 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):
@@ -105,7 +99,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(blocks=((4, 4), (4, 4)))
+        kw = dict(fine_block=(4, 4), base_block=(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)
@@ -124,7 +118,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(blocks=((4, 4), (4, 4)))
+        kw = dict(fine_block=(4, 4), base_block=(4, 4))
         leaf = s.fine_kind == ACTIVE
         support = s.fine_layout((4, 4)).owned
         outer = support & (s.fine_kind == OUTER_GHOST)
@@ -170,7 +164,7 @@ class MultipatchBodiesGPU(unittest.TestCase):
 
     def run_steps(self, context, cfg, bounds, coarse, fine, switch=None, **kw):
         out = []
-        with FixedPatchStepper(cfg, context=context, **levels(bounds, **kw)) as backend:
+        with FixedPatchStepper(cfg, bounds, context=context, **kw) as backend:
             backend.initialize(coarse, fine)
             for step in self.STEPS:
                 if switch is not None and step == 8:
@@ -188,7 +182,8 @@ 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, blocks=((3, 4), (8, 4)), objects=[circle()])
+                kw = dict(coarse_boundary=mode, fine_block=(8, 4), base_block=(3, 4),
+                          fine_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):
@@ -207,17 +202,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(blocks=((4, 4), (4, 4)))
+        kw = dict(fine_block=(4, 4), base_block=(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,
-                                     objects=[body(8., 8.)], jets=jets, **kw)
+                                     fine_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,
-                                     objects=[body(8., 8., 5), body(24., 8., 5)],
+                                     fine_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,16 +27,7 @@ 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)
-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()
+KW = dict(fine_block=(4, 4), base_block=(4, 4), child=L2, child_block=(4, 4))
 
 
 def child_body(spec=None):
@@ -61,29 +52,26 @@ 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, context=context,
-                                 **levels(objects=((), (), (child_body(),)), jets=((), (), (child_jet(),))))
+        return FixedPatchStepper(CFG, L1, context=context, child_objects=[child_body()],
+                                 child_jets=[child_jet()], **KW)
 
 
 class ThreeLevelCPU(unittest.TestCase):
     def test_rejections_before_cuda(self):
         context = object()  # never activated: every case must fail first
-        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))))
+        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)))
         for kw in bad:
             with self.subTest(kw=kw), self.assertRaises(ValueError):
-                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(),)))):
+                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()])):
             with self.assertRaises(ValueError):
-                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))))
+                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)
 
 
     def test_four_halo_bfl_steps_and_fresh_middle_feeds(self):
@@ -168,7 +156,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, context=ctx, **KW) as run:
+        with CudaContext(0) as ctx, FixedPatchStepper(CFG, L1, 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)
@@ -181,7 +169,7 @@ 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, context=ctx, **KW) as run:
+        with CudaContext(0) as ctx, FixedPatchStepper(CFG, L1, 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]
@@ -217,9 +205,8 @@ 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, context=ctx, coarse_boundary=mode,
-                **levels(blocks=(B, B, block), objects=((), (), (child_body(spec),)),
-                         jets=((), (), tuple(jets)))) as run:
+                cfg, L1, context=ctx, coarse_boundary=mode, child_objects=[child_body(spec)],
+                child_jets=jets, **dict(KW, child_block=block)) 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):
-        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:
+        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:
             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,8 +26,9 @@ FP64_ATOL = 3e-6
 
 
 def make(ctx):
-    return FixedPatchStepper(CFG, context=ctx, refinements=[PARENTS, A],
-                             level_blocks=[(4, 4)] * 3)
+    return FixedPatchStepper(CFG, PARENTS, context=ctx, base_block=(4, 4),
+                             fine_block=(4, 4), child_block=(4, 4),
+                             child_rectangles_in_parent=A)
 
 
 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,context=ctx,refinements=[(8,6,24,18)],level_blocks=[(4,4)]*2,
-                        level_objects=[(),(obj,)],level_jets=[(),(whole,)],observe_raw_exchange=True) as run:
+                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:
                     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,8 +153,7 @@ 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(context=ctx,refinements=[(8,6,24,18)],level_blocks=[(4,4)]*2,
-                              level_objects=[(),(obj,)],level_jets=[(),(j,)])
+                settings=dict(bounds=(8,6,24,18),context=ctx,base_block=(4,4),fine_block=(4,4),fine_objects=(obj,),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))
@@ -187,9 +186,8 @@ 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,context=ctx,refinements=[parents,a],
-                level_blocks=[(4,4)]*3,level_objects=[(),(),(obj,)],level_jets=[(),(),(j,)],
-                observe_raw_exchange=True))
+            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._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,8 +22,9 @@ B=((20,16,36,40),)
 
 
 def make(ctx):
-    return FixedPatchStepper(CFG,context=ctx,refinements=[PARENTS,A],
-                             level_blocks=[(4,4)]*3)
+    return FixedPatchStepper(CFG,PARENTS,context=ctx,base_block=(4,4),
+                             fine_block=(4,4),child_block=(4,4),
+                             child_rectangles_in_parent=A)
 
 
 def initial(run):

tests/validation/run_dfg_re20.py

@@ -126,8 +126,8 @@ def run_case(a, h, case):
             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,)],
+                cfg, rr.PLACEMENTS["wide"], context=context, coarse_boundary="channel",
+                base_block=rr.BASE_BLOCK, fine_block=rr.FINE_BLOCK, fine_objects=(body,),
                 observe_raw_exchange=True, observe_window_exchange=True))
             from CelerisLab.lbm.channel import channel_masks
             from CelerisLab.lbm.refinement import ACTIVE

tests/validation/run_jet_circle_cv.py

@@ -279,11 +279,12 @@ 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,
-            refinements=[c["bounds"]],
-            level_blocks=[(8, 8), (8, 8)],
-            level_objects=[(), (obj,)],
-            level_jets=[(), tuple(jets)],
+            base_block=(8, 8),
+            fine_block=(8, 8),
+            fine_objects=[obj],
+            jets=jets,
             observe_raw_exchange=True,
         ) as amr,
     ):

tests/validation/run_refined_dfg.py

@@ -383,12 +383,13 @@ def run_case(a, h, periodic, case):
             backend = stack.enter_context(
                 FixedPatchStepper(
                     cfg,
+                    PLACEMENTS[a.placement],
                     context=context,
                     coarse_boundary="channel",
-                    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)],
+                    base_block=BASE_BLOCK,
+                    fine_block=FINE_BLOCK,
+                    fine_objects=(body,),
+                    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 {}),
                 )

tests/validation/taylor_green/test_projected_pairs.py

@@ -249,41 +249,19 @@ class ProjectedPairsTests(unittest.TestCase):
         self.assertFalse(assess(incomplete)["comparisons"][3]["passed"])
 
     def test_pinned_build_sources_match_host_generated_inputs(self):
-        """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
+        from CelerisLab.cuda import compiler_v2
 
-        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)
+        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)
 
     def test_source_preflight_requires_clean_exact_revision_and_tree(self):
         import oracle

tests/validation/time_refined_dfg.py

@@ -173,9 +173,8 @@ 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, context=context, coarse_boundary="channel", refinements=[r.BOUNDS],
-            level_blocks=[r.BASE_BLOCK, r.FINE_BLOCK], level_objects=[(), (body,)],
-            observe_raw_exchange=True))
+            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))
         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, context=context, coarse_boundary="channel", refinements=[r.BOUNDS],
-            level_blocks=[r.BASE_BLOCK, r.FINE_BLOCK], level_objects=[(), (body,)], observe_raw_exchange=True,
+            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,
             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)]