research: VIV driver — fail closed on any non-finite recorded output

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

docs/development/src_coherence.md

@@ -1,15 +1,29 @@
 # src coherence stage (from main 914906e)
 
 Owner direction 2026-10-28: unified style over historical compatibility, inside `src`.
-No numerical change is intended; kernel text was moved verbatim (PTX/bit parity is
-pending a GPU run).
+No numerical change is intended. Stages after 5f69eba: 65339ab (one boundary/Jet record,
+host preparation helpers, pair accessors), c5a3cdb (`block_collide` families), 442979b
+(exchange metadata on the boundary record). All three received parent/candidate
+generated-config and host/GPU equality checks. Only the collision slice ran a new PTX
+matrix (a4fc46d, whose src is identical to c5a3cdb); 65339ab and 442979b reused prior
+PTX evidence after unchanged CUDA/compiler and generated-input checks. In the collision
+matrix, only the two channel-root SRT modules changed PTX (omega load/clamp moved);
+their parent/candidate outputs matched exactly in the exercised cases.
 
 ## Conventions (defined once, in `lbm.hierarchy.Hierarchy`)
 
 - Level `l = 0 .. depth-1` (0 = root). Pair `hierarchy.pairs[l-1]` builds level `l`.
 - Array shapes are `(ny, nx)`; rectangles, origins and centers are `(x, y)`.
-- FixedPatchStepper bodies/Jets use the root EDGE frame (root cell `i` = `[i, i+1)`);
-  uniform `Simulation` objects use the integer-NODE frame. Nothing converts between them.
+- Coarse/fine name the two members (parent/child) of ANY pair, not absolute L0/L1:
+  `fine_shape`, `fine_origin`, `prepare_fine_boundary`, `blocked_fine_tables`,
+  `blocked_coarse_tables`.
+- Refined levels (l >= 1) of FixedPatchStepper place bodies/Jets in the root EDGE frame
+  (root cell `i` = `[i, i+1)`). Uniform `Simulation` objects AND the static circles of a
+  single-level (depth 1) FixedPatchStepper use the integer-NODE frame. Nothing converts
+  between them. This concerns input placement only: within the EDGE frame, refined
+  levels still project root-edge coordinates to fine-local cell indices with the
+  half-cell shift (`prepare_fine_boundary`: `(v - o)/dx - .5`; hits back via `fine_hits_to_edge_frame`
+  `origin + dx*(hit + .5)`), using `edge_frame`/`fine_origin` below.
 - `Hierarchy.edge_frame(l)` -> `((ox, oy), h)`: cell `(i, j)` of level `l`'s allocated
   array (ghost band and holes included) covers `[o + h*i, o + h*(i+1))`. A scalar
   FixedPatch2D window carries the half-coarse-cell ghost shift `bounds - HALO/2` in
@@ -20,17 +34,25 @@ pending a GPU run).
 
 ## FixedPatchStepper migration map
 
+Historical: left column = names at or after 914906e that no longer exist; right column =
+the current name.
+
 | removed / changed | replacement |
 |---|---|
 | `patch`, `child_patch` | `hierarchy.pairs[0]`, `hierarchy.pairs[1]` |
 | `base_layout`, `fine_layout`, `child_layout` (stepper properties) | `layouts[0]`, `layouts[1]`, `layouts[2]` (`FixedPatchSet2D.fine_layout(block)` is unchanged) |
 | `fine_boundary`, `child_boundary` | `boundaries[1]`, `boundaries[2]` (host geometry; `None` = no bodies/Jets) |
 | `_jets` | `jets[1]` |
-| `_boundaries[l]` (dict from L1, level keys) | `_level_boundaries[l]` (list from L0, `None` at L0 unless single-level static circles) |
-| `_fine_tables`, `_actuated`, `_jet_cache` | `_level_boundaries[1].tables` / `.actuated` / `.jet_cache` |
+| `_boundaries[l]` (dict from L1, level keys), `_level_boundaries` (b4eacf7), `_device_boundaries` (65339ab) | `_boundaries[l]`: the one combined per-level `_BlockedBoundary` record (list from L0, `None` at L0 unless single-level static circles): host geometry `.host`, `.tables`, `.jets`, `.jet_cache`, exchange `.spacing`/`.centers`, device handles |
+| `_fine_tables`, `_actuated`, `_jet_cache` | `_boundaries[1].tables` / `.actuated` / `.jet_cache` |
+| `_level_jets` (private) | `_boundaries[l].jets` (public `jets[l]` is derived) |
+| `_exchange_spacing[l]`, `_exchange_centers[l]` (private) | `_boundaries[l].spacing` (h_l, root units) / `.centers` (root edge frame); only levels with a boundary report exchange |
+| `_level_blocks` (private) | `_layouts[l].block` |
+| `_remap_cfg` (private) | `_root_cfg` |
+| `refinement_tables.blocked_base_tables` (also via `refinement_stepper`) | `refinement_tables.blocked_coarse_tables` (import from `refinement_tables`; the stepper no longer re-exports) |
 | `_fine_leaf`, `_child_leaf` | `leaf[1]`, `leaf[2]` (copies) |
 | `_middle_kind` | `hierarchy.kinds[1]` |
-| `_boundary_hosts` (private) | `_boundaries` (private list behind `boundaries`) |
+| `_boundary_hosts` (private) | public `boundaries[l]` = `_boundaries[l].host` (host geometry only; the private record is the combined one) |
 | `remap_child(rects, child_objects=...)` | `remap_finest(rects, objects=...)` |
 | `set_jet_amplitude(i, a)` (implicit level 1) | `set_jet_amplitude(i, a, level=l)` (required) |
 | `read_raw_exchange(b)` / `read_exchange_window(b)` (implicit level 1) | `read_raw_exchange(b, level=l)` / `read_exchange_window(b, level=l)` (required) |
@@ -45,18 +67,50 @@ the replaced level relative to its parent).
 
 ## Responsibilities
 
-- Construction: `_check_levels` (inputs), `_kernel_configs` (per-level kernel configs),
-  `Hierarchy.edge_frame` (geometry frames), `_build_level` / `_build_pair` / `_upload_boundary`
-  (device). Observation reads share `_observed_boundary` + `_summarize`.
+- Construction: `_check_levels` (inputs); host preparation, numpy only, before any device
+  work: `_level_layouts`, `_kernel_configs` (per-level kernel configs), `_host_boundaries`,
+  `_prepare_blocked_boundary`, `_fluid_leaves`; `Hierarchy.edge_frame` (geometry frames);
+  then `_planned_bytes()` (memory preflight) and `_build_level` / `_build_pair` /
+  `_upload_boundary` (device). A construction failure closes the partial instance.
+  Observation reads share `_observed_boundary` + `_summarize`.
+- Pair accessors: `rectangles` (a scalar FixedPatch2D gives `(bounds,)`) and
+  `fine_level_offset` (`(0, 0)` for a set) on both pair types; the types and their
+  storage frames stay distinct.
+- Per-level state, by lifecycle:
+  - Fixed plan structure (built once in `__init__`, read by `_planned_bytes()` before any
+    allocation, readable after `close()` through the public views): `_layouts`,
+    `_level_cfgs`, `_leaf`, `_state`, and in each `_boundaries[l]` record the host geometry
+    `.host`, link tables, `.jet_cache`, `.spacing`, `.centers`.
+  - Mutable boundary data in the same record: `.jets` and `.tables["surface_velocity"]`,
+    changed only by `set_jet_amplitude` (host prepare -> device write -> commit; an uncertain
+    write fails the instance); device handles (`.flags`, `.device`, `.kernel`, ledgers) are
+    set by `_upload_boundary` and cleared by `close()`.
+  - Runtime: `_levels` (`_Level`) holds compiled modules, kernels, the two population
+    buffers (swapped per step) and transfer/kind handles. It is created only after the
+    memory preflight and cleared on `close()`; `.layout`/`.boundary` reference the plan for
+    `_hop`. Folding the plan into `_Level` would mix both lifecycles in one record.
 - Blocked kernels: `operators/block_collision.cuh block_collide` is the one collision extension
-  point of `BlockStep` and `BlockChannelStep`. MRT shares one path; SRT keeps its two explicit
-  arithmetic variants (level: unclamped, no-force, shifted direct-momentum feq; channel: uniform
-  OneStep arithmetic with omega clamp), selected by `BLOCK_ROOT_BC`.
+  point of `BlockStep` and `BlockChannelStep`. It computes rho/u once and reads no runtime
+  force (`d_params.fx/fy`). Two compile-time families: uniform-equivalent (MRT on any level,
+  SRT on the channel root `BLOCK_ROOT_BC == 2`): clamped omega, `collide_mrt_no_force` or
+  `collide_srt` with its zero Guo term; and level SRT (`BLOCK_ROOT_BC` 0/1): unclamped omega,
+  `collide_srt_no_force`, with the shifted direct-momentum feq under `USE_DDF_SHIFTING`.
+  The uniform path keeps `collide_dispatch` (runtime force, LES, TRT).
 - `explosion="projected_linear"` stays an optional capability; `"constant"` stays default.
 
 ## Not done / open
 
+- Two per-level config derivations stay distinct on purpose: `Hierarchy.configs`
+  (`_middle_config`, wall fields reset on every refined level, geometry validation) vs.
+  `FixedPatchStepper._kernel_configs` (reset only on level 1 of a slot_reservoir root;
+  generates each level's config.h). Folding them needs a config.h/PTX diff gate.
 - No single-`Simulation` facade over FixedPatchStepper: the two own different state models
-  (uniform node frame, body sync, forcing, checkpoint vs. immutable blocked hierarchy, edge
+  (uniform node frame, body sync, forcing, checkpoint vs. fixed blocked hierarchy, edge
   frame, remap-as-new-instance). A facade would only forward calls; no concrete need found.
+- Deferred, each needing its own evidence first: sharing the channel-root cell/donor selector
+  with OneStep (the closure functions are already shared; a shared selector changes uniform
+  signatures and needs a uniform PTX gate); merging FixedPatch2D into FixedPatchSet2D
+  (different origin, ghost band, layout size and addressing; needs a geometry/parity table and
+  would need tiles that are not grid-aligned); one `clamp_omega` helper for uniform and
+  blocked paths (it lives in the hash-pinned LES header `turbulence_smag.cuh`).
 - Tests, examples and validation callers are migrated separately (pinned runners stay frozen).

docs/moving_bodies.md

@@ -0,0 +1,403 @@
+# Moving circles on the uniform dense grid — S0 DRAFT contract
+
+Status: **DRAFT, not frozen** for the force RHS, kernels and physical
+acceptance. Host-only S1 (state, validation, guards; no kernels) is on branch
+`claude/moving-bodies-s1`; positive-step runs with a moving body are rejected
+until S2. When S2 lifts that guard, `run(n > 1)` with a pending target must
+raise `ValueError` before mutation (4a), and `run(0)` must not upload a motion
+record or consume the pending target. Base: main
+`b2034f87e03656c5f67ab88562a70d6c338b9918`. Scope approved by Frank (via TRON,
+thread C0C78UFKGRM/1793335557.460509): uniform-grid prescribed translation and
+single-cylinder free transverse VIV. Out of scope: refinement/blocked path,
+pinball assembly, EsoPull, shifted storage, D3Q19, DynamisLab, manuscript.
+
+Items marked **OPEN** are undecided; items marked **CANDIDATE** are proposals.
+
+## 1. Working standards
+
+- Lean, readable kernels in the existing idiom (`kernels/step/*.cu`, shared
+  `apply_bouzidi_link`); Python manages state and validation. No framework,
+  no compatibility layer, no new dependency.
+- The existing stationary/rotating path (`CurvedBoundaryKernel`, host link
+  packing, `sync_bodies`) stays byte-identical. Moving circles are a separate,
+  explicitly enabled path with its own numerical contract (this file).
+- One authority per state (section 4). No per-step full-domain download,
+  rebuild or recompile; local swept-band device updates are in scope.
+- Small slices (S1 host spec, S2 kernels without motion, S3 motion + refill,
+  S4 oscillator); each reviewed SHA is frozen before targeted runs.
+- Tolerances, benchmark parameters and the force composition used in the
+  structural RHS are declared before the runs that judge them. Fresh vs carried
+  evidence and existing baseline failures are recorded; smoke, convergence and
+  physical validation are reported separately. No undeclared force filtering.
+- External code (waLBerla GPLv3, FluidX3D source-available, VIVSIM) is a
+  concept reference only; nothing is copied.
+
+## 2. Current update order (main b2034f87, dense double buffer)
+
+`Simulation.run(n)`: upload host action (if `upload_act`) → zero obs (if
+`zero_obs`) → `DoubleBufferStepper.step(n)` → download obs → synchronize.
+
+Per step (`lbm/stepper.py` `step`):
+
+1. `CurvedBoundaryKernel` over the packed cut links: reads post-collision
+   `G^{n-1}` at `k_f` (and `k_f - c` for q < 1/2), adds the moving-wall term
+   (`Uw = -omega*ry, Vw = omega*rx` + optional per-link `surface_velocity`),
+   writes the reflected value into slot `opp` of the solid node `k_s`, and
+   accumulates `c(f_toward + f_reflected)` and torque into `obs`.
+2. `OneStep` pull/collide `ddf → temp` (collides FLUID cells; obstacle
+   interiors flagged `SOLID|OBSTACLE|BC_CURVED`), then buffer swap.
+3. `ForceRegionKernel`, `SensorKernel`.
+
+Raw `obs` force is the force **on the body** (link derivation in section 6;
+DFG docs report `Cd > 0` from it). Translation `vx, vy` in the action slots is
+currently stored but not read by any curved kernel.
+
+## 3. Proposed update order (moving circles)
+
+Per step n, before the existing step 1:
+
+- **M1 `MovingCircleCells`** (one thread per swept-band cell): classify with the
+  analytic predicates `in_old = inside(x, c^{n-1})`, `in_new = inside(x, c^n)`;
+  write `flag = in_new ? SOLID|OBSTACLE|BC_CURVED : FLUID`; covered cells
+  (`!in_old && in_new`) accumulate `node_cover`; fresh cells
+  (`in_old && !in_new`) are refilled (section 5) and accumulate `node_uncover`.
+- **M2 `MovingCircleLinks`** (one thread per band cell × direction): compute
+  q and hit point at `c^n` (FP64, same formula as `find_circle_intersection`,
+  same clamp [0.01, 1] and fallback classes as `CircleGeometry.build_cut_links`),
+  skip non-cut slots, call `apply_bouzidi_link` with
+  `u_w = (vx^n, vy^n) + omega × (hit - c^n)`, accumulate the link terms
+  (section 6).
+
+Then the unchanged steps 1–3. M2 must follow M1 because a q < 1/2 link may
+read `k_f - c` at a fresh cell.
+
+Exclusivity: moving bodies are never packed into the `CurvedBoundaryKernel`
+link arrays (`n_curved` covers stationary/rotating bodies only), so M2 is the
+sole writer of their solid slots and the sole source of their link force; no
+double write and no double `obs`. The stationary path's packing, launch and
+arithmetic are unchanged.
+
+Completeness: after M1, every (fluid cell, direction) pair in the swept band
+whose neighbour is `in_new` is visited by M2. This holds because the band is
+the circle AABB at the declared extreme offsets plus 2 cells, and the host
+rejects per-step moves with `|c_new - c_old|_2 > 0.5` and bands touching the domain edge,
+other bodies, sensors or force regions. Each M2 write lands at
+`index_f(k_s, opp)`, the exact slot the adjacent fluid cell pulls, so every
+fluid pull from a solid slot of a moving body is BFL-written.
+
+## 4. State, masks and authority
+
+- Host `ObjectState` is the single authority for one displacement
+  `offset = (dx, dy)` relative to the immutable base centre, plus
+  `velocity = (vx, vy)` and `omega`. Centres are always derived
+  (`c = base + offset`), never stored as independent state. Velocity is
+  supplied, never differenced from offsets.
+- `c^{n-1}` is derived from the last *committed* offset, `c^n` from the
+  offset being applied. Per body a small device **motion record**
+  `{c^{n-1}, c^n, u^n}` is uploaded with the action buffer; action-slot meaning
+  is not overloaded. (CANDIDATE layout; OPEN.)
+- Fail-closed commit: an enqueued upload is not a completed one. The new
+  offset becomes the committed offset only after the stream synchronization
+  that follows the upload and the step group succeeds. Any upload or later
+  synchronization failure leaves host and device motion state unknown: the
+  simulation is marked unusable (existing `_require_usable` lifecycle), and
+  no API may report a host/device match afterwards.
+
+### 4a. Applied geometry vs pending target (CANDIDATE, narrow API)
+
+Two distinct facts per moving body, one authority each:
+`applied` = offset of the geometry last applied on the device, changed only
+by a successful commit; `pending` = an optional requested target
+`(offset, velocity)`, changed only by `set_body`. A setter returning proves
+nothing about the device. The uploaded motion record is derived each run:
+`{c_old = base + applied, c_new = base + (pending.offset if pending is not
+None else applied), u = pending.velocity if pending is not None else 0}`
+(explicit `is not None`, so a zero offset or zero velocity is a valid target). Translational wall velocity exists only on a step
+that advances to a pending target; a hold step (`c_old == c_new`) has
+translational `u_w = 0` (spin unchanged), so a non-advancing circle never
+acts as a sliding wall. One moving step consumes one target.
+
+| situation | contract |
+|---|---|
+| `set_body` target | validated against `applied` (box, `|Δc|_2 ≤ 0.5` cell, speed bound) before mutation; on success replaces `pending` |
+| target with `offset == applied` and nonzero translational velocity | `ValueError` before mutation (would be a hold step with sliding wall velocity) |
+| two setters before a run | last wins; each validated against `applied`, not against the earlier pending target |
+| `run(1)` with pending | upload record, step, sync; on success `applied := pending.offset` (applied stores only the offset), `pending := None` |
+| `run(1)` without pending | hold: `c_old == c_new`, no cover/uncover, links at `applied`, translational `u_w = 0` |
+| `run(n > 1)` with pending | `ValueError` before mutation (single-step-only moving API) |
+| `run(n > 1)`, no pending | allowed: hold for n steps, translational `u_w = 0`, spin only |
+| `run(0)` | no motion upload, no commit; `pending` kept |
+| validation failure | `ValueError`, nothing mutated, retry allowed |
+| upload/launch/sync failure after enqueue | simulation unusable; `applied` and `pending` are not reported as device truth; no re-entry |
+
+Displacement norm: Euclidean per-step `|c_new - c_old|_2 ≤ 0.5` lattice
+cells (declared bound; CANDIDATE value).
+- Membership is never read from a flag that M1 is overwriting. The host
+  verifies at `add_body` that every swept-band cell has base flag FLUID, so the
+  band's base state is known without reading `flag_gpu`.
+- `flag_gpu` is the device authority for band cells while a moving body
+  exists; the host flag mirror is stale for those cells and is not used.
+
+## 5. Refill of fresh cells (CANDIDATE)
+
+- Stored time level: everything in `ddf` at M1 is post-collision `G^{n-1}`; the
+  refill writes all 9 slots at that same level, so every slot of a fresh cell
+  that the following pull or a q < 1/2 link may read is defined. Donor values
+  are read at the same level `G^{n-1}` (old-fluid cells, untouched by M1).
+- `G = f_eq(rho_bar, u_w(x)) + G_neq(x_d)` where `rho_bar` = mean density of
+  valid neighbours and `x_d` = the valid neighbour best aligned with the outward
+  normal `(x - c^n)/|x - c^n|`; `G_neq(x_d) = G(x_d) - f_eq(rho(x_d), u(x_d))`.
+- Valid donor: inside the domain, base FLUID, `!in_old && !in_new`.
+- Sets: `V(x)` = the 8 lattice neighbours that are in the domain, base FLUID
+  and `!in_old && !in_new`. `rho_bar` = mean density over `V(x)`. NEQ donor =
+  the member of `V(x)` with the largest `c·n` among those with `c·n > 0`
+  (`n` = outward normal at `c^n`); none → NEQ part 0 (counted as `eq_only`).
+- `V(x)` empty → no defined density: the cell is refilled with
+  `f_eq(RHO, u_w(x))` so all 9 slots stay defined, the `no_donor` counter is
+  incremented, and the host treats `no_donor > 0` after the run as a hard
+  failure (simulation unusable). It is not a silent fallback.
+- Non-emptiness is NOT assumed from the displacement bound. Geometry
+  assumptions: circle radius `R ≥ R_min` (CANDIDATE 4 cells), `|Δc|_2 ≤ 0.5`.
+  S1 includes a host enumeration over sub-cell centre positions and
+  displacement directions on a fine grid within these bounds that checks every
+  fresh cell has nonempty `V(x)`; it is evidence for those bounds, not a
+  proof, and runtime keeps the counter.
+- Fresh cells only write, old-fluid cells only read: no race.
+
+## 6. Discrete momentum ledger (to review; not a selection)
+
+Notation per step n: `F^n` = fluid set after M1; links L = (k_f, dir) at
+`c^n`; `f_t = G_dir(k_f)`; `f_r` = value M2 writes into `(k_s, opp)`.
+Fluid momentum before streaming:
+
+    P_pre = P(G^{n-1} on F^{n-1}) - sum_cover sum_q c_q G_q + sum_fresh sum_q c_q G_q^refill
+
+During the pull, every population entering a fluid cell comes from a fluid
+cell except the BFL-written slots; `f_t` leaves into solid and is never pulled
+by a fluid cell. Collision conserves momentum for SRT/MRT without forcing.
+Hence, in exact arithmetic,
+
+    P_fluid^n - P_fluid^{n-1} = - sum_L c (f_t + f_r) - I_cover + I_uncover + B^n + S^n
+
+with `B^n` external-boundary momentum (inlet/outlet/walls) and `S^n` forcing.
+So **`F_CME + node_cover - node_uncover`** is, by construction, the body
+momentum that balances the fluid ledger of *this* storage/update order.
+Branch details: halfway fallback links (`f_r = f_t` plus the moving-wall
+addend) and q < 1/2 links enter the link accumulators exactly like Bouzidi
+links, with `c(f_t + f_r)`. In the q < 1/2 branch `G_dir(k_f - c)` is an
+interpolation input only, not a fluid-momentum sink. Pull-slot completeness
+is the section 3 contract. The scope excludes `ddf_shifting` (rejected), so
+the unshifted force form applies. The balance is a fluid-side identity; it
+does not close body+fluid globally in a domain with inlet/outlet/forcing
+(`B^n`, `S^n`), and the solid-slot writes are boundary plumbing, not physical
+streaming packets.
+
+The Wen 2014 relative-velocity link sum adds `-u_w (f_t - f_r) = u_w dm` per
+link (`dm = f_r - f_t`, the BFL mass defect). That term has no counterpart in
+the fluid ledger above. Peng et al. (PRE 2016 submitted, §II.C, Eqs. 20–21)
+show for a straight wall that the bounce-back error accumulated before a node
+switch cancels with the switching momentum. Whether `sum u_w dm` approximates
+`I_cover - I_uncover` here (implying double counting if both are used) is
+**NOT demonstrated** for this implementation: it needs a bounded experiment
+(prescribed slow translation; cumulative `sum u_w dm` vs cumulative node
+impulse, per crossing cycle). Momentum conservation is also distinct from
+Galilean invariance, which G4 tests.
+
+External reference note: waLBerla's covered/refill momentum switches
+(`conserveMomentumWhenMapping`, `conserveMomentumWhenRestoring`) default to
+false (Astra), so their existence does not establish which combination its
+validation cases use. A bookkeeping identity alone does not select the
+physical force either.
+
+Mass ledger (same proposition status). With `M^n` the active-fluid mass
+summed over `F^n` at the sample point,
+
+    M^n - M^{n-1} = sum_L dm - m_cover + m_uncover + B_m^n
+
+(`dm = f_r - f_t`, `m_cover`/`m_uncover` = summed `sum_q G` of covered and
+refilled cells, `B_m` external-boundary mass flux; collision conserves mass).
+
+Sample point (fixed): post-collision storage after the swap of step n, i.e.
+the same level `G` that M1 reads at step n+1; both ledgers are evaluated
+there. Ledger residual = left side minus every listed term. Zero mass flux
+does not imply zero momentum exchange: closed walls (bounce-back, free-slip)
+still exchange momentum, so "no inlet/outlet, no forcing" does NOT give
+`B = S = 0`. The rounding-level momentum check therefore uses either a
+momentum-isolated case (periodic in both directions, no forcing) or counts
+the wall impulse explicitly as part of `B`; the mass check may use a closed
+domain. Otherwise `B`, `S` are not separately measured, the residual is
+reported as their estimate and is not a pass criterion.
+Both ledger identities are **proposed, unreviewed**: they are not called
+proven until review and a host replay agree.
+
+Accumulators (separate, per body, per step window; all reported raw):
+`link_cme`, `link_wen_correction` (= `sum u_w dm`), `node_cover`,
+`node_uncover` (momentum), `link_dm`, `mass_cover`, `mass_uncover` (mass),
+plus counters of crossings, `eq_only` and `no_donor`, so force spikes can be
+correlated with crossing events.
+
+**OPEN — structural RHS candidates (not selected):**
+(a) `link_cme + node_cover - node_uncover`; (b) `link_cme + link_wen_correction`
+(no node impulse). Selection follows the ledger review and the bounded
+experiment above.
+
+## 7. Rejection semantics
+
+- Before any mutation (`ValueError`, no device write, host state unchanged):
+  moving body with EsoPull, `ddf_shifting`, refinement/blocked stepper, D3Q19,
+  Jets or surface actuation on that body, `sync_bodies`/checkpoint while a
+  moving body exists; swept band touching domain boundaries, other bodies,
+  sensors or force regions; `set_body` offset outside the declared box,
+  per-step move `|c_new - c_old|_2 > 0.5` cell (the single bound of 4a), speed
+above the declared bound.
+- Host motion state is committed only after the upload *and* the following
+  synchronization of the step group succeed (section 4); any failure in
+  between makes the simulation unusable and no host/device match is reported.
+- The only soft device diagnostic is `eq_only` (NEQ part omitted, counted).
+  `no_donor > 0` is a hard failure: it is checked after the step group and
+  before the step is committed or any results are returned; on failure the
+  simulation is unusable (section 5).
+
+## 8. Coupling and units (CANDIDATE)
+
+Lattice units, rho0 = 1, per unit span. Step n uses `c^n`, `u^n` and yields
+`F^n`. Weak staggered update
+`v^{n+1} = v^n + (F^n - c v^n - k y^n)/m`, `y^{n+1} = y^n + v^{n+1}`
+(Newmark is an alternative; OPEN). Mass, natural frequency, damping and `U_r`
+follow the chosen benchmark's own definitions. Leontini, Thompson & Hourigan,
+JFS 22:857–864 (2006), DOI 10.1016/j.jfluidstructs.2006.04.003, as read by
+Grok from the full text: Re = 200, `m* = m/m_f = 10` (m_f displaced fluid
+mass), `zeta = c/c_crit = 0.01`, `U* = U/(f_N D)` with `f_N` the natural
+frequency in fluid assuming an inviscid added-mass term; the paper also
+reports frequencies normalised by the in-vacuo value. Our runner derives `k`
+from that in-fluid `f_N` (added-mass coefficient as stated in the paper) and
+never from the vacuum formula; `c_crit` convention and the original
+comparison data are still OPEN. Stability/convergence is gated by fluid
+time/space refinement of the coupled run, not assumed from the mass ratio;
+structural sub-stepping does not substitute for fluid-coupling time
+refinement.
+
+## 9. Gates (map; tolerances OPEN, declared before runs)
+
+G0 zero-motion regression, split: G0a populations of the existing stationary
+and spin-only path bitwise unchanged vs main; G0b moving path at rest (with
+and without spin) reproduces the stationary link set and populations; force
+estimator outputs are compared separately from populations, since the moving
+path adds estimator terms · G1 raw force sign (`Cd > 0`) ·
+G2/G3 prescribed translation incl. ledger accumulators · G4 matched
+relative-flow (moving body in quiescent fluid vs fixed body in uniform
+stream), optional boost, co-moving sanity · G5 Dütsch et al. 1998 Re 100,
+KC 5 in-line (needs 2-component offset) · G6 Leontini free transverse VIV
+(m* = 10, ζ = 0.01, Re = 200, definitions pending) · G7 time-level contract
+frozen before the oscillator.
+
+Candidate pre-run engineering tolerances (Astra; for discussion, not
+literature limits or measured outcomes; norms, averaging window, convergence
+protocol and physical reference uncertainty must be fixed before adoption, and
+none may be weakened after seeing runs):
+normalised momentum-ledger residual FP32 ≤ 5e-6 (FP64 oracle ≤ 2e-13);
+10-cycle relative mass drift ≤ 1e-4; prescribed-motion grid study D = 40/80/160
+with change in the force-response amplitude (observed quantity and norm to be
+named before adoption) ≤ 2 % and phase ≤ 2°; enlarged domain
+amplitude change ≤ 1 %; matched-frame force RMS ≤ 2 % and improving with
+refinement; VIV refinement amplitude ≤ 2 %, frequency ≤ 1 %, five consecutive
+cycle amplitudes within 1 %.
+
+Conditions stated before any moving kernel (S2/S3), not only at the
+oscillator stage: (i) fluid refinement protocol for the prescribed-motion
+gates — diffusive scaling (dx halved, dt quartered, physical Re, KC and
+amplitude fixed), the observed quantities and norms named in advance;
+(ii) Galilean constant-boost G4-boost requires boundaries compatible with a
+uniform boost (periodic in the boost direction or boundaries that move with
+it); if the current dense configuration cannot represent that, G4-boost is
+recorded as not runnable rather than replaced by another test.
+
+## 10. Proposed minimal S1 scope (host only; held until cross-review)
+
+- `add_body` moving-circle declaration: base centre, declared offset box
+  (x/y), speed bound; swept-band computation and band-placement validation.
+- `set_body(offset=..., velocity=..., omega=...)` validation before mutation
+  (section 7), applied/pending state and the run-transition table of 4a,
+  and the fail-closed commit of section 4.
+- Host enumeration check of refill-donor non-emptiness (section 5) within
+  the declared `R_min` and displacement bound.
+- Motion-record host layout and upload path; exclusion of moving bodies from
+  the packed `CurvedBoundaryKernel` links (stationary packing unchanged).
+- Explicit rejections for unsupported modes.
+- S1 execution guard: until the S2 moving kernels exist, any `run(n)` with
+  `n > 0` (and `step(n)`) while a moving body is declared raises before any
+  device mutation (no action/motion upload, no launch), including the
+  nominal at-rest case, because the moving body is excluded from the static
+  packed links and nothing else would apply its boundary. `run(0)` keeps its
+  existing behaviour. Simulations without moving bodies are unaffected.
+  Preparation/commit helpers are exercised by host tests with mocks; no
+  operational unsupported stepping path is exposed.
+- Host tests (Astra runs them): rejection-before-mutation, band geometry,
+  derived centres, fail-closed behaviour under an injected upload/sync
+  failure, stationary packing byte-identical, execution guard (positive-step
+  run with a moving body raises before upload/launch; stationary unaffected). No kernels, no GPU runs, no
+  y-only VIV oscillator (S4), no force RHS.
+
+## 10a. S2a — fixed-centre moving links (implemented on `claude/moving-bodies-s2a`)
+
+- `MovingCircleLinks` (`kernels/step/moving_circle.cu`, dense double buffer)
+  runs before the static curved kernel when moving tables are attached: one
+  thread per (band cell, direction) slot, FP64 `{cx, cy, r}` at the applied
+  offset, q/hit/clamp/fallback in the order of `build_cut_links` with explicit
+  `_rn` operations (host/device bit parity is a test target, not a claim),
+  then the unchanged `apply_bouzidi_link` with spin wall velocity only.
+- `obs` keeps its contract (CME and torque from `apply_bouzidi_link`). A
+  separate per-body diagnostics buffer receives `cme`, the Wen term
+  `u_w * dm`, `link_dm` and the link count from the same link outputs; it is
+  zeroed with obs and only exists when moving circles do. Optional
+  `slot_debug` (q, fallback, rx, ry per slot) is null in production.
+- Displacement is disabled: `set_body(offset=..., vx=..., vy=...)` on a
+  translating circle raises before mutation, before and after initialize, so
+  no pending target can exist and `run`/`step`/direct `LBMStepper.step` all
+  hold at the applied geometry. `initialize()` never reads or consumes
+  `pending`. `stepper.blocked` is cleared only when the moving tables are
+  attached (initialize, stepper recreation).
+- S3 rule (not implemented): once targets are re-enabled, any pending target
+  must raise in `run`/`step` and in `LBMStepper.step` alike before any launch,
+  swap or step-count change; zero-step calls never consume it.
+- Not validation: a pass of the zero-motion gates says nothing about moving
+  geometry. GPU parity (G0a/G0b) is unverified until run on a free device.
+
+## 10b. S3a — translation with refill (implemented on `claude/moving-bodies-s3a`)
+
+- Targets are enabled again (4a). `set_body(offset, vx, vy)` validates against
+  `applied`; a pending target blocks direct `LBMStepper.step(n > 0)` before any
+  launch; `run(n > 1)` with a target raises before mutation; `run(0)` never
+  consumes it; `run(1)` advances it.
+- Per advancing step, on one stream: upload the immutable FP64 records
+  `{c_old, c_new, r, u}` (only when changed) → M1 `MovingCircleCells` → M2
+  `MovingCircleLinks` (`u_w = u + omega x r` at the hit point) → static curved →
+  `OneStep` → swap → force region → sensors → synchronize → read `no_donor` →
+  commit (`applied := pending.offset`) → upload the hold record
+  (`c_old = c_new = applied`, `u = 0`) → synchronize. Hold steps launch no M1
+  and are the S2a path. Commit point: the advancing step-group sync plus a
+  zero `no_donor` count; a failure before it commits nothing. The hold-record
+  upload stages the next step after the commit: if it (or its sync) fails,
+  `applied` already matches the moved device geometry, but the Simulation is
+  still marked failed and refuses reuse. Any failure from the first upload on
+  marks the circles and the Simulation failed.
+- M1 classifies with the old/new predicates only (no flag reads). Covered
+  cells add raw `sum c G` / `sum G` of `G^{n-1}` to `cover_*` and become
+  `SOLID|OBSTACLE|BC_CURVED`. Uncovered cells: donors are fluid in both
+  geometries; `rho_bar` over all of them; NEQ from the donor with the largest
+  unnormalized `c·n > 0` (first direction wins ties, `n` from `c_new`);
+  `u_w` (translation + spin) at the cell position; all nine slots written;
+  raw `sum c G`, `sum G` of the refill go to `uncover_*`; no donor →
+  `f_eq(RHO, u_w)` plus `no_donor` (hard failure before commit); no NEQ donor
+  → `eq_only`.
+- `get_flags()` reads the device flags back (synchronized) when moving tables
+  exist; otherwise it returns the host array as before.
+- Not selected: the structural force. CME+node vs CME+Wen stay OPEN until X1
+  and G4 have run.
+
+## 11. Remaining decisions
+
+Structural RHS (section 6); motion-record layout; refill details; coupling
+scheme; benchmark definitions and digitized reference data; tolerances;
+execution/resource feasibility (Astra, pending); whether a quiescent domain
+(zero inlet velocity) runs cleanly in the current config.

docs/validation/duetsch_re100_kc5_protocol.md

@@ -0,0 +1,242 @@
+# Prescribed in-line oscillation: Dütsch Re 100 / KC 5 — PROTOCOL (review draft)
+
+Production under test: `claude/moving-bodies-s3a` at `653fd997` (unchanged).
+Primary source: Dütsch, Durst, Becker & Lienhart, J. Fluid Mech. 360, 249–271
+(1998), doi:10.1017/S002211209800860X. Source facts below were verified from
+the paper (Slack C0C78UFKGRM/1793335557.460509, TRON 1793356317–1793356724,
+page numbers given). Everything marked ENGINEERING is our a-priori choice,
+fixed before any benchmark result; it is not a paper uncertainty.
+
+## 0. Source facts (with provenance)
+
+- Motion (p.251): `x_c = -A sin(2 pi f t)`; paper phase phi = 0 is the mean
+  position moving left at `-U_max`. KC = 5 gives `A/D = 5/(2 pi) = .7957747`.
+- Start (p.254): the computation starts a quarter cycle before phi = 0, at
+  `x = +A` with zero velocity (phi = -90°).
+- Experiment (pp.252, 254–256): D = 10 mm, L = .65 m, tank 1.2 × .7 × .7 m,
+  rho = 998.2 kg/m³, mu = 1.00818 mPa·s, Re = 100, KC = 5; LDA with 1° phase
+  bins and 200 samples per node and degree; no quantified experimental error
+  recovered.
+- Profiles (Fig. 8, p.260): LAB-FIXED stations x1 = -6, 0, 6, 12 mm, i.e.
+  x/D = -.6, 0, .6, 1.2, at phi = 180°, 210°, 330°; u1 and u2 versus x2.
+- Original numerics (§2): cylinder-fixed accelerated frame (`x_i = x~_i +
+  x_i,s`, outer velocity `-u_c`, inertial source `-rho du_c/dt`), O-grid outer
+  diameter 120 D, illustrated 192 circumferential nodes, 720 steps/cycle; tank
+  influence neglected.
+- Force fit (verified): force ON THE BODY per unit span
+  `F1 = -0.5 rho D c_d xdot|xdot| - (pi/4) rho D² c_i xddot` with `x = -A sin phi`;
+  normalized by `0.5 rho D U_max²`: `c_d cos phi |cos phi| - (pi² c_i / KC) sin phi`.
+  `c_i` is the ADDED-MASS coefficient; total inertia `c_m = c_i + 1` is a
+  different convention and is not used.
+- Table 1 computed Morison coefficients (c_d, c_i) for 96 / 192 / 384 perimeter
+  nodes at dt/T = 1/720: (2.15, 1.42), (2.10, 1.45), (2.09, 1.45). These are
+  fitted computed coefficients, not a measured force and not an error bar. The
+  paper states the coarse direct force is closer to the finest direct force
+  than the Morison fit, but gives no waveform-error norm.
+
+## 1. Kinematics, time levels and phase bookkeeping
+
+- Our run is in the LAB frame: the circle moves, the fluid is initially at rest
+  (rho = 1, u = 0); no fictitious inertial source is added. Lab-fixed profile
+  stations are compared directly (no frame transform needed).
+- Geometry reference: the circle's base centre is the INITIAL position, the
+  right turning point `x_mean + A`; the declared offset box is x in [-2A, 0],
+  y = 0. The body is never teleported: offset 0 at s = 0.
+- Elapsed time s = n dt from the start. Analytic position relative to the mean
+  `x(s) = A cos(2 pi s / T)`, velocity `u(s) = -U_max sin(2 pi s / T)`; paper
+  phase `phi = (360 s / T - 90) mod 360`.
+- Step n moves from s_n to s_{n+1}: target offset `x(s_{n+1}) - A`; wall
+  velocity = the single-step interval average `(x(s_{n+1}) - x(s_n)) / dt`
+  (analytic, supplied, not differenced from stored offsets). The per-step
+  force of `run(1)` / `read_force(normalize=True)` is attributed to that same
+  interval.
+- Paper cycle k = 1, 2, ... spans `s in [T/4 + (k-1) T, T/4 + k T)`; the
+  initial quarter cycle (s < T/4) is cycle 0 and never scored. Every logged
+  sample carries (k, phi). Example: phi = 330° of cycle 1 is at s = 7T/6.
+- T is chosen as a multiple of 360 native steps, so phi = 180°, 210°, 330° fall
+  exactly on step boundaries and 1° bins hold whole steps.
+
+## 2. Resolution ladder (uniform grid, fixed Mach), ENGINEERING
+
+| D (cells) | T (steps) | U_max = 5 D / T | nu_lat = U_max D / 100 | tau (3 nu + .5) |
+|---|---|---|---|---|
+| 20 | 2,520 | .039683 | .0079365 | .5238 |
+| 40 | 5,040 | .039683 | .015873 | .5476 |
+| 80 | 10,080 | .039683 | .031746 | .5952 |
+
+- MRT collision. Mach fixed (Ma ≈ .0687); a separate Mach check at D = 40
+  with T = 10,080 (U_max halved, tau .5238) is labelled MACH and not part of the
+  ladder.
+- Per-step displacement ≤ U_max ≈ .04 cells (bound .5).
+
+## 3. Domain and boundaries (finite-domain reproduction), ENGINEERING
+
+The paper's 120 D O-grid is not reproduced; our rectangular box is a
+reproduction choice. The native dense channel (zero-velocity inlet,
+pressure-NEQ outlet, bounce-back y walls) is a finite-domain approximation,
+not infinite still fluid. The G4 x-periodic / y-slip override is a separately
+labelled sensitivity family and never mixed into the primary ladder. The
+runner owner confirms the zero-velocity inlet is supported as configured
+before any run.
+
+Order (closed):
+1. Smoke + ledger at D = 40, native BC, 40 D × 40 D (finite box).
+2. Domain sensitivity before any literature statement: same BC, D = 40,
+   L ∈ {20 D, 40 D, 60 D}; record the change in Fx waveform, c_d, c_i and the
+   scored profiles.
+3. Optional labelled BC sensitivity: one x-periodic / y-slip case at D = 40,
+   40 D; reported as a delta against native only.
+4. Only after step 2 passes: the D = 20 / 40 / 80 ladder on the chosen BC and
+   domain.
+
+Domain criteria (a-priori, ENGINEERING; separate from the Table 1 bands): the
+reference is 60 D at D = 40; 20 D and 40 D are compared with it on the common
+window (same cycle indices) and the common valid point set:
+- `RMS(Fx_L - Fx_60D) / F_ref ≤ .01` over the window (per-step waveforms);
+- `|c_d,L - c_d,60D| ≤ .02` and `|c_i,L - c_i,60D| ≤ .02` (window fits);
+- `max(|u_L - u_60D|, |v_L - v_60D|)` over common points `/ U_max ≤ .02`.
+The chosen domain is the smallest of {20 D, 40 D} passing all three; if
+neither passes, the domain verdict is FAIL and the ladder is not run.
+- Memory/runtime: measured by a smoke run before each new size; measurements
+  inform scheduling only, never tolerances.
+
+## 4. Observables (raw, per step, preserved)
+
+- `obs` Fx, Fy (CME) and all `MovingTables.DIAG` columns per step: CME,
+  CME + Wen and CME + cover − uncover reported separately; crossings, eq_only,
+  no_donor (no_donor must stay 0).
+- Velocity u1, u2 on lab-fixed lines x/D = -.6, 0, .6, 1.2 (relative to the
+  mean cylinder position) versus x2, at the scored phases. Lab frame
+  throughout; if any body-frame intermediate is used, position and velocity
+  are both transformed.
+- Sampling cost: raw force and diagnostics counters are read every step
+  (small). Flags and velocity fields are read only at the steps inside the
+  three prescribed 1° phase bins (φ ∈ [180, 181), [210, 211), [330, 331),
+  T/360 steps each) of the cycles needed for settling and scoring, plus
+  sparse declared audits; never every step.
+- Point set: fixed x2/D sample points per station, identical on every grid
+  and declared before running; per point the interpolation weights are
+  recorded; the scored set is the common subset valid on all compared grids
+  at that phase, with its coverage recorded; point sets never change silently.
+- Interpolation: bilinear on cell centres using only cells that are fluid at
+  that step (flags read back via `get_flags()`); weights renormalized over the
+  fluid corners; a sample with no fluid corner, or inside the circle at that
+  step, is MISSING and logged, never set to zero. One-degree bins keep valid
+  fluid weights. The numerical average need not mimic the 200 LDA samples.
+
+## 5. Periodic-regime gate (every case, before any statistic), ENGINEERING
+
+Ledger rows, same schema for every case: dx, dt (in D and T units); U_max, nu,
+Re, KC (lattice and nondimensional); T (native steps, nondimensional);
+duration (native steps, s/T, paper cycles); discard; accepted window; cycle
+ledger; verdict PASS or FAIL-CLOSED.
+
+Scales: `F_ref = 0.5 rho U_max² D`, `U_ref = U_max`; never a near-zero mean
+force as denominator.
+
+- Settling metrics, cycle k−1 → k, phase-aligned by step within the cycle
+  (error budget: one fifth of the tightest comparison band of the same
+  observable, so residual drift cannot obscure a domain or resolution pass):
+  - waveform `RMS(dFx) / F_ref ≤ .002`;
+  - `|d mean|`, `|d RMS|`, `|d peak|` of Fx, each `/ F_ref ≤ .002`;
+  - scored profiles (fluid-only) `|du| / U_max`, `|dv| / U_max ≤ .004` at the
+    three phases × four stations;
+  - Morison `|dc_d| ≤ .004`, `|dc_i| ≤ .004`.
+- Conservative comparison: every domain and cross-resolution criterion is
+  applied to `|difference| + drift_a + drift_b`, where drift is the largest
+  pair metric of that case inside the scored window.
+- These constants (protocol.py SETTLE, DOMAIN, CROSS) are frozen before any
+  physical result; no post-result loosening.
+- Rule: all metrics for N = 3 consecutive whole paper cycles → settled.
+
+Cycle indexing (worked example, binding). Paper cycle k spans
+`s/T in [0.25 + (k-1), 0.25 + k)`, i.e. φ0 → φ0; cycle 0 is the startup
+quarter `s/T < 0.25`. Pair p_k compares cycle k−1 with cycle k.
+
+- Discarded: cycles 0, 1, 2 (s/T < 2.25). Pairs are formed only between
+  cycles ≥ 3, so the first pair is p_4 (3 → 4).
+- Settling proof for a case: three consecutive pairs all pass; K_case = the
+  first cycle AFTER the last of those pairs. Earliest possible: pairs 3→4,
+  4→5, 5→6 pass ⇒ K_case = 7.
+- Common window: K = max over compared cases of K_case; every case scores the
+  same three complete cycles K, K+1, K+2 and must stay settled continuously:
+  every pair p_j for j = its own K_case .. K+2 must pass (including the wait
+  between its local settle and the common K), otherwise that case is
+  FAIL-CLOSED.
+- Cap: cycle k ends at s/T = k + 0.25, so the last admissible scored cycle is
+  11 (ends 11.25 T) and K ≤ 9. Examples: K = 7 → cycles 7–9 (ends 9.25 T);
+  K = 9 → cycles 9–11 (ends 11.25 T); any K_case > 9 → FAIL-CLOSED.
+- Run length: cases run sequentially; each runs to the end of cycle 11
+  (11.25 T: 28,350 / 56,700 / 113,400 steps for D = 20 / 40 / 80; the 12 T
+  cap is 30,240 / 60,480 / 120,960) and keeps its scalar and profile data, then
+  the common K is chosen. An unscorable partial cycle 12 is not run.
+- Profiles: an empty common valid point set is FAIL-CLOSED; coverage, missing
+  points, bin time weights (equal per step) and the original fluid corner
+  weight sums are reported.
+- The earlier G4 timing check (100 coarse vs 400 diffusive-fine steps) was a
+  matched short transient, not periodic evidence.
+
+## 6. Metrics and per-observable criteria (ENGINEERING, not paper uncertainty)
+
+Three criteria stay separate; the G4 same-grid 2% target is not reused.
+
+- Morison fit: least squares of the normalized per-step force over the
+  accepted window against the two basis functions, each averaged over its
+  single step interval, fixed paper phase, no free phase shift; c_d and
+  added-mass c_i per force candidate.
+- Cross-resolution: every scored observable's error against its reference must
+  decrease monotonically D20 → D40 → D80 (necessary, not sufficient).
+- Against Table 1 (computed fit, not measured force): at D ≥ 40,
+  `|c_d - 2.09| ≤ .15` and `|c_i - 1.45| ≤ .15`, and monotone.
+- Force waveform against its own Morison reconstruction: residual reported
+  only, no pass/fail (the paper reports fit error above its coarse numerical
+  error and gives no percentage).
+- Fig. 8 LDA profiles: QUALITATIVE ONLY. The original figure could not be
+  accessed or rendered; no digitized ordinates and no measurement uncertainty
+  exist. After a digitization with figure number, tool and uncertainty note:
+  `RMS(|du|, |dv|) / U_max ≤ .15` at D ≥ 40 and monotone; LDA error stays
+  UNQUANTIFIED and is stated with every comparison.
+- Fy sanity (not literature): `|mean Fy| / F_ref ≤ .05` in the accepted window.
+- No force candidate is selected for an RHS by these results; each is
+  reported.
+
+Superseded (earlier drafts): the velocity-jump start, guessed stations and
+phases, "C_M" for c_i, and the uncertainty wording; cycles 4–6, the 30-cycle
+cap and the earlier tolerance numbers.
+
+## 6b. Amendment: duration-only diagnostic (declared before the run)
+
+- Background: the D40 domain stage (20 D / 40 D / 60 D, 11.25 T each) failed
+  settling in every box (no window). That result stands.
+- One declared diagnostic: D = 40, 60 D box, from rest, through cycle 31
+  (157,500 steps = 31.25 T). Physics, settling tolerances, bands and
+  production unchanged; three proof pairs plus a continuously settled
+  K..K+2 with K ≤ 29.
+- Invocation: `run.py --D 40 --box 60 --cycles 31` then
+  `analyze.py --horizon 31`; the JSON records `horizon` and `horizon_mode`.
+  Without `--horizon` the analyzer is the original cap-11 scoring.
+- Failure at 31.25 T: report and stop; no automatic extension, matrix, D80
+  or VIV. A pass shows only that this box settles with time; it is not a
+  domain or literature pass.
+- The phase-locked high-frequency force component (residual of the mean
+  waveform after its Morison reconstruction, above 20/T) is reported per case
+  as unqualified; it needs separate qualification before any Morison or
+  literature statement.
+
+## 6a. Driver unit tests (CPU, before any run)
+
+- Phase/cycle arithmetic: s ↔ (k, φ), cycle boundaries, bin membership,
+  window/cap selection for K = 7, 8, 9, 10.
+- Solid-aware interpolation: all-fluid, partial-solid (renormalized weights),
+  no-fluid (MISSING) and inside-circle cases.
+- Interval-averaged Morison bases and the two-parameter fit on a synthetic
+  waveform with known c_d, c_i.
+
+## 7. Order and scope
+
+Protocol review → one-shot Dütsch driver (analytic `set_body` + `run(1)`,
+cycle ledger, lab-fixed fluid-only station sampler, two-basis Morison fit; no
+reusable runner until the first smoke shows shared helpers pay off; no
+production change) → smoke (memory/runtime) → domain sensitivity →
+ladder → physical-validation verdict → only then single-cylinder VIV. No AMR,
+no pinball, no manuscript work.

docs/validation/leontini_viv_predeclare.md

@@ -0,0 +1,146 @@
+# Single-cylinder free transverse VIV — PREDECLARATION (agreed at 1abfe954)
+
+Stage goal (Yanqi via TRON): prescribed Dütsch case + one VIV case both
+"reasonably close" to a predeclared reference = development-stage success, then
+cleanup. The frozen Dütsch settling failure (0/28 pairs at 31.25 T) stays on
+record. Practical closeness is not a formal validation claim. Agreed by
+Claude, Grok and Astra at 1abfe954; implemented in research/viv_re150/.
+
+## Primary case (one case)
+
+- Reference: Qin et al., arXiv:2003.12186, §4.5 (maximum amplitudes at
+  U_red = 3…8: .075, .551, .520, .459, .370, .084 D in the text, as quoted by
+  Grok/Astra) for the target; setup from Kolahdouz et al., arXiv:2003.12046,
+  §4.1.1 (read 2026-10-30 via ar5iv): U* sweep 3–8 with amplitudes only in
+  Fig. 3(b) (no text values), domain [−30, 45] D × [−30, 30] D with the
+  cylinder "initially at rest and centered at the origin", uniform inflow on
+  the left, zero normal traction and tangential velocity at the outlet, zero
+  normal velocity and tangential traction on y = ±30 D; no initial
+  perturbation and no total simulated time stated.
+- Re = 150, U* = U/(f_n D) = 5, m/(rho D²) = 2 (m* = rho_s/rho_f = 8/pi),
+  zeta = 0, f_n = sqrt(K/M)/(2 pi) in vacuum. Transverse DOF only.
+- Target: A_max = .520 D (peak amplitude, Qin §4.5). Multi-code reference
+  (Kolahdouz ILE-NS, Qin II-LBM). Reference uncertainty: not stated; the
+  spread between codes is not quantified here.
+
+## Our reproduction (declared differences)
+
+- D = 40, MRT, dense double buffer; U_lat = .04, nu = U D / 150 = .010667,
+  tau = .532.
+- Box [0, 75 D] × [0, 60 D] (3000 × 2400) with the cylinder centre at
+  (30 D, 30 D): 30 D from the inlet, 45 D from the outlet, centred in y
+  (Kolahdouz §4.1.1). Qin's own box ([0, 40 D] × [0, 20 D], centre (10 D, 10 D))
+  is the target paper's, not ours.
+- Boundaries: uniform Zou–He inlet, NEQ outlet (rho = 1, backflow clamp),
+  free-slip y walls. Kolahdouz uses a traction outlet: declared mismatch.
+- Qin's own VIV box (40 D × 20 D) is not used; a smaller box would be a
+  declared cost cut with more blockage, not "the same case".
+- Allowed y offset ±1.5 D; peak transverse speed ≈ .026 cells/step (bound .5).
+
+## Structure, coupling and time levels
+
+- M = 2 rho D² (lattice rho = 1), f_n = U/(5 D) = 2e-4 per step, K = M (2 pi f_n)²,
+  zeta = 0. Equation of motion uses the structural mass only.
+- Per-step sequence (binding): known (y_n, v_n) and the last completed force
+  F_last → candidate (y_{n+1}, v_{n+1}) by the update below → `set_body` with
+  offset y_{n+1} and wall velocity (y_{n+1} − y_n)/dt with lattice dt = 1 (interval average; a
+  different quantity from the endpoint v_{n+1}) → `run(1)` → only on success:
+  commit y_n ← y_{n+1}, v_n ← v_{n+1}, F_last ← force read for that step (failure
+  semantics: see the window section). Whether the
+  measured force represents the interval's start, midpoint or end is not
+  asserted; the one-step lag is declared, not corrected.
+- Initial state: fluid initialized by `Simulation.initialize()` to the uniform
+  inlet velocity U everywhere (the code's actual initial condition; the
+  source gives no fluid initial condition), body at rest at the declared
+  centre, y_0 = v_0 = 0, F_last = 0.
+- Structural update over one step with constant force F = F_last (exact
+  harmonic solution, zeta = 0): y_e = F/K, w = 2 pi f_n,
+  y^{n+1} = y_e + (y^n − y_e) cos w + (v^n / w) sin w,
+  v^{n+1} = −(y^n − y_e) w sin w + v^n cos w.
+- Primary RHS (single coupled run): CME + Wen (theory-motivated stage choice:
+  Galilean-invariant link correction; G4 could not separate it from CME).
+  Raw CME and the node terms are logged as diagnostics of that trajectory;
+  the node impulse is not added. A run with CME-only as RHS is a separate,
+  labelled sensitivity case, decided separately. Neither RHS is claimed proven.
+
+## Window, failure semantics, acceptance (engineering, development stage)
+
+Clocks: lattice dt = 1; T_n = 1/f_n = 5,000 steps. Response cycles are
+counted from the measured signal; T_n is used only for the seed check.
+
+- Start: body at rest at the declared centre (the only initial condition the
+  source states), y_0 = v_0 = 0, F_last = 0; fluid uniform at U (as
+  initialized by the code; our choice, not the source's).
+- Seed (conditional, one time, declared): at step n_seed = 200,000
+  (t = 40 T_n), if max |y_t| over states t in [175,000, 200,000) is below
+  .01 D, then
+  in that step, after the check and before the candidate update, set
+  v_n += 2 pi f_n (.01 D) (F_last unchanged) and proceed normally; the
+  geometry moves continuously through the update. `seed_applied`, step and
+  size are logged; never repeated or changed.
+- Cycle boundary (discrete, fixed threshold): step n closes a cycle iff
+  y[n−1] < 0 and y[n] ≥ 0 (upward crossing of the equilibrium y = 0);
+  y_0 = 0 is not a crossing; no boundary fewer than 2 samples after the
+  previous one.
+- Per cycle: positive peak max(y), negative peak min(y), mean offset
+  (time-mean of y), amplitude a = (max − min)/2 (the only amplitude used for
+  settling).
+- Eligibility: a completed cycle is counted (toward the settling proof and
+  the 150 budget) only if a ≥ .01 D; ineligible cycles are logged, never
+  counted, so near-zero motion cannot settle or pass.
+- Seed reset: if the seed is applied, the incomplete cycle at n_seed is
+  discarded and the settling proof starts empty; the 150 budget and the
+  native cap keep running from t = 0.
+- Windows: adjacent non-overlapping blocks of 10 counted cycles;
+  W_k = mean amplitude of block k (W_k ≥ .01 D by eligibility, so no zero
+  denominator).
+- Block rule R(k), k ≥ 2: |W_k − W_{k−1}| / W_{k−1} ≤ .02 and, for the last 5
+  cycles of block k, max_i |a_i − ā_5| / ā_5 ≤ .01 with ā_5 their mean.
+- Settled at the end of block k when R(k) holds. Scoring = the next two
+  blocks k+1, k+2 (20 counted cycles); both must satisfy R. If one fails,
+  that block becomes the new reference: the settling proof restarts from it
+  and continues within the caps (deterministic; no other retry). A pass
+  requires one complete stable 20-cycle scoring window before a cap; success
+  stops after its 20th cycle.
+- Caps (both always armed; the first hit wins and is reported): 150 counted
+  response cycles including the 20 scoring cycles (so settling must complete
+  by counted cycle 130) → `timeout_response_cycles`; 1,000,000 native steps
+  (200 T_n) → `timeout_native_steps`.
+- Durations reported: total steps, t U/D, counted and uncounted cycles,
+  measured T_resp (mean boundary interval); N_resp / f_n only as a labelled
+  secondary clock.
+- Growth is not a failure: initial VIV growth is physical. Checked every
+  step (FAIL-CLOSED): non-finite y, v or force; |y| > 1.0 D (the tighter
+  gate; about twice the largest reference peak, .551 D) or the ±1.5 D offset
+  box reached; no_donor > 0; a failed upload/step/sync.
+- On any failure the driver stops and reports; it never retries,
+  sub-iterates, retunes the RHS or parameters, or starts a second case.
+  Device-side failures (upload/step/sync, no_donor) mark the Simulation
+  failed through the existing lifecycle; input rejection by `set_body`
+  (offset box, speed, step bound) and timeouts end the driver run without
+  that mark.
+- Target metric over the 20 scoring cycles: A_max/D = max over those cycles
+  of max(|positive peak|, |negative peak|) about y = 0 (Qin reports "maximum
+  amplitude"); also reported: mean amplitude, positive and negative peaks,
+  mean offset, f/f_n (no target).
+- Acceptance: |A_max − .520 D| ≤ .052 D (±10 %), stated together with the
+  boundary mismatch and the absence of a quoted reference uncertainty. A pass
+  is "reasonably close" for the stage: not grid convergence, not force-RHS
+  validation, and it does not touch the Dütsch settling record.
+- Driver reads force and the small diagnostics per step only; field reads
+  only at a few predeclared snapshots. No runtime claim: there is no measured
+  timing on this grid yet (the D40/60D receipt is a different, sampled case).
+
+## Plumbing smoke (not the benchmark)
+
+`run.py --smoke --D 20 --steps 32` runs exactly 32 steps of the same BC
+package and coupling loop (no seed, settle or score; never scored);
+status `smoke_ok` iff every step is finite, no_donor stays 0 and the body
+stays inside its box. The physical case is only the default invocation
+(`--D 40`, no `--smoke`).
+
+## Queued, not in this stage
+
+Leontini et al. JFS 22 (2006) Re = 200, m* = 10, zeta = 0.01, U* = 5.0 (needs
+the owners' c_crit convention and a Fig. 1 digitization). D80, literature
+claims, pinball, AMR, BC redesign, threshold changes.

research/duetsch_re100_kc5/analyze.py

@@ -0,0 +1,224 @@
+"""Ledger, settling, common window and Morison fits for Dütsch cases (numpy only).
+
+    python analyze.py CASE_DIR [CASE_DIR ...]   # cases of one comparison
+"""
+from __future__ import annotations
+
+import json
+import math
+import os
+import sys
+
+import numpy as np
+
+import protocol as p
+
+
+def load(path):
+    with open(os.path.join(path, "case.json")) as fh:
+        case = json.load(fh)
+    raw = np.load(os.path.join(path, "raw.npz"))
+    case.update(force=raw["force"], diag=raw["diag"], samples=raw["samples"], path=path)
+    return case
+
+
+def cycle_slices(case):
+    """{k: slice of step indices m whose interval [m, m+1] lies in paper cycle k}."""
+    T, q = case["T"], case["T"] // 4
+    return {k: slice(q + (k - 1) * T, q + k * T)
+            for k in range(1, case["cycles"] + 1) if q + k * T <= len(case["force"])}
+
+
+def coarse_wave(case, k, r=1):
+    """Normalized per-step Fx of cycle k, averaged over r native steps per sample."""
+    f = case["force"][cycle_slices(case)[k], 0] / case["F_ref"]
+    return f.reshape(-1, r).mean(axis=1)
+
+
+def profiles(case, k):
+    """{(phase, station, x2): (u, v)} of cycle k, time-averaged over the bin steps
+    (equal weights per step) from valid samples only."""
+    s = case["samples"]
+    acc = {}
+    for _, _, ph, xs, x2, u, v, _ in s[s[:, 1] == k]:
+        if not (math.isnan(u) or math.isnan(v)):
+            acc.setdefault((ph, xs, x2), []).append((u, v))
+    return {key: tuple(np.mean(vals, axis=0)) for key, vals in acc.items()}
+
+
+def coverage(case, k):
+    s = case["samples"]
+    rows = s[s[:, 1] == k]
+    missing = int(np.isnan(rows[:, 5]).sum())
+    return dict(samples=len(rows), missing=missing, bin_steps=len(set(rows[:, 0])),
+                min_fluid_weight=float(np.nanmin(np.where(rows[:, 7] > 0, rows[:, 7], np.nan)))
+                if len(rows) else math.nan)
+
+
+def common_points(dicts):
+    keys = set(dicts[0])
+    for d in dicts[1:]:
+        keys &= set(d)
+    return sorted(keys)
+
+
+def candidates(case):
+    """Per-step Fx of each reported force candidate (none is selected as an RHS)."""
+    d = case["diag"]
+    return {"cme": d[:, 0], "cme_wen": d[:, 0] + d[:, 2], "cme_node": d[:, 0] + d[:, 6] - d[:, 9]}
+
+
+def cycle_morison(case, ks, fx=None):
+    T = case["T"]
+    fx = case["force"][:, 0] if fx is None else fx
+    idx = np.concatenate([np.arange(len(fx))[cycle_slices(case)[k]] for k in ks])
+    b = np.array([p.interval_bases(int(m), T) for m in idx])
+    return p.morison_fit(fx[idx] / case["F_ref"], b[:, 0], b[:, 1])
+
+
+def pair_metrics(case, k):
+    a, b = coarse_wave(case, k - 1), coarse_wave(case, k)
+    pa, pb = profiles(case, k - 1), profiles(case, k)
+    common = common_points([pa, pb])  # empty set -> inf -> pair fails (FAIL-CLOSED)
+    prof = max((max(abs(pa[c][0] - pb[c][0]), abs(pa[c][1] - pb[c][1])) for c in common),
+               default=math.inf) / case["u_max"]
+    (cd0, ci0), (cd1, ci1) = cycle_morison(case, [k - 1]), cycle_morison(case, [k])
+    m = dict(wave=float(np.sqrt(np.mean((b - a) ** 2))),
+             scalar=float(max(abs(b.mean() - a.mean()), abs(np.sqrt(np.mean(b * b)) - np.sqrt(np.mean(a * a))),
+                              abs(np.abs(b).max() - np.abs(a).max()))),
+             profile=float(prof), morison=float(max(abs(cd1 - cd0), abs(ci1 - ci0))),
+             profile_points=len(common))
+    m["pass"] = all(m[key] <= tol for key, tol in p.SETTLE.items())
+    return m
+
+
+def hf_component(case, ks):
+    """Unqualified phase-locked high-frequency force component (reported, not checked):
+    mean CME waveform over cycles ks minus its two-basis Morison reconstruction,
+    split at 20/T; cycle differencing does not remove it."""
+    T = case["T"]
+    if not ks:
+        return dict(label="unqualified phase-locked HF force component, not checked", cycles=[])
+    mean = np.mean([coarse_wave(case, k) for k in ks], axis=0)
+    c_d, c_i = cycle_morison(case, ks)
+    q = T // 4
+    b = np.array([p.interval_bases(q + m, T) for m in range(T)])
+    resid = mean - (c_d * b[:, 0] - (np.pi ** 2 * c_i / p.KC) * b[:, 1])
+    spec = np.fft.rfft(resid)
+    high = spec.copy()
+    high[:21] = 0.0
+    energy = one_sided_energy(spec, T)
+    return dict(label="unqualified phase-locked HF force component, not checked",
+                cycles=list(ks), hf_rms_over_F_ref=float(np.sqrt(np.mean(np.fft.irfft(high, T) ** 2))),
+                hf_energy_fraction=float(energy[21:].sum() / max(energy.sum(), 1e-300)),
+                top_harmonics=[int(j) for j in np.argsort(energy[21:])[::-1][:5] + 21])
+
+
+def one_sided_energy(spec, n):
+    """Parseval-consistent energy per rfft bin: interior bins count twice, DC and
+    (for even n) Nyquist once."""
+    w = np.full(len(spec), 2.0)
+    w[0] = 1.0
+    if n % 2 == 0:
+        w[-1] = 1.0
+    return w * np.abs(spec) ** 2
+
+
+def ledger(case, horizon=p.LAST_SCORED):
+    ks = [k for k in sorted(cycle_slices(case)) if k <= horizon]
+    pairs = {k: pair_metrics(case, k) for k in ks if k - 1 > p.DISCARD}
+    return pairs, p.settled_start({k: v["pass"] for k, v in pairs.items()})
+
+
+def main(argv):
+    """argv: CASE_DIR ... [--reference CASE_DIR]; the reference defaults to the
+    largest D (resolution ladder); domain runs name the 60 D case explicitly."""
+    ref_path = None
+    if "--reference" in argv:
+        i = argv.index("--reference")
+        ref_path, argv = argv[i + 1], argv[:i] + argv[i + 2:]
+    horizon, explicit = p.LAST_SCORED, "--horizon" in argv  # default: original scoring
+    if explicit:
+        i = argv.index("--horizon")
+        horizon, argv = int(argv[i + 1]), argv[:i] + argv[i + 2:]
+    paths = argv if ref_path is None or ref_path in argv else argv + [ref_path]
+    cases = [load(path) for path in paths]
+    for c in cases if explicit else ():  # complete cycles 1..N; never pad or shift
+        if c["cycles"] < horizon or len(c["force"]) < round((0.25 + horizon) * c["T"]):
+            raise SystemExit(f"{c['path']}: run does not contain complete cycles 1..{horizon}")
+    ledgers = [ledger(c, horizon) for c in cases]
+    window = p.common_window([k for _, k in ledgers],
+                             [{k: v["pass"] for k, v in pairs.items()} for pairs, _ in ledgers],
+                             last_scored=horizon)
+    report = dict(window=window, horizon=horizon, cases=[],
+                  horizon_mode="default-cap" if horizon == p.LAST_SCORED else
+                  "extended duration diagnostic (declared amendment, D40/60D only)")
+    for case, (pairs, k_case) in zip(cases, ledgers):
+        limit = round((0.25 + horizon) * case["T"])  # events inside the horizon only
+        entry = {key: case[key] for key in ("path", "D", "T", "box", "n", "u_max", "nu", "tau", "A",
+                                            "steps", "cycles", "F_ref")}
+        entry.update(duration_T=case["steps"] / case["T"], K_case=k_case, pairs=pairs,
+                     coverage={k: coverage(case, k) for k in sorted(cycle_slices(case))
+                               if p.DISCARD < k <= horizon},
+                     no_donor=float(case["diag"][:limit, 13].sum()),
+                     eq_only=float(case["diag"][:limit, 12].sum()),
+                     hf_component=hf_component(case, window or tuple(
+                         [k for k in sorted(cycle_slices(case)) if p.DISCARD < k <= horizon][-3:])))
+        if window:  # residual drift inside the scored window (pairs K..K+2)
+            entry["window_drift"] = {key: max(pairs[k][key] for k in window)
+                                     for key in ("wave", "profile", "morison")}
+            entry["morison_obs"] = cycle_morison(case, window)
+            entry["morison"] = {name: cycle_morison(case, window, fx)
+                                for name, fx in candidates(case).items()}
+            entry["fy_mean"] = float(np.mean(np.concatenate(
+                [case["force"][cycle_slices(case)[k], 1] for k in window])) / case["F_ref"])
+        report["cases"].append(entry)
+    if window and len(cases) > 1:  # cross-resolution waveform against the finest case
+        fine = (next(c for c in cases if c["path"] == ref_path) if ref_path
+                else max(cases, key=lambda c: c["D"]))
+        report["reference"] = fine["path"]
+        ref = np.concatenate([coarse_wave(fine, k) for k in window])
+        for case, entry in zip(cases, report["cases"]):
+            r = fine["T"] // case["T"]  # 1 for equal-D domain runs
+            ref_c = ref.reshape(-1, r).mean(axis=1)
+            own = np.concatenate([coarse_wave(case, k) for k in window])
+            entry["wave_rms_vs_finest"] = float(np.sqrt(np.mean((own - ref_c) ** 2)))
+        for k in window:  # profiles on the common valid point set of all compared cases
+            per_case = [profiles(c, k) for c in cases]
+            keys = common_points(per_case)
+            report.setdefault("common_profile_points", {})[k] = len(keys)
+            if not keys:
+                report["verdict"] = "FAIL-CLOSED: empty common profile set"
+                continue
+            fine_p = per_case[cases.index(fine)]
+            for case, entry, prof in zip(cases, report["cases"], per_case):
+                d = max(max(abs(prof[key][0] - fine_p[key][0]), abs(prof[key][1] - fine_p[key][1]))
+                        for key in keys) / case["u_max"]
+                entry.setdefault("profile_max_vs_finest", {})[k] = float(d)
+        ref_entry = report["cases"][cases.index(fine)]
+        for entry in report["cases"]:  # conservative: difference + both drifts
+            if entry is ref_entry:
+                continue
+            drift = {key: entry["window_drift"][key] + ref_entry["window_drift"][key]
+                     for key in ("wave", "profile", "morison")}
+            dm = max(abs(entry["morison"]["cme"][0] - ref_entry["morison"]["cme"][0]),
+                     abs(entry["morison"]["cme"][1] - ref_entry["morison"]["cme"][1]))
+            entry["conservative"] = dict(
+                wave=entry["wave_rms_vs_finest"] + drift["wave"],
+                profile=max(entry.get("profile_max_vs_finest", {}).values(), default=math.inf)
+                + drift["profile"],
+                morison_cme=dm + drift["morison"])
+            bands = p.DOMAIN if ref_path else p.CROSS  # domain runs name the 60 D reference
+            c = entry["conservative"]
+            entry["pass_scope"] = (
+                ("wave+profile+morison_cme" if "morison" in bands else "wave+profile")
+                + ": conservative (|diff| + both window drifts) on the obs/CME Fx waveform and "
+                "simulation-to-simulation profiles; CME+Wen and CME+node reported, not checked; "
+                "no Fig. 8 / literature verdict")
+            entry["pass"] = (c["wave"] <= bands["wave"] and c["profile"] <= bands["profile"]
+                             and ("morison" not in bands or c["morison_cme"] <= bands["morison"]))
+    print(json.dumps(report, indent=1, default=float))
+
+
+if __name__ == "__main__":
+    main(sys.argv[1:])

research/duetsch_re100_kc5/capture.py

@@ -0,0 +1,52 @@
+"""Research-only field capture helpers for the Dütsch driver (numpy only).
+
+Snapshots are taken at the first step of the scored phase bins in declared
+cycles: a near-body crop at full resolution and a fluid-weighted block-mean
+downsample of the whole domain; plus scalar trends at every bin read.
+"""
+from __future__ import annotations
+
+import numpy as np
+
+CROP_X, CROP_Y, DOWN = 8.0, 6.0, 8      # crop half-widths in D; downsample factor
+STRIP = 2.0                             # inlet/outlet strip width in D
+
+
+def crop_box(xm, ym, D, n, body_extent):
+    """Inclusive integer crop (x0, x1, y0, y1); fail closed if it leaves the domain
+    interior or does not contain the body's swept extent (x_lo, x_hi, y_lo, y_hi)."""
+    x0, x1 = int(round(xm - CROP_X * D)), int(round(xm + CROP_X * D))
+    y0, y1 = int(round(ym - CROP_Y * D)), int(round(ym + CROP_Y * D))
+    if x0 < 1 or y0 < 1 or x1 > n - 2 or y1 > n - 2:
+        raise ValueError("capture crop leaves the domain interior")
+    bx0, bx1, by0, by1 = body_extent
+    if not (x0 <= bx0 and bx1 <= x1 and y0 <= by0 and by1 <= y1):
+        raise ValueError("capture crop clips the body's swept extent")
+    return x0, x1, y0, y1
+
+
+def downsample(field, fluid, f=DOWN):
+    """Fluid-weighted f x f block means; NaN where a block has no fluid cell."""
+    ny, nx = field.shape
+    h, w = ny // f, nx // f
+    v = np.where(fluid, field, 0.0)[:h * f, :w * f].reshape(h, f, w, f).sum(axis=(1, 3))
+    c = fluid[:h * f, :w * f].reshape(h, f, w, f).sum(axis=(1, 3))
+    with np.errstate(invalid="ignore", divide="ignore"):
+        return np.where(c > 0, v / np.maximum(c, 1), np.nan).astype(np.float32)
+
+
+def fluid_mean(field, fluid, xs=slice(None), ys=slice(None)):
+    m = fluid[ys, xs]
+    return float(field[ys, xs][m].mean()) if m.any() else float("nan")
+
+
+def trend_row(s, k, ph, rho, ux, uy, fluid, D, n, crop):
+    """(s, cycle, phase, <rho>, <ux>, <uy>, <rho>_crop, inlet <rho>, <ux>, outlet <rho>, <ux>)."""
+    w = max(1, int(round(STRIP * D)))
+    x0, x1, y0, y1 = crop
+    cx, cy = slice(x0, x1 + 1), slice(y0, y1 + 1)
+    inlet, outlet, inner = slice(1, 1 + w), slice(n - 1 - w, n - 1), slice(1, n - 1)
+    return (s, k, ph, fluid_mean(rho, fluid), fluid_mean(ux, fluid), fluid_mean(uy, fluid),
+            fluid_mean(rho, fluid, cx, cy),
+            fluid_mean(rho, fluid, inlet, inner), fluid_mean(ux, fluid, inlet, inner),
+            fluid_mean(rho, fluid, outlet, inner), fluid_mean(ux, fluid, outlet, inner))

research/duetsch_re100_kc5/protocol.py

@@ -0,0 +1,129 @@
+"""Pure arithmetic of the Dütsch Re 100 / KC 5 protocol (numpy only, no CUDA).
+
+docs/validation/duetsch_re100_kc5_protocol.md is the contract: elapsed time s
+in native steps from the right turning point, paper phase
+phi = (360 s / T - 90) mod 360, paper cycle k >= 1 spans
+s/T in [0.25 + (k-1), 0.25 + k); cycle 0 is the startup quarter.
+"""
+from __future__ import annotations
+
+import math
+
+import numpy as np
+
+KC, RE = 5.0, 100.0
+STATIONS = (-0.6, 0.0, 0.6, 1.2)          # lab-fixed x/D from the mean position
+PHASES = (180.0, 210.0, 330.0)            # degrees; bins [p, p + 1)
+X2 = tuple(np.round(np.arange(-1.5, 1.5001, 0.05), 10))  # declared x2/D points
+DISCARD, LAST_SCORED, N_PAIRS = 2, 11, 3   # cycles 0..2 discarded, cap 12 T
+# A-priori bands. Settling is one fifth of the tightest comparison band of the
+# same observable, and comparisons add both cases' residual window drift
+# (conservative: |difference| + drift_a + drift_b <= band).
+DOMAIN = dict(wave=.01, morison=.02, profile=.02)  # vs the 60 D reference at D = 40
+CROSS = dict(wave=.02, profile=.02)                # D40 vs D80
+SETTLE = dict(wave=.002, scalar=.002, profile=.004, morison=.004)
+
+
+def amplitude(D):
+    return KC * D / (2.0 * math.pi)
+
+
+def position(s, T, A):
+    """Centre position relative to the mean: x = A cos(2 pi s / T), evaluated from the
+    native step modulo T so every cycle repeats the same geometry bit for bit."""
+    return A * math.cos(2.0 * math.pi * (int(s) % T) / T)
+
+
+def phase(s, T):
+    """Paper phase in degrees from the native step modulo T (the unwrapped cycle
+    index is kept separately by ``cycle``)."""
+    return (360.0 * (int(s) % T) / T - 90.0) % 360.0
+
+
+def cycle(s, T):
+    return 0 if s < 0.25 * T else int(math.floor(s / T - 0.25)) + 1
+
+
+def bin_phase(s, T):
+    """Scored phase whose 1-degree bin [p, p + 1) contains time s, else None."""
+    phi = phase(s, T)
+    for p in PHASES:
+        if p <= phi < p + 1.0:
+            return p
+    return None
+
+
+def interval_bases(n, T, points=16):
+    """Single-step interval averages over [n, n + 1] of cos phi |cos phi| and sin phi
+    (phi in radians = 2 pi s / T - pi / 2), Gauss-Legendre quadrature."""
+    x, w = np.polynomial.legendre.leggauss(points)
+    s = int(n) % T + 0.5 * (x + 1.0)  # periodic: identical bases in every cycle
+    phi = 2.0 * np.pi * s / T - 0.5 * np.pi
+    c = np.cos(phi)
+    return 0.5 * float(np.sum(w * c * np.abs(c))), 0.5 * float(np.sum(w * np.sin(phi)))
+
+
+def morison_fit(f_norm, b1, b2):
+    """Least squares of f_norm = c_d b1 - (pi^2 c_i / KC) b2 (paper phase, no shift)."""
+    a = np.column_stack([b1, -(math.pi ** 2 / KC) * np.asarray(b2)])
+    (c_d, c_i), *_ = np.linalg.lstsq(a, np.asarray(f_norm), rcond=None)
+    return float(c_d), float(c_i)
+
+
+def inside_circle(x, y, centre, radius):
+    return (x - centre[0]) ** 2 + (y - centre[1]) ** 2 < radius * radius
+
+
+def sample(field, fluid, x, y, centre, radius):
+    """Profile sample: MISSING if the exact point is inside the body (checked before
+    any corner interpolation), else ``bilinear_fluid``."""
+    if inside_circle(x, y, centre, radius):
+        return None, (), 0.0
+    return bilinear_fluid(field, fluid, x, y)
+
+
+def bilinear_fluid(field, fluid, x, y):
+    """Bilinear value at (x, y) in cell coordinates from fluid cells only.
+
+    Returns (value, ((ix, iy, weight), ...), fluid_weight_sum) with weights
+    renormalized over the fluid corners and the original sum of their bilinear
+    weights, or (None, (), 0.0) when no corner is fluid (MISSING, never zero).
+    """
+    x0, y0 = int(math.floor(x)), int(math.floor(y))
+    fx, fy = x - x0, y - y0
+    corners = ((x0, y0, (1 - fx) * (1 - fy)), (x0 + 1, y0, fx * (1 - fy)),
+               (x0, y0 + 1, (1 - fx) * fy), (x0 + 1, y0 + 1, fx * fy))
+    used = [(i, j, w) for i, j, w in corners if w > 0.0 and fluid[j, i]]
+    total = sum(w for _, _, w in used)
+    if total == 0.0:
+        return None, (), 0.0
+    used = tuple((i, j, w / total) for i, j, w in used)
+    return float(sum(w * field[j, i] for i, j, w in used)), used, float(total)
+
+
+def settled_start(pairs_pass):
+    """K_case from {k: pass of pair (k-1 -> k)}: first cycle after three consecutive
+    passing pairs among cycles > DISCARD, or None."""
+    run = 0
+    for k in sorted(pairs_pass):
+        if k - 1 <= DISCARD:
+            continue
+        run = run + 1 if pairs_pass[k] else 0
+        if run == N_PAIRS:
+            return k + 1
+    return None
+
+
+def common_window(k_cases, pairs_pass_by_case, last_scored=LAST_SCORED):
+    """Common scored cycles K..K+2 with K = max K_case, or None (FAIL-CLOSED) when a
+    case never settled, K + 2 exceeds the cap, or a case fails any pair p_j for
+    j = its own K_case .. K + 2 (continuous settling, including the wait for K)."""
+    if any(k is None for k in k_cases):
+        return None
+    K = max(k_cases)
+    if K + 2 > last_scored:
+        return None
+    for k_case, pairs in zip(k_cases, pairs_pass_by_case):
+        if not all(pairs.get(j, False) for j in range(k_case, K + 3)):
+            return None
+    return K, K + 1, K + 2

research/duetsch_re100_kc5/run.py

@@ -0,0 +1,133 @@
+"""One-shot Dütsch Re 100 / KC 5 prescribed in-line oscillation driver.
+
+Protocol: docs/validation/duetsch_re100_kc5_protocol.md. Native dense channel
+(zero-velocity inlet, pressure-NEQ outlet, bounce-back y walls), MRT, lab frame.
+Writes raw per-step force/diagnostics and phase-bin profile samples to --out
+(outside the checkout); analysis is in analyze.py. No production change.
+
+    python run.py --D 40 --box 40 --cycles 11 --out /path/outside/checkout
+"""
+from __future__ import annotations
+
+import argparse
+import json
+import math
+import os
+import tempfile
+
+import numpy as np
+
+import protocol as p
+
+
+TEMPLATE = os.path.join(os.path.dirname(os.path.abspath(__file__)), "..", "..", "src",
+                        "CelerisLab", "configs", "config_lbm.json")
+
+
+def config(D, box, T):
+    """Shipped config_lbm.json (full parser schema) with only the case fields
+    overridden: grid, still fluid (velocity 0), nu from Re, MRT, uniform inlet
+    profile, bounce-back y walls. Inlet scheme and outlet mode stay the
+    template's (zou_he_local, neq_extrap) and are recorded in case.json."""
+    u = p.KC * D / T
+    n = int(round(box * D))
+    with open(TEMPLATE) as fh:
+        cfg = json.load(fh)
+    cfg["grid"].update(lattice_model="D2Q9", nx=n, ny=n, nz=1)
+    cfg["physics"].update(data_type="FP32", viscosity=u * D / p.RE, velocity=0.0, rho=1.0)
+    cfg["method"].update(collision="MRT", streaming="double_buffer", store_precision="FP32",
+                         ddf_shifting=False, y_wall_bc="bounce_back")
+    cfg["method"]["les"]["enabled"] = False
+    cfg["method"]["inlet"]["profile"] = "uniform"
+    return cfg, u, n
+
+
+def main():
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--D", type=int, required=True)
+    ap.add_argument("--box", type=float, default=40.0, help="box side in D")
+    ap.add_argument("--T", type=int, default=None, help="period in steps (multiple of 360)")
+    ap.add_argument("--cycles", type=int, default=p.LAST_SCORED, help="run to the end of this cycle")
+    ap.add_argument("--out", required=True)
+    ap.add_argument("--dump-fields", default="",
+                    help="research capture: comma-separated paper cycles, e.g. 11,21,30,31")
+    args = ap.parse_args()
+    D, T = args.D, args.T or 126 * args.D
+    if T % 360:
+        raise SystemExit("T must be a multiple of 360 steps")
+    if os.path.realpath(args.out).startswith(os.path.realpath(os.path.join(os.path.dirname(__file__), "..", ".."))):
+        raise SystemExit("--out must be outside the checkout")
+    cfg, u_max, n = config(D, args.box, T)
+    A, R = p.amplitude(D), 0.5 * D
+    xm = ym = 0.5 * (n - 1)
+    os.makedirs(args.out, exist_ok=True)
+
+    from CelerisLab.body.motion import Translation
+    from CelerisLab.simulation import Simulation
+    with tempfile.NamedTemporaryFile("w", suffix=".json", delete=False) as fh:
+        json.dump(cfg, fh)
+    sim = Simulation(fh.name)
+    try:
+        body = sim.add_body("circle", (xm + A, ym), R,
+                            translation=Translation((-2.0 * A, 0.0, 0.0, 0.0), 1.01 * u_max))
+        sim.initialize()
+        dump = {int(c) for c in args.dump_fields.split(",") if c}
+        if dump:
+            import capture
+            if max(dump) > args.cycles or min(dump) <= p.DISCARD:
+                raise SystemExit("--dump-fields cycles must lie in the scored range of this run")
+            crop = capture.crop_box(xm, ym, D, n, (xm - A - R, xm + A + R, ym - R, ym + R))
+            trend = []
+        steps = int(round((0.25 + args.cycles) * T))
+        force = np.zeros((steps, 2)); diag = []; samples = []
+        for m in range(steps):
+            x0, x1 = p.position(m, T, A), p.position(m + 1, T, A)
+            sim.set_body(body, offset=(max(x1 - A, -2.0 * A), 0.0), vx=x1 - x0)
+            sim.run(1)
+            force[m] = sim.read_force(body)[:2]
+            diag.append(sim.moving_diagnostics()[0])
+            s = m + 1  # state after the step
+            ph = p.bin_phase(s, T)
+            if ph is not None and p.cycle(s, T) > p.DISCARD:
+                fields = sim.get_macroscopic()
+                fluid = (sim.get_flags().reshape(n, n) & 1) != 0
+                ux, uy = (np.asarray(fields[k]).reshape(n, n) for k in ("ux", "uy"))
+                if dump:  # research capture only; the default path is unchanged
+                    rho = np.asarray(fields["rho"]).reshape(n, n)
+                    trend.append(capture.trend_row(s, p.cycle(s, T), ph, rho, ux, uy, fluid, D, n, crop))
+                    if p.cycle(s, T) in dump and p.phase(s, T) == ph:  # first step of the bin
+                        x0c, x1c, y0c, y1c = crop
+                        cy_, cx_ = slice(y0c, y1c + 1), slice(x0c, x1c + 1)
+                        np.savez_compressed(
+                            os.path.join(args.out, f"fields_c{p.cycle(s, T)}_p{int(ph)}.npz"),
+                            s=s, cycle=p.cycle(s, T), phase=ph, crop=np.asarray(crop),
+                            force=force[m], centre=np.asarray((xm + p.position(s, T, A), ym)),
+                            ux=ux[cy_, cx_].astype(np.float32), uy=uy[cy_, cx_].astype(np.float32),
+                            rho=rho[cy_, cx_].astype(np.float32), fluid=fluid[cy_, cx_].astype(np.uint8),
+                            down_ux=capture.downsample(ux, fluid), down_uy=capture.downsample(uy, fluid),
+                            down_rho=capture.downsample(rho, fluid))
+                centre = (xm + p.position(s, T, A), ym)  # analytic centre after the step
+                for xs in p.STATIONS:
+                    for x2 in p.X2:
+                        x, y = xm + xs * D, ym + x2 * D
+                        u, wu, wsum = p.sample(ux, fluid, x, y, centre, R)
+                        v, _, _ = p.sample(uy, fluid, x, y, centre, R)
+                        samples.append((s, p.cycle(s, T), ph, xs, x2,
+                                        math.nan if u is None else u, math.nan if v is None else v,
+                                        wsum, json.dumps(wu)))
+        if dump:
+            np.save(os.path.join(args.out, "trend.npy"), np.asarray(trend, float))
+        np.savez(os.path.join(args.out, "raw.npz"), force=force, diag=np.asarray(diag),
+                 samples=np.asarray([r[:8] for r in samples], float))
+        with open(os.path.join(args.out, "case.json"), "w") as fh:
+            json.dump(dict(D=D, T=T, box=args.box, n=n, u_max=u_max, nu=cfg["physics"]["viscosity"],
+                           tau=3 * cfg["physics"]["viscosity"] + 0.5, A=A, steps=steps,
+                           cycles=args.cycles, F_ref=0.5 * u_max ** 2 * D,
+                           inlet=cfg["method"]["inlet"], outlet=cfg["method"]["outlet"],
+                           weights=[r[8] for r in samples]), fh)
+    finally:
+        sim.close()
+
+
+if __name__ == "__main__":
+    main()

research/viv_re150/run.py

@@ -0,0 +1,134 @@
+"""One-shot predeclared VIV run (docs/validation/leontini_viv_predeclare.md).
+
+Re 150, U* = 5, m/(rho D^2) = 2, zeta = 0, D = 40 on a 75 D x 60 D native
+channel (uniform inlet, NEQ outlet, free-slip y walls), cylinder at
+(30 D, 30 D). Structural RHS = CME + Wen (diagnostics: raw CME, node terms).
+Research driver only: Simulation API, no production change.
+
+    PYTHONPATH=src python research/viv_re150/run.py --out /path/outside/checkout
+"""
+from __future__ import annotations
+
+import argparse
+import json
+import math
+import os
+import tempfile
+
+import numpy as np
+
+import viv
+
+D, U = 40, 0.04
+TEMPLATE = os.path.join(os.path.dirname(os.path.abspath(__file__)), "..", "..", "src",
+                        "CelerisLab", "configs", "config_lbm.json")
+
+
+def config(D=D, U=U):
+    nx, ny = 75 * D, 60 * D
+    with open(TEMPLATE) as fh:
+        cfg = json.load(fh)
+    cfg["grid"].update(lattice_model="D2Q9", nx=nx, ny=ny, nz=1)
+    cfg["physics"].update(data_type="FP32", viscosity=U * D / viv.RE, velocity=U, rho=1.0)
+    cfg["method"].update(collision="MRT", streaming="double_buffer", store_precision="FP32",
+                         ddf_shifting=False, y_wall_bc="free_slip")
+    cfg["method"]["les"]["enabled"] = False
+    cfg["method"]["inlet"]["profile"] = "uniform"
+    return cfg, nx, ny
+
+
+def main():
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--out", required=True)
+    ap.add_argument("--D", type=int, default=D)
+    ap.add_argument("--smoke", action="store_true",
+                    help="plumbing smoke only: exactly --steps steps (<= 64), never scored")
+    ap.add_argument("--steps", type=int, default=0)
+    args = ap.parse_args()
+    if args.smoke != bool(args.steps) or args.steps > 64 or args.steps < 0:
+        raise SystemExit("--smoke requires --steps N with 1 <= N <= 64 (and only with --smoke)")
+    if args.D != D and not args.smoke:
+        raise SystemExit("the physical case is predeclared at D = 40")
+    if os.path.realpath(args.out).startswith(os.path.realpath(os.path.join(os.path.dirname(__file__), "..", ".."))):
+        raise SystemExit("--out must be outside the checkout")
+    os.makedirs(args.out, exist_ok=True)
+    for name in ("report.json", "steps.npy"):  # never overwrite earlier evidence
+        if os.path.exists(os.path.join(args.out, name)):
+            raise SystemExit(f"{args.out} already holds {name}; use a fresh output directory")
+    Dn = args.D
+    cfg, nx, ny = config(Dn)
+    s, cyc = viv.Structure(Dn, U), viv.Cycles(Dn)
+    T_n = 1.0 / s.f_n
+    step_cap, n_seed = int(round(viv.STEP_CAP_TN * T_n)), int(round(viv.SEED_TN * T_n))
+    quiet = int(round(5 * T_n))
+    if args.smoke:
+        step_cap = args.steps
+    centre = (30.0 * Dn, 0.5 * (ny - 1))
+
+    from CelerisLab.body.motion import Translation
+    from CelerisLab.simulation import Simulation
+    with tempfile.NamedTemporaryFile("w", suffix=".json", delete=False) as fh:
+        json.dump(cfg, fh)
+    sim = Simulation(fh.name)
+    rows, status, seed_step, f_last = [], "running", None, 0.0
+    try:
+        body = sim.add_body("circle", centre, 0.5 * Dn,
+                            translation=Translation((0.0, 0.0, -1.5 * Dn, 1.5 * Dn), 0.5))
+        sim.initialize()
+        for n in range(step_cap):
+            # rows[k] holds the state y_{k+1}; the declared window is y_t, t in [n_seed - 5 T_n, n_seed)
+            if not args.smoke and n == n_seed and \
+                    max(abs(r[1]) for r in rows[n_seed - quiet - 1:n_seed - 1]) < viv.ELIGIBLE * Dn:
+                s.v += s.seed_velocity(Dn)         # after the quiet check, before the update
+                seed_step = n
+                cyc.discard_current()
+                cyc.reset_proof()
+            y1, v1 = s.candidate(f_last)
+            if not (math.isfinite(y1) and math.isfinite(v1)) or abs(y1) > viv.Y_LIMIT * Dn:
+                status = "fail_bound_or_nonfinite"
+                break
+            try:
+                sim.set_body(body, offset=(0.0, y1), vy=y1 - s.y)
+            except ValueError as exc:              # input rejection: stop, no retry
+                status = f"rejected: {exc}"
+                break
+            sim.run(1)                             # device failures raise and mark the Simulation failed
+            d = sim.moving_diagnostics()[0]
+            f_obs = sim.read_force(body)[1]
+            f_rhs = float(d[1] + d[3])             # CME + Wen (y)
+            row = (n, y1, v1, f_rhs, float(d[1]), float(d[3]), float(f_obs),
+                   float(d[7] - d[10]), float(d[13]))
+            if not viv.finite(row):                # every recorded output, not only the RHS
+                status = "fail_nonfinite_output"
+                break
+            s.y, s.v, f_last = y1, v1, f_rhs       # commit only after a successful step
+            rows.append(row)
+            if args.smoke:
+                continue                           # no settle / score in the plumbing smoke
+            if cyc.update(s.y) == "scored":
+                status = "scored"
+                break
+            if cyc.counted >= viv.CYCLE_CAP:
+                status = "timeout_response_cycles"
+                break
+        else:
+            status = ("smoke_ok" if all(r[8] == 0.0 for r in rows) else "smoke_no_donor") \
+                if args.smoke else "timeout_native_steps"
+    except BaseException as exc:
+        status = f"error: {exc!r}"
+        raise
+    finally:
+        np.save(os.path.join(args.out, "steps.npy"), np.asarray(rows, float))
+        report = dict(status=status, smoke=args.smoke, D=Dn, steps=len(rows), t_U_over_D=len(rows) * U / Dn, T_n=T_n,
+                      seed_step=seed_step, counted_cycles=cyc.counted, cycles=cyc.log,
+                      columns=["n", "y", "v", "F_rhs_cme_wen", "F_cme", "F_wen", "F_obs",
+                               "node_cover_minus_uncover_y", "no_donor"])
+        if status == "scored" and not args.smoke:
+            report["score"] = viv.score(cyc.window)
+        with open(os.path.join(args.out, "report.json"), "w") as fh:
+            json.dump(report, fh, indent=1)
+        sim.close()
+
+
+if __name__ == "__main__":
+    main()

research/viv_re150/viv.py

@@ -0,0 +1,126 @@
+"""Pure logic of the predeclared Re 150 free transverse VIV case (stdlib only).
+
+Contract: docs/validation/leontini_viv_predeclare.md. Oscillator update with a
+constant force over one lattice step, response-cycle tracking with fixed y = 0
+crossings, eligibility, 10-cycle blocks, settling/scoring with deterministic
+restart, and the conditional seed.
+"""
+from __future__ import annotations
+
+import math
+
+RE, U_STAR, MASS_NUMBER = 150.0, 5.0, 2.0      # m/(rho D^2) = 2 -> m* = 8/pi
+TARGET_A, BAND = 0.520, 0.10                   # Qin §4.5 maximum amplitude, ±10 %
+ELIGIBLE, BLOCK, LAST5 = 0.01, 10, 5           # in D; cycles; cycles
+WINDOW_TOL, LAST5_TOL = 0.02, 0.01
+CYCLE_CAP, STEP_CAP_TN, SEED_TN, Y_LIMIT = 150, 200, 40, 1.0
+
+
+def finite(values):
+    """True iff every recorded value is finite (fail-closed output check)."""
+    return all(math.isfinite(float(v)) for v in values)
+
+
+class Structure:
+    """Undamped spring-mass, exact harmonic update over dt = 1 with constant force."""
+
+    def __init__(self, D, u):
+        self.m = MASS_NUMBER * D * D           # lattice rho = 1
+        self.f_n = u / (U_STAR * D)
+        self.w = 2.0 * math.pi * self.f_n
+        self.k = self.m * self.w * self.w
+        self.y = self.v = 0.0
+
+    def candidate(self, force):
+        ye = force / self.k
+        c, s = math.cos(self.w), math.sin(self.w)
+        return (ye + (self.y - ye) * c + self.v / self.w * s,
+                -(self.y - ye) * self.w * s + self.v * c)
+
+    def seed_velocity(self, D):
+        return self.w * ELIGIBLE * D           # amplitude ~ .01 D
+
+
+class Cycles:
+    """Upward y = 0 crossings, eligibility, blocks, settle/score with restart."""
+
+    def __init__(self, D):
+        self.D = D
+        self.reset_proof()
+        self.counted = 0
+        self.log = []                          # one dict per completed cycle
+        self._cur = None
+        self._prev = 0.0
+        self._since = 0
+
+    def reset_proof(self):
+        self.amps = []                         # eligible amplitudes since proof start
+        self.blocks = []                       # block means
+        self.settled_at = None                 # block index where R held
+        self.window = None                     # completed scoring cycles (list of dicts)
+
+    def discard_current(self):
+        self._cur, self._since = None, 0
+
+    def rule(self, k):
+        """R(k): block-mean change ≤ 2 % and last-five spread ≤ 1 %."""
+        if k < 1:
+            return False
+        w0, w1 = self.blocks[k - 1], self.blocks[k]
+        last = self.amps[(k + 1) * BLOCK - LAST5:(k + 1) * BLOCK]
+        a5 = sum(last) / LAST5
+        return abs(w1 - w0) / w0 <= WINDOW_TOL and max(abs(a - a5) for a in last) / a5 <= LAST5_TOL
+
+    def update(self, y):
+        """Feed y after a committed step; returns 'scored' when a stable window completes."""
+        out = None
+        crossing = self._prev < 0.0 <= y and self._since >= 2
+        if crossing:
+            if self._cur is not None:
+                out = self._close()
+            self._cur, self._since = dict(ymax=y, ymin=y, ysum=0.0, n=0), 0
+        if self._cur is not None:
+            c = self._cur
+            c["ymax"], c["ymin"] = max(c["ymax"], y), min(c["ymin"], y)
+            c["ysum"] += y
+            c["n"] += 1
+        self._since += 1
+        self._prev = y
+        return out
+
+    def _close(self):
+        c = self._cur
+        a = 0.5 * (c["ymax"] - c["ymin"]) / self.D
+        cyc = dict(ymax=c["ymax"] / self.D, ymin=c["ymin"] / self.D,
+                   offset=c["ysum"] / c["n"] / self.D, amp=a, steps=c["n"], eligible=a >= ELIGIBLE)
+        self.log.append(cyc)
+        if not cyc["eligible"]:
+            return None
+        self.counted += 1
+        self.amps.append(a)
+        if self.settled_at is not None:
+            self.window.append(cyc)
+        if len(self.amps) % BLOCK:
+            return None
+        self.blocks.append(sum(self.amps[-BLOCK:]) / BLOCK)
+        k = len(self.blocks) - 1
+        if self.settled_at is None:
+            if self.rule(k):
+                self.settled_at, self.window = k, []
+            return None
+        if not self.rule(k):                   # scoring block failed: restart from it
+            amps, blocks = self.amps[-BLOCK:], self.blocks[-1:]
+            self.reset_proof()
+            self.amps, self.blocks = amps, blocks
+            return None
+        return "scored" if k == self.settled_at + 2 else None
+
+
+def score(window):
+    """Target metric and reported quantities over the 20 scoring cycles."""
+    a_max = max(max(abs(c["ymax"]), abs(c["ymin"])) for c in window)
+    return dict(A_max=a_max, mean_amp=sum(c["amp"] for c in window) / len(window),
+                ymax=max(c["ymax"] for c in window), ymin=min(c["ymin"] for c in window),
+                offset=sum(c["offset"] for c in window) / len(window),
+                T_resp=sum(c["steps"] for c in window) / len(window),
+                within_band=abs(a_max - TARGET_A) <= BAND * TARGET_A)

src/CelerisLab/body/manager.py

@@ -405,7 +405,9 @@ class ObjectManager:
 
         for obj in objects:
             obj.validate_surface_actuation(self.cfg)
-            if (not obj.is_sensor) and (not obj.is_force_region) and hasattr(obj, 'get_curved_list'):
+            # Moving circles never enter the packed static links (docs/moving_bodies.md, section 3).
+            if (not obj.is_sensor) and (not obj.is_force_region) and obj.motion is None \
+                    and hasattr(obj, 'get_curved_list'):
                 (
                     fluid_idx, dirs, q_vals, body_ids,
                     rx_vals, ry_vals, rz_vals, fallback_vals

src/CelerisLab/body/motion.py

@@ -0,0 +1,259 @@
+# CelerisLab/body/motion.py
+"""Host state of a translating circle (dense double buffer, D2Q9 only).
+
+Contract: docs/moving_bodies.md, sections 4a, 5 and 7. The applied geometry
+(last successfully committed offset) and an optional pending target are two
+separate facts; centres are always ``base + offset``. Nothing here touches
+the device: the caller supplies upload/step/synchronize callables.
+"""
+from __future__ import annotations
+
+import math
+from dataclasses import dataclass
+from typing import Callable, Optional, Tuple
+
+MAX_STEP = 0.5     # Euclidean centre displacement per step (lattice cells)
+BAND_MARGIN = 2    # cells added around the swept circle AABB
+R_MIN = 4.0        # smallest radius covered by the donor enumeration
+
+Vec = Tuple[float, float]
+
+# D2Q9 non-rest directions, paired order of lbm/descriptors.
+_CX = (1, -1, 0, 0, 1, -1, 1, -1)
+_CY = (0, 0, 1, -1, 1, -1, -1, 1)
+
+
+def _vec(value, name: str) -> Vec:
+    x, y = (float(v) for v in value)
+    if not (math.isfinite(x) and math.isfinite(y)):
+        raise ValueError(f"{name} must be finite, got {value!r}")
+    return x, y
+
+
+@dataclass(frozen=True)
+class Translation:
+    """Declared motion envelope: offset box (x_min, x_max, y_min, y_max) and speed bound."""
+    offset_box: Tuple[float, float, float, float]
+    max_speed: float
+
+    def __post_init__(self):
+        box = tuple(float(v) for v in self.offset_box)
+        speed = float(self.max_speed)
+        if len(box) != 4 or not all(math.isfinite(v) for v in box + (speed,)):
+            raise ValueError("offset_box needs 4 finite values and max_speed must be finite")
+        if not (box[0] <= 0.0 <= box[1] and box[2] <= 0.0 <= box[3]):
+            raise ValueError("offset_box must contain the zero offset")
+        if not 0.0 < speed <= MAX_STEP:
+            raise ValueError(f"max_speed must satisfy 0 < max_speed <= {MAX_STEP}")
+        object.__setattr__(self, "offset_box", box)
+        object.__setattr__(self, "max_speed", speed)
+
+
+class MovingCircle:
+    """Applied offset, pending target and the fail-closed commit of one circle."""
+
+    def __init__(self, base, radius: float, translation: Translation,
+                 on_failure: Optional[Callable[[BaseException], None]] = None):
+        self.base = _vec(base, "base")
+        self.radius = float(radius)
+        if not (math.isfinite(self.radius) and self.radius >= R_MIN):
+            raise ValueError(f"moving circle radius must be >= {R_MIN}")
+        if not isinstance(translation, Translation):
+            raise TypeError("translation must be a Translation")
+        self.translation = translation
+        self.applied: Vec = (0.0, 0.0)
+        self.pending: Optional[Tuple[Vec, Vec]] = None
+        self.failed = False
+        self._on_failure = on_failure  # owner lifecycle hook (Simulation marks itself failed)
+
+    def centre(self, offset: Vec) -> Vec:
+        return self.base[0] + offset[0], self.base[1] + offset[1]
+
+    def check_target(self, offset, velocity) -> Tuple[Vec, Vec]:
+        """Validate a target against ``applied``; returns it, mutates nothing."""
+        self._require_usable()
+        offset, velocity = _vec(offset, "offset"), _vec(velocity, "velocity")
+        x0, x1, y0, y1 = self.translation.offset_box
+        if not (x0 <= offset[0] <= x1 and y0 <= offset[1] <= y1):
+            raise ValueError(f"offset {offset} outside declared box {self.translation.offset_box}")
+        step = math.hypot(offset[0] - self.applied[0], offset[1] - self.applied[1])
+        if step > MAX_STEP:
+            raise ValueError(f"per-step displacement {step} exceeds {MAX_STEP}")
+        if math.hypot(*velocity) > self.translation.max_speed:
+            raise ValueError(f"velocity {velocity} exceeds max_speed {self.translation.max_speed}")
+        if offset == self.applied and velocity != (0.0, 0.0):
+            raise ValueError("a hold target (offset == applied) needs zero translational velocity")
+        return offset, velocity
+
+    def set_target(self, offset, velocity) -> None:
+        """Replace the pending target (last wins); invalid input keeps the old one."""
+        self.pending = self.check_target(offset, velocity)
+
+    def record(self) -> Tuple[Vec, Vec, Vec]:
+        """Motion record (c_old, c_new, u) for the next single moving step."""
+        self._require_usable()
+        if self.pending is not None:
+            offset, velocity = self.pending
+        else:
+            offset, velocity = self.applied, (0.0, 0.0)
+        return self.centre(self.applied), self.centre(offset), velocity
+
+    def advance(self, upload: Callable, step: Callable, synchronize: Callable) -> None:
+        """Upload the record, run one step and commit only after a successful sync.
+
+        Any exception after the upload is attempted leaves device geometry
+        unknown: the circle is marked failed, never reports applied state again,
+        and the owner's ``on_failure`` hook is called.
+        """
+        record = self.record()
+        try:
+            upload(record)
+            step()
+            synchronize()
+        except BaseException as exc:
+            self.fail(exc)
+            raise
+        self.commit()
+
+    def commit(self) -> None:
+        """Applied := pending offset (caller has synchronized a successful step)."""
+        if self.pending is not None:
+            self.applied = self.pending[0]
+            self.pending = None
+
+    def fail(self, exc: BaseException) -> None:
+        self.failed = True
+        if self._on_failure is not None:
+            self._on_failure(exc)
+
+    def band(self) -> Tuple[int, int, int, int]:
+        """Inclusive cell AABB (x0, x1, y0, y1) swept over the whole offset box."""
+        x0, x1, y0, y1 = self.translation.offset_box
+        r, m = self.radius, BAND_MARGIN
+        return (math.floor(self.base[0] + x0 - r) - m, math.ceil(self.base[0] + x1 + r) + m,
+                math.floor(self.base[1] + y0 - r) - m, math.ceil(self.base[1] + y1 + r) + m)
+
+    def _require_usable(self) -> None:
+        if self.failed:
+            raise RuntimeError("moving circle state is unknown after a failed step")
+
+
+def _inside(x: int, y: int, c: Vec, r: float) -> bool:
+    return (x - c[0]) ** 2 + (y - c[1]) ** 2 < r * r
+
+
+def empty_donor_cells(c_old: Vec, c_new: Vec, radius: float):
+    """Fresh cells (inside c_old, outside c_new) whose donor set V(x) is empty."""
+    r = radius
+    xs = range(math.floor(min(c_old[0], c_new[0]) - r) - 1, math.ceil(max(c_old[0], c_new[0]) + r) + 2)
+    ys = range(math.floor(min(c_old[1], c_new[1]) - r) - 1, math.ceil(max(c_old[1], c_new[1]) + r) + 2)
+    empty = []
+    for x in xs:
+        for y in ys:
+            if not _inside(x, y, c_old, r) or _inside(x, y, c_new, r):
+                continue
+            if not any(not _inside(x + cx, y + cy, c_old, r) and not _inside(x + cx, y + cy, c_new, r)
+                       for cx, cy in zip(_CX, _CY)):
+                empty.append((x, y))
+    return empty
+
+
+def enumerate_donor_coverage(radius: float, positions: int = 8, directions: int = 16) -> int:
+    """Count empty-donor fresh cells over sub-cell centres and moves of MAX_STEP.
+
+    Evidence for the declared bounds (radius >= R_MIN, |dc|_2 <= MAX_STEP),
+    not a proof; the device keeps the ``no_donor`` counter.
+    """
+    count = 0
+    for i in range(positions):
+        for j in range(positions):
+            c_old = (i / positions, j / positions)
+            for k in range(directions):
+                a = 2.0 * math.pi * k / directions
+                c_new = (c_old[0] + MAX_STEP * math.cos(a), c_old[1] + MAX_STEP * math.sin(a))
+                count += len(empty_donor_cells(c_old, c_new, radius))
+    return count
+
+
+class MovingTables:
+    """Device tables of the translating circles for ``MovingCircleCells``/``Links``.
+
+    Band table (x0, y0, width, height, first_slot, body_id), immutable per run;
+    FP64 motion records (cx_old, cy_old, cx_new, cy_new, r, ux, uy), immutable
+    during a step; separate per-body diagnostics (see DIAG), never part of obs.
+    """
+    BAND_INTS, GEOM_DOUBLES = 6, 7
+    DIAG = ("cme_x", "cme_y", "wen_x", "wen_y", "link_dm", "n_links",
+            "cover_jx", "cover_jy", "cover_mass", "uncover_jx", "uncover_jy", "uncover_mass",
+            "eq_only", "no_donor", "crossings")
+    DIAG_FLOATS = len(DIAG)
+
+    def __init__(self, band, geom):
+        import numpy as np
+        import pycuda.driver as cuda
+        self.n_bodies = len(geom) // self.GEOM_DOUBLES
+        self.n_slots = int(band[-2] + band[-4] * band[-3] * 8)
+        self.n_cells = self.n_slots // 8
+        self.geom = list(geom)  # host copy of the uploaded records
+        self.band_gpu = cuda.to_device(np.asarray(band, np.int32))
+        self.geom_gpu = cuda.to_device(np.asarray(geom, np.float64))
+        # Owned page-locked staging for async record uploads (allocated once).
+        self._geom_pinned = cuda.pagelocked_empty(len(geom), np.float64)
+        self.diag_gpu = cuda.to_device(np.zeros(self.n_bodies * self.DIAG_FLOATS, np.float32))
+        self.slot_debug_gpu = None  # optional (q, fallback, rx, ry) per slot for parity tests
+
+    @classmethod
+    def upload(cls, circles):
+        """``circles``: [(MovingCircle, body_id)] in body order."""
+        return cls(*cls.layout(circles))
+
+    @staticmethod
+    def layout(circles):
+        """Host band and geometry tables (no device work)."""
+        band, geom, first = [], [], 0
+        for motion, body_id in circles:
+            x0, x1, y0, y1 = motion.band()
+            width, height = x1 - x0 + 1, y1 - y0 + 1
+            band += [x0, y0, width, height, first, int(body_id)]
+            first += width * height * 8
+            geom += MovingTables.record(motion)
+        return band, geom
+
+    @staticmethod
+    def record(motion):
+        """FP64 motion record of one circle for its next step (hold if nothing pending)."""
+        c_old, c_new, u = motion.record()
+        return [*c_old, *c_new, motion.radius, *u]
+
+    def upload_records(self, circles, stream) -> bool:
+        """Enqueue the records when they differ from the device copy; True if advancing.
+
+        The pinned staging buffer is reused: callers synchronize ``stream`` after every
+        upload before the next one (``Simulation._run_moving``), so no copy is in flight
+        when it is rewritten. ``geom`` is the enqueued state; after a failed copy the
+        owner is failed and the cache is never trusted again.
+        """
+        import pycuda.driver as cuda
+        geom = [v for motion, _ in circles for v in self.record(motion)]
+        if geom != self.geom:
+            self._geom_pinned[:] = geom
+            cuda.memcpy_htod_async(self.geom_gpu, self._geom_pinned, stream)
+            self.geom = geom
+        g = self.GEOM_DOUBLES
+        return any(geom[i:i + 2] != geom[i + 2:i + 4] for i in range(0, len(geom), g))
+
+    def enable_slot_debug(self):
+        import numpy as np
+        import pycuda.driver as cuda
+        self.slot_debug_gpu = cuda.to_device(np.zeros(self.n_slots * 4, np.float32))
+
+    def zero_diagnostics(self, stream):
+        import pycuda.driver as cuda
+        cuda.memset_d32_async(self.diag_gpu, 0, self.n_bodies * self.DIAG_FLOATS, stream)
+
+    def download_diagnostics(self):
+        import numpy as np
+        import pycuda.driver as cuda
+        out = np.empty((self.n_bodies, self.DIAG_FLOATS), np.float32)
+        cuda.memcpy_dtoh(out, self.diag_gpu)
+        return out

src/CelerisLab/body/objects.py

@@ -140,6 +140,7 @@ class SimObject:
             x=float(c[0]), y=float(c[1]), z=z0,
         )
         self.control = ObjectControl()
+        self.motion = None  # MovingCircle for a translating circle (body/motion.py)
 
     @property
     def surface_actuation(self):

src/CelerisLab/lbm/blocks.py

@@ -18,7 +18,9 @@ all NQ populations from an owner core cell into another block's halo cell.
 Periodicity is explicit per axis. Across a non-periodic exterior, halo cells
 have no pair and stay invalid (NaN after packing): no wrap, no implied BC.
 
-Preparation only: no CUDA, no boundary conditions, no refinement levels.
+Host preparation plus thin launch wrappers: no CUDA import or allocation
+(launch_* only call kernels and device tables the caller owns), no refinement
+levels; level_config and static_circle_tables prepare one blocked level.
 The one-cell same-level halo here is unrelated to the four-cell fine halo of
 the fixed refinement patch (refinement.py), whose semantics are unchanged.
 

src/CelerisLab/lbm/hierarchy.py

@@ -26,8 +26,15 @@ class Hierarchy:
     Index conventions (the single place they are defined):
       * level l = 0 (root) .. depth-1; pair index p = l-1 builds level l from level p;
       * array shapes are (ny, nx); rectangles, origins and centers are (x, y);
-      * the root EDGE frame puts root cell i at [i, i+1), centers at i+.5 (bodies and
-        Jets of FixedPatchStepper); uniform Simulation objects use the integer-NODE frame;
+      * a pair's two members are its COARSE (parent) and FINE (child) level, for every
+        pair, not absolute L0/L1 (fine_shape, fine_origin, blocked_coarse_tables, ...);
+      * refined levels (l >= 1) of FixedPatchStepper place bodies and Jets in the root
+        EDGE frame: root cell i covers [i, i+1), center i+.5. Uniform Simulation objects
+        and the static circles of a single-level (depth 1) FixedPatchStepper use the
+        integer-NODE frame (CircleGeometry.build_flag_mask at integer (x, y)). No
+        NODE<->EDGE input conversion is applied anywhere; inside the EDGE frame, refined
+        boundaries still project root-edge coordinates to fine-local cell indices with
+        the half-cell shift ((v - o)/h - .5, see prepare_fine_boundary);
       * cell (i, j) of level l's allocated array (ghost band and holes included) covers
         [o + h*i, o + h*(i+1)) with (o, h) = edge_frame(l); level 1 is built from the root
         config directly, deeper levels from their parent frame (see edge_frame).
@@ -115,7 +122,7 @@ def leaf_partition(h):
     for level, kind in enumerate(h.kinds):
         if level:
             pair = h.pairs[level - 1]
-            py, px = (pair.fine_level_offset if isinstance(pair, FixedPatch2D) else (0, 0))
+            py, px = pair.fine_level_offset
             oy, ox = 2 * oy + py, 2 * ox + px
         y, x = np.nonzero(kind == ACTIVE)
         np.add.at(base, ((y + oy) >> level, (x + ox) >> level), 0.25 ** level)

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

@@ -1,10 +1,12 @@
+// CelerisLab – operators/block_collision.cuh
 // The one collision extension point of the blocked steps (BlockStep, BlockChannelStep).
-// Compile-time selection: MRT shares one path; SRT keeps two explicit arithmetic
-// variants, unchanged from their original kernels:
-//   level  (BLOCK_ROOT_BC 0/1): unclamped omega, no forcing term; shifted storage
-//          builds feq from the retained drho and direct momentum (numerical variant);
-//   channel (BLOCK_ROOT_BC 2): matches uniform OneStep SRT arithmetic (omega clamp,
-//          zero Guo forcing term, compute_feq).
+// Unforced by construction: FixedPatchStepper supplies no forcing, so no runtime
+// d_params.fx/fy branch is read. Two compile-time arithmetic families:
+//   uniform-equivalent (MRT on every level; SRT on the channel root, BLOCK_ROOT_BC 2):
+//          collide_dispatch arithmetic with zero force and no LES: clamped omega;
+//          SRT keeps the zero Guo forcing term of collide_srt;
+//   level SRT (BLOCK_ROOT_BC 0/1): unclamped omega, no forcing term; shifted storage
+//          builds feq from the retained drho and direct momentum (numerical variant).
 #ifndef CELERIS_OPERATORS_BLOCK_COLLISION_CUH
 #define CELERIS_OPERATORS_BLOCK_COLLISION_CUH
 
@@ -16,17 +18,19 @@
 #endif
 
 __device__ __forceinline__ void block_collide(float* f) {
-#if COLLISION_MODEL == 2
-    collide_mrt_unforced(f);
-#elif BLOCK_ROOT_BC == 2
-    float rho, ux, uy, feq[NQ], forcing[NQ] = {};
+    float rho, ux, uy;
     compute_rho_u(f, rho, ux, uy);
-    compute_feq(rho, ux, uy, feq);
+#if COLLISION_MODEL == 2 || BLOCK_ROOT_BC == 2
     const float omega = fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega));
+  #if COLLISION_MODEL == 2
+    collide_mrt_no_force(f, rho, ux, uy, omega);
+  #else
+    float feq[NQ], forcing[NQ] = {};
+    compute_feq(rho, ux, uy, feq);
     collide_srt(f, feq, forcing, omega);
+  #endif
 #else
-    float rho, ux, uy, feq[NQ];
-    compute_rho_u(f, rho, ux, uy);
+    float feq[NQ];
 #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];

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/step/moving_circle.cu

@@ -0,0 +1,183 @@
+// Translating circles (docs/moving_bodies.md): M1 MovingCircleCells reclassifies
+// band cells and refills uncovered ones; M2 MovingCircleLinks applies the cut links.
+// Geometry comes from an immutable FP64 record per body
+// {cx_old, cy_old, cx_new, cy_new, r, ux, uy}; the M2 operation order mirrors
+// CircleGeometry.build_cut_links / find_circle_intersection with explicit _rn
+// operations (bit parity with the host builder is a test target, not a claim).
+// Moving circles are never in the packed static links, so this kernel is the
+// only writer of their solid slots.
+#ifndef CELERIS_MOVING_CIRCLE_CU
+#define CELERIS_MOVING_CIRCLE_CU
+
+#if DIM == 2
+constexpr int MOVING_BAND_INTS = 6;   // x0, y0, width, height, first_slot, body_id
+constexpr int MOVING_GEOM_DOUBLES = 7;
+// cme_x, cme_y, wen_x, wen_y, link_dm, n_links, cover_jx, cover_jy, cover_mass,
+// uncover_jx, uncover_jy, uncover_mass, eq_only, no_donor, crossings
+constexpr int MOVING_DIAG_FLOATS = 15;
+
+__device__ __forceinline__ double mc_sq(double v) { return __dmul_rn(v, v); }
+
+__device__ __forceinline__ bool mc_inside(int x, int y, double cx, double cy, double r)
+{
+    return __dadd_rn(mc_sq(__dsub_rn((double)x, cx)), mc_sq(__dsub_rn((double)y, cy))) < mc_sq(r);
+}
+
+// Smallest root t in [0, 1] of |p + t c - centre| = r (find_circle_intersection order).
+__device__ __forceinline__ bool mc_hit(int x, int y, int ex, int ey,
+                                       double cx, double cy, double r, double& t)
+{
+    const double dx = (double)ex, dy = (double)ey;
+    const double px = __dsub_rn((double)x, cx), py = __dsub_rn((double)y, cy);
+    const double a = __dadd_rn(mc_sq(dx), mc_sq(dy));
+    const double b = __dmul_rn(2.0, __dadd_rn(__dmul_rn(dx, px), __dmul_rn(dy, py)));
+    const double c = __dsub_rn(__dadd_rn(mc_sq(px), mc_sq(py)), mc_sq(r));
+    const double det = __dsub_rn(mc_sq(b), __dmul_rn(__dmul_rn(4.0, a), c));
+    if (det < 0.0) return false;
+    const double s = __dsqrt_rn(det), a2 = __dmul_rn(2.0, a);
+    const double t1 = __ddiv_rn(__dsub_rn(-b, s), a2);
+    const double t2 = __ddiv_rn(__dadd_rn(-b, s), a2);
+    if (0.0 <= t1 && t1 <= 1.0) { t = t1; return true; }
+    if (0.0 <= t2 && t2 <= 1.0) { t = t2; return true; }
+    return false;
+}
+
+__global__ void MovingCircleLinks(
+    fpxx* fi, const uint16_t* flag,
+    const int* band, const double* geom,
+    unsigned int n_bodies, unsigned int n_slots,
+    const float* action, float* obs, float* diag, float* slot_debug)
+{
+    const unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
+    if (tid >= n_slots) return;
+    unsigned int m = 0;
+    while (m + 1 < n_bodies && tid >= (unsigned int)band[(m + 1) * MOVING_BAND_INTS + 4]) ++m;
+    const int* b = band + m * MOVING_BAND_INTS;
+    const unsigned int local = tid - (unsigned int)b[4];
+    const unsigned int dir = local % 8u + 1u;
+    const int x = b[0] + (int)((local / 8u) % (unsigned int)b[2]);
+    const int y = b[1] + (int)((local / 8u) / (unsigned int)b[2]);
+    const double* g = geom + MOVING_GEOM_DOUBLES * m;
+    const double cx = g[2], cy = g[3], r = g[4];
+    const int ex = d_cx[dir], ey = d_cy[dir];
+    float* debug = slot_debug == nullptr ? nullptr : slot_debug + 4 * tid;
+    if (debug != nullptr) debug[0] = -1.0f;
+
+    double t;
+    if (mc_inside(x, y, cx, cy, r) || !mc_inside(x + ex, y + ey, cx, cy, r)
+            || !mc_hit(x, y, ex, ey, cx, cy, r, t))
+        return;
+    const double hx = __dadd_rn((double)x, __dmul_rn(t, (double)ex));
+    const double hy = __dadd_rn((double)y, __dmul_rn(t, (double)ey));
+    const double hdx = __dsub_rn(hx, (double)x), hdy = __dsub_rn(hy, (double)y);
+    const double norm = __dsqrt_rn((double)(ex * ex + ey * ey));
+    double qd = __ddiv_rn(__dsqrt_rn(__dadd_rn(mc_sq(hdx), mc_sq(hdy))), norm);
+    qd = fmax(0.01, fmin(qd, 1.0));
+    const float q = (float)qd;
+    const float rx = (float)__dsub_rn(hx, cx), ry = (float)__dsub_rn(hy, cy);
+
+    // Same two-stage order as build_cut_links + pack_cut_links_to_soa: the fallback
+    // class is chosen from the clamped FP64 q (q < 1/2 needs a fluid donor at x - c);
+    // only then are q, rx, ry rounded to FP32 for the BFL branch and spin velocity.
+    unsigned char fallback = CURVED_FALLBACK_BOUZIDI;
+    if (qd < 0.5) {
+        const int xd = x - ex, yd = y - ey;
+        if (xd < 0 || xd >= (int)NX || yd < 0 || yd >= (int)NY || mc_inside(xd, yd, cx, cy, r)
+                || (flag[linear_index((unsigned int)xd, (unsigned int)yd)] & FLAG_FLUID) == 0)
+            fallback = CURVED_FALLBACK_HALFWAY;
+    }
+    if (debug != nullptr) { debug[0] = q; debug[1] = (float)fallback; debug[2] = rx; debug[3] = ry; }
+
+    const int bid = b[5];
+    const float omega = action_omega(action, bid);
+    const float Uw = (float)g[5] - omega * ry, Vw = (float)g[6] + omega * rx;  // u + omega x r
+    float exchange[3], dm;
+    apply_bouzidi_link(dir, q, linear_index((unsigned int)x, (unsigned int)y), fi,
+                       Uw, Vw, rx, ry, fallback, obs, bid, exchange, &dm);
+    if (diag != nullptr) {
+        float* d = diag + m * MOVING_DIAG_FLOATS;
+        atomicAdd(&d[0], exchange[0]);
+        atomicAdd(&d[1], exchange[1]);
+        atomicAdd(&d[2], Uw * dm);   // Wen 2014 correction -u_w (f_t - f_r)
+        atomicAdd(&d[3], Vw * dm);
+        atomicAdd(&d[4], dm);
+        atomicAdd(&d[5], 1.0f);
+    }
+}
+// M1: one thread per band cell, launched only on steps with c_old != c_new.
+// Membership comes from the immutable old/new predicates, never from flags, so
+// classification, refill and the flag write of one cell cannot race with another.
+// ddf holds post-collision G^{n-1}; refilled cells get all nine slots at that level.
+__global__ void MovingCircleCells(
+    fpxx* fi, uint16_t* flag, const int* band, const double* geom,
+    unsigned int n_bodies, unsigned int n_cells, const float* action, float* diag)
+{
+    const unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
+    if (tid >= n_cells) return;
+    unsigned int m = 0;
+    while (m + 1 < n_bodies && tid >= (unsigned int)band[(m + 1) * MOVING_BAND_INTS + 4] / 8u) ++m;
+    const int* b = band + m * MOVING_BAND_INTS;
+    const unsigned int local = tid - (unsigned int)b[4] / 8u;
+    const int x = b[0] + (int)(local % (unsigned int)b[2]);
+    const int y = b[1] + (int)(local / (unsigned int)b[2]);
+    const double* g = geom + MOVING_GEOM_DOUBLES * m;
+    const bool in_old = mc_inside(x, y, g[0], g[1], g[4]);
+    const bool in_new = mc_inside(x, y, g[2], g[3], g[4]);
+    if (in_old == in_new) return;
+    const unsigned long k = linear_index((unsigned int)x, (unsigned int)y);
+    float* d = diag + m * MOVING_DIAG_FLOATS;
+    float rho = 0.0f, jx = 0.0f, jy = 0.0f;
+    if (in_new) {  // covered: its fluid content leaves the active set
+        for (unsigned int i = 0; i < NQ; ++i) {
+            const float f = load_ddf(fi, index_f(k, i));
+            rho += f; jx += (float)d_cx[i] * f; jy += (float)d_cy[i] * f;
+        }
+        flag[k] = FLAG_SOLID | FLAG_OBSTACLE | FLAG_BC_CURVED;
+        atomicAdd(&d[6], jx); atomicAdd(&d[7], jy); atomicAdd(&d[8], rho);
+        atomicAdd(&d[14], 1.0f);
+        return;
+    }
+    // Uncovered: donors are fluid in BOTH geometries; NEQ from the neighbour with the
+    // largest unnormalized c.n > 0 (first direction wins ties), n from c_new.
+    const float nx = (float)__dsub_rn((double)x, g[2]), ny = (float)__dsub_rn((double)y, g[3]);
+    float rho_sum = 0.0f, best_dot = 0.0f;
+    int n_valid = 0, best = -1;
+    for (unsigned int i = 1; i < NQ; ++i) {
+        const int xn = x + d_cx[i], yn = y + d_cy[i];
+        if (mc_inside(xn, yn, g[0], g[1], g[4]) || mc_inside(xn, yn, g[2], g[3], g[4])) continue;
+        const unsigned long kn = linear_index((unsigned int)xn, (unsigned int)yn);
+        for (unsigned int q = 0; q < NQ; ++q) rho_sum += load_ddf(fi, index_f(kn, q));
+        ++n_valid;
+        const float dot = (float)d_cx[i] * nx + (float)d_cy[i] * ny;
+        if (dot > best_dot) { best_dot = dot; best = (int)i; }
+    }
+    const float omega = action_omega(action, b[5]);
+    const float uwx = (float)g[5] - omega * ny, uwy = (float)g[6] + omega * nx;  // at the cell
+    float feq[NQ], neq[NQ];
+    compute_feq(n_valid > 0 ? rho_sum / (float)n_valid : (float)RHO, uwx, uwy, feq);
+    for (unsigned int i = 0; i < NQ; ++i) neq[i] = 0.0f;
+    if (best >= 0) {
+        const unsigned long kd = linear_index((unsigned int)(x + d_cx[best]), (unsigned int)(y + d_cy[best]));
+        float gd[NQ], rd = 0.0f, mx = 0.0f, my = 0.0f, fd[NQ];
+        for (unsigned int i = 0; i < NQ; ++i) {
+            gd[i] = load_ddf(fi, index_f(kd, i));
+            rd += gd[i]; mx += (float)d_cx[i] * gd[i]; my += (float)d_cy[i] * gd[i];
+        }
+        compute_feq(rd, mx / rd, my / rd, fd);
+        for (unsigned int i = 0; i < NQ; ++i) neq[i] = gd[i] - fd[i];
+    } else {
+        atomicAdd(&d[12], 1.0f);
+    }
+    if (n_valid == 0) atomicAdd(&d[13], 1.0f);  // hard failure, checked before commit
+    for (unsigned int i = 0; i < NQ; ++i) {
+        const float f = feq[i] + neq[i];
+        store_ddf(fi, index_f(k, i), f);
+        rho += f; jx += (float)d_cx[i] * f; jy += (float)d_cy[i] * f;
+    }
+    flag[k] = FLAG_FLUID;
+    atomicAdd(&d[9], jx); atomicAdd(&d[10], jy); atomicAdd(&d[11], rho);
+    atomicAdd(&d[14], 1.0f);
+}
+#endif  // DIM == 2
+
+#endif  // CELERIS_MOVING_CIRCLE_CU

src/CelerisLab/lbm/kernels/uniform.cu

@@ -71,6 +71,7 @@ extern "C"
 #include "step/one_step_double.cu"
 #include "step/one_step_esopull.cu"
 #include "step/aux_kernels.cu"
+#include "step/moving_circle.cu"
 #include "step/esopull_macro.cu"
 #include "step/esopull_sync.cu"
 

src/CelerisLab/lbm/refinement.py

@@ -120,6 +120,11 @@ class FixedPatch2D:
         (an index offset, not an (x, y) origin)."""
         return 2 * self.bounds[1] - HALO, 2 * self.bounds[0] - HALO
 
+    @property
+    def rectangles(self) -> tuple:
+        """The one parent rectangle, as a sequence like FixedPatchSet2D.rectangles."""
+        return (self.bounds,)
+
     @property
     def covered(self) -> np.ndarray:
         """Dense coarse mask of cells replaced by the patch (not leaves)."""
@@ -225,11 +230,18 @@ class FixedPatchSet2D:
     def bounds(self) -> tuple:
         return tuple(p.bounds for p in self.patches)
 
+    rectangles = bounds
+
     @property
     def fine_origin(self) -> tuple[float, float]:
         """Lower edge of the level-wide fine grid in coarse physical units."""
         return 0.0, 0.0
 
+    @property
+    def fine_level_offset(self) -> tuple[int, int]:
+        """The level-wide grid itself: zero (y, x) index offset."""
+        return 0, 0
+
     @property
     def covered(self) -> np.ndarray:
         mask = np.zeros(self.coarse_shape, bool)

src/CelerisLab/lbm/refinement_stepper.py

@@ -22,7 +22,7 @@ from .channel import channel_masks, validate_channel_config, validate_channel_tr
 from ..body.jets import jet_link_velocity
 from .hierarchy import build_hierarchy
 from .refinement import ACTIVE, HALO, FixedPatch2D, FixedPatchSet2D, prepare_fine_boundary
-from .refinement_tables import (blocked_base_tables, blocked_fine_tables,  # noqa: F401 (re-export)
+from .refinement_tables import (blocked_coarse_tables, blocked_fine_tables,
                                 blocked_kind, projected_explosion_tables, _prepare_blocked_boundary,
                                 _BlockedBoundary,
                                 BLOCKED_BOUNDARY_DEVICE_COLUMNS)
@@ -66,8 +66,12 @@ class _Level:
 def _kernel_configs(cfg, pairs, wall_aware):
     """Per-level kernel LBMConfig (window shape, viscosity doubling per level).
 
-    Geometry configs live in Hierarchy.configs; level 1 of a slot_reservoir root
-    replaces the root wall fields by the internal bounce-back/velocity ghost setup.
+    Deliberately distinct from Hierarchy.configs (refinement._middle_config), which
+    resets the wall fields on EVERY refined level for geometry validation. Here only
+    level 1 of a slot_reservoir root resets them (internal bounce-back/velocity ghost
+    setup); other levels inherit their parent's fields. These configs generate each
+    level's compiled config.h, so folding the two derivations is a numerical-parity
+    change (config.h/PTX diff gate), not a style one.
     """
     configs = [replace(cfg)]
     for index, pair in enumerate(pairs):
@@ -107,6 +111,57 @@ def _check_levels(refinements, blocks, objects, jets):
             raise ValueError("bodies/Jets are supported only on the finest level of a nested hierarchy")
 
 
+# Host preparation (numpy only, before any device work); each returns one entry per level.
+
+def _level_layouts(cfg, pairs, level_blocks, wall_aware):
+    """Blocked storage layout of every level, checked against the launch limits."""
+    layouts = []
+    for level, block in enumerate(level_blocks):
+        if level == 0:
+            layout = BlockLayout2D.build(cfg.ny, cfg.nx, *block, periodic=(True, not wall_aware))
+        elif isinstance(pairs[level - 1], FixedPatchSet2D):
+            layout = pairs[level - 1].fine_layout(block)
+        else:
+            layout = BlockLayout2D.build(*pairs[level - 1].fine_shape, *block,
+                                         periodic=(False, False))
+        if layout.nblocks > MAX_GRID_DIM or layout.local_shape[0] > MAX_GRID_DIM:
+            raise ValueError(f"level {level} blocks exceed launch limits")
+        layouts.append(layout)
+    return layouts
+
+
+def _host_boundaries(hierarchy, level_cfgs, level_objects, level_jets, wall_aware, shifted):
+    """Host boundary geometry of every refined level in the root edge frame (None at L0)."""
+    hosts = [None]
+    for level in range(1, hierarchy.depth):
+        pair = hierarchy.pairs[level - 1]
+        if level == 1:  # the root config carries the wall-aware boundary fields
+            host = prepare_fine_boundary(level_cfgs[0], pair, level_objects[1],
+                                         wall_aware=wall_aware, jets=level_jets[1],
+                                         _shifted_blocked=shifted)
+        else:
+            origin, spacing = hierarchy.edge_frame(level - 1)
+            host = prepare_fine_boundary(level_cfgs[level - 1], pair, level_objects[level],
+                                         jets=level_jets[level], parent_origin=origin,
+                                         parent_spacing=spacing, _shifted_blocked=shifted)
+        hosts.append(host)
+    return hosts
+
+
+def _fluid_leaves(hierarchy, boundaries):
+    """ACTIVE cells of every level that no static circle or body makes solid."""
+    leaves = [kind == ACTIVE for kind in hierarchy.kinds]
+    for level, b in enumerate(boundaries):
+        if b is None:
+            continue
+        if b.host is None:  # single level: static root circles
+            leaves[level] &= ~b.tables["solid"]
+        else:
+            fine_shape = hierarchy.pairs[level - 1].fine_shape
+            leaves[level] &= (b.host.flags.reshape(fine_shape) & SOLID) == 0
+    return leaves
+
+
 class FixedPatchStepper:
     """D2Q9 SRT/MRT FP32 fixed 2:1 refinement hierarchy, built from level lists.
 
@@ -119,7 +174,8 @@ class FixedPatchStepper:
     supported).
     ``refinements=[]`` is one blocked level: periodic SRT root without LES,
     Jets or exchange observation; its ``level_objects[0]`` are stationary
-    passive circles (static BFL).
+    passive circles (static BFL) in the integer-NODE frame, like Simulation
+    objects (see Hierarchy for frames).
 
     ``coarse_boundary`` is ``periodic`` (default) or ``slot_reservoir``:
     a bottom on-site slot and top fixed-density reservoir from ``cfg``.
@@ -131,7 +187,7 @@ class FixedPatchStepper:
     Periodic mode does not apply the wall/slot boundary fields in ``cfg``.
     Outer walls remain coarse; fine halo donors must be ordinary fluid.
 
-    Bodies are fixed circles in root (L0) physical coordinates; velocities are
+    Bodies of refined levels 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
@@ -241,40 +297,15 @@ class FixedPatchStepper:
         self.hierarchy = build_hierarchy(cfg, refinements, wall_aware=self._wall_aware)
         pairs = self.hierarchy.pairs
         if self._wall_aware:
-            rects = (pairs[0].bounds,) if isinstance(pairs[0], FixedPatch2D) else pairs[0].bounds
-            if any(r[1] < 4 or r[3] > cfg.ny - 4 for r in rects):
+            if any(r[1] < 4 or r[3] > cfg.ny - 4 for r in pairs[0].rectangles):
                 raise ValueError("slot_reservoir requires y0 >= 4 and y1 <= ny-4")
-        self._layouts = []
-        for level, block in enumerate(level_blocks):
-            if level == 0:
-                layout = BlockLayout2D.build(cfg.ny, cfg.nx, *block,
-                                             periodic=(True, not self._wall_aware))
-            elif isinstance(pairs[level - 1], FixedPatchSet2D):
-                layout = pairs[level - 1].fine_layout(block)
-            else:
-                layout = BlockLayout2D.build(*pairs[level - 1].fine_shape, *block,
-                                             periodic=(False, False))
-            if layout.nblocks > MAX_GRID_DIM or layout.local_shape[0] > MAX_GRID_DIM:
-                raise ValueError(f"level {level} blocks exceed launch limits")
-            self._layouts.append(layout)
+        self._layouts = _level_layouts(cfg, pairs, level_blocks, self._wall_aware)
         if self._channel:
-            base = blocked_base_tables(pairs[0], self._layouts[0], np.argsort(pairs[0].coalesce_coarse))
+            base = blocked_coarse_tables(pairs[0], self._layouts[0], np.argsort(pairs[0].coalesce_coarse))
             validate_channel_transfers(pairs[0], self._layouts[0], base[1])
         self._level_cfgs = _kernel_configs(cfg, pairs, self._wall_aware)
-        # Host boundary geometry per level in the root edge frame (Hierarchy.edge_frame).
-        self._boundaries = [None]
-        for level in range(1, depth):
-            pair = pairs[level - 1]
-            if level == 1:  # the root config carries the wall-aware boundary fields
-                host = prepare_fine_boundary(cfg, pair, level_objects[1],
-                                             wall_aware=self._wall_aware, jets=level_jets[1],
-                                             _shifted_blocked=self._ddf_shifting)
-            else:
-                origin, spacing = self.hierarchy.edge_frame(level - 1)
-                host = prepare_fine_boundary(self._level_cfgs[level - 1], pair, level_objects[level],
-                                             jets=level_jets[level], parent_origin=origin,
-                                             parent_spacing=spacing, _shifted_blocked=self._ddf_shifting)
-            self._boundaries.append(host)
+        hosts = _host_boundaries(self.hierarchy, self._level_cfgs, level_objects, level_jets,
+                                 self._wall_aware, self._ddf_shifting)
         if not isinstance(observe_raw_exchange, bool):
             raise ValueError("observe_raw_exchange must be bool")
         self._observe_raw_exchange = observe_raw_exchange
@@ -286,32 +317,25 @@ class FixedPatchStepper:
         self._window_valid = False
         self._window_start = None
         self._exchange_valid = False
-        self._exchange_spacing = tuple(self.hierarchy.edge_frame(l)[1] for l in range(depth))
-        self._exchange_centers = tuple(tuple(tuple(o.geometry.center) for o in objects)
-                                       for objects in level_objects)
-        self._level_jets = level_jets
         root = None
         if depth == 1 and level_objects[0]:
             tables = static_circle_tables(self._layouts[0], level_objects[0])
             root = _BlockedBoundary(None, tables, (), False)
-        self._level_boundaries = [root] + [
-            _prepare_blocked_boundary(self._boundaries[level], self._layouts[level],
+        # The one boundary/Jet record per level; host geometry and Jets are views of it.
+        self._boundaries = [root] + [
+            _prepare_blocked_boundary(hosts[level], self._layouts[level],
                                       self._level_cfgs[level], level_jets[level],
                                       _shifted_blocked=self._ddf_shifting)
             for level in range(1, depth)]
-        self._leaf = [kind == ACTIVE for kind in self.hierarchy.kinds]
-        if root is not None:
-            self._leaf[0] &= ~root.tables["solid"]
-        for level in range(1, depth):
-            host = self._boundaries[level]
-            if host is not None:
-                self._leaf[level] &= (host.flags.reshape(pairs[level - 1].fine_shape) & SOLID) == 0
+        for level, b in enumerate(self._boundaries):
+            if b is not None:  # exchange scaling (dt_l = h_l) and root-edge-frame moment centers
+                b.spacing = self.hierarchy.edge_frame(level)[1]
+                b.centers = tuple(tuple(o.geometry.center) for o in level_objects[level])
+        self._leaf = _fluid_leaves(self.hierarchy, self._boundaries)
         self._state = [mask.copy() for mask in self._leaf]
         if self._channel:
             self._leaf[0] &= channel_masks((cfg.ny, cfg.nx))[0]
-        # Private constructor inputs for remap_finest.
-        self._remap_cfg = replace(cfg)
-        self._level_blocks = level_blocks
+        self._root_cfg = replace(cfg)  # private constructor input for remap_finest
         self.ctx = context
         self._lock = RLock()
         self._status = "new"
@@ -366,7 +390,7 @@ class FixedPatchStepper:
         for index, layout in enumerate(self._layouts):
             total += 2 * 4 * int(np.prod(layout.blocked_shape))
             total += layout.halo_dst.nbytes + layout.halo_src.nbytes
-            b = self._level_boundaries[index]
+            b = self._boundaries[index]
             if b is not None:
                 total += sum(np.asarray(v).nbytes for k, v in b.tables.items()
                              if k == "flags" or k in BLOCKED_BOUNDARY_DEVICE_COLUMNS)
@@ -389,19 +413,19 @@ class FixedPatchStepper:
         grid = ((config.nx + self._block[0] - 1) // self._block[0], config.ny, layout.nblocks)
         level = _Level(shape, module, self._allocate(size), self._allocate(size),
                        module.get_function(name), grid, layout)
-        level.boundary = self._level_boundaries[index]
+        level.boundary = self._boundaries[index]
         level.halo_copy = module.get_function("BlockHaloCopy")
         if len(layout.halo_dst):
             level.halo_dst = self._upload(layout.halo_dst)
             level.halo_src = self._upload(layout.halo_src)
         if index == 0 and self._channel:
-            cfg = self._remap_cfg
+            cfg = self._root_cfg
             ptr, size = module.get_global("d_block_channel")
             if size != struct.calcsize("=IIf"):
                 raise RuntimeError("Block channel host/device layout mismatch")
             cuda.memcpy_htod(ptr, struct.pack("=IIf", cfg.nx, cfg.ny, cfg.velocity))
         if index == 0 and self._wall_aware:
-            cfg = self._remap_cfg
+            cfg = self._root_cfg
             ptr, size = module.get_global("d_block_wall")
             if size != struct.calcsize("=IIIf"):
                 raise RuntimeError("Block wall host/device layout mismatch")
@@ -415,11 +439,11 @@ class FixedPatchStepper:
         order = np.argsort(pair.coalesce_coarse)
         ghosts, children, child_kind = blocked_fine_tables(pair, child.layout, order)
         if index == 0:
-            donors, targets, parent_kind = blocked_base_tables(pair, parent.layout, order,
+            donors, targets, parent_kind = blocked_coarse_tables(pair, parent.layout, order,
                                                                self._wall_aware)
         else:
             # Nested pairs keep inherited INNER/OUTER and sparse holes of the parent.
-            donors, targets, parent_kind = blocked_base_tables(
+            donors, targets, parent_kind = blocked_coarse_tables(
                 pair, parent.layout, order, composed_kind=self.hierarchy.kinds[index])
         result = _Pair(child.module.get_function("BlockExplode"),
                        parent.module.get_function("BlockCoalesce"),
@@ -551,19 +575,17 @@ class FixedPatchStepper:
             pairs = self.hierarchy.pairs
             if (self._wall_aware or self._channel or self.hierarchy.depth != 3
                     or not isinstance(pairs[1], FixedPatchSet2D)
-                    or self._boundaries[1] is not None
-                    or any(b is not None for b in self._level_boundaries[:2])):
+                    or any(b is not None for b in self._boundaries[:2])):
                 raise ValueError("remap_finest requires a periodic three-level hierarchy with a "
                                  "level-2 set and no level-0/1 bodies")
-            cfg = self._remap_cfg
+            cfg = self._root_cfg
             objects = tuple(objects)
-            actual = self._level_boundaries[2]
+            actual = self._boundaries[2]
             jets = () if actual is None else actual.jets
             old = RemapTopology.from_config(cfg, pairs[0], pairs[1].bounds,
                                            child_objects=objects, child_jets=jets,
                                            _shifted_blocked=self._ddf_shifting)
-            if (not same_physical_boundary(old.boundary, self._boundaries[2])
-                    or (self._boundaries[2] is not None and actual is None)):
+            if not same_physical_boundary(old.boundary, None if actual is None else actual.host):
                 raise ValueError("objects must reproduce the original root-edge-frame level-2 boundary")
             if actual is not None:
                 velocity = actual.tables.get("surface_velocity")
@@ -580,7 +602,7 @@ class FixedPatchStepper:
                     or not np.array_equal(old.surface_velocity, new.surface_velocity)):
                 raise ValueError("remap requires the identical pinned physical boundary and BFL support")
             # Check sparse block alignment/launch size before any device read/allocation.
-            layout = new.child.fine_layout(self._level_blocks[2])
+            layout = new.child.fine_layout(self._layouts[2].block)
             if layout.nblocks > MAX_GRID_DIM or layout.local_shape[0] > MAX_GRID_DIM:
                 raise ValueError("level 2 blocks exceed launch limits")
             with self.ctx.activate():
@@ -598,7 +620,7 @@ class FixedPatchStepper:
                 phase="postcollision_g", **options)
             candidate = type(self)(
                 cfg, context=self.ctx, refinements=[pairs[0].bounds, new.child.bounds],
-                level_blocks=list(self._level_blocks),
+                level_blocks=[layout.block for layout in self._layouts],
                 level_objects=[(), (), objects], level_jets=[(), (), jets],
                 observe_raw_exchange=self._observe_raw_exchange,
                 observe_window_exchange=self._observe_window, explosion=self._explosion)
@@ -617,7 +639,7 @@ class FixedPatchStepper:
             return candidate
 
     def set_jet_amplitude(self, index, amplitude, *, level):
-        """Replace Jet ``index`` of ``level_jets[level]`` between steps; no geometry/table rebuild.
+        """Replace Jet ``index`` of level ``level`` (``jets[level]``) between steps; no table rebuild.
 
         Invalid input raises before any change. A device/sync error marks the
         instance failed (as step/initialize do); host state commits only after
@@ -628,27 +650,24 @@ class FixedPatchStepper:
             if (isinstance(level, bool) or not isinstance(level, Integral)
                     or not 1 <= level < self.hierarchy.depth):
                 raise ValueError("Jet level must name an existing refined level (1 .. depth-1)")
-            boundary = self._level_boundaries[level]
-            jets = list(self._level_jets[level])
+            boundary = self._boundaries[level]
+            # Host preparation: validate and build the new column without touching state.
+            jets = list(() if boundary is None else boundary.jets)
             jets[index] = jets[index].with_amplitude(amplitude)  # validates first
-            velocity = None
-            if boundary is not None and boundary.jet_cache is not None:
-                shape, jet_id = boundary.jet_cache
-                velocity = boundary.tables["surface_velocity"].copy()
-                driven = jet_id >= 0
-                velocity[driven] = jet_link_velocity(jets, shape, jet_id)[driven]
+            shape, jet_id = boundary.jet_cache
+            velocity = boundary.tables["surface_velocity"].copy()
+            driven = jet_id >= 0
+            velocity[driven] = jet_link_velocity(jets, shape, jet_id)[driven]
+            # Device action: an uncertain write fails the instance closed.
             try:
-                if velocity is not None:
-                    self._stream.synchronize()  # no in-flight step reads the column
-                    cuda.memcpy_htod(boundary.device["surface_velocity"], velocity)
+                self._stream.synchronize()  # no in-flight step reads the column
+                cuda.memcpy_htod(boundary.device["surface_velocity"], velocity)
             except BaseException:
                 self._status = "failed"
                 raise
-            if velocity is not None:
-                boundary.tables["surface_velocity"] = velocity
-            if boundary is not None:
-                boundary.jets = tuple(jets)
-            self._level_jets[level] = tuple(jets)
+            # Commit: the one boundary record holds both the column and the Jets.
+            boundary.tables["surface_velocity"] = velocity
+            boundary.jets = tuple(jets)
 
     def _blocked_boundary(self, level, buffer, **kw):
         b = level.boundary
@@ -691,7 +710,7 @@ class FixedPatchStepper:
         b = level.boundary
         if b is not None and b.last_raw is not None:
             count = len(b.tables["cell"])
-            h = self._exchange_spacing[index]  # acoustic scaling: dt_l = h_l
+            h = b.spacing  # acoustic scaling: dt_l = h_l
             b.accumulate(b.last_raw, b.interval_integral, np.uint32(count),
                          np.float64(h), np.float64(h),
                          grid=((count + self._block[0] - 1) // self._block[0], 1), **kw)
@@ -735,7 +754,7 @@ class FixedPatchStepper:
                 self._exchange_valid = False
                 for _ in range(n):
                     if self._observe_raw_exchange:  # stream-ordered: holds the last interval only
-                        for b in self._level_boundaries:
+                        for b in self._boundaries:
                             if b is not None:
                                 cuda.memset_d8_async(b.interval_integral, 0,
                                                      len(b.tables["cell"]) * 32, self._stream)
@@ -758,7 +777,7 @@ class FixedPatchStepper:
                 raise RuntimeError("window exchange observation is disabled")
             self._window_valid = False
             try:
-                for b in self._level_boundaries:
+                for b in self._boundaries:
                     if b is not None and b.window_integral is not None:
                         cuda.memset_d8(b.window_integral, 0, len(b.tables["cell"]) * 32)
             except BaseException:
@@ -771,7 +790,7 @@ class FixedPatchStepper:
         if (isinstance(level, bool) or not isinstance(level, Integral)
                 or not 0 <= level < len(self._levels)):
             raise ValueError("level must name an existing level")
-        b = self._level_boundaries[level]
+        b = self._boundaries[level]
         if (b is None or isinstance(body_id, bool) or not isinstance(body_id, Integral)
                 or not 0 <= body_id < len(b.host.objects)):
             raise ValueError("body_id must name a body on this level")
@@ -779,9 +798,9 @@ class FixedPatchStepper:
 
     def _summarize(self, b, rows, level, body_id, window, **kw):
         return summarize_exchange(rows, b.tables["body_id"], b.tables["aperture"],
-                                  body_id=body_id, level=level, h=self._exchange_spacing[level],
+                                  body_id=body_id, level=level, h=b.spacing,
                                   time_end=self._step_count, window=window,
-                                  reference_center=self._exchange_centers[level][body_id], **kw)
+                                  reference_center=b.centers[body_id], **kw)
 
     def read_exchange_window(self, body_id, *, level):
         """Non-destructive read of [window start, current step count] in base-lattice units."""
@@ -868,7 +887,7 @@ class FixedPatchStepper:
                         self._stream.synchronize()
                 except Exception as exc:
                     error = exc
-                for b in getattr(self, "_level_boundaries", ()):
+                for b in getattr(self, "_boundaries", ()):
                     if b is not None:
                         b.flags = b.device = b.kernel = None
                         b.last_raw = b.interval_integral = b.accumulate = None
@@ -896,9 +915,20 @@ class FixedPatchStepper:
                 raise
 
     # Read-only per-level state, indexed by level (0 = root). Pairs: hierarchy.pairs[l-1].
-    layouts = property(lambda self: tuple(self._layouts))
-    boundaries = property(lambda self: tuple(self._boundaries),
-                          doc="Host boundary geometry per level (None: no bodies/Jets).")
-    jets = property(lambda self: tuple(self._level_jets))
-    leaf = property(lambda self: tuple(m.copy() for m in self._leaf),
-                    doc="Per-level (ny, nx) masks of public (non-NaN) download cells.")
+    @property
+    def layouts(self):
+        return tuple(self._layouts)
+
+    @property
+    def boundaries(self):
+        """Host boundary geometry per level (None: no bodies/Jets on a refined level)."""
+        return tuple(None if b is None else b.host for b in self._boundaries)
+
+    @property
+    def jets(self):
+        return tuple(() if b is None else b.jets for b in self._boundaries)
+
+    @property
+    def leaf(self):
+        """Per-level (ny, nx) masks of public (non-NaN) download cells."""
+        return tuple(m.copy() for m in self._leaf)

src/CelerisLab/lbm/refinement_tables.py

@@ -36,7 +36,7 @@ def blocked_kind(layout, dense_kind):
     kind[owner[owned]] = np.asarray(dense_kind).ravel()[owned]
     return kind
 
-def blocked_base_tables(patch, layout, order, wall_aware=False, *, composed_kind=None):
+def blocked_coarse_tables(patch, layout, order, wall_aware=False, *, composed_kind=None):
     """Parent donor cells and receiver-local arrival slots, in transfer row order.
 
     Stage arrivals only after child reads and parent halo gather, immediately
@@ -130,6 +130,8 @@ class _BlockedBoundary:
     interval_integral: object = None
     accumulate: object = None
     window_integral: object = None  # optional observation-window FP64 ledger
+    spacing: float = None  # level spacing h_l in root units (exchange scaling)
+    centers: tuple = ()    # body reference centers in the root edge frame
 
 
 def _prepare_blocked_boundary(boundary, layout, cfg, jets, *, _shifted_blocked=False):

src/CelerisLab/lbm/stepper.py

@@ -20,6 +20,10 @@ class LBMStepper:
         self.cfg = cfg
 
         self._esopull = (cfg.streaming == "esopull")
+        # Set by the owner when a step would be wrong (translating circles
+        # excluded from the packed links before the moving kernels exist).
+        self.blocked: str | None = None
+        self.moving = None  # MovingTables attached by the owner (translating circles)
 
         if self._esopull:
             self.step_fn = module.get_function("EsoPullStep")
@@ -35,6 +39,8 @@ class LBMStepper:
         self.apply_esopull_curved_fn = module.get_function("ApplyEsoPullCurvedKernel")
         self.sensor_fn = module.get_function("SensorKernel")
         self.force_region_fn = module.get_function("ForceRegionKernel")
+        self.moving_fn = None if self._esopull else module.get_function("MovingCircleLinks")
+        self.moving_cells_fn = None if self._esopull else module.get_function("MovingCircleCells")
 
         tpb = cfg.threads_per_block
         self.block = (tpb, 1, 1)
@@ -70,6 +76,13 @@ class LBMStepper:
         stream: cuda.Stream | None = None,
     ):
         """Advance *n* time steps."""
+        if n > 0 and self.blocked is not None:
+            raise RuntimeError(self.blocked)
+        self._step(n, action_gpu=action_gpu, obs_gpu=obs_gpu, stream=stream)
+
+    def _step(self, n, *, action_gpu, obs_gpu, stream=None, advance=False):
+        """Unguarded steps; ``advance`` runs one step that moves translating circles."""
+        assert not advance or n == 1
         f = self.field
         launch_kw = {}
         if stream is not None:
@@ -81,6 +94,8 @@ class LBMStepper:
                     f.ddf_gpu, self.boundary_values_gpu, np.uint64(self._step_count),
                     block=self.block, grid=self.boundary_grid, **launch_kw,
                 )
+            if self.moving is not None:
+                self._launch_moving(action_gpu, obs_gpu, advance, **launch_kw)
             self._launch_curved(action_gpu, obs_gpu, **launch_kw)
 
             if self._esopull:
@@ -106,6 +121,21 @@ class LBMStepper:
             # Any completed GPU step invalidates host DDF mirror.
             f.invalidate_host_ddf_cache()
 
+    def _launch_moving(self, action_gpu, obs_gpu, advance, **launch_kw):
+        m, tpb = self.moving, self.cfg.threads_per_block
+        if advance:  # M1 completes before M2 on the same stream
+            self.moving_cells_fn(
+                self.field.ddf_gpu, self.field.flag_gpu, m.band_gpu, m.geom_gpu,
+                np.uint32(m.n_bodies), np.uint32(m.n_cells), action_gpu, m.diag_gpu,
+                block=(tpb, 1, 1), grid=((m.n_cells + tpb - 1) // tpb, 1, 1), **launch_kw,
+            )
+        self.moving_fn(
+            self.field.ddf_gpu, self.field.flag_gpu, m.band_gpu, m.geom_gpu,
+            np.uint32(m.n_bodies), np.uint32(m.n_slots), action_gpu, obs_gpu,
+            m.diag_gpu, m.slot_debug_gpu if m.slot_debug_gpu is not None else np.uintp(0),
+            block=(tpb, 1, 1), grid=((m.n_slots + tpb - 1) // tpb, 1, 1), **launch_kw,
+        )
+
     def _launch_curved(self, action_gpu, obs_gpu, **launch_kw):
         f = self.field
         if f.n_curved == 0:

src/CelerisLab/simulation.py

@@ -26,6 +26,13 @@ import pycuda.driver as cuda
 from .config import LBMConfig, BodyConfig, load_lbm_config, load_body_config
 from .cuda.context import CudaContext, context_bound
 from .cuda import compiler
+from .lbm.descriptors import FLUID
+
+# Translating circles step only through their moving tables (fixed centre in S2a).
+PENDING_BLOCKED = "a pending translation target is advanced only by Simulation.run(1)"
+MOVING_BLOCKED = ("translating circles need the moving kernels and their tables (attached "
+                  "at initialize); "
+                  "positive-step runs are rejected")
 from .lbm.field import LBMField
 from .lbm.stepper import LBMStepper
 from .body.objects import SimObject
@@ -97,7 +104,8 @@ class Simulation:
 
     # -- Object management ---------------------------------------------------
     def add_body(self, type: str = "circle", center=None,
-                 radius: float = 0.0, *, surface_actuation=None, **kwargs) -> int:
+                 radius: float = 0.0, *, surface_actuation=None, translation=None,
+                 **kwargs) -> int:
         """Add a simulation body. Returns body_id.
 
         Before ``initialize()`` the object is added to the formal registry
@@ -113,6 +121,9 @@ class Simulation:
                 angle, half_width (radians), normal_velocity (positive outward).
                 D2Q9/FP32/unshifted/DB only; set before initialization or stage
                 a new body. No runtime motion or amplitude setter.
+            translation: ``body.motion.Translation`` declaring a translating
+                circle (before initialize only; docs/moving_bodies.md). Until
+                the moving kernels exist, positive-step runs are rejected.
             **kwargs: passed to the geometry constructor.
 
         Returns:
@@ -123,6 +134,10 @@ class Simulation:
         type_lower = str(type).strip().lower()
         if surface_actuation is not None and type_lower != "circle":
             raise ValueError("Surface actuation is only supported on circles")
+        if translation is not None and (type_lower != "circle" or surface_actuation is not None
+                                        or self._initialized):
+            raise ValueError("translation needs a circle without surface actuation, "
+                             "added before initialize()")
 
         if type_lower == "circle":
             from .body.geometry.circle import CircleGeometry
@@ -131,6 +146,11 @@ class Simulation:
             obj = SimObject(obj_id=-1, geometry=geom,
                             center=center, radius=radius,
                             is_sensor=False, surface_actuation=surface_actuation)
+            if translation is not None:
+                from .body.motion import MovingCircle
+                obj.motion = MovingCircle((cx, cy), radius, translation,
+                                          on_failure=self._fail_moving)
+                self.stepper.blocked = MOVING_BLOCKED
         elif type_lower == "sensor":
             from .body.geometry.circle import CircleGeometry
             cx, cy = float(center[0]), float(center[1])
@@ -190,6 +210,78 @@ class Simulation:
                 raise
         self.bodies.jets[index] = jet  # later geometry syncs keep the new amplitude
 
+    def _fail_moving(self, exc: BaseException) -> None:
+        # Device geometry is unknown after a failed moving step: same lifecycle as Jets.
+        self._failed = f"moving circle step failed: {exc!r}"
+        self.stepper.blocked = self._failed  # direct stepper calls are refused as well
+
+    def _attach_moving(self) -> None:
+        # Direct stepper calls stay blocked until the moving tables exist.
+        tables = getattr(self, "_moving_tables", None)
+        self.stepper.moving = tables
+        moving = self._moving_bodies()
+        if moving and tables is None:
+            self.stepper.blocked = MOVING_BLOCKED
+        elif any(obj.motion.pending is not None for obj in moving):
+            self.stepper.blocked = PENDING_BLOCKED  # direct steps never consume a target
+        else:
+            self.stepper.blocked = None
+
+    def _build_moving_tables(self):
+        moving = self._moving_bodies()
+        if not moving:
+            return None
+        from .body.motion import MovingTables
+        return MovingTables.upload([(obj.motion, obj.obj_id) for obj in moving])
+
+    @context_bound
+    def moving_diagnostics(self) -> np.ndarray:
+        """Per translating circle (in body order), columns ``MovingTables.DIAG``,
+        accumulated since obs were last zeroed (diagnostics only, not a force RHS)."""
+        self._require_usable()
+        tables = self.stepper.moving
+        if tables is None:
+            return np.zeros((0, 6), np.float32)
+        self.stream.synchronize()
+        return tables.download_diagnostics()
+
+    def _moving_bodies(self):
+        return [obj for obj in self.bodies.objects if obj.motion is not None]
+
+    def _reject_moving_bodies(self, operation: str) -> None:
+        if self._moving_bodies():
+            raise ValueError(f"{operation} is not supported with translating circles")
+
+    def _validate_moving_bodies(self) -> None:
+        """Reject unsupported modes and swept-band placement before any mutation."""
+        moving = self._moving_bodies()
+        if not moving:
+            return
+        cfg = self.lbm_cfg
+        if (not cfg.is_d2q9 or cfg.streaming != "double_buffer" or cfg.ddf_shifting
+                or cfg.store_precision != "FP32"):
+            raise ValueError("translating circles need D2Q9/FP32/unshifted/double_buffer")
+        if self.bodies.jets:
+            raise ValueError("Jets are not supported together with translating circles")
+        nx, ny = self.field.nx, self.field.ny
+        base = self.field.build_channel_flags().reshape(ny, nx)
+        for obj in moving:
+            x0, x1, y0, y1 = obj.motion.band()
+            if x0 < 1 or y0 < 1 or x1 > nx - 2 or y1 > ny - 2:
+                raise ValueError(f"swept band {(x0, x1, y0, y1)} touches the domain edge")
+            if not np.all((base[y0:y1 + 1, x0:x1 + 1] & FLUID) != 0):
+                raise ValueError(f"swept band {(x0, x1, y0, y1)} contains non-fluid base cells")
+            for other in self.bodies.objects:
+                if other is obj:
+                    continue
+                (cx, cy), r = other.center[:2], float(other.radius)
+                if other.motion is not None:
+                    ox0, ox1, oy0, oy1 = other.motion.band()
+                else:
+                    ox0, ox1, oy0, oy1 = cx - r - 2, cx + r + 2, cy - r - 2, cy + r + 2
+                if ox0 <= x1 and x0 <= ox1 and oy0 <= y1 and y0 <= oy1:
+                    raise ValueError(f"swept band {(x0, x1, y0, y1)} overlaps body {other.obj_id}")
+
     def _require_usable(self):
         failed = getattr(self, "_failed", None)
         if failed is not None:
@@ -260,7 +352,7 @@ class Simulation:
 
     # -- Runtime control -----------------------------------------------------
     def set_body(self, id: int, omega: float = None,
-                 vx: float = None, vy: float = None) -> None:
+                 vx: float = None, vy: float = None, offset=None) -> None:
         """Set runtime body state (host action array only, no H2D).
 
         The GPU action buffer is updated at the next ``run()`` call.
@@ -270,8 +362,21 @@ class Simulation:
             omega: rotation speed (lattice units).  ``None`` = no change.
             vx: translation velocity x (future).  ``None`` = no change.
             vy: translation velocity y (future).  ``None`` = no change.
+            offset: target centre offset of a translating circle; with vx/vy
+                (default 0) it replaces the pending target after validation
+                against the applied offset (docs/moving_bodies.md, 4a). Only
+                ``run(1)`` advances it; direct stepper calls are blocked meanwhile.
         """
         obj = self.bodies.get(id)
+        motion = obj.motion
+        if motion is not None:
+            if offset is None and (vx is not None or vy is not None):
+                raise ValueError("a translating circle takes vx/vy only with an offset")
+            target = None if offset is None else motion.check_target(
+                offset, (0.0 if vx is None else vx, 0.0 if vy is None else vy))
+            vx = vy = None  # the motion target, not ObjectState, owns translation
+        elif offset is not None:
+            raise ValueError(f"body {id} was not declared with a translation")
         proposed = dict(vars(obj.state))
         proposed.update({k: float(v) for k, v in (("omega", omega), ("vx", vx), ("vy", vy))
                          if v is not None})
@@ -287,6 +392,9 @@ class Simulation:
         if vy is not None:
             obj.state.vy = float(vy)
             changed = True
+        if motion is not None and target is not None:
+            motion.pending = target
+            self._attach_moving()
         if changed:
             self.bodies._refresh_action_from_objects()
 
@@ -443,6 +551,7 @@ class Simulation:
             self.field, self._module, self.lbm_cfg,
         )
         self.stepper._step_count = _prev_step_count
+        self._attach_moving()
 
         self._assert_object_count_contract(expected_count=expected_n_objects)
 
@@ -468,6 +577,8 @@ class Simulation:
         # Validate the complete geometry before flags, populations or modules
         # change. Compact-list construction publishes no outputs on failure.
         next_flags = self.bodies.build_flags(self.field.build_channel_flags())
+        self._validate_moving_bodies()
+        self._attach_moving()
         if self.bodies.jets and self.bodies.count == 0:
             raise ValueError("Jets need bodies with cut links")
         self.bodies.build_compact_lists_for(self.bodies.objects, next_flags)
@@ -484,6 +595,8 @@ class Simulation:
             self.bodies.sync_to_gpu(self.field, rebuild_flags=False)
         else:
             self.bodies.ensure_packed_buffers()
+        self._moving_tables = self._build_moving_tables()
+        self._attach_moving()
         self._assert_runtime_contracts()
         self._initialized = True
 
@@ -499,6 +612,7 @@ class Simulation:
         self._require_usable()
         if not self._initialized:
             raise RuntimeError("Call initialize() before sync_bodies()")
+        self._reject_moving_bodies("sync_bodies()")
 
         self._assert_allocation_contract()
         # 1. Synchronize GPU
@@ -599,6 +713,13 @@ class Simulation:
         checkpoint_interval = operator.index(checkpoint_interval)
         if steps < 0 or checkpoint_interval < 0:
             raise ValueError("steps and checkpoint_interval must be nonnegative")
+        moving = self._moving_bodies()
+        if steps > 0 and moving and self.stepper.moving is None:
+            raise RuntimeError(MOVING_BLOCKED)
+        if moving and checkpoint_interval > 0:
+            raise ValueError("checkpoints are not supported with translating circles")
+        if steps > 1 and any(obj.motion.pending is not None for obj in moving):
+            raise ValueError("a pending translation target needs run(1)")
         self._assert_runtime_contracts()
         # Discard any uncommitted body edits before stepping.
         if self.bodies.has_pending_edit():
@@ -615,6 +736,8 @@ class Simulation:
         # Zero obs segments before step group
         if zero_obs:
             self.bodies.zero_obs_async(stream)
+            if self.stepper.moving is not None:
+                self.stepper.moving.zero_diagnostics(stream)
 
         if checkpoint_interval > 0:
             done = 0
@@ -631,6 +754,9 @@ class Simulation:
                 if done < steps or done % checkpoint_interval == 0:
                     stream.synchronize()
                     self.save_checkpoint()
+        elif self.stepper.moving is not None and steps > 0:
+            self._run_moving(steps, stream, sync_obs)
+            return
         else:
             self.stepper.step(
                 steps,
@@ -649,6 +775,34 @@ class Simulation:
 
         stream.synchronize()
 
+    def _run_moving(self, steps, stream, sync_obs):
+        """Steps with translating circles: hold, or one advancing step that commits
+        each pending target only after synchronization and a zero ``no_donor`` count
+        (docs/moving_bodies.md 4a, 5). Any failure leaves the Simulation unusable."""
+        tables = self.stepper.moving
+        circles = [(obj.motion, obj.obj_id) for obj in self._moving_bodies()]
+        try:
+            advancing = tables.upload_records(circles, stream)
+            self.stepper._step(steps, action_gpu=self.bodies.action_gpu,
+                               obs_gpu=self.bodies.obs_gpu, stream=stream, advance=advancing)
+            self.bodies.increment_obs_steps(steps)
+            if sync_obs:
+                self.bodies.download_obs_full_async(stream)
+            stream.synchronize()
+            if advancing:
+                no_donor = tables.download_diagnostics()[:, tables.DIAG.index("no_donor")]
+                if np.any(no_donor > 0):
+                    raise RuntimeError("translating circle refill found no donor cell")
+                for motion, _ in circles:
+                    motion.commit()
+                tables.upload_records(circles, stream)  # first hold: old = new = applied, u = 0
+                stream.synchronize()
+                self._attach_moving()
+        except BaseException as exc:
+            for motion, _ in circles:  # device geometry/state unknown after any failure
+                motion.fail(exc)
+            raise
+
     def step(self, n: int = 1):
         """Advance *n* steps (convenience for interactive use)."""
         self.run(n)
@@ -666,8 +820,16 @@ class Simulation:
         self.field.download_ddf(step_id=self.stepper.step_count)
         return self.field.ddf.copy()
 
+    @context_bound
     def get_flags(self) -> np.ndarray:
-        return self.field.flag.copy()
+        """Flag array; read back from the device when translating circles may have moved."""
+        if getattr(self.stepper, "moving", None) is None:
+            return self.field.flag.copy()
+        self._require_usable()
+        self.stream.synchronize()
+        flags = np.empty_like(self.field.flag)
+        cuda.memcpy_dtoh(flags, self.field.flag_gpu)
+        return flags
 
     # -- Runtime parameter updates -------------------------------------------
     @context_bound
@@ -698,6 +860,7 @@ class Simulation:
         """Save full simulation state to an HDF5 file."""
         self._require_usable()
         from .common.checkpoint import save_checkpoint
+        self._reject_moving_bodies("save_checkpoint()")
         if self.bodies.jets:
             raise NotImplementedError("Jets are not checkpointed yet")
         return save_checkpoint(self.field, self.stepper, self.lbm_cfg,
@@ -708,6 +871,7 @@ class Simulation:
         """Restore simulation state from an HDF5 checkpoint."""
         self._require_usable()
         from .common.checkpoint import load_checkpoint
+        self._reject_moving_bodies("load_checkpoint()")
         if not self._initialized:
             self.initialize()
         load_checkpoint(path, self.field, self.stepper, self.lbm_cfg,
@@ -726,6 +890,7 @@ class Simulation:
                 cuda.Context.synchronize()
                 # Release owned GPU resources while their context is current.
                 self.stepper = None
+                self._moving_tables = None
                 self.bodies = None
                 self.field = None
                 self._module = None

tests/unit/test_channel_root.py

@@ -7,7 +7,7 @@ from CelerisLab.config import LBMConfig
 from CelerisLab.lbm.blocks import BlockLayout2D
 from CelerisLab.lbm.channel import channel_masks, validate_channel_config, validate_channel_transfers
 from CelerisLab.lbm.refinement import FixedPatch2D
-from CelerisLab.lbm.refinement_tables import blocked_base_tables
+from CelerisLab.lbm.refinement_tables import blocked_coarse_tables
 from .test_numerical_operators import compile_run
 
 
@@ -31,7 +31,7 @@ def transfer(bounds, block=(4, 6)):
     cfg = LBMConfig(nx=24, ny=20, viscosity=.1)
     pair = FixedPatch2D.from_config(cfg, bounds)
     layout = BlockLayout2D.build(20, 24, *block)
-    targets = blocked_base_tables(pair, layout, np.argsort(pair.coalesce_coarse))[1]
+    targets = blocked_coarse_tables(pair, layout, np.argsort(pair.coalesce_coarse))[1]
     return pair, layout, targets
 
 

tests/unit/test_duetsch_protocol.py

@@ -0,0 +1,191 @@
+"""CPU tests of the Dütsch protocol arithmetic (docs/validation/duetsch_re100_kc5_protocol.md §6a)."""
+import importlib.util
+import math
+from pathlib import Path
+
+import numpy as np
+import pytest
+
+_spec = importlib.util.spec_from_file_location(
+    "duetsch_protocol", Path(__file__).resolve().parents[2] / "research/duetsch_re100_kc5/protocol.py")
+p = importlib.util.module_from_spec(_spec)
+_spec.loader.exec_module(p)
+
+T = 2520
+
+
+def test_phase_and_cycle_bookkeeping():
+    assert p.phase(0, T) == 270.0 and p.cycle(0, T) == 0          # +A, at rest, phi = -90
+    assert p.phase(T // 4, T) == 0.0 and p.cycle(T // 4, T) == 1  # mean position, u = -U_max
+    assert p.position(0, T, 1.0) == 1.0
+    assert p.cycle(T // 4 - 1, T) == 0 and p.cycle(T // 4 + T - 1, T) == 1
+    assert p.cycle(T // 4 + T, T) == 2
+    assert p.phase(T // 6, T) == pytest.approx(330.0) and p.cycle(T // 6, T) == 0  # startup quarter
+    assert p.phase(7 * T // 6, T) == pytest.approx(330.0) and p.cycle(7 * T // 6, T) == 1
+    for ph, frac in ((180.0, .75), (210.0, 5 / 6)):
+        s = round(frac * T) + T
+        assert p.phase(s, T) == pytest.approx(ph) and p.bin_phase(s, T) == ph
+    assert p.bin_phase(round(.75 * T) + T - 1, T) is None
+    assert sum(p.bin_phase(s, T) == 180.0 for s in range(T, 2 * T)) == T // 360
+
+
+def test_settling_and_common_window():
+    passes = {k: True for k in range(3, 12)}
+    assert p.settled_start(passes) == 7                          # pairs 3->4, 4->5, 5->6
+    late = {**passes, 5: False}
+    assert p.settled_start(late) == 9                            # pairs 6->7, 7->8, 8->9
+    assert p.common_window([7, 9], [passes, passes]) == (9, 10, 11)
+    assert p.common_window([7, 10], [passes, passes]) is None    # K + 2 = 12 beyond the cap
+    assert p.common_window([7, None], [passes, passes]) is None
+    assert p.common_window([7, 7], [passes, {**passes, 8: False}]) is None  # must stay settled
+    # a case settled early must keep passing while waiting for the common K
+    assert p.common_window([7, 9], [{**passes, 8: False}, passes]) is None
+    long = {k: True for k in range(3, 32)}
+    assert p.common_window([10], [long]) is None                 # default cap 11
+    assert p.common_window([29], [long], last_scored=31) == (29, 30, 31)
+    assert p.common_window([30], [long], last_scored=31) is None
+
+
+def test_solid_aware_bilinear():
+    field = np.arange(16.0).reshape(4, 4)
+    fluid = np.ones((4, 4), bool)
+    value, used, total = p.bilinear_fluid(field, fluid, 1.25, 1.5)
+    assert total == pytest.approx(1.0) and value == pytest.approx(0.75 * 0.5 * (5 + 9) + 0.25 * 0.5 * (6 + 10))
+    fluid[1, 1] = False                                           # one solid corner
+    value, used, total = p.bilinear_fluid(field, fluid, 1.25, 1.5)
+    assert total == pytest.approx(1.0 - 0.75 * 0.5)
+    assert all(not (i == 1 and j == 1) for i, j, _ in used)
+    assert sum(w for *_, w in used) == pytest.approx(1.0)
+    fluid[:] = False
+    assert p.bilinear_fluid(field, fluid, 1.25, 1.5) == (None, (), 0.0)
+
+
+def test_interval_bases_and_morison_fit_recover_known_coefficients():
+    b = np.array([p.interval_bases(n, T) for n in range(T // 4, T // 4 + 3 * T)])
+    phi = 2 * np.pi * (np.arange(T // 4, T // 4 + 3 * T) + .5) / T - np.pi / 2
+    assert np.allclose(b[:, 1], np.sin(phi), atol=1e-5)          # midpoint ~ interval mean
+    f = 2.09 * b[:, 0] - (math.pi ** 2 * 1.45 / p.KC) * b[:, 1]
+    assert p.morison_fit(f, b[:, 0], b[:, 1]) == pytest.approx((2.09, 1.45), abs=1e-9)
+
+
+def test_geometry_repeats_exactly_across_cycles():
+    A = p.amplitude(20)
+    for m in (0, 1, T // 4, T // 2 - 1, 7 * T // 12):
+        for c in (1, 5, 11):
+            assert p.position(m + c * T, T, A) == p.position(m, T, A)
+            assert p.phase(m + c * T, T) == p.phase(m, T)
+            v0 = p.position(m + 1, T, A) - p.position(m, T, A)
+            assert p.position(m + 1 + c * T, T, A) - p.position(m + c * T, T, A) == v0
+    assert p.cycle(T // 4 + 5 * T, T) == 6  # unwrapped index stays separate
+    assert p.interval_bases(123 + 7 * T, T) == p.interval_bases(123, T)
+
+
+def test_exact_point_inside_body_is_missing_even_with_fluid_corners():
+    field = np.zeros((6, 6))
+    fluid = np.ones((6, 6), bool)
+    fluid[3, 3] = False                         # only one corner solid
+    centre, radius = (3.0, 3.0), 0.3
+    value, used, total = p.sample(field, fluid, 3.2, 3.1, centre, radius)   # inside the circle
+    assert value is None and used == () and total == 0.0
+    value, used, total = p.sample(field, fluid, 3.35, 3.0, centre, radius)  # just outside
+    assert value is not None and total > 0.0
+
+
+def test_generated_config_passes_the_real_loader(tmp_path):
+    import json
+    import sys
+    sys.path.insert(0, str(Path(__file__).resolve().parents[2] / "research/duetsch_re100_kc5"))
+    try:
+        import run
+    finally:
+        sys.path.pop(0)
+    from CelerisLab.config import load_lbm_config
+    cfg_dict, u_max, n = run.config(20, 10.0, 2520)
+    path = tmp_path / "config_lbm.json"
+    path.write_text(json.dumps(cfg_dict))
+    cfg = load_lbm_config(str(path))
+    assert (cfg.nx, cfg.ny) == (n, n) == (200, 200)
+    assert cfg.collision == "MRT" and cfg.streaming == "double_buffer" and not cfg.ddf_shifting
+    assert cfg.velocity == 0.0 and cfg.viscosity == pytest.approx(u_max * 20 / 100)
+    assert cfg.y_wall_bc == "bounce_back"
+
+
+def _analyze():
+    import sys
+    sys.path.insert(0, str(Path(__file__).resolve().parents[2] / "research/duetsch_re100_kc5"))
+    try:
+        import analyze
+    finally:
+        sys.path.pop(0)
+    return analyze
+
+
+def test_hf_energy_is_parseval_weighted():
+    a = _analyze()
+    n = 360
+    t = np.arange(n)
+    x = 1.0 + np.cos(2 * np.pi * 5 * t / n) + 2.0 * np.cos(2 * np.pi * 30 * t / n) + 0.5 * np.cos(np.pi * t)
+    e = a.one_sided_energy(np.fft.rfft(x), n) / n
+    assert e.sum() == pytest.approx(np.sum(x ** 2))               # Parseval
+    expected = (2.0 ** 2 / 2 + 0.5 ** 2) / np.mean(x ** 2)       # bins >= 21 (30 and Nyquist)
+    assert e[21:].sum() / e.sum() == pytest.approx(expected)
+
+
+def test_horizon_excludes_future_coverage_and_events(monkeypatch):
+    a = _analyze()
+    T, horizon = 360, 4
+    steps = round((0.25 + 6) * T)                                 # cycles 1..6 recorded
+    diag = np.zeros((steps, 15))
+    diag[round((0.25 + horizon) * T):, 13] = 1.0                  # synthetic future no_donor
+    rows = [(s, p.cycle(s, T), 180.0, 0.0, 0.0, 0.1, 0.0, 1.0) for s in range(steps)
+            if p.bin_phase(s, T) is not None and p.cycle(s, T) > p.DISCARD]
+    case = dict(path="c", D=20, T=T, box=10, n=200, u_max=.04, nu=.008, tau=.524, A=1.0,
+                steps=steps, cycles=6, F_ref=1.0, force=np.zeros((steps, 2)), diag=diag,
+                samples=np.asarray(rows, float))
+    monkeypatch.setattr(a, "load", lambda path: case)
+    out = []
+    monkeypatch.setattr("builtins.print", lambda text: out.append(text))
+    a.main(["c", "--horizon", str(horizon)])
+    import json
+    entry = json.loads(out[0])["cases"][0]
+    assert sorted(int(k) for k in entry["coverage"]) == [3, 4]   # nothing from cycles 5, 6
+    assert entry["no_donor"] == 0.0
+
+
+
+def _capture():
+    spec = importlib.util.spec_from_file_location(
+        "duetsch_capture", Path(__file__).resolve().parents[2] / "research/duetsch_re100_kc5/capture.py")
+    mod = importlib.util.module_from_spec(spec)
+    spec.loader.exec_module(mod)
+    return mod
+
+
+def test_capture_crop_downsample_and_trends():
+    c = _capture()
+    n, D = 2400, 40
+    xm = ym = 0.5 * (n - 1)
+    A, R = p.amplitude(D), 20.0
+    box = c.crop_box(xm, ym, D, n, (xm - A - R, xm + A + R, ym - R, ym + R))
+    assert box == (880, 1520, 960, 1440) and box[1] - box[0] + 1 == 641 and box[3] - box[2] + 1 == 481
+    with pytest.raises(ValueError):
+        c.crop_box(xm, ym, D, 600, (0, 1, 0, 1))                    # leaves the domain
+    with pytest.raises(ValueError):
+        c.crop_box(xm, ym, D, n, (xm - 9 * D, xm, ym, ym))           # clips the swept body
+    field = np.arange(64.0).reshape(8, 8)
+    fluid = np.ones((8, 8), bool)
+    fluid[:4, :4] = False
+    down = c.downsample(field, fluid, 4)
+    assert np.isnan(down[0, 0]) and down[1, 1] == field[4:, 4:].mean()
+    rho = np.ones((n, n)); rho[:, 1:81] = 1.01
+    row = c.trend_row(10, 3, 180.0, rho, np.zeros((n, n)), np.zeros((n, n)),
+                      np.ones((n, n), bool), D, n, box)
+    assert row[7] == pytest.approx(1.01) and row[9] == pytest.approx(1.0)
+
+
+def test_dump_steps_are_the_first_step_of_each_bin():
+    T = 5040
+    for k in (11, 21, 30, 31):
+        starts = [s for s in range(T // 4 + (k - 1) * T, T // 4 + k * T)
+                  if p.bin_phase(s, T) is not None and p.phase(s, T) == p.bin_phase(s, T)]
+        assert len(starts) == 3 and all(p.cycle(s, T) == k for s in starts)

tests/unit/test_jet_update_failure.py

@@ -102,7 +102,7 @@ def test_simulation_before_initialize_commits_host_only(monkeypatch):
 
 
 def stepper_state(run):
-    boundary = run._level_boundaries[2]
+    boundary = run._boundaries[2]
     return boundary.tables["surface_velocity"].copy(), boundary.jets, run.jets[2]
 
 
@@ -116,8 +116,8 @@ def test_stepper_invalid_input_writes_nothing_and_stays_ready():
             with pytest.raises(error):
                 run.set_jet_amplitude(*args, **kw)
         module.cuda.memcpy_htod.assert_not_called()
-        np.testing.assert_array_equal(run._level_boundaries[2].tables["surface_velocity"], velocity)
-        assert (run._level_boundaries[2].jets, run.jets[2]) == (jets, level_jets)
+        np.testing.assert_array_equal(run._boundaries[2].tables["surface_velocity"], velocity)
+        assert (run._boundaries[2].jets, run.jets[2]) == (jets, level_jets)
         assert run._status == "ready"
         run.set_jet_amplitude(0, .017, level=2)
         assert run.jets[2][0].amplitude == .017
@@ -136,8 +136,8 @@ def test_stepper_device_failure_keeps_host_and_blocks_reuse(where):
             run._stream.synchronize.side_effect = RuntimeError("upload failed")
         with pytest.raises(RuntimeError, match="upload failed"):
             run.set_jet_amplitude(0, .017, level=2)
-        np.testing.assert_array_equal(run._level_boundaries[2].tables["surface_velocity"], velocity)
-        assert (run._level_boundaries[2].jets, run.jets[2]) == (jets, level_jets)
+        np.testing.assert_array_equal(run._boundaries[2].tables["surface_velocity"], velocity)
+        assert (run._boundaries[2].jets, run.jets[2]) == (jets, level_jets)
         assert run._status == "failed"
         module.cuda.memcpy_htod.side_effect = None
         run._stream.synchronize.side_effect = None

tests/unit/test_moving_bodies.py

@@ -0,0 +1,536 @@
+"""Translating-circle host contract, S1 (docs/moving_bodies.md 4a, 5, 7).
+
+Host doubles only: no kernels exist yet, so every positive-step run with a
+moving body must be rejected before any device write.
+"""
+from types import SimpleNamespace
+
+import pytest
+
+from CelerisLab.body.motion import (MAX_STEP, R_MIN, MovingCircle, Translation,
+                                    enumerate_donor_coverage)
+
+BAD = (float("nan"), float("inf"))
+
+
+def circle(box=(-2., 2., -3., 3.), speed=.1):
+    return MovingCircle((50., 40.), 10., Translation(box, speed))
+
+
+def test_translation_and_radius_validation():
+    for box in ((0., 1., -1., 1.)[:3], (.5, 1., -1., 1.), (-1., 1., BAD[0], 1.)):
+        with pytest.raises(ValueError):
+            Translation(box, .1)
+    for speed in (0., MAX_STEP + 1e-9) + BAD:
+        with pytest.raises(ValueError):
+            Translation((0., 0., -1., 1.), speed)
+    with pytest.raises(ValueError):
+        MovingCircle((0., 0.), R_MIN - .1, Translation((0., 0., -1., 1.), .1))
+
+
+def test_targets_validate_against_applied_and_last_wins():
+    c = circle()
+    c.set_target((0., .3), (0., .05))
+    c.set_target((0., .5), (0., .05))  # second setter replaces, checked vs applied (0, 0)
+    assert c.pending == ((0., .5), (0., .05))
+    for offset, velocity in (((0., .51), (0., 0.)),           # Euclidean bound
+                             ((.36, .36), (0., 0.)),          # |dc|_2 = .509
+                             ((0., 3.1), (0., 0.)),           # outside box
+                             ((0., .1), (0., .11)),           # speed bound
+                             ((0., 0.), (0., .01)),           # hold with sliding velocity
+                             ((BAD[0], 0.), (0., 0.)), ((0., 0.), (0., BAD[1]))):
+        with pytest.raises(ValueError):
+            c.set_target(offset, velocity)
+        assert c.pending == ((0., .5), (0., .05))  # invalid replacement keeps the old target
+    c.set_target((.35, .35), (0., 0.))  # |dc|_2 = .495 inside the bound
+    assert c.applied == (0., 0.)
+
+
+def test_record_commit_and_hold():
+    c = circle()
+    assert c.record() == ((50., 40.), (50., 40.), (0., 0.))  # hold: no translation velocity
+    c.set_target((0., .25), (0., .05))
+    assert c.record() == ((50., 40.), (50., 40.25), (0., .05))
+    uploads = []
+    c.advance(uploads.append, lambda: None, lambda: None)
+    assert uploads == [((50., 40.), (50., 40.25), (0., .05))]
+    assert c.applied == (0., .25) and c.pending is None  # offset-only commit
+    c.advance(uploads.append, lambda: None, lambda: None)  # no new target: hold step
+    assert uploads[-1] == ((50., 40.25), (50., 40.25), (0., 0.)) and c.applied == (0., .25)
+    c.set_target((0., 0.), (0., -.05))  # return to zero offset is a real target
+    assert c.pending == ((0., 0.), (0., -.05))
+    c.advance(uploads.append, lambda: None, lambda: None)
+    assert c.applied == (0., 0.) and c.pending is None
+
+
+@pytest.mark.parametrize("stage", ("upload", "step", "synchronize"))
+def test_failure_after_upload_attempt_is_fail_closed(stage):
+    c = circle()
+    c.set_target((0., .25), (0., .05))
+
+    def call(name):
+        def run(*_):
+            if name == stage:
+                raise RuntimeError(f"{name} failed")
+        return run
+    with pytest.raises(RuntimeError, match=f"{stage} failed"):
+        c.advance(call("upload"), call("step"), call("synchronize"))
+    assert c.failed and c.applied == (0., 0.)
+    for later in (c.record, lambda: c.set_target((0., .1), (0., 0.)),
+                  lambda: c.advance(*(lambda *_: None,) * 3)):
+        with pytest.raises(RuntimeError, match="unknown"):
+            later()
+
+
+def test_band_covers_offset_box_with_margin():
+    assert circle().band() == (50 - 2 - 10 - 2, 50 + 2 + 10 + 2, 40 - 3 - 10 - 2, 40 + 3 + 10 + 2)
+
+
+@pytest.mark.parametrize("radius", (R_MIN, 4.3, 7.5, 12.0))
+def test_refill_donor_sets_nonempty_within_bounds(radius):
+    # Evidence for R >= R_MIN and |dc|_2 <= MAX_STEP on a fixed sample, not a proof.
+    assert enumerate_donor_coverage(radius) == 0
+
+
+pytest.importorskip("pycuda.driver")
+from CelerisLab import simulation as module  # noqa: E402
+from tests.unit.test_simulation_lifecycle import Owner, make_simulation  # noqa: E402
+
+
+class ClosableOwner(Owner):
+    """Owner double that refuses activation after close, like CudaContext.activate."""
+    def activate(self):
+        if self.is_closed:
+            raise RuntimeError("CUDA context is closed")
+        return super().activate()
+
+
+def moving_simulation(monkeypatch, *, moving=True):
+    sim = make_simulation(monkeypatch)
+    obj = SimpleNamespace(motion=circle() if moving else None, obj_id=0,
+                          center=(50., 40.), radius=10.)
+    writes = []
+    sim._failed, sim._initialized = None, True
+    sim.bodies = SimpleNamespace(objects=[obj], count=1, jets=[], action_gpu=None, obs_gpu=None,
+                                 has_pending_edit=lambda: False,
+                                 _upload_action_async=lambda s: writes.append("action"),
+                                 zero_obs_async=lambda s: writes.append("obs"),
+                                 increment_obs_steps=lambda n: writes.append("count"),
+                                 download_obs_full_async=lambda s: writes.append("download"))
+    sim.stepper = SimpleNamespace(step=lambda n, **k: writes.append(("step", n)), blocked=None,
+                                  moving=None)
+    sim._assert_runtime_contracts = lambda: None
+    return sim, writes
+
+
+STREAM = SimpleNamespace(synchronize=lambda: None)
+
+
+def test_positive_steps_rejected_before_device_writes(monkeypatch):
+    sim, writes = moving_simulation(monkeypatch)
+    sim.bodies.objects[0].motion.set_target((0., .1), (0., .01))
+    sim.bodies.objects[0].motion.applied = (0., 0.)
+    for call in (lambda: sim.run(1, stream=STREAM), lambda: sim.run(5, stream=STREAM),
+                 lambda: sim.step(), lambda: sim.run(3, stream=STREAM, checkpoint_interval=1)):
+        with pytest.raises(RuntimeError, match="moving kernels"):
+            call()
+    with pytest.raises(RuntimeError, match="moving kernels"):  # nominal at rest as well
+        sim.bodies.objects[0].motion.pending = None
+        sim.run(1)
+    sim.bodies.objects[0].state = SimpleNamespace(omega=.02)  # spin does not lift the guard
+    with pytest.raises(RuntimeError, match="moving kernels"):
+        sim.run(2, stream=STREAM)
+    assert writes == [] and sim.bodies.objects[0].motion.applied == (0., 0.)
+
+
+def test_run_zero_keeps_pending_and_static_runs_proceed(monkeypatch):
+    sim, writes = moving_simulation(monkeypatch)
+    motion = sim.bodies.objects[0].motion
+    motion.set_target((0., .1), (0., .01))
+    sim.run(0, stream=STREAM)  # existing action/obs work only; no motion upload or commit
+    assert motion.pending == ((0., .1), (0., .01)) and motion.applied == (0., 0.)
+    assert ("step", 0) in writes
+    static, writes = moving_simulation(monkeypatch, moving=False)
+    static.save_checkpoint = lambda *a: writes.append("checkpoint")
+    static.run(3, stream=STREAM)
+    static.run(4, stream=STREAM, checkpoint_interval=2)
+    assert writes.count(("step", 3)) == 1 and writes.count(("step", 2)) == 2
+
+
+def test_direct_stepper_entry_is_guarded():
+    from CelerisLab.lbm.stepper import LBMStepper
+    launches = []
+    module = SimpleNamespace(get_function=lambda name: lambda *a, **k: launches.append(name))
+    cfg = SimpleNamespace(streaming="double_buffer", threads_per_block=128, is_d2q9=True,
+                          nx=64, ny=32, nz=1)
+    field = SimpleNamespace(ddf_gpu="ddf", temp_gpu="temp", n_curved=0)
+    stepper = LBMStepper(field, module, cfg)
+    assert stepper.blocked is None
+    stepper.blocked = "translating circles need the moving kernels"
+    for n in (1, 3):
+        with pytest.raises(RuntimeError, match="moving kernels"):
+            stepper.step(n, action_gpu=None, obs_gpu=None)
+    assert launches == [] and (field.ddf_gpu, field.temp_gpu) == ("ddf", "temp")
+    assert stepper.step_count == 0
+    stepper.step(0, action_gpu=None, obs_gpu=None)  # zero steps stay a no-op
+    assert launches == [] and stepper.step_count == 0
+
+
+def test_moving_body_rejects_topology_and_checkpoint(monkeypatch):
+    sim, _ = moving_simulation(monkeypatch)
+    for call in (sim.sync_bodies, sim.save_checkpoint, lambda: sim.load_checkpoint("x.h5")):
+        with pytest.raises(ValueError, match="translating circles"):
+            call()
+
+
+# -- Real-object routing, initialize rejections, packing, lifecycle -------------
+import numpy as np  # noqa: E402
+
+from CelerisLab.body.geometry.circle import CircleGeometry  # noqa: E402
+from CelerisLab.body.manager import ObjectManager  # noqa: E402
+from CelerisLab.body.objects import SimObject  # noqa: E402
+from CelerisLab.lbm.descriptors import FLUID, SOLID  # noqa: E402
+
+
+class Cfg:
+    dim, is_d3q19, is_d2q9 = 2, False, True
+    streaming, ddf_shifting, store_precision, y_wall_bc = "double_buffer", False, "FP32", "bounce_back"
+
+
+def body(x, y, r, translation=None, **kw):
+    obj = SimObject(obj_id=-1, geometry=CircleGeometry(x, y, r), center=(x, y), radius=r, **kw)
+    if translation is not None:
+        obj.motion = MovingCircle((x, y), r, translation)
+    return obj
+
+
+def routed_simulation(monkeypatch):
+    sim = make_simulation(monkeypatch)
+    sim._failed, sim._initialized, sim.lbm_cfg = None, False, Cfg()
+    sim.stepper = SimpleNamespace(blocked=None)
+    sim.bodies = ObjectManager(nx=120, ny=80, nz=1, nq=9, cfg=Cfg())
+    return sim
+
+
+def test_add_body_and_set_body_route_targets(monkeypatch):
+    sim = routed_simulation(monkeypatch)
+    bid = sim.add_body("circle", (50., 40.), 10., translation=Translation((-2., 2., -3., 3.), .1))
+    still = sim.add_body("circle", (100., 40.), 5.)
+    motion = sim.bodies.get(bid).motion
+    sim._moving_tables = "tables"
+    sim._attach_moving()
+    assert sim.stepper.blocked is None
+    sim.set_body(bid, offset=(0., .3), vy=.05)
+    assert sim.stepper.blocked == module.PENDING_BLOCKED  # direct steps never consume a target
+    sim.set_body(bid, offset=(0., .4), vy=.05)  # last wins
+    for kwargs in (dict(offset=(0., .6)), dict(vy=.05), dict(vx=.05),
+                   dict(offset=(0., .1), vy=.2), dict(offset=(0., 0.), vy=.01)):
+        with pytest.raises(ValueError):
+            sim.set_body(bid, **kwargs)
+        assert motion.pending == ((0., .4), (0., .05))  # invalid input keeps the target
+    sim.set_body(bid, omega=.01)  # omega-only keeps the pending target
+    obj = sim.bodies.get(bid)
+    assert obj.state.omega == .01 and (obj.state.vx, obj.state.vy) == (0., 0.)
+    assert motion.pending == ((0., .4), (0., .05)) and motion.applied == (0., 0.)
+    with pytest.raises(ValueError, match="not declared"):
+        sim.set_body(still, offset=(0., .1))
+    sim.set_body(still, vx=.01)  # stationary API unchanged
+    assert sim.bodies.get(still).state.vx == .01 and sim.bodies.get(still).motion is None
+    for kwargs in (dict(type="sensor"), dict(surface_actuation=dict(angle=0., half_width=.2,
+                                                                    normal_velocity=.01))):
+        kind = kwargs.pop("type", "circle")
+        with pytest.raises(ValueError, match="translation"):
+            sim.add_body(kind, (20., 20.), 5., translation=Translation((0., 0., 0., 0.), .1),
+                         **kwargs)
+
+
+class Tables:
+    """MovingTables double: records uploads, steps and diagnostics reads."""
+    DIAG = ("no_donor",)
+
+    def __init__(self, log, no_donor=0., fail=None):
+        self.log, self.no_donor, self.fail = log, no_donor, fail
+
+    def upload_records(self, circles, stream):
+        records = [m.record() for m, _ in circles]
+        self.log.append(("upload", records))
+        uploads = sum(1 for w in self.log if isinstance(w, tuple) and w[0] == "upload")
+        if self.fail == "upload" or (self.fail == "hold_upload" and uploads == 2):
+            raise RuntimeError("upload failed")
+        return any(r[0] != r[1] for r in records)
+
+    def zero_diagnostics(self, stream):
+        self.log.append("zero_diag")
+
+    def download_diagnostics(self):
+        return np.array([[self.no_donor]], np.float32)
+
+
+def advancing_simulation(monkeypatch, **tables):
+    sim, writes = moving_simulation(monkeypatch)
+    sim.bodies.objects[0].motion = circle()
+    sim.stepper = SimpleNamespace(
+        moving=Tables(writes, **tables), blocked=None,
+        _step=lambda n, advance=False, **k: writes.append(("step", n, advance)))
+    from CelerisLab.lbm.stepper import LBMStepper
+    sim.stepper.step = LBMStepper.step.__get__(sim.stepper)  # real public entry
+    sim._moving_tables = sim.stepper.moving
+    return sim, writes, sim.bodies.objects[0].motion
+
+
+def test_run_one_advances_and_commits_after_sync(monkeypatch):
+    sim, writes, motion = advancing_simulation(monkeypatch)
+    motion.set_target((0., .25), (0., .05))
+    with pytest.raises(ValueError, match="run\\(1\\)"):  # multi-step with a target
+        sim.run(2, stream=STREAM)
+    assert writes == [] and motion.pending is not None
+    sim.run(0, stream=STREAM)  # zero steps never consume the target
+    assert motion.pending is not None
+    sim.run(1, stream=STREAM)
+    uploads = [w for w in writes if isinstance(w, tuple) and w[0] == "upload"]
+    assert uploads[0][1] == [((50., 40.), (50., 40.25), (0., .05))]
+    assert uploads[1][1] == [((50., 40.25), (50., 40.25), (0., 0.))]  # first hold after move
+    assert ("step", 1, True) in writes
+    assert motion.applied == (0., .25) and motion.pending is None
+    writes.clear()
+    sim.run(3, stream=STREAM)  # hold steps: no advance
+    assert ("step", 3, False) in writes and motion.applied == (0., .25)
+
+
+@pytest.mark.parametrize("failure", ("upload", "first_sync", "no_donor", "hold_upload",
+                                     "hold_sync"))
+def test_failed_advance_fails_simulation(monkeypatch, failure):
+    kwargs = dict(no_donor=1.) if failure == "no_donor" else dict(fail=failure)
+    sim, writes, motion = advancing_simulation(monkeypatch, **kwargs)
+    motion._on_failure = sim._fail_moving
+    motion.set_target((0., .25), (0., .05))
+    syncs = []
+
+    def synchronize():  # stands in for an asynchronous failure surfacing at sync
+        syncs.append(1)
+        if (failure, len(syncs)) in (("first_sync", 1), ("hold_sync", 2)):
+            raise RuntimeError("sync failed")
+    with pytest.raises(RuntimeError):
+        sim.run(1, stream=SimpleNamespace(synchronize=synchronize))
+    # Before the step-group sync and no_donor check nothing is committed; the hold
+    # record is staged after the commit (applied already matches the moved device).
+    committed = failure in ("hold_upload", "hold_sync")
+    assert motion.applied == ((0., .25) if committed else (0., 0.))
+    assert (motion.pending is None) == committed
+    assert motion.failed and sim._failed is not None
+    assert sim.stepper.blocked == sim._failed  # direct stepper reuse refused too
+    with pytest.raises(RuntimeError, match="Simulation failed"):
+        sim.run(1, stream=STREAM)
+
+
+def test_failed_hold_step_blocks_the_direct_stepper(monkeypatch):
+    # No target is pending, so only the failure latch can refuse direct steps.
+    from CelerisLab.lbm.stepper import LBMStepper
+    sim, writes, motion = advancing_simulation(monkeypatch, fail="upload")
+    motion._on_failure = sim._fail_moving
+    launches = []
+    kernels = SimpleNamespace(get_function=lambda name: lambda *a, **k: launches.append(name))
+    cfg = SimpleNamespace(streaming="double_buffer", threads_per_block=128, is_d2q9=True,
+                          nx=64, ny=32, nz=1)
+    field = SimpleNamespace(ddf_gpu="ddf", temp_gpu="temp", flag_gpu="flag", n_curved=0)
+    stepper = LBMStepper(field, kernels, cfg)
+    stepper.moving = sim.stepper.moving
+    sim.stepper = stepper
+    assert motion.pending is None and stepper.blocked is None
+    with pytest.raises(RuntimeError, match="upload failed"):
+        sim.run(2, stream=STREAM)  # hold step
+    assert sim._failed is not None and stepper.blocked == sim._failed
+    with pytest.raises(RuntimeError, match="moving circle step failed"):
+        stepper.step(1, action_gpu=None, obs_gpu=None)
+    assert launches == [] and stepper.step_count == 0
+    assert (field.ddf_gpu, field.temp_gpu) == ("ddf", "temp")
+
+
+def test_moving_tables_layout_and_attachment(monkeypatch):
+    from CelerisLab.body.motion import MovingTables
+    a = MovingCircle((50., 40.), 10., Translation((-1., 1., -3., 3.), .1))
+    b = MovingCircle((90.25, 30.5), 5., Translation((0., 0., 0., 0.), .1))
+    band, geom = MovingTables.layout([(a, 2), (b, 5)])
+    wa, ha = 63 - 37 + 1, 55 - 25 + 1
+    assert band == [37, 25, wa, ha, 0, 2, 83, 23, 16, 16, wa * ha * 8, 5]
+    assert geom == [50., 40., 50., 40., 10., 0., 0., 90.25, 30.5, 90.25, 30.5, 5., 0., 0.]  # hold
+    sim = routed_simulation(monkeypatch)
+    sim.add_body("circle", (50., 40.), 10., translation=Translation((0., 0., 0., 0.), .1))
+    sim._moving_tables = "tables"
+    sim._attach_moving()
+    assert sim.stepper.moving == "tables" and sim.stepper.blocked is None
+    sim._moving_tables = None
+    sim._attach_moving()
+    assert sim.stepper.moving is None and sim.stepper.blocked is not None
+
+
+class CountingOwner(Owner):
+    """Owner double that counts nested activations like CudaContext."""
+    def __init__(self):
+        super().__init__()
+        self.active = 0
+
+    def activate(self):
+        from contextlib import contextmanager
+
+        @contextmanager
+        def scope():
+            self.active += 1
+            try:
+                yield self
+            finally:
+                self.active -= 1
+        return scope()
+
+
+def test_public_diagnostics_keep_owner_active_through_sync_and_download(monkeypatch):
+    sim = make_simulation(monkeypatch)
+    sim.ctx, sim._failed = CountingOwner(), None
+    seen = []
+    sim._stream = SimpleNamespace(synchronize=lambda: seen.append(("sync", sim.ctx.active)))
+    expected = np.arange(6, dtype=np.float32).reshape(1, 6)
+
+    def download():
+        seen.append(("download", sim.ctx.active))
+        return expected
+    sim.stepper = SimpleNamespace(moving=SimpleNamespace(download_diagnostics=download))
+    assert sim.moving_diagnostics() is expected
+    assert [name for name, _ in seen] == ["sync", "download"]
+    assert all(active > 0 for _, active in seen) and sim.ctx.active == 0
+
+
+def test_stepper_launches_moving_links_before_static_curved():
+    from CelerisLab.lbm.stepper import LBMStepper
+    launches = []
+    module = SimpleNamespace(get_function=lambda name: lambda *a, **k: launches.append((name, a, k)))
+    cfg = SimpleNamespace(streaming="double_buffer", threads_per_block=128, is_d2q9=True,
+                          nx=64, ny=32, nz=1)
+    curved = SimpleNamespace(**{f"{name}_gpu": name for name in (
+        "fluid_idx", "dir", "q", "rx", "ry", "rz", "fallback_class", "body_id")},
+        surface_velocity=None, surface_velocity_gpu=None, link_exchange_gpu=None)
+    field = SimpleNamespace(ddf_gpu="ddf", temp_gpu="temp", flag_gpu="flag", n_curved=5,
+                            curved=curved, invalidate_host_ddf_cache=lambda: None)
+    stepper = LBMStepper(field, module, cfg)
+    stepper._launch_force_region = stepper._launch_sensor = lambda *a, **k: None
+    stepper.moving = SimpleNamespace(band_gpu="band", geom_gpu="geom", diag_gpu="diag",
+                                     slot_debug_gpu=None, n_bodies=1, n_slots=300)
+    stepper.step(2, action_gpu="act", obs_gpu="obs")
+    names = [name for name, _, _ in launches]
+    assert names == ["MovingCircleLinks", "CurvedBoundaryKernel", "OneStep"] * 2
+    assert stepper.step_count == 2
+    assert launches[1][1][0] == "ddf" and launches[3][1][0] == "temp"  # swap between steps
+    args, kw = launches[0][1], launches[0][2]
+    assert args[:4] == ("ddf", "flag", "band", "geom") and int(args[5]) == 300
+    assert kw["grid"] == (3, 1, 1) and kw["block"] == (128, 1, 1)
+
+
+def test_recompile_installs_a_guarded_stepper(monkeypatch):
+    from CelerisLab.lbm.stepper import LBMStepper
+    sim = routed_simulation(monkeypatch)
+    sim.add_body("circle", (50., 40.), 10., translation=Translation((0., 0., -1., 1.), .1))
+    launches = []
+    kernel_module = SimpleNamespace(get_function=lambda name: lambda *a, **k: launches.append(name))
+    monkeypatch.setattr(module.compiler, "build_module", lambda *a, **k: kernel_module)
+    sim.lbm_cfg = SimpleNamespace(validate_runtime=lambda: None, viscosity=.01, omega=1.,
+                                  streaming="double_buffer", threads_per_block=128,
+                                  is_d2q9=True, nx=120, ny=80, nz=1)
+    sim._compiled_viscosity = .01
+    sim._assert_allocation_contract = lambda: None
+    sim._resolve_compile_arch = lambda: "sm_70"
+    sim._assert_object_count_contract = lambda **k: None
+    sim.field = SimpleNamespace(module=None, invalidate_module_constant_cache=lambda: None,
+                                _upload_params=lambda: None, _runtime_params={},
+                                ddf_gpu="ddf", temp_gpu="temp", n_curved=0)
+    sim.stepper = SimpleNamespace(_step_count=7, blocked=None)
+    sim._recompile()
+    assert isinstance(sim.stepper, LBMStepper) and sim.stepper.step_count == 7
+    with pytest.raises(RuntimeError, match="moving kernels"):
+        sim.stepper.step(1, action_gpu=None, obs_gpu=None)
+    assert launches == [] and (sim.field.ddf_gpu, sim.field.temp_gpu) == ("ddf", "temp")
+    assert sim.stepper.step_count == 7
+
+
+def validate(monkeypatch, objects, flags=None, **cfg):
+    sim = routed_simulation(monkeypatch)
+    for obj in objects:
+        sim.bodies.add(obj)
+    for name, value in cfg.items():
+        setattr(sim.lbm_cfg, name, value)
+    base = np.full(120 * 80, FLUID, np.uint16) if flags is None else flags
+    sim.field = SimpleNamespace(nx=120, ny=80, build_channel_flags=lambda: base.copy())
+    sim._validate_moving_bodies()
+
+
+def test_initialize_rejects_band_and_modes(monkeypatch):
+    t = Translation((-1., 1., -3., 3.), .1)
+    validate(monkeypatch, [body(50., 40., 10., t), body(100., 40., 5.)])  # accepted
+    walls = np.full(120 * 80, FLUID, np.uint16)
+    walls[60 + 30 * 120] = SOLID
+    cases = [([body(10., 40., 10., t)], None, {}, "domain edge"),
+             ([body(50., 40., 10., t)], walls, {}, "non-fluid"),
+             ([body(50., 40., 10., t), body(68., 40., 4.)], None, {}, "overlaps"),
+             ([body(50., 40., 10., t), body(50., 60., 3., is_sensor=True)], None, {}, "overlaps"),
+             ([body(50., 40., 10., t), body(64., 40., 3., is_force_region=True)], None, {}, "overlaps"),
+             ([body(50., 40., 10., t)], None, {"streaming": "esopull"}, "double_buffer"),
+             ([body(50., 40., 10., t)], None, {"ddf_shifting": True}, "unshifted"),
+             ([body(50., 40., 10., t)], None, {"is_d2q9": False}, "D2Q9")]
+    for objects, flags, cfg, match in cases:
+        with pytest.raises(ValueError, match=match):
+            validate(monkeypatch, objects, flags, **cfg)
+
+
+def test_initialize_validates_before_mutation(monkeypatch):
+    sim = routed_simulation(monkeypatch)
+    sim.bodies.add(body(10., 40., 10., Translation((0., 0., 0., 0.), .1)))
+    uploads = []
+    sim.field = SimpleNamespace(nx=120, ny=80, build_channel_flags=lambda: np.full(9600, FLUID, np.uint16),
+                                upload_flags=lambda: uploads.append("flags"))
+    sim._assert_allocation_contract = lambda: None
+    with pytest.raises(ValueError, match="domain edge"):
+        sim.initialize()
+    assert uploads == [] and not sim._initialized
+
+
+def test_moving_circle_excluded_from_static_links():
+    def manager(*objects):
+        mgr = ObjectManager(nx=120, ny=80, nz=1, nq=9, cfg=Cfg())
+        ids = [mgr.add(obj) for obj in objects]
+        return mgr, ids
+    reference_mgr, _ = manager(body(100., 40., 5.), body(20., 40., 6.))
+    mixed_mgr, ids = manager(body(100., 40., 5.), body(20., 40., 6.),
+                             body(55., 40., 10., Translation((0., 0., -1., 1.), .1)))
+    reference = reference_mgr.build_compact_lists_for(reference_mgr.objects)
+    mixed = mixed_mgr.build_compact_lists_for(mixed_mgr.objects)
+    for a, b in zip(reference[:8], mixed[:8]):  # static link columns byte-identical
+        assert np.asarray(a).tobytes() == np.asarray(b).tobytes()
+    assert len(reference[0]) > 0 and ids[2] not in set(np.asarray(mixed[3]).tolist())
+
+
+def test_failed_moving_step_marks_simulation_failed(monkeypatch):
+    sim = routed_simulation(monkeypatch)
+    sim.ctx = ClosableOwner()
+    bid = sim.add_body("circle", (50., 40.), 10., translation=Translation((0., 0., -1., 1.), .1))
+    motion = sim.bodies.get(bid).motion
+    motion.set_target((0., .2), (0., .05))
+
+    def fail(*_):
+        raise RuntimeError("sync failed")
+    with pytest.raises(RuntimeError, match="sync failed"):
+        motion.advance(lambda *_: None, lambda: None, fail)
+    assert sim._failed is not None and motion.applied == (0., 0.)
+    for call in (lambda: sim.run(0), lambda: sim.run(1), sim.step, sim.sync_bodies,
+                 sim.save_checkpoint, lambda: sim.load_checkpoint("x.h5")):
+        with pytest.raises(RuntimeError, match="Simulation failed"):
+            call()
+    with pytest.raises(RuntimeError, match="unknown"):  # no retry on the failed circle
+        motion.advance(lambda *_: None, lambda: None, lambda: None)
+    with sim.ctx.activate():  # a refused close keeps the failed state and resources
+        with pytest.raises(RuntimeError, match="active CUDA"):
+            sim.close()
+    assert sim._failed is not None and sim.bodies is not None and not sim._closed
+    sim.close()  # retry outside the activation succeeds
+    assert sim._closed and sim._failed is not None
+    with pytest.raises(RuntimeError, match="CUDA context is closed"):  # no re-entry after close
+        sim.run(1)
+    with pytest.raises(RuntimeError, match="unknown"):  # the failed circle stays refused
+        motion.record()

tests/unit/test_multiple_pinned_remap.py

@@ -96,11 +96,11 @@ class MultiplePinnedRemap(unittest.TestCase):
         before = run.download()
         clone = run.remap_finest((PINS+MOVED+ADDED)[::-1], objects=bodies)
         try:
-            self.assertEqual(tuple(j.amplitude for j in clone._level_boundaries[2].jets), (.017, -.011))
-            np.testing.assert_array_equal(clone._level_boundaries[2].tables['surface_velocity'],
-                                          run._level_boundaries[2].tables['surface_velocity'])
+            self.assertEqual(tuple(j.amplitude for j in clone._boundaries[2].jets), (.017, -.011))
+            np.testing.assert_array_equal(clone._boundaries[2].tables['surface_velocity'],
+                                          run._boundaries[2].tables['surface_velocity'])
             clone.set_jet_amplitude(1, .005, level=2)
-            self.assertEqual(run._level_boundaries[2].jets[1].amplitude, -.011)
+            self.assertEqual(run._boundaries[2].jets[1].amplitude, -.011)
             clone.step()
             self.assertEqual((run.step_count, clone.step_count), (7, 8))
             for got, want in zip(run.download(), before):

tests/unit/test_pinned_body_remap.py

@@ -62,9 +62,9 @@ class PinnedBodyRemap(unittest.TestCase):
         run.set_jet_amplitude(0,.017,level=2)
         before=run.download()
         clone=run.remap_finest(B,objects=objects)
-        self.assertEqual(clone._level_boundaries[2].jets[0].amplitude,.017)
-        np.testing.assert_array_equal(clone._level_boundaries[2].tables['surface_velocity'],
-                                      run._level_boundaries[2].tables['surface_velocity'])
+        self.assertEqual(clone._boundaries[2].jets[0].amplitude,.017)
+        np.testing.assert_array_equal(clone._boundaries[2].tables['surface_velocity'],
+                                      run._boundaries[2].tables['surface_velocity'])
         self.assertTrue(same_physical_boundary(run.boundaries[2],clone.boundaries[2]))
         clone.step()
         self.assertEqual((run.step_count,clone.step_count),(7,8))
@@ -74,8 +74,8 @@ class PinnedBodyRemap(unittest.TestCase):
     def test_static_body_without_actuation(self):
         run,module,allocations,storage,objects=ready(jets=())
         clone=run.remap_finest(A,objects=objects)
-        self.assertFalse(clone._level_boundaries[2].actuated)
-        self.assertNotIn('surface_velocity',clone._level_boundaries[2].tables)
+        self.assertFalse(clone._boundaries[2].actuated)
+        self.assertNotIn('surface_velocity',clone._boundaries[2].tables)
         for a,b in zip(run.download(),clone.download()): np.testing.assert_array_equal(a,b)
         clone.close();run.close()
 

tests/unit/test_raw_exchange.py

@@ -85,7 +85,7 @@ class RawExchange(unittest.TestCase):
 
     def test_four_level_schedule_publication_and_fixed_storage(self):
         run, host, calls, arrays = backend()
-        b = run._level_boundaries[3]
+        b = run._boundaries[3]
         rows = len(b.tables['cell'])
         self.assertEqual(b.last_raw.size, rows * 16)
         self.assertEqual(b.interval_integral.size, rows * 32)
@@ -128,7 +128,7 @@ class RawExchange(unittest.TestCase):
 
     def test_disabled_has_nulls_and_no_observation_allocation_or_launch(self):
         run, host, calls, _ = backend(False)
-        b = run._level_boundaries[3]
+        b = run._boundaries[3]
         self.assertIs(b.last_raw, b.interval_integral)
         self.assertIs(b.interval_integral, None)
         run.step()

tests/unit/test_refinement_mrt.py

@@ -1,5 +1,6 @@
 """Host guards and source contract for MRT reuse in refined/blocked stepping (no CUDA)."""
 from pathlib import Path
+import re
 import unittest
 
 from CelerisLab.config import LBMConfig
@@ -32,22 +33,25 @@ class RefinementMRTHost(unittest.TestCase):
                      "collide_srt_no_force(f, feq, d_params.omega);"):
             self.assertIn(text, block_collision)
         self.assertIn("collide_srt(f, feq, forcing, omega);", block_collision)
-        self.assertIn("#elif BLOCK_ROOT_BC == 2", block_collision)
+        self.assertIn("#if COLLISION_MODEL == 2 || BLOCK_ROOT_BC == 2", block_collision)
         self.assertEqual(channel.count("block_collide(f);"), 1)
-        self.assertNotIn("collide_mrt_unforced(f);", channel)
+        # Executable text only: the header comment explains the absent force read.
+        code = re.sub(r"//[^\n]*", "", channel + "\n" + block_collision)
+        self.assertIn("d_params.omega", code)
+        for text in ("collide_mrt_unforced", "collide_dispatch(", "d_params.fx", "d_params.fy"):
+            self.assertNotIn(text, code)
         for source in (blocks,):
             self.assertIn("(COLLISION_MODEL != 0 && COLLISION_MODEL != 2)", source)
             self.assertIn("USE_LES", source)
             self.assertIn("STORE_PRECISION != 0", source)
             self.assertIn('#if COLLISION_MODEL == 2\n#include "operators/collision_mrt.cuh"\n#endif', source)
-        # 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 + "\n" + block_collision):
-            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")
+        # One unforced MRT call, in block_collide only; the steps never collide directly.
+        self.assertEqual(block_collision.count("collide_mrt_no_force(f, rho, ux, uy, omega);"), 1)
+        self.assertEqual(block_collision.count("compute_rho_u(f, rho, ux, uy);"), 1)
+        self.assertEqual(block_collision.count("fmaxf(OMEGA_COLLISION_MIN, d_params.omega)"), 1)
+        for source in (blocks, channel):
+            self.assertNotIn("collide_mrt", source)
+            self.assertNotIn("collide_srt", source)
 
 
 if __name__ == "__main__": unittest.main()

tests/unit/test_refinement_three_level_multipatch.py

@@ -141,7 +141,7 @@ class Gate(unittest.TestCase):
     def test_sparse_parent_tables_preserve_inherited_kinds(self):
         child,kind=r.nest_patch_set(CFG,self.parent,CHILDREN)
         layout=self.parent.fine_layout((4,4));order=np.argsort(child.coalesce_coarse)
-        donors,targets,bkind=host.blocked_base_tables(child,layout,order,composed_kind=kind)
+        donors,targets,bkind=host.blocked_coarse_tables(child,layout,order,composed_kind=kind)
         self.assertTrue((donors < layout.hole).all())
         self.assertTrue((targets < np.prod(layout.blocked_shape)).all())
         np.testing.assert_array_equal(bkind[layout.owner[layout.owned]],kind[layout.owned])
@@ -158,10 +158,10 @@ class Gate(unittest.TestCase):
         hole=int(np.flatnonzero(~layout.owned.ravel())[0])
         bad=child.ghost_coarse.copy();bad[0]=hole
         with self.assertRaisesRegex(ValueError,'owned uncovered'):
-            host.blocked_base_tables(replace(child,ghost_coarse=bad),layout,order,composed_kind=kind)
+            host.blocked_coarse_tables(replace(child,ghost_coarse=bad),layout,order,composed_kind=kind)
         bad=child.coalesce_coarse.copy();bad[0]=hole
         with self.assertRaisesRegex(ValueError,'owned uncovered'):
-            host.blocked_base_tables(replace(child,coalesce_coarse=bad),layout,order,composed_kind=kind)
+            host.blocked_coarse_tables(replace(child,coalesce_coarse=bad),layout,order,composed_kind=kind)
 
 
 def bodies_and_jets(speed=.002):

tests/unit/test_runtime_child_remap.py

@@ -1,5 +1,6 @@
 """CPU mocks execute production remap lifecycle; no GPU numerical claim."""
 from dataclasses import replace
+from types import SimpleNamespace
 import unittest
 from unittest.mock import patch
 
@@ -69,10 +70,11 @@ class RuntimeChildRemap(unittest.TestCase):
         run._status = 'ready'
         for bad in ([], [(40,16,48,40)], [(16,16,32,24),(16,27,32,55)]):
             with self.assertRaises(ValueError): run.remap_finest(bad)
-        hosts, pairs = run._boundaries, run.hierarchy.pairs
+        records, pairs = run._boundaries, run.hierarchy.pairs
         for attribute, value in (('_wall_aware', True),
-                                 ('_boundaries', [*hosts[:2], object()]),        # child boundary
-                                 ('_boundaries', [hosts[0], object(), hosts[2]]),  # fine boundary
+                                 # level-2 record whose host geometry differs from objects=()
+                                 ('_boundaries', [*records[:2], SimpleNamespace(host=object(), jets=())]),
+                                 ('_boundaries', [records[0], object(), records[2]]),  # level-1 body
                                  ('hierarchy', replace(run.hierarchy, pairs=(pairs[0], None)))):
             with patch.object(run, attribute, value):
                 with self.assertRaises(ValueError): run.remap_finest(B)

tests/unit/test_shifted_jets.py

@@ -67,7 +67,7 @@ class ShiftedJets(unittest.TestCase):
     def test_switch_upload_only_then_raw_remap_keeps_live_velocity(self):
         run,module,allocations,storage,objects=ready();clone=None
         try:
-            boundary=run._level_boundaries[2]
+            boundary=run._boundaries[2]
             before=run._download_raw()
             original=boundary.tables['surface_velocity'].copy()
             module.cuda.memcpy_htod.reset_mock()
@@ -79,7 +79,7 @@ class ShiftedJets(unittest.TestCase):
             for a,b in zip(before.values,run._download_raw().values):self.assertEqual(a.tobytes(),b.tobytes())
             clone=run.remap_finest(B,objects=objects)
             self.assertEqual(clone.jets[2][0].amplitude,-jet().amplitude)
-            np.testing.assert_array_equal(clone._level_boundaries[2].tables['surface_velocity'],
+            np.testing.assert_array_equal(clone._boundaries[2].tables['surface_velocity'],
                                           boundary.tables['surface_velocity'])
             self.assertEqual(clone.step_count,7)
             saved=run._download_raw()
@@ -92,12 +92,12 @@ class ShiftedJets(unittest.TestCase):
     def test_invalid_amplitude_does_not_upload_or_commit(self):
         run,module,_,_,_=ready()
         try:
-            before=run._level_boundaries[2].tables['surface_velocity'].copy()
+            before=run._boundaries[2].tables['surface_velocity'].copy()
             module.cuda.memcpy_htod.reset_mock()
             for amplitude in (float('nan'),float('inf'),1.):
                 with self.assertRaises(ValueError):run.set_jet_amplitude(0,amplitude,level=2)
             module.cuda.memcpy_htod.assert_not_called()
-            np.testing.assert_array_equal(before,run._level_boundaries[2].tables['surface_velocity'])
+            np.testing.assert_array_equal(before,run._boundaries[2].tables['surface_velocity'])
         finally:run.close()
 
 

tests/unit/test_single_level.py

@@ -73,7 +73,7 @@ class SingleLevel(unittest.TestCase):
 
     def test_circle_solids_are_not_state(self):
         run, _, _ = backend((circle(),))
-        solid = run._level_boundaries[0].tables["solid"]
+        solid = run._boundaries[0].tables["solid"]
         self.assertTrue(solid.any())
         np.testing.assert_array_equal(run._state[0], ~solid)
 

tests/unit/test_state_regressions.py

@@ -294,8 +294,11 @@ def test_passive_overlap_failure_keeps_live_state(device, entry, second_x):
     sim = SimpleNamespace(bodies=bodies, field=field, lbm_cfg=field.cfg,
                           _initialized=entry == "runtime_edit",
                           _assert_allocation_contract=lambda: None,
-                          stream=SimpleNamespace(synchronize=device.synchronize))
-    sim._require_usable = Simulation._require_usable.__get__(sim)
+                          stream=SimpleNamespace(synchronize=device.synchronize),
+                          stepper=SimpleNamespace(blocked=None))
+    for name in ("_require_usable", "_moving_bodies", "_reject_moving_bodies",
+                 "_validate_moving_bodies", "_attach_moving"):
+        setattr(sim, name, getattr(Simulation, name).__get__(sim))
     with pytest.raises(ValueError, match="active fluid receivers|unique cut-link writers"):
         if entry == "sync":
             bodies.sync_to_gpu(field)

tests/unit/test_union_body_support.py

@@ -37,7 +37,7 @@ class UnionBodySupport(unittest.TestCase):
             run.set_jet_amplitude(0,.017,level=2)
             clone=run.remap_finest(((16,16,24,40),(24,16,32,40),A[1]),objects=objects)
             try:
-                self.assertEqual(clone._level_boundaries[2].jets[0].amplitude,.017)
+                self.assertEqual(clone._boundaries[2].jets[0].amplitude,.017)
                 self.assertTrue(same_physical_boundary(run.boundaries[2],clone.boundaries[2]))
                 for a,b in zip(run.download(),clone.download()):np.testing.assert_array_equal(a,b)
             finally:clone.close()

tests/unit/test_viv_logic.py

@@ -0,0 +1,96 @@
+"""CPU tests of the predeclared VIV logic (research/viv_re150/viv.py; stdlib only)."""
+import importlib.util
+import math
+from pathlib import Path
+
+_spec = importlib.util.spec_from_file_location(
+    "viv_logic", Path(__file__).resolve().parents[2] / "research/viv_re150/viv.py")
+viv = importlib.util.module_from_spec(_spec)
+_spec.loader.exec_module(viv)
+
+D, U = 40, 0.04
+
+
+def _feed(cyc, amps, period=100):
+    """Feed sine cycles of the given amplitudes (in D); return the first non-None result."""
+    n = 0
+    for a in amps:
+        for i in range(period):
+            out = cyc.update(a * D * math.sin(2 * math.pi * (i + 1) / period))
+            n += 1
+            if out:
+                return out, n
+    return None, n
+
+
+def test_structure_exact_update_conserves_free_oscillation_and_equilibrium():
+    s = viv.Structure(D, U)
+    assert math.isclose(1 / s.f_n, 5000.0) and math.isclose(s.m, 2 * D * D)
+    s.y, s.v = 0.3 * D, 0.0
+    e0 = 0.5 * s.k * s.y ** 2
+    for _ in range(12345):
+        s.y, s.v = s.candidate(0.0)
+    assert math.isclose(0.5 * s.k * s.y ** 2 + 0.5 * s.m * s.v ** 2, e0, rel_tol=1e-9)
+    s.y, s.v = 2.0 / s.k, 0.0
+    assert all(math.isclose(a, b, abs_tol=1e-12) for a, b in zip(s.candidate(2.0), (2.0 / s.k, 0.0)))
+    assert math.isclose(s.seed_velocity(D) / s.w, viv.ELIGIBLE * D)
+
+
+def test_constant_amplitude_settles_then_scores_twenty_cycles():
+    cyc = viv.Cycles(D)
+    out, _ = _feed(cyc, [0.5] * 60)
+    assert out == "scored" and len(cyc.window) == 20 and cyc.settled_at == 1
+    r = viv.score(cyc.window)
+    assert math.isclose(r["A_max"], 0.5, rel_tol=1e-3) and abs(r["offset"]) < 1e-3
+    assert r["T_resp"] == 100 and r["within_band"]          # |.500 - .520| <= .052
+
+
+def test_small_cycles_are_not_counted_and_y0_is_not_a_crossing():
+    cyc = viv.Cycles(D)
+    assert cyc.update(0.0) is None and cyc.log == []
+    out, _ = _feed(cyc, [0.005] * 50)
+    assert out is None and cyc.counted == 0 and cyc.log and not any(c["eligible"] for c in cyc.log)
+
+
+def test_scoring_failure_restarts_from_the_failing_block():
+    cyc = viv.Cycles(D)
+    amps = [0.5] * 20 + [0.55] * 10 + [0.55] * 40      # settle at block 1, block 2 jumps 10 %
+    out, _ = _feed(cyc, amps)
+    assert out == "scored"
+    assert all(math.isclose(c["amp"], 0.55, rel_tol=1e-3) for c in cyc.window)
+
+
+def test_seed_reset_discards_incomplete_cycle_and_proof():
+    cyc = viv.Cycles(D)
+    _feed(cyc, [0.5] * 15)
+    cyc.update(0.2 * D)                                  # inside an incomplete cycle
+    counted = cyc.counted
+    cyc.discard_current()
+    cyc.reset_proof()
+    assert cyc.amps == [] and cyc.blocks == [] and cyc.counted == counted
+
+
+def test_viv_config_passes_the_real_loader(tmp_path):
+    import json
+    import sys
+    sys.path.insert(0, str(Path(__file__).resolve().parents[2] / "research/viv_re150"))
+    try:
+        import run
+    finally:
+        sys.path.pop(0)
+    from CelerisLab.config import load_lbm_config
+    cfg_dict, nx, ny = run.config()
+    path = tmp_path / "config_lbm.json"
+    path.write_text(json.dumps(cfg_dict))
+    cfg = load_lbm_config(str(path))
+    assert (cfg.nx, cfg.ny) == (nx, ny) == (3000, 2400)
+    assert cfg.collision == "MRT" and cfg.y_wall_bc == "free_slip" and cfg.velocity == .04
+    assert math.isclose(cfg.viscosity, .04 * 40 / 150)
+
+
+def test_finite_rejects_any_nonfinite_recorded_output():
+    row = (3, .1, .2, .3, .4, .5, .6, .7, 0.0)
+    assert viv.finite(row)
+    for i in range(len(row)):
+        for bad in (float("nan"), float("inf"), -float("inf")):
+            assert not viv.finite(row[:i] + (bad,) + row[i + 1:])

tests/validation/multigrid/test_block_stepper.py

@@ -240,7 +240,7 @@ class SingleLevelGPU(unittest.TestCase):
         with CudaContext(0) as context:
             for split in SPLITS:
                 with single(context=context, block=split, objects=[circle()]) as backend:
-                    tables = backend._level_boundaries[0].tables
+                    tables = backend._boundaries[0].tables
                     for key, values, dtype in (("q", [l.q for l in links], np.float32),
                                                ("dir", [l.dir for l in links], np.uint8),
                                                ("fallback", [l.fallback for l in links], np.uint8)):

tests/validation/multigrid/test_exchange_window.py

@@ -59,7 +59,7 @@ class ExchangeWindowCPU(unittest.TestCase):
             stack.enter_context(patch.object(cuda, "memcpy_dtoh",
                                              side_effect=lambda out, src: out.fill(0)))
             backend = build(m)
-            b = backend._level_boundaries[1]
+            b = backend._boundaries[1]
             self.assertIsNotNone(b.window_integral)
             with self.assertRaises(RuntimeError):
                 backend.start_exchange_window()  # not initialized
@@ -96,7 +96,7 @@ class ExchangeWindowCPU(unittest.TestCase):
         with mocked_backend() as m, ExitStack() as stack:
             stack.enter_context(patch.object(cuda, "memset_d8_async"))
             backend = build(m, window=False)
-            self.assertIsNone(backend._level_boundaries[1].window_integral)
+            self.assertIsNone(backend._boundaries[1].window_integral)
             backend.initialize(*[np.full(l.shape, 1 / 9, np.float32) for l in backend._levels])
             backend.step(1)
             self.assertEqual(len(accumulations(m)), 2)
@@ -120,7 +120,7 @@ class ExchangeWindowGPU(unittest.TestCase):
                 windowed.step(2)
                 plain.step(2)
                 windowed.start_exchange_window()
-                b = windowed._level_boundaries[1]
+                b = windowed._boundaries[1]
                 count = len(b.tables["cell"])
                 factors = np.array([.25, .25, .125, .25])  # h*dt, h*h*dt, h*h for level-1 h=.5
                 scale = np.zeros(4)
@@ -212,7 +212,7 @@ class ExchangeWindowShiftedMRTGPU(unittest.TestCase):
                     run.step(2)
                     if run is not plain:
                         run.start_exchange_window()
-                b = windowed._level_boundaries[1]
+                b = windowed._boundaries[1]
                 count = len(b.tables["cell"])
                 record = dict(scale=np.zeros(4), substeps=0)
                 audit_substeps(windowed, b, count, record)
@@ -228,7 +228,7 @@ class ExchangeWindowShiftedMRTGPU(unittest.TestCase):
                     same_bits(self, a, e)
                 rows = [np.empty((count, 4), np.float64) for _ in range(2)]
                 cuda.memcpy_dtoh(rows[0], b.window_integral)
-                cuda.memcpy_dtoh(rows[1], split._level_boundaries[1].window_integral)
+                cuda.memcpy_dtoh(rows[1], split._boundaries[1].window_integral)
                 np.testing.assert_array_equal(rows[0].view(np.uint64), rows[1].view(np.uint64))
 
                 out = windowed.read_exchange_window(0, level=1)

tests/validation/multigrid/test_multiple_pinned_remap_gpu.py

@@ -57,7 +57,7 @@ def make(context):
 def topology(run, rectangles, objects):
     return RemapTopology.from_config(
         CFG, run.hierarchy.pairs[0], rectangles, child_objects=objects,
-        child_jets=run._level_boundaries[2].jets,
+        child_jets=run._boundaries[2].jets,
     )
 
 
@@ -74,7 +74,7 @@ class MultiplePinnedRemapGPU(unittest.TestCase):
     def change_amplitudes(self, run, amplitudes):
         # Each jet selects a different body. Updating one must preserve the
         # other body's current per-link velocities, including its driven links.
-        boundary = run._level_boundaries[2]
+        boundary = run._boundaries[2]
         body_ids = run.boundaries[2].curved_host[3]
         for index, amplitude in enumerate(amplitudes):
             before = boundary.tables['surface_velocity'].copy()
@@ -91,7 +91,7 @@ class MultiplePinnedRemapGPU(unittest.TestCase):
         self.assertTrue(same_physical_boundary(new.boundaries[2], mesh.boundary))
         # Compare physical cut-link rows, never layout-dependent blocked addresses.
         for run in (old, new):
-            boundary = run._level_boundaries[2]
+            boundary = run._boundaries[2]
             self.assertEqual(tuple(jet.amplitude for jet in boundary.jets), amplitudes)
             exact(self, boundary.tables['surface_velocity'], mesh.surface_velocity,
                   'latest physical per-link velocities')

tests/validation/multigrid/test_pinned_body_remap_gpu.py

@@ -39,9 +39,9 @@ class PinnedBodyRemapGPU(unittest.TestCase):
             exact(self,getattr(after.boundaries[2],name),getattr(before.boundaries[2],name),name)
         for a,b in zip(after.boundaries[2].curved_host,before.boundaries[2].curved_host):
             exact(self,a,b,'physical cut links')
-        exact(self,after._level_boundaries[2].tables['surface_velocity'],
-              before._level_boundaries[2].tables['surface_velocity'],'current per-link velocities')
-        self.assertEqual(after._level_boundaries[2].jets[0].amplitude,.017)
+        exact(self,after._boundaries[2].tables['surface_velocity'],
+              before._boundaries[2].tables['surface_velocity'],'current per-link velocities')
+        self.assertEqual(after._boundaries[2].jets[0].amplitude,.017)
 
     def test_noop_next_step_bitwise(self):
         with CudaContext(0) as ctx, ExitStack() as owned:
@@ -63,7 +63,7 @@ class PinnedBodyRemapGPU(unittest.TestCase):
             current.set_jet_amplitude(0,.017,level=2)
             for rects in (B,A):
                 previous=current;before=previous.download();count=previous.step_count
-                jets=previous._level_boundaries[2].jets
+                jets=previous._boundaries[2].jets
                 old=RemapTopology.from_config(CFG,previous.hierarchy.pairs[0],previous.hierarchy.pairs[1].bounds,
                                              child_objects=objects,child_jets=jets)
                 new=RemapTopology.from_config(CFG,previous.hierarchy.pairs[0],rects,

tests/validation/multigrid/test_raw_exchange_gpu.py

@@ -105,7 +105,7 @@ class RawExchangeGPU(unittest.TestCase):
                     arrays = [state(l.shape) for l in measured._levels]
                     measured.initialize(*arrays)
                     plain.initialize(*arrays)
-                    b = measured._level_boundaries[1]
+                    b = measured._boundaries[1]
                     count = len(b.tables['cell'])
                     kernel = measured._levels[1].kernel
                     samples = []
@@ -158,7 +158,7 @@ class RawExchangeGPU(unittest.TestCase):
                  FixedPatchStepper(cfg, context=ctx, **settings) as plain:
                 arrays = [state(l.shape, ux=0.) for l in run._levels]
                 run.initialize(*arrays); plain.initialize(*arrays)
-                b = run._level_boundaries[3]
+                b = run._boundaries[3]
                 count = len(b.tables['cell'])
                 kernel = run._levels[3].kernel
                 samples, independent_samples = [], []
@@ -213,7 +213,7 @@ class RawExchangeGPU(unittest.TestCase):
                  FixedPatchStepper(CFG, context=ctx, **fine((16, 20)), observe_raw_exchange=True) as amr:
                 amr.initialize(*(state(l.shape) for l in amr._levels))
                 uniform.initialize(state((9, 48, 64)))
-                b = amr._level_boundaries[1]
+                b = amr._boundaries[1]
                 count = len(b.tables['cell'])
                 links = obj.geometry.build_cut_links(64, 48)
                 np.testing.assert_array_equal([v.dir for v in links], b.tables['dir'])
@@ -221,7 +221,7 @@ class RawExchangeGPU(unittest.TestCase):
                 # Same fine-grid force arms after translation to the same center: the
                 # single level's static BFL gets the AMR Jet column and force arms; the raw
                 # output pointer is substituted at launch (no stepper-side accumulation).
-                root = uniform._level_boundaries[0]
+                root = uniform._boundaries[0]
                 velocity = uniform._upload(b.tables['surface_velocity'])
                 root.device.update(surface_velocity=velocity, rx=uniform._upload(b.tables['rx']),
                                    ry=uniform._upload(b.tables['ry']))

tests/validation/multigrid/test_refinement_blocked_base.py

@@ -24,7 +24,8 @@ from CelerisLab.lbm import refinement_stepper
 from CelerisLab.lbm.blocks import BlockLayout2D
 from CelerisLab.lbm.descriptors import D2Q9_EX, D2Q9_EY
 from CelerisLab.lbm.refinement import ACTIVE, OUTER_GHOST, FixedPatch2D
-from CelerisLab.lbm.refinement_stepper import FixedPatchStepper, blocked_base_tables
+from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+from CelerisLab.lbm.refinement_tables import blocked_coarse_tables
 from test_refinement_blocked_fine import NX, NY, BOUNDS, config, state, circle, actuation
 
 BASE_BLOCKS = ((24, 32), (12, 8), (3, 4), (1, 1))
@@ -108,7 +109,7 @@ class BlockedBaseCPU(unittest.TestCase):
                 leaves[[0, -1]] = False
             for by, bx in BASE_BLOCKS:
                 layout = BlockLayout2D.build(NY, NX, by, bx, periodic=(True, mode == "periodic"))
-                donors, targets, kind = blocked_base_tables(p, layout, order, mode != "periodic")
+                donors, targets, kind = blocked_coarse_tables(p, layout, order, mode != "periodic")
                 ly, lx = layout.local_shape
                 L = ly * lx
                 expected, source_cells = [], []
@@ -138,7 +139,7 @@ class BlockedBaseCPU(unittest.TestCase):
         p = FixedPatch2D.from_config(config("periodic"), BOUNDS)
         order = np.argsort(p.coalesce_coarse)
         layout = BlockLayout2D.build(NY, NX, 3, 4)
-        donors, targets, _ = blocked_base_tables(p, layout, order)
+        donors, targets, _ = blocked_coarse_tables(p, layout, order)
         dense = np.arange(9 * NY * NX, dtype=np.float32).reshape(9, NY, NX)
         initial = layout.to_blocked(dense)
         arrivals = np.arange(len(targets), dtype=np.float32) + 10000
@@ -178,7 +179,7 @@ class BlockedBaseCPU(unittest.TestCase):
         # With one parent block, all eight grazing sources are owner cores.
         # Each is read for four fine ghost children: late explosion corrupts32.
         single = BlockLayout2D.build(NY, NX, NY, NX)
-        donors, targets, _ = blocked_base_tables(p, single, order)
+        donors, targets, _ = blocked_coarse_tables(p, single, order)
         initial = single.to_blocked(dense)
         staged = initial.copy()
         staged.ravel()[targets.astype(np.intp)] = arrivals

tests/validation/multigrid/test_refinement_blocked_fine.py

@@ -21,7 +21,8 @@ from CelerisLab.lbm.descriptors import D2Q9_EX, D2Q9_EY
 from CelerisLab.lbm.descriptors import SOLID
 from CelerisLab.lbm.refinement import (ACTIVE, INNER_GHOST, OUTER_GHOST, FixedPatch2D,
                                        prepare_fine_boundary)
-from CelerisLab.lbm.refinement_stepper import FixedPatchStepper, blocked_fine_tables
+from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+from CelerisLab.lbm.refinement_tables import blocked_fine_tables
 
 NX, NY, BOUNDS, STEPS = 32, 24, (8, 6, 24, 18), 8
 FINE_BLOCKS = ((32, 40), (16, 20), (16, 10))  # fine window 32x40: 1x1, 2x2, 4x2
@@ -186,8 +187,8 @@ class BlockedFineCPU(unittest.TestCase):
             backend, calls = self.mocked_backend([circle(spec)])
             c, f = backend._levels
             coarse, cur, spare = c.current, f.current, f.spare
-            flags, kind = backend._level_boundaries[1].flags, backend._levels[1].kind
-            velocity = backend._level_boundaries[1].device.get("surface_velocity")
+            flags, kind = backend._boundaries[1].flags, backend._levels[1].kind
+            velocity = backend._boundaries[1].device.get("surface_velocity")
             self.assertEqual(velocity is None, spec is None)
             backend.step(1)
             backend.close()
@@ -212,8 +213,8 @@ class BlockedFineCPU(unittest.TestCase):
     def test_actuation_without_velocity_buffer_fails_before_bfl(self):
         for speed in (0.0, .002):  # any spec, even zero amplitude
             backend, calls = self.mocked_backend([circle(actuation(speed))])
-            self.assertIsNotNone(backend._level_boundaries[1].tables["surface_velocity"])
-            del backend._level_boundaries[1].device["surface_velocity"]
+            self.assertIsNotNone(backend._boundaries[1].tables["surface_velocity"])
+            del backend._boundaries[1].device["surface_velocity"]
             with self.assertRaisesRegex(RuntimeError, "Surface velocity device buffer is missing"):
                 backend.step(1)
             names = [name for name, _ in calls]

tests/validation/multigrid/test_refinement_jets.py

@@ -62,8 +62,8 @@ class JetStepperCPU(unittest.TestCase):
         C = (16.0, 12.0)
         whole = Jet(Sector(C, .1, math.pi, 3.), Radial(C, inward=False), amplitude=0.)
         backend, _ = mocked([circle()], [whole], FINE_BLOCKS[1])
-        _, jet_id = backend._level_boundaries[1].jet_cache
-        self.assertGreater(len(np.unique(backend._level_boundaries[1].tables["block"][jet_id == 0])), 1)
+        _, jet_id = backend._boundaries[1].jet_cache
+        self.assertGreater(len(np.unique(backend._boundaries[1].tables["block"][jet_id == 0])), 1)
         backend.close()
 
     def test_fine_ownership_uses_remapped_physical_owner(self):
@@ -76,10 +76,10 @@ class JetStepperCPU(unittest.TestCase):
         source = SimObject(99, geometry=CircleGeometry(*C, 2.15), center=C, radius=2.15)
         whole = Jet(Sector(C, .1, math.pi, 3.), Radial(C, inward=False), amplitude=0.)
         backend, _ = mocked([source], [whole], FINE_BLOCKS[1])
-        _, jet_id = backend._level_boundaries[1].jet_cache
+        _, jet_id = backend._boundaries[1].jet_cache
         owners = backend.boundaries[1].curved_host[3][jet_id == 0]
         self.assertEqual(np.unique(owners).tolist(), [0])
-        self.assertGreater(len(np.unique(backend._level_boundaries[1].tables["block"][jet_id == 0])), 1)
+        self.assertGreater(len(np.unique(backend._boundaries[1].tables["block"][jet_id == 0])), 1)
         self.assertEqual(source.obj_id, 99)
         backend.close()
         cfg = config("periodic")
@@ -97,8 +97,8 @@ class JetStepperCPU(unittest.TestCase):
     def test_blocked_jet_velocity_matches_legacy_rows(self):
         legacy, _ = mocked([circle(actuation(SPEED))], (), FINE_BLOCKS[1])
         driven, _ = mocked([circle()], [jet()], FINE_BLOCKS[1])
-        a = legacy._level_boundaries[1].tables["surface_velocity"]
-        b = driven._level_boundaries[1].tables["surface_velocity"]
+        a = legacy._boundaries[1].tables["surface_velocity"]
+        b = driven._boundaries[1].tables["surface_velocity"]
         diff = np.abs(a - b).max()
         print(f"fine jet vs legacy velocity rows: max |dU| = {diff:.3e} (exact={diff == 0})")
         self.assertLessEqual(diff, 4 * np.finfo(np.float32).eps * SPEED)
@@ -106,8 +106,8 @@ class JetStepperCPU(unittest.TestCase):
 
     def test_blocked_amplitude_update_rewrites_only_velocity_column(self):
         backend, htod = mocked([circle()], [jet()], FINE_BLOCKS[1])
-        before = backend._level_boundaries[1].tables["surface_velocity"].copy()
-        device = backend._level_boundaries[1].device["surface_velocity"]
+        before = backend._boundaries[1].tables["surface_velocity"].copy()
+        device = backend._boundaries[1].device["surface_velocity"]
         with patch.object(cuda, "memcpy_htod", htod):
             htod.reset_mock()
             backend.set_jet_amplitude(0, -SPEED, level=1)
@@ -115,7 +115,7 @@ class JetStepperCPU(unittest.TestCase):
             self.assertEqual(htod.call_count, 1)
             self.assertIs(ptr, device)
             np.testing.assert_allclose(sent, -before, rtol=1e-6, atol=0)
-            np.testing.assert_array_equal(backend._level_boundaries[1].tables["surface_velocity"], sent)
+            np.testing.assert_array_equal(backend._boundaries[1].tables["surface_velocity"], sent)
             with self.assertRaises(ValueError):
                 backend.set_jet_amplitude(0, 1.5, level=1)
             self.assertEqual(htod.call_count, 1)
@@ -133,7 +133,7 @@ class JetStepperCPU(unittest.TestCase):
             with self.assertRaises(IndexError):
                 backend.set_jet_amplitude(1, -SPEED, level=1)
             self.assertEqual(backend._status, status)
-            before = backend._level_boundaries[1].tables["surface_velocity"].copy()
+            before = backend._boundaries[1].tables["surface_velocity"].copy()
             boom = Mock(side_effect=RuntimeError("copy failed"))
             with patch.object(cuda, "memcpy_htod", boom), patch.object(cuda, "memcpy_htod_async", boom), \
                     patch.object(cuda, "pagelocked_empty", side_effect=lambda shape, dtype: np.empty(shape, dtype)):
@@ -141,7 +141,7 @@ class JetStepperCPU(unittest.TestCase):
                     backend.set_jet_amplitude(0, -SPEED, level=1)
             self.assertEqual(backend._status, "failed", fine_block)
             self.assertEqual(backend.jets[1][0].amplitude, SPEED)
-            np.testing.assert_array_equal(backend._level_boundaries[1].tables["surface_velocity"], before)
+            np.testing.assert_array_equal(backend._boundaries[1].tables["surface_velocity"], before)
             with self.assertRaisesRegex(RuntimeError, "failed"):
                 backend.step(1)
             with self.assertRaisesRegex(RuntimeError, "failed"):

tests/validation/multigrid/test_shifted_jets_gpu.py

@@ -102,7 +102,7 @@ class Ledger:
 
     def step(self):
         run=self.run;self.samples.clear();run.step()
-        b=run._level_boundaries[self.index];h=run._exchange_spacing[self.index]
+        b=run._boundaries[self.index];h=b.spacing
         self.test.assertEqual(len(self.samples),2**self.index)
         expected=np.zeros_like(self.samples[0],dtype=np.float64)
         weights=np.array([h*h,h*h,h*h*h,h*h])
@@ -140,7 +140,7 @@ class ShiftedJetsGPU(unittest.TestCase):
                         level_objects=[(),(obj,)],level_jets=[(),(whole,)],observe_raw_exchange=True) as run:
                     run._initialize_raw(initial(run,rho,u=.013,v=0.))
                     expected=run._download_raw();ledger=Ledger(self,run,1)
-                    self.assertTrue(run._level_boundaries[1].tables['aperture'].all())
+                    self.assertTrue(run._boundaries[1].tables['aperture'].all())
                     for _ in range(4):ledger.step()
                     error=max(float(abs(a[:,leaf]-b[:,leaf]).max()) for a,b,leaf in
                               zip(run._download_raw().values,expected.values,run.leaf))
@@ -201,7 +201,7 @@ class ShiftedJetsGPU(unittest.TestCase):
             self.exact_raw(before,old._download_raw())
             self.exact_raw(_StoredPopulations(expected.g,True),new._download_raw())
             self.assertEqual(new.jets[2][0].amplitude,-.001)
-            np.testing.assert_array_equal(new._level_boundaries[2].tables['surface_velocity'],old._level_boundaries[2].tables['surface_velocity'])
+            np.testing.assert_array_equal(new._boundaries[2].tables['surface_velocity'],old._boundaries[2].tables['surface_velocity'])
             newledger=Ledger(self,new,2);ledger.step();newledger.step()
             self.assertEqual((old.step_count,new.step_count),(count+1,count+1))
             ledger.report('L2 source');newledger.report('L2 remapped')

tests/validation/run_refined_dfg.py

@@ -81,7 +81,7 @@ def public_jets(center, radius, scenario):
     )
 
 
-def jet_interval(handle, scenario, begin, n, diameter, *, level=None):
+def jet_interval(handle, scenario, begin, n, diameter, level=None):
     """Apply a reversal only at a completed interval boundary, before advancing."""
     if scenario != "none":
         boundaries = [round(t * diameter / SPEED) for t in (250, 260, 270, 280, 290, 300, 305)]
@@ -92,9 +92,10 @@ def jet_interval(handle, scenario, begin, n, diameter, *, level=None):
     if begin < switch:
         return min(n, switch - begin), None
     if begin == switch:
-        target = {} if level is None else dict(level=level)
-        handle.set_jet_amplitude(0, -JET_PEAK, **target)
-        handle.set_jet_amplitude(1, JET_PEAK, **target)
+        # FixedPatchStepper needs an explicit Jet level; Simulation has one level.
+        kw = {} if level is None else {"level": level}
+        handle.set_jet_amplitude(0, -JET_PEAK, **kw)
+        handle.set_jet_amplitude(1, JET_PEAK, **kw)
         return n, dict(step=begin, tstar=300.0, amplitudes=[-JET_PEAK, JET_PEAK])
     return n, None
 
@@ -257,7 +258,7 @@ def host_geometry(source, h, collision="SRT", jet_scenario="none", placement="wi
         validate_channel_config,
         validate_channel_transfers,
     )
-    from CelerisLab.lbm.refinement_tables import blocked_base_tables
+    from CelerisLab.lbm import refinement_tables
     from CelerisLab.body.geometry.circle import CircleGeometry
     from CelerisLab.body.objects import SimObject
 
@@ -269,7 +270,15 @@ def host_geometry(source, h, collision="SRT", jet_scenario="none", placement="wi
     pair = FixedPatch2D.from_config(cfg, PLACEMENTS[placement])
     layout = BlockLayout2D.build(cfg.ny, cfg.nx, *BASE_BLOCK)
     validate_channel_config(cfg, [BASE_BLOCK, FINE_BLOCK])
-    targets = blocked_base_tables(pair, layout, np.argsort(pair.coalesce_coarse))[1]
+    # Select by the imported source's own API, not by scenario: the pinned historical
+    # source (verify_source) defines blocked_base_tables, the current checkout (host
+    # tests, public Jet scenarios) blocked_coarse_tables. Exactly one must exist.
+    names = [n for n in ("blocked_base_tables", "blocked_coarse_tables")
+             if hasattr(refinement_tables, n)]
+    if len(names) != 1:
+        raise RuntimeError(f"Unexpected coarse transfer table API: {names}")
+    coarse_tables = getattr(refinement_tables, names[0])
+    targets = coarse_tables(pair, layout, np.argsort(pair.coalesce_coarse))[1]
     validate_channel_transfers(pair, layout, targets)
     BlockLayout2D.build(*pair.fine_shape, *FINE_BLOCK, periodic=(False, False))
     body = SimObject(