docs: drop 'rigorous' wording for the window tolerance (no threshold change)

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

README.md

@@ -663,6 +663,17 @@ for step in range(steps):
 streak.render("streakline.png")
 ```
 
+## Validation campaigns
+
+- [Periodic Taylor–Green AMR](docs/validation/taylor_green.md): analytic periodic-decay pilot and frozen baseline-relative assessment.
+- [Matched refined DFG](docs/validation/refined_dfg.md): frozen SRT uniform/refined comparison, requiring channel GPU gates before execution.
+- [DFG 2D-2 Re100](docs/validation/dfg_re100.md): frozen periodicity and official-reference comparisons, with a declared noncanonical outlet.
+- [Uniform LES shear decay](docs/validation/uniform_les_flow.md): independent continuum-oracle flow diagnostic, baseline and isolated corrected source variants.
+- [Smagorinsky helper comparison](docs/validation/les_helpers.md): pinned baseline versus an output-local two-coefficient candidate, without production edits.
+- [Jet-circle control volume](docs/validation/jet_circle_cv.md): bounded AMR versus uniform-grid conservation diagnostic.
+
+The runners are in `tests/validation`, their CPU tests in `tests/unit`, and compact reference/source guards in `tests/validation/data`. Current Re100/Jet source guards explicitly pin the reviewed AMR-channel experiment; historical d79 evidence remains tied to its original checkpoint commits. New runs record actual Git HEAD separately from source origin. These campaigns preserve pinned production code; successful CPU checks alone do not establish GPU or physical validation. Keep generated results outside tracked files.
+
 ## Citation
 
 ```bibtex

docs/development/refinement_mrt.md

@@ -0,0 +1,34 @@
+# MRT reuse in blocked/refined stepping (candidate)
+
+Base `9805383`. Opt-in by `LBMConfig(collision="MRT")`; SRT remains the default and its
+kernel statements, including the shifted retained-delta/direct-momentum path, are unchanged
+under the alternate compile branch.
+
+## What changes
+- `blocks.cu` (`BlockStep`), `step/block_channel.cu` (`BlockChannelStep`) and `refinement.cu`
+  (`patch_collide`) gain a compile-time `COLLISION_MODEL == 2` branch that repeats the uniform
+  `collide_dispatch` sequence with zero forcing: `compute_rho_u`, omega clamp to
+  `[OMEGA_COLLISION_MIN, OMEGA_COLLISION_MAX]`, `collide_mrt_no_force`.
+- Kernel and host guards accept SRT or MRT only. TRT, LES, non-FP32 storage, non-double-buffer
+  streaming, and shifted dense/mixed refinement remain rejected.
+- Transfers, subcycling, BFL, raw exchange ledgers and channel closures are untouched. The
+  existing pressure-NEQ outlet selects its MRT branch (unknowns-only NEQ extrapolation) from
+  `COLLISION_MODEL`, so MRT channel parity must use uniform MRT as reference.
+
+## Contracts (existing operator, not redefined here)
+- Shifted storage uses the existing MRT representation (equilibrium moments from `rho - 1`),
+  not the SRT retained-delta/direct-momentum precision. Mass and momentum are conserved
+  moments (relaxation rate 0).
+- Kinetic rates are the operator's fixed `s_e = s_eps = s_q = 1.2` on every level; the shear
+  rate is the existing per-level omega. Under acoustic 2:1 scaling the physical bulk viscosity
+  therefore halves on each finer level. No rate rescaling is applied.
+- No LES, no external forcing (`d_params` force fields stay zero).
+- Shifted Jets stay SRT-only: `body/jets.py` keeps its existing guard. This extension does not
+  provide shifted MRT Jets.
+- `PeriodicBlockStepper` (uniform blocked stepper) remains SRT-only.
+
+## Evidence status
+CPU: `tests/unit/test_refinement_mrt.py` (host guards, source contract). GPU:
+`tests/validation/multigrid/test_refinement_mrt_gpu.py` (uniform-MRT channel parity, AMR MRT
+smoke) and unchanged SRT GPU tests. Strict parity may be subject to cross-kernel FMA
+contraction differences already documented for SRT. Not a DFG/benchmark validation.

docs/development/root_channel_candidate.md

@@ -0,0 +1,22 @@
+# Blocked AMR root channel increment
+
+Opt in with `FixedPatchStepper(..., coarse_boundary="channel", base_block=(by,bx), fine_block=(fy,fx))`, or the existing level-list form. Every level must be blocked. The supported boundary is the existing uniform D2Q9/SRT/FP32 double-buffer channel: local Zou-He west inlet, pressure-NEQ east outlet, stationary halfway bounce-back y walls. Uniform and parabolic inlet profiles and shifted/unshifted storage are supported; the existing configuration guards still reject MRT/LES and incompatible storage. This does not change the external DFG MRT reference.
+
+`step/block_channel.cu` is a small root-only kernel. Fine kernels, transfer kernels and recursive hop scheduling are unchanged. Global coordinates and the global profile height/velocity are separate from block-local storage strides. Corner ownership and collision phases match `OneStep`:
+
+- y-edge ghosts (including corners): pull, pair swap, no collision
+- west ghosts at interior y: Zou-He reconstruction then SRT collision
+- east ghosts at interior y: pressure NEQ from the immutable left input node, no collision
+- physical fluid beside y walls: same-node opposite-input correction then collision
+
+The root uses the uniform SRT equilibrium arithmetic, including the existing shifted rho-minus-one path, to preserve matched-root semantics. Fine `BlockStep` retains its existing shifted direct-momentum/retained-density arithmetic. This increment is not a migration of fine arithmetic.
+
+Root halo topology stays wrapped like the uniform kernel, solely to reproduce its complete evolved ghost state as well as physical populations. It does not make the physical inlet/outlet periodic. Evolved root state includes boundary ghosts, while public fluid-leaf downloads exclude them. Initialization must provide finite positive-density root ghost states. Private raw snapshots preserve them exactly for same-representation continuation; public leaf-only downloads are not restart states. Channel remapping remains explicitly unsupported.
+
+`channel.py` checks the actual explosion donors and restriction receivers are uncovered bulk fluid and checks actual receiver-local restriction slots against pressure-donor and wall-correction reads. No arbitrary large padding or block-aligned patch restrictions are added. Keep the physical domain dimensions exact: choose legal root block divisors (including full width) instead of padding a benchmark domain. Stationary circles/Jets retain the existing finest-level and transfer-footprint rules.
+
+## Gates
+
+- `tests/unit/test_channel_root.py`: host semantics, minimal footprint, overlap rejection, raw-state/public-leaf separation; actual production uniform and blocked kernels under a CPU launch shim, five non-equilibrium steps, all corners, profiles, representations, y seams and a one-cell-wide block outlet donor.
+- `tests/validation/multigrid/test_channel_root_gpu.py`: exact uniform-versus-blocked full-state GPU parity; AMR finite-state and block-split invariance with/without a stationary circle and velocity Jet; amplitude update and raw snapshot round-trip.
+- GPU tests require explicit device admission by the caller. CPU-header algebra is not CUDA execution or external physical validation. The current increment does not establish DFG drag/lift/pressure acceptance or a performance speedup; those need matched SRT controls and a separate benchmark run.

docs/shifted_blocked_candidate.md

@@ -0,0 +1,61 @@
+# Opt-in shifted blocked candidate
+
+This candidate is based on fbe7377. The default remains unshifted. Enable with
+`LBMConfig(ddf_shifting=True)` only on the supported routes below.
+
+| Route | Candidate support |
+|---|---|
+| PeriodicBlockStepper, FP32 D2Q9 SRT double-buffer, no LES | shifted or unshifted |
+| FixedPatchStepper, every level blocked and same representation | shifted or unshifted |
+| Passive stationary circular BFL boundary and raw exchange | supported candidate scope |
+| Synchronized three-level child remap | private raw shifted snapshot, exact overlap/copy |
+| Dense root/refinement or mixed representation | rejected for shifted |
+| Slot/reservoir, surface actuation, velocity Jet | existing public shifted exclusion retained |
+| Other collision/lattice/storage/streaming/LES combinations | existing exclusions retained |
+
+The D2Q9 equilibrium helper accepts a retained density offset only in shifted
+compilation. BlockStep computes delta=sum(g) before adding the unit density;
+velocity uses the existing macro calculation. The existing SRT collision and
+all unshifted expressions remain. No global precision type or duplicate
+collision implementation is introduced. Ordinary shifted Simulation callers
+still use the existing rho−1 overload; this candidate does not change their
+numerical method.
+
+Public initialize receives physical populations and performs shifted subtraction
+in FP64 before the one FP32 storage cast. Public shifted download returns FP64
+physical populations with exact-rational D2Q9 background weights. Internal raw
+snapshots retain FP32 g and a representation tag; runtime remap never decodes
+and re-encodes overlap. Exact mass is interpreted as 1+sum(g), and momentum as
+sum(c*g); analyses should use raw values for tiny perturbation budgets. The
+FP32 BFL force constants have their own arithmetic rounding; no exact global
+conservation claim follows from shifting.
+
+Planar helper formulas include their shifted reference-density offsets and have
+actual-header host tests. Their public shifted route remains excluded until a
+separate test-gated migration. Next velocity-Jet work must validate physical
+wall density, incoming/outgoing populations, raw force/mass, ownership and
+transfer interactions before relaxing that restriction.
+
+Evidence so far: selected-cell CUDA replay favored retained-delta shifting;
+this is not full-solver validation. Host algebra/storage/remap tests and staged
+GPU continuation, three-grid cylinder physics, raw budgets, Couette helper
+stress and measured runtime are separate gates. Original failed experiments
+remain evidence; no acceptance threshold is relaxed.
+
+## Direct-input-momentum equilibrium candidate
+
+The shifted BlockStep now constructs the linear equilibrium term from input
+momentum j, without reconstructing rho*u. For opposite axis pairs the odd term
+is j/3; for diagonal pairs it is (jx +/- jy)/12, evaluated with explicit
+round-to-nearest division. The even term is
+w*[delta + 4.5*(c dot j)*(c dot u) - 1.5*(j dot u)].
+This is the same second-order D2Q9 equilibrium in exact arithmetic. Original SRT,
+storage, boundaries, and transfers are unchanged. No collision residual or force
+is corrected afterward. Existing four/five-argument equilibrium APIs are retained;
+only shifted blocked collision calls the new six-argument momentum helper.
+
+The all-fluid read-only CUDA replay reduced signed x-momentum error by 120–2285x
+at five sampled states, with lower mass/momentum and second-moment RMS errors.
+Signed mass bias was larger (~1e-12/cell/hop), so full evolving mass budgets remain
+mandatory. This local evidence does not establish physical convergence or speed.
+The unshifted default and its arithmetic path remain unchanged.

docs/validation/dfg_re100.md

@@ -0,0 +1,60 @@
+# DFG 2D-2 Re100 GPU campaign (validation-only)
+
+## Current experimental source provenance
+
+This development branch now includes the reviewed AMR-channel source changes originating in `5338864748ac67fe2544487c059ba149ebf916da` and `47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5`. The current full-source manifest pins the latter reviewed source snapshot, including the new channel files. This is an explicit new experiment; neither source review nor host tests transfer historical d79 GPU/flow acceptance to it. The channel GPU gate has not yet run at integration time.
+
+New run JSON records the source-origin pin plus `checkout_revision` (actual Git HEAD), `checkout_dirty`, and the experimental-source label. The origin pin is not claimed to be the run's checkout commit. Numerical definitions and thresholds are unchanged. Runs require a Git checkout for this runtime identity; source hash checks remain active.
+
+Historical d79 Re100/Jet execution stays reproducible at `69e9b78de9483a5ff21e2d0c74311d240bfaf88c`; the d79 LES diagnostic checkpoint is `20ae6770924df15f31305aad56e44a54a3dd4d88`. Their manifests and existing result files are not rewritten. Use those commits to reproduce or interpret old d79 evidence. The current AMR-channel experiment requires its own GPU and applicable flow evidence.
+
+## Scope and locked definitions
+
+Explicitly pinned AMR-channel experimental production snapshot (see provenance below). Uniform D2Q9/MRT, shifted FP32, double buffer, no LES, stationary cylinder and no-slip walls. This is not an AMR, jet, or full canonical-outlet certification.
+
+Physical geometry: [0,2.2] × [0,0.41], center (0.2,0.2), radius .05, rho=1, nu=.001, inlet mean=1 and peak=1.5. D is lattice diameter; dx=.1/D, dt=U_lattice*dx, nu_lattice=U_lattice*D/100. U_lattice=.01 gives peak inlet Mach .02598. Existing local Zou–He inlet is selected in validation config only. Existing pressure-NEQ outlet is unclamped: **it does not impose the canonical vector do-nothing condition nu*d_n(u)-p*n=0**. The existing production channel-flow initial condition is used, not zero velocity. Periodic eligibility is delayed to t*=250 (25 physical seconds), with fixed end t*=300 (30 seconds).
+
+Official source: https://wwwold.mathematik.tu-dortmund.de/~featflow/en/benchmarks/cfdbenchmarking/flow/dfg_benchmark2_re100.html
+
+The 20141124 official newer reference archive has higher-order line-integration force data and supersedes the older example table for this comparison. The compact tests/validation/data/dfg_re100/reference_cycle.json preserves the last complete lift-minimum cycle from lv6/dt4. Reference Cd=[3.164266,3.227392], Cl=[-1.021286,.986571], pressure difference=[2.420371,2.507013], St=.301841. These are derived sampled extrema, not exact analytical values. Official mean=(min+max)/2 and amplitude=max-min; separate time averages are also recorded.
+
+## Exact commands and startup ownership
+
+Use the existing CUDA/PyCUDA environment; no package installation or production edits required. From the repository root, use a new task-owned output path:
+
+    export PYTHONPATH="$PWD/src${PYTHONPATH:+:$PYTHONPATH}"
+    python -m pytest -q tests/unit/test_re100.py
+    CUDA_VISIBLE_DEVICES=0 python tests/validation/run_dfg_re100.py --diameter 20 --speed .01 --mode pilot --output output/validation/pilot_d20_u001 > pilot.log 2>&1
+
+Pilot is t*=5 (10,000 steps), with dense .02-t* sampling to calibrate diagnostic cost. It cannot pass periodicity or external benchmark validation. Root launch owner reviews health, memory, timing and stdout before full runs; the full pair is not launched automatically.
+
+    CUDA_VISIBLE_DEVICES=0 python tests/validation/run_dfg_re100.py --diameter 40 --speed .01 --mode full --output output/validation/full_d40_u001 > d40.log 2>&1
+    CUDA_VISIBLE_DEVICES=1 python tests/validation/run_dfg_re100.py --diameter 80 --speed .01 --mode full --output output/validation/full_d80_u001 > d80.log 2>&1
+
+These are separate invocations, not a parallel shell script. A single root launch owner serializes startup/compilation across all Horsea work. Do not launch the next process until the previous log contains an explicit `READY {...}` line. READY comes after the first physical sample, including Simulation initialize, sim.run, read_force and get_macroscopic, so lazy runtime/diagnostic paths are exercised. That sample remains in the physical trajectory/history. No hidden warmup/reset, no later new diagnostic kernel. Do not treat a process merely existing as READY. Keep separate output directories and retain exit codes/logs. GPU assignment is chosen by root after checking existing workloads; shown indices are examples of physical device selection.
+
+## Measurement, frozen criteria and outcomes
+
+Force uses advancing link-momentum exchange, accumulated in batches <=100 steps, combined in FP64; Cd/Cl=2F/(U_lattice² D). Force timestamp is the batch midpoint. Pressure uses fluid-only h/2h extrapolation to cylinder front/back points and physical conversion 1/(3 U_lattice²), at each sample endpoint. For phase comparisons endpoint pressure is interpolated to the same cycle phase. Both raw times are saved.
+
+Before t*=240 sample width=.5; afterwards .02 (about 166 samples per lift cycle). Pilot uses dense sampling throughout. Extract complete lift-minimum-to-minimum cycles using quadratic local minimum locations, accepting periods 2–5 t*. Latest five complete cycles must be consecutive with the last ending no more than 5 t* before the final observation, and must start at or after t*=250. Eligibility requires lift amplitude >=.2, period range/mean <=.005, Cd/Cl/pressure min and max ranges <=.005 absolute, cycle-mean mass range/mean <=1e-4, and each cycle's mean mass-flux imbalance <=.005. Instantaneous flux equality is not required in this unsteady compressible LBM diagnostic.
+
+Reference gates: Cd extrema relative errors <=2%, Cl extrema absolute errors <=.03, pressure extrema relative errors <=2%, St relative error <=2%. These are predeclared engineering tolerances, not official uncertainty intervals. Phase-aligned waveform RMS errors, extrema, midpoint means, time means, amplitudes, outlet residual, sampled density bounds and Mach are all reported. Waveform RMS is diagnostic (no retroactive gate). No threshold changes after GPU results. At cap, failed periodicity stays incomplete. Invalid/nonfinite state, rho outside [.5,1.5], or sampled max Mach>.2 causes a safety stop; this is a gross-health guard, not a Mach-independence claim.
+
+A case_reference_pass requires both periodicity and reference gates. Every raw case keeps suite_validated=false. The two-grid comparison must report each case and errors; matched-Mach D40/D80 is evidence of resolution sensitivity only, not an asymptotic convergence order. If either fails or finer errors increase materially, report the failure instead of adjusting thresholds. No full canonical/AMR/Mach-independent validation claim is allowed.
+
+## Resource and calibration plan
+
+D20: 441×84=37,044 nodes, pilot 10,000 steps. D40: 881×166=146,246 nodes, 1,200,000 full steps. D80: 1761×330=581,130 nodes, 2,400,000 full steps. Main double-buffer DDF arrays alone cost 72 bytes/node: 2.7/10.5/41.8 MB decimal. Flags, masks, macro host arrays, body buffers and CUDA context/JIT memory add overhead; root verifies actual free GPU/host memory and disk before launch. Conservative per-process reservation: 1 GiB VRAM and 2 GiB host RAM, excluding shared/compiler overhead; measurements take precedence.
+
+Historical Re20 D80 U=.01 step throughput was ~6082 MLUPS (45.9 seconds/480k steps). At equal throughput D80 Re100 full advancing steps would take ~229 seconds, excluding dense full-field downloads/reduction and compile overhead. D40 estimate must be calibrated, not assumed proportional. Pilot records startup_seconds, step_seconds, measurement_seconds, mlups. For each full grid, calibrate a same-grid pilot if D20 cannot predict cost adequately. Full has about 3480 records (480 coarse+3000 dense); use measured seconds/step × full steps plus seconds/record ×3480 plus startup. D20-to-D80 extrapolation is only a rough planning upper/lower guide, not an ETA commitment.
+
+Only geometry.npz, one final_field.npz, history.csv, summary.json and stdout/stderr are persisted. No per-step or repeated full-field snapshots. D80 final float64 rho/ux/uy plus mask are <15 MB uncompressed; history/summary several MB, compact versioned reference/source guards <50 KiB. Budget <100 MB per case and <500 MB whole campaign including logs and optional externally retained archives, safely below 2 GiB; root still checks baseline Horsea usage. Retain sparse evidence only.
+
+## CPU validation and remaining execution
+
+Ten tests cover geometry/units/config, pressure scaling, periodic and nonperiodic synthetic histories, immutable production hashes, and self-reference comparison. CUDA execution is not available in the preparation environment and has not been claimed. Run pilot first on Horsea; do not confuse successful CPU tests with GPU readiness.
+
+## Reference provenance and reproduction
+
+Only the compact derived cycle is versioned. Its source URL, official ZIP SHA-256, column selection, and cycle-selection description are embedded in reference_cycle.json. Raw archives and GPU results are not repository inputs. To reproduce it, obtain the exact archive identified there, verify its SHA-256, and extract logs/results_cc2d_bench2_referencenew/bdforces_q2_lv6_dt4 and pointvalues_q2_lv6_dt4. Use physical time column 2 multiplied by 10, drag/lift columns 4/5, and pressure columns 7 minus 12 (one-based columns). Build rows with mass=flux_in=flux_out=1, then call periodic.extract(rows, min_time=0) and select the last complete cycle. The source guard is tests/validation/data/dfg_re100/source_manifest.json; it verifies the current reviewed production snapshot rather than equating its origin with the validation branch HEAD.

docs/validation/jet_circle_cv.md

@@ -0,0 +1,115 @@
+# Jet-circle closed-control-volume pilot
+
+## Current experimental source provenance
+
+This development branch now includes the reviewed AMR-channel source changes originating in `5338864748ac67fe2544487c059ba149ebf916da` and `47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5`. The current full-source manifest pins the latter reviewed source snapshot, including the new channel files. This is an explicit new experiment; neither source review nor host tests transfer historical d79 GPU/flow acceptance to it. The channel GPU gate has not yet run at integration time.
+
+New run JSON records the source-origin pin plus `checkout_revision` (actual Git HEAD), `checkout_dirty`, and the experimental-source label. The origin pin is not claimed to be the run's checkout commit. Numerical definitions and thresholds are unchanged. Runs require a Git checkout for this runtime identity; source hash checks remain active.
+
+Historical d79 Re100/Jet execution stays reproducible at `69e9b78de9483a5ff21e2d0c74311d240bfaf88c`; the d79 LES diagnostic checkpoint is `20ae6770924df15f31305aad56e44a54a3dd4d88`. Their manifests and existing result files are not rewritten. Use those commits to reproduce or interpret old d79 evidence. The current AMR-channel experiment requires its own GPU and applicable flow evidence.
+
+Status: CPU preparation only. No GPU run or production change. Current source origin: reviewed AMR-channel snapshot 47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5; this new experiment has no inherited d79 GPU acceptance. tests/validation/data/jet_circle_cv/source_manifest.json checks all tracked source file hashes before importing GPU code.
+
+## Execute only after root-owned GPU serialization
+
+From the repository root, with REPO set to the repository root and an absent output directory:
+
+```sh
+PYTHONPATH=src python -m pytest -q tests/unit/test_jet_circle_cv.py
+python tests/validation/run_jet_circle_cv.py --repo "$REPO" --output "$OUTPUT" --plan
+CUDA_VISIBLE_DEVICES="$VERIFIED_IDLE_GPU" python tests/validation/run_jet_circle_cv.py --repo "$REPO" --output "$OUTPUT"
+```
+
+The runner does not select, reserve, or contact a GPU machine. It requires one explicit CUDA_VISIBLE_DEVICES value, at least 1 GiB free, and refuses an existing output directory. It imports existing PyCUDA; no installation. Default pilot: 800 base steps, 8 output windows, at most 1800 seconds after READY. READY is printed only after both steppers compile, execute one base/two fine warmup hops, exercise download/raw paths, and reset initial state/time. A pre-READY silence may be compilation, not stepping progress. Root owns any overall startup timeout and all launch serialization.
+
+## Physical/numerical case
+
+Periodic root box 192 by 128, center (96,64) in root cell-edge coordinates, D=24.084, U0=.03, nu=.0072252, Re_D=100, beta=Vj/U0=.1, rho0=1. Unforced D2Q9 SRT FP32, no DDF shifting, no LES. The slightly noninteger diameter was selected solely to avoid q clipping at all three planned scales. Circle is stationary. Two uniform radial sectors at pi/2 and 3pi/2, half-width pi/12: top blows outward +.003, bottom sucks inward -.003; remaining wall has zero velocity. This periodic transient is NOT DFG, infinite-stream drag, or statistical high-Re validation. End time is t*=t U0/D approximately .9965.
+
+AMR bounds (48,32,144,96). Finest dx=dt=.5 in root lattice units, nu_fine=2 nu, finest D=48.168. Full-domain fine reference is 384 by 256 with exactly the same finest spacing, timestep, viscosity, collision and storage. Root tau=.5216756, fine tau=.5433512, Mach approximately .05196.
+
+AMR uses actual public FixedPatchStepper Jet construction and raw observation. The uniform blocked reference reuses the existing test_raw_exchange_gpu.py optional-input validation adapter: production Jet-derived velocity/force-arm arrays feed the unchanged blocked curved kernel. Its normal public PeriodicBlockStepper API is passive-only. This does not claim there is no uniform Jet route anywhere: Simulation.add_jet exists, but the same-blocked-kernel adapter removes a backend/storage confound for this comparison. No new public backend is introduced. Runtime asserts matching directions, q, fallback and true hit coordinates after frame conversion; both sectors must resolve links. Unclamped q/no fallback is required.
+
+## Independent measurements and signs
+
+Two closed rectangles: (72,44,120,84) and (64,40,128,88), wholly within the finest patch. Every contour face has two active-fluid stencil layers. Capture physical postcollision populations at t=0 and every completed base interval. Collision conserves density/momentum; these are valid endpoint moments, NOT precollision stress estimates.
+
+For each contour, M and P sum fluid-cell rho and rho*u times h^2, excluding the staircase solid. Face midpoint quadrature integrates rho*u dot n and [rho*u*u+p*I-mu*(grad u+grad u.T)] dot n. Velocity gradients use centered differences in root-length units, mu=rho*nu, p=(rho-1)/3. Constant pressure subtraction integrates to zero on a closed rectangle. No raw exchange enters these field quantities.
+
+Trapezoidal time integration samples every base step over the SAME reported [step_start,step_end] as cumulative raw impulses/masses. For each window:
+
+- mass into fluid through body surface = delta M + integral outer mass flux
+- raw boundary reaction target = -delta P - integral outer momentum-minus-traction flux
+
+Raw rows are [Jx,Jy,angular impulse,mass into fluid]; each uniform fine hop scales by [.25,.25,.125,.25], exactly as the AMR interval ledger. Both fine hops are accumulated. AMR observation time endpoints and the public raw reduction are checked. Reports separate total, wall, aperture, blowing and suction raw rows. Coefficients divide force by .5*rho0*U0^2*D. These are momentum reactions including aperture transport; neither raw aperture nor total raw reaction is labelled solid-only hydrodynamic traction.
+
+Independent aperture mass estimates integrate measured rho*u dot e_r over true arcs using 64-point Gauss quadrature and nearest active-fluid samples. No target velocity or raw row enters this estimate. Reports include maximum sampling distance, actual signed raw mass rates and nominal signed volume flux ±2*(pi/12)*R*Vj. Near-wall nearest-node sampling is only a first-order diagnostic, not certified wall extrapolation. Nominal volume flux is not asserted to equal measured mass flow at finite density variation.
+
+Inventory near the circle is a staircase-cell approximation; its geometric error and face/time quadrature errors can appear in CV residuals. Two-contour spread and grid refinement are essential. No exact continuum force value, solid traction recovery, jet hardware load, or physics acceptance is claimed. Output status is deliberately PILOT_DIAGNOSTIC_ONLY even on successful execution.
+
+## Scope, refinement, resources
+
+Default one pilot only. Estimated device envelope 128 MiB, required free 1 GiB; bounded JSON only, conservatively under 2 MiB, no population snapshots or per-step logs. Every base step downloads both fields for independent quadrature, so PCIe/CPU diagnostic cost is intentional. Six CPU tests cover descriptor/source identity, uniform flux cancellation, affine convective divergence, quadratic viscous sign/coefficient, matched geometry/Jet selection and diffusive scaling.
+
+Optional --scale 2 or 4 doubles/quadruples all lengths and divides U,V by that scale, holds lattice nu fixed, and scales steps/windows by scale^2: physical Re and t* match, Mach decreases. Source and geometry guards repeat. Device estimate scales as scale^2; memory gate requires twice the estimate or 1 GiB, whichever is greater. Scale 4 is not part of default authorization/readiness. --re 400 requires --allow-re400 and a separate reviewed root launch decision after Re100 evidence; no automatic promotion.
+
+Compare AMR/reference velocity error, signed reaction coefficients, each contour's independently measured reaction and residual, contour spread, signed mass budgets, and measured aperture flux. Freeze numerical tolerances only after agreeing a scientifically appropriate convergence criterion; this preparation imposes only structural/finite-state guards. A completed pilot cannot establish convergence.
+
+Follow-on design retained: porous Couette crossflow with u=U*expm1(Re_V*y/H)/expm1(Re_V), v=V and unequal exact wall shear. It requires validation-only planar geometry and is intentionally not implemented here.
+
+### Amplitude, geometry, and high-Re limits
+
+Both Jet amplitudes are constant throughout the run and every fine substep; geometry is fixed. There is no runtime switch, per-step forcing schedule, or moving-wall velocity in this campaign. Construction/packing happens once, with identical uploaded amplitudes in the uniform adapter. This avoids any substep timing ambiguity from amplitude changes.
+
+Every case records actual tau_root and tau_fine. For this fixed s=1 geometry and speed, Re400 would give nu=.0018063, tau_root=.5054189 and tau_fine=.5108378, close to the SRT limit .5. The implemented Re choices are a bounded campaign menu, not a universal physical Reynolds-number ceiling. Higher-Re evidence may require a larger initial resolved geometry or a different justified grid/relaxation design before launch. Diffusive refinement at constant lattice nu improves spatial/Mach resolution but does not move tau farther from .5; simply selecting scale=2 does not cure a poor tau. No all-Re, LES or three-dimensional realism claim follows from this pilot.
+
+The source guard checks production file content against the current reviewed source origin, not validation-branch HEAD. This runner is validation-only; generated results and frozen execution-package archives are not committed.
+
+## Measurement-only continuation (frozen design before execution)
+
+The original 800-step pilot and its audit remain historical evidence, not an acceptance target. The following adds diagnostics only; it neither modifies production kernels nor adjusts velocity commands to improve measured flux. Use the canonical validation Git branch, with the same source-manifest check and checkout provenance record. No standalone replacement source archive is needed.
+
+Run the pair serially under root-owned GPU allocation, from the checked-out repository:
+
+```sh
+python tests/unit/test_jet_circle_cv.py -v
+python tests/unit/test_jet_cv_measurements.py -v
+CUDA_VISIBLE_DEVICES="$VERIFIED_IDLE_GPU" python tests/validation/run_jet_circle_cv.py --repo "$PWD" --output "$NEW_OUTPUT_S1" --continuation --steps 3200 --scale 1
+CUDA_VISIBLE_DEVICES="$VERIFIED_IDLE_GPU" python tests/validation/run_jet_circle_cv.py --repo "$PWD" --output "$NEW_OUTPUT_S2" --continuation --steps 3200 --scale 2
+python tests/validation/analyze_jet_cv_continuation.py "$NEW_OUTPUT_S1" "$NEW_OUTPUT_S2" --output "$COMPARISON_JSON"
+```
+
+Both output directories must be absent. `--plan` prints configuration without CUDA. `--continuation` deliberately accepts only Re100, scale1/2, and 3200 reference steps: a bounded campaign, not a physical Re ceiling. Actual steps are 3200 and 12800, reaching the same t*=3.986048829. Physical circle/box/contours and beta remain fixed; all root lengths grow by s, U and jet speed shrink by s, lattice nu remains .0072252. Both runs retain tau_root=.5216756 and tau_fine=.5433512, with Mach .0519615/s. This diffusive refinement changes the finite-Mach/compressibility error as well as grid resolution; it is not an isolated spatial-error order study.
+
+### Units, cadence, and reconstructibility
+
+Reference length/time are one scale-1 root lattice length/time; density reference is one. Convert scale-s root quantities by: length /s, time /s², velocity times s, mass /s², momentum and impulse /s, force times s, angular impulse /s², pressure times s². Volume and mass rates are unchanged. The nominal signed aperture volume rate is ±.0189155293672 in reference units at every scale. A mass flux divided by rho_ref times that rate is a reference-density diagnostic, not an actual-volume-flow ratio. Saved Q and M remain separate.
+
+Exact raw observations accumulate every completed base interval, including both fine substeps. Field downloads occur only at shared reference times 0..100 at dt=1, then 100..3200 at dt=4: 876 samples per run, compared with 3201/12801 possible every-base samples. Every 100-reference-time window boundary is included. Trapezoidal integration uses the actual variable time gap, with the raw ledger covering exactly the same endpoints. No missing sample is silently treated as unit dt. Sampling is denser during abrupt startup; it is not a guarantee of resolving all high-frequency acoustic modes. Offline coarsening can expose some time-quadrature sensitivity but cannot rule out unresolved-frequency aliasing.
+
+`samples_s*.npz` stores at every sample: density minima/maxima/max perturbation and max speed for uniform fluid, AMR root leaves and AMR fine leaves; composite full-domain mass; CV endpoint inventory and flux; separate advective, pressure and negative-viscous-traction momentum fluxes; nearest-fluid mass proxy; extrapolation ray primitive values; linear/quadratic reconstructed Q and M; and cumulative raw wall/aperture/total/blowing/suction rows. Associated JSON identifies axes, units, conversion factors, hashes and snapshot metadata. Composite AMR mass sums root leaves at area1 and fine leaves at area.25, then divides by s². Covered cells are excluded. Global periodic mass change minus total raw body mass is reported, so a transfer/interface mass error cannot hide behind the two contours lying wholly within the fine patch.
+
+Five compressed primitive snapshots at reference times 0,100,800,1600,3200 retain rho, reference velocity, origin and spacing for uniform, AMR root and AMR fine grids. Inactive/solid entries are NaN. These are full primitive grids at sparse times, not downsampled images or distributions. They permit independent spatial quadrature, field inspection, masks and inventory reconstruction. Per-sample split flux/inventory arrays allow independent temporal integration without rerunning GPU work. FP32 saved primitive snapshots introduce their own rounding; FP64 sampled diagnostic/ray arrays remain in the sample file.
+
+### Independent aperture estimators and limitations
+
+Thirty-two Gauss points per true aperture arc, with rays at distances 1.5,2.5,3.5 finest-cell widths into the fluid, sample rho,u,v by bilinear interpolation using only active fluid nodes. No commanded wall value or raw population enters reconstruction. Linear and quadratic outward-normal extrapolations separately reconstruct rho and u; integrate Q=integral(u.n ds) and M=integral(rho*u.n ds). Retain ray positions, normals, weights, primitive samples and coefficients, making every estimate reproducible. The nearest-fluid 64-point mass proxy is preserved for continuity with the pilot and remains explicitly a proxy.
+
+Neither estimator is an exact wall value or validated wall model. Extrapolation amplifies noise and may yield unphysical density; do not clamp it to force agreement. Sector endpoints contain nonsmooth prescribed velocity changes, and interpolation may mix their neighborhood. Linear-versus-quadratic spread, ray density extrema, resolution changes and analytic manufactured tests constrain interpretation, but their agreement is not proof of low bias. Two grids alone cannot establish asymptotic convergence. Nominal geometry/profile flux is exact for the commanded uniform radial profile; comparison with reconstructed Q diagnoses estimator plus numerical-boundary error, not automatically a Jet defect. Raw reaction remains momentum transport plus traction; no nozzle-load or solid-only force claim.
+
+### Predeclared comparison criteria
+
+These criteria are fixed before the continuation runs. There is no new physical acceptance threshold and no tolerance fitted to the pilot:
+
+1. Structural validity: exact source/provenance, matched geometry/parameters, unclamped q/no fallback, both apertures resolved, finite positive active-node density and finite velocity at every saved field sample, complete scheduled samples, contiguous matching raw intervals, and all declared files/hashes. Runtime does not promise positivity between saved samples.
+2. Report every window's signed and gross quantities, maximum sampled errors, and early (reference t<=800) versus later (t>800) norms separately. Normalize mass residual by rho_ref*abs(Q_command)*window_duration, and reaction residual by .5*rho_ref*U_ref²*D_ref*duration. Avoid dividing by a near-zero net mass or instantaneous force.
+3. Recompute contour residuals using the complete sampled grid and every second/fourth sample, always retaining endpoints. Save their changes as a temporal-sensitivity proxy, not an error bound. Compare both contours and both backends; show mass, Cx and Cy components separately.
+4. A directionally resolved reduction in a component's RMS CV residual is flagged only if coarse-minus-fine RMS exceeds twice the largest observed temporal coarsening change. This is an evidence-of-trend flag, never a physics PASS. If a component plateaus or grows, decompose temporal/spatial/staircase/boundary effects before proposing production changes.
+5. Compare reconstructed Q/command and M/(rho_ref*command) separately, at identical physical windows, retaining both extrapolations. Regard refinement improvement as persuasive only when it exceeds estimator spread and temporal variation. Inspect density and global mass alongside it. Agreement of AMR and uniform ledgers alone cannot validate their shared boundary treatment.
+6. Four D/U in a periodic unforced domain remains transient consistency evidence. Two scales do not prove order, stable shedding, infinite-stream drag, all-Re capability, or readiness to promote Re400. Formal physical acceptance needs an additional reference/error budget and appropriate time/domain evidence.
+
+### Runtime and storage budget
+
+Actual earlier scale-1 pilot cost was 52.3884s launch-to-exit, 46.1414s after READY for800 steps with full-field diagnostics every base step. A naive same-diagnostic3200-step scale-1 extrapolation is about210s; the new876 field samples reduce transfer count but add ray diagnostics and compression. Use conservative planning envelopes210–500s for scale1 and900–3200s for scale2, not timing guarantees. The scale2 conservative envelope accounts for4x cells and4x root steps; sparse sampling mitigates but does not eliminate that cost. Post-READY safety caps are1200s and5400s respectively, replacing the unsuitable fixed1800s assumption. Root owns compilation/startup timeout and serial launch decisions.
+
+Device-envelope estimates remain128/512MiB for s1/s2, at least1GiB free (or twice estimate). Output estimate is under64MiB per run including five primitive snapshots and876 sample records; no all-time populations are saved. Plan/metadata and sampled files are rewritten at each100-reference-time window to retain bounded partial evidence if interrupted. READY still follows compiled warmup and state/time reset, before measured stepping begins.

docs/validation/les_helpers.md

@@ -0,0 +1,40 @@
+# Smagorinsky baseline/candidate helper diagnostic
+
+This validation-only diagnostic compares exact d79 production headers with a temporary candidate. It does not modify production source. The candidate changes only the D2Q9 and D3Q19 closed-form factors from 18 to 18 sqrt(2), plus comments defining |S| = sqrt(2 S:S) and self-consistent effective tau. The small reviewed patch and expected header hashes are in `tests/validation/les/data`; no copied headers, solver trees, raw results, or binary fixtures are versioned.
+
+## Prepare and verify on the CPU
+
+From the repository root, choose a new output directory:
+
+```sh
+python tests/validation/les/prepare.py --output-dir output/les-host/candidate
+python tests/validation/les/test_candidate.py --candidate-dir output/les-host/candidate --output output/les-host/host_results.json
+```
+
+Preparation checks both actual production header hashes against d79, copies only `descriptors.cuh` and `turbulence_smag.cuh` into output-local baseline/patched include trees, applies the reviewed patch only to the latter, and verifies all four resulting hashes. It refuses an existing candidate directory. The host test verifies candidate hashes again and refuses an existing evidence file. The Git branch HEAD may include validation commits; header content, rather than HEAD identity, is the pin.
+
+The independent host oracle constructs second-order Hermite populations and solves the independently derived quadratic by bisection:
+
+    tau * (tau - tau0) = 9 Cs^2 ||Pi_neq|| / (sqrt(2) rho)
+
+There are **736 parameter cases**, covering D2Q9/D3Q19; Cs=0,.08,.16,.24; tau0=.6,.8,1.2; rho=.7,1,1.3; equilibrium, isotropic stress, pure shear, and trace-free diagonal extension; all coordinate planes and both signs; density-scaled Pi and separate fixed-Pi/nonunit-rho probes. The test checks quadratic residuals, closed/default-three/converged-200 iterations, bit-identical unchanged iterative output, and an additional explicit cap-active manufactured probe. Existing g++ and Python stdlib are sufficient; no installation is attempted. Compilation products are temporary. Previously reviewed host maxima were omega error 1.394302e-7 and quadratic residual 1.201202e-7; this is not CUDA fast-math evidence.
+
+## Authorized GPU diagnostic
+
+Use the existing working PyCUDA/NumPy environment and an explicitly selected device, through the same serialized GPU launch owner as the other campaigns:
+
+```sh
+CUDA_VISIBLE_DEVICES="$VERIFIED_IDLE_GPU" python tests/validation/les/run_gpu.py --run-gpu --device 0 --output-dir output/les-gpu
+```
+
+Device 0 is the logical device after CUDA visibility filtering. The output directory must not exist. The runner materializes and verifies the isolated include trees under `output/les-gpu/candidate`, then runs baseline and patched processes sequentially. Each process compiles one SourceModule at a time synchronously and synchronizes each probe. There are no background jobs, worker pools, or parallel compiler calls. No machine contact, device selection, reservation, or environment setup is performed by the runner.
+
+Each version has exactly **48 scalar GPU cases**, 96 total: D2Q9/D3Q19, Cs=.16/0, equilibrium/shear/isotropic/cap-active manufactured stress, and closed/default-three/converged-200 branches. This is the frozen GPU matrix, not the expanded 736-case host matrix. Outputs `baseline.json` and `patched.json` contain raw scalar omega values and device name. Completion means only the diagnostic executed; an independent comparison against the declared host oracle is still required. A failed run leaves partial evidence for inspection and must not be overwritten.
+
+## Interpretation and limits
+
+Full-stress versus deviatoric-stress selection, iterative viscosity cap, iteration count/relaxation, omega guards, collision dispatch, and AMR guards remain unchanged. Closed and iterative paths intentionally remain different for isotropic or cap-active stresses. No settings or Cs defaults change. High stress is a helper diagnostic, not an admissible flow-state claim.
+
+The corrected scalar formula is justified only under the stated stress-strain estimate. MRT full-trace stress does not generally share one shear relaxation rate, so its bulk coupling remains uncertified. Force-corrected non-equilibrium stress and Guo-force interaction also remain uncertified. These manufactured tests are unforced and do not cover full TRT/MRT collision, streaming modes, shifted storage, walls/Jets, flow solutions, or 3D turbulence. Existing guard statements are preserved, not newly comprehensively tested; the main matrix is guard-inactive. No AMR LES support is added or enabled.
+
+Production integration and removal of this temporary candidate fixture are separate decisions after GPU and applicable flow evidence. CPU success alone does not approve a production change.

docs/validation/refined_dfg.md

@@ -0,0 +1,70 @@
+# Matched SRT uniform / refined DFG Re100 preparation
+
+Status: CPU-preflight checked; **not GPU executed or approved to execute**. No physical acceptance or speedup claim. No production source changes. Root owns GPU admission and integration.
+
+## Frozen setup
+
+Canonical source snapshot: integrated commit `cd3a657f9d3833e812c3cca173525f165b48710a`. Production origin remains candidate `47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5` (on `5338864`, on `69e9b78`); production files are byte-identical. The integrated Re100 helper header/PIN records that origin explicitly. Both commit identities are retained in the manifest. `tests/validation/data/refined_dfg/source_manifest.json` hashes all production source files plus the existing DFG geometry/pressure helper, periodicity helper and official reference cycle. An integrated checkout may have a different commit but must have identical pinned file contents. No MRT external-reference run is reused as the matched uniform control.
+
+Both cases: SRT, shifted FP32, double buffer, no LES, no applied force, stationary cylinder, local Zou–He parabolic inlet, halfway y walls, identical pressure-NEQ outlet with **all nine nonequilibrium components damped by 0.5**, backflow clamp disabled. This outlet is not the exact canonical DFG do-nothing condition. Raw cylinder reaction observation is a diagnostic, not applied forcing.
+
+Physical units use D=0.1 m, mean U=1 m/s, ν=0.001 m²/s, density reference=1; domain 2.2×0.41 m, cylinder center (0.2,0.2) m. Uniform D80: 1761×330, dx=0.00125 m, dt=0.0000125 s, ν_lattice=.008. Refined root D40: 881×166, dx=.0025 m, dt=.000025 s, ν_lattice=.004; fine dx/dt/ν match uniform D80. All lattice mean speeds are .01.
+
+Fixed half-open patch **(40,24,320,140)** in root edge coordinates maps to x=[.09875,.79875] m, y=[.0575,.3475] m. The body and about six diameters of downstream wake are fine; wall layers and far wake remain coarse. Bounds are frozen before outcomes, not tuned afterward. Geometry-only leaf-count and nominal collision-work ratios to uniform are ~.419 and ~.321. These omit real launch/transfer/halo/measurement costs and are not measured speedups. The narrow root block height is legal but potentially expensive; do not infer speed from leaf count.
+
+Exact root dimensions are tiled by (by,bx)=(2,881); no padding. Fine allocated shape (240,568), including four-cell halos, uses (8,8). `host_geometry` calls production channel transfer-overlap checks and complete fine cut-link/donor/transfer guards before any CUDA import.
+
+### Half-cell conversion
+
+Uniform integer nodes have physical coordinates (i*dx, (j-.5)*dx). Root edge coordinates therefore map to ((X-.5)*dx_root, (Y-1)*dx_root). Body edge center=(80.5,81), radius=20; wall edges Y=1,165; source planes X=.5,880.5. Fine local body nodes are `(body_edge - fine_origin)/.5 - .5`.
+
+Body and physical wall planes match exactly. The axes have different phases. With fine origin O=(38,22):
+
+- x_f = [O_x+(j+.5)/2-.5] dx40 = (j+75.5) dx80, versus x_uniform = i dx80. Fine x samples have a half-fine-cell (quarter-root-cell) offset.
+- y_f = [O_y+(k+.5)/2-1] dx40 = (k+42.5) dx80, versus y_uniform = (l-.5) dx80. The y samples align exactly for l=k+43.
+
+This difference follows from the inlet's ghost-node plane versus the half-way y-wall planes. The host phase test checks both mappings. Compare interpolated observables, not pointwise arrays or bitwise equality; x-phase error, the different fine collision arithmetic, wall resolution, and coarse/fine transfer effects all contribute to discrepancies. No field-difference acceptance is claimed by this runner. Initial fluid state samples the same analytic parabolic equilibrium as uniform production initialization; root boundary ghosts start at rest. It is not asserted to be bitwise identical across different grids/arithmetic.
+
+## Observations and gates
+
+- Every AMR base step is individually followed by `read_raw_exchange(..., window='last_base_step')`: two fine substeps, root-lattice per-unit-depth impulse integrated over [n,n+1]. `step(n)` alone retains only its last base interval, so it cannot support averaging the whole batch.
+- Uniform `Simulation.run(batch)` / `read_force` gives the mean over that batch; multiply by batch before accumulating. Cd/Cl divide the interval mean reaction by 0.5 U_lattice² D_lattice in each case. Force unit is ρ dx³/dt² per depth, so fine-uniform lattice force is twice root force for the same physical load.
+- Force timestamps are observation-interval midpoints. Pressure uses physical decoded postcollision density at the interval end, extrapolated from fluid-only points one/two **fine** cells outside the cylinder using existing `pressure_probes`. Pressure difference is normalized by ρ_ref U_physical². Existing periodic helpers align pressure onto force cycles.
+- Public AMR downloads exclude nonleaves. Fixed geometry masks, not state-dependent finiteness masks, check all fluid populations and sum ρ dx_level²; covered root cells, body, boundaries and fine halos are excluded. Corrupt NaN fluid cannot disappear from the mass check. Mass is per-unit-depth ρ_ref·m². Uniform uses its physical fluid flags and dx². Different sampling of the circular boundary means initial discrete fluid volume can differ slightly; compare temporal stability, not identical absolute mass.
+- Flux samples root/uniform x-index 2 and -3, summed with dx and physical velocity conversion. Those planes differ by grid-scale offsets; cycle-average inlet/outlet balance is the gate, not direct pointwise flux equality.
+- Dense measurement width Δt*=.02 in the pilot and for t*≥240, otherwise .5. Pilot ends t*=5; full ends t*=300. Same existing five-cycle periodicity, mass, flux, and reference gates are imported without modification. Both official-reference diagnostics and refined-versus-uniform cycle comparisons are written. `matched_case_pass` refers only to matched SRT agreement, not certification of canonical DFG or the broader suite.
+- Safety stop: nonfinite state/reaction, Mach>.2, or density outside [.5,1.5]. GPU durations include required exchange reads and synchronization in stepping; full wall time also includes host measurements. Do not report pure-kernel throughput from these numbers.
+
+## Commands
+
+From the repository root, CPU only:
+
+```sh
+PYTHONPATH=src python tests/unit/test_refined_dfg.py
+python tests/validation/run_refined_dfg.py --mode plan --output /tmp/refined-dfg-plan-new
+```
+
+Fourteen CPU tests are retained in tests/unit/test_refined_dfg.py. Plan JSON is configuration output only; no preparation result files are committed. The initial attempt to use pytest was unavailable in this interpreter; tests use standard-library unittest and NumPy instead.
+
+Only after root channel GPU gates pass and root explicitly admits an idle physical GPU, root may run:
+
+```sh
+CUDA_VISIBLE_DEVICES=<admitted-device> timeout --signal=TERM --kill-after=30s <caller-hard-budget>s python tests/validation/run_refined_dfg.py --mode pilot --gpu-gates-approved --max-wall-seconds <caller-soft-budget> --output <new-pilot-directory>
+CUDA_VISIBLE_DEVICES=<admitted-device> timeout --signal=TERM --kill-after=30s <caller-hard-budget>s python tests/validation/run_refined_dfg.py --mode full --gpu-gates-approved --max-wall-seconds <caller-soft-budget> --pilot-results <completed-pilot-directory> --output <new-full-directory>
+```
+
+`--gpu-gates-approved` is a caller attestation, not proof that GPU tests ran. Full mode additionally requires both matched healthy pilots through t*=5 with identical source hashes, bounds and runner hash. No GPU/Mac/remote operation has been performed for this preparation. Pilot/full outputs must be new/empty. Uniform runs first; an unhealthy case stops the campaign. A full run that fails periodic/matched gates exits nonzero.
+
+## Bounded execution and receipt validation
+
+No pilot throughput is known yet; the caller must set a finite positive `--max-wall-seconds` rather than treating an invented estimate as calibration. After valid pilots, the full-run calibrated budget is 3 × sum_over_cases[startup_seconds + 60 × (pilot_wall_seconds − startup_seconds)]. The full run uses the smaller of that estimate and the caller's explicit ceiling. This uses total measured pilot evolution/observation cost and the full/pilot physical-time ratio (300/5); 3× is a declared safety margin, not a runtime guarantee. The plan records both budgets.
+
+The campaign deadline begins before the first case initialization. Cooperative checks surround initialization, every AMR root step or uniform ≤100-step batch, and measurements. On observed expiry, the runner stops, writes a nonpassing `wall_time_limit` receipt, and exits 3. A blocking CUDA call cannot be interrupted by these checks: the external `timeout` watchdog is required, with a hard budget chosen slightly above the soft budget to permit cleanup. Hard termination may leave partial CSV data and no final receipt; that never qualifies as a completed pilot.
+
+Each case emits `READY` with strict finite JSON and saves `ready.json` only after its first healthy complete observation, after initialization, actual stepping, force readback, macro/pressure/mass/flux paths and CSV flush all succeeded. No early initialization-only READY is emitted.
+
+Full-run pilot validation rejects nonfinite JSON constants and numeric overflow, wrong case/mode/status, nonboolean health, missing/duplicate/out-of-order rows, wrong step counts or observation times, missing numeric observations, health-limit violations, absent/mismatched READY receipts and invalid timing calibration. It requires exactly 250 observations spanning the full dense t*=5 pilot, source/runner hashes and frozen bounds. Tests include the reviewer's exact NaN-endpoint regression and malformed receipt variants. Synthetic test fixtures exist only in temporary directories and are not simulation results.
+
+## Integrated source import guard
+
+The runner's source defaults to this repository root. Before geometry or GPU work, it verifies the pinned source/helper/reference contents and requires the imported CelerisLab package to be exactly `<source>/src/CelerisLab/__init__.py`. An installed DynamisLab copy or another cached checkout is rejected. Integration changes paths/readability and this explicit import assertion; geometry, numerical definitions, gates, and receipt validation remain frozen. Any new path-formatted runner has its own hash, so old preparation pilot receipts cannot authorize a full run of this integrated runner.

docs/validation/refined_dfg_mrt.md

@@ -0,0 +1,71 @@
+# Matched MRT DFG pilot preparation
+
+Preparation only: this test-only extension reuses the existing runner and existing MRT
+operator. GPU admission awaits reviewed MRT capability results; no pilot/full execution
+is implied by CPU preflight. No production, historical SRT manifest, reference data,
+periodic gate, geometry, initialization or observation algorithm is changed here.
+
+## Selection and provenance
+
+Use `--collision MRT` for **both** uniform D80 and refined root D40/fine D80. There are
+no independently selectable collision models per case. Default remains SRT.
+
+MRT verifies `data/refined_dfg/source_manifest_mrt.json`, pinned to reviewed numerical
+source snapshot `3ea131e92c057843034f11a1c574301a2430dfce`, including the shared geometry,
+pressure/cycle helpers and official reference. Test-only commits may change HEAD without
+changing this numerical source pin. Plan, READY, case summaries and comparison separately
+record `checkout_revision` and `checkout_dirty` using existing source-provenance helper;
+the runner hash remains separate. Collision is recorded in every receipt, and full-mode
+validation rejects missing or different collision labels in plan, either summary or either
+READY receipt, in addition to existing source/runner/health/time checks.
+
+The historical SRT manifest remains byte-for-byte unchanged. SRT on the MRT candidate
+checkout deliberately fails that source pin; use a matching historical checkout via
+`--source` for SRT. No fallback or automatic repinning is permitted. Historical SRT
+failure/diagnostic evidence remains attached to its original source and runner hashes;
+it cannot authorize an MRT full run. The existing failed-observation step-reporting
+undercount is not changed in this narrowly scoped selection patch.
+
+## Frozen physical comparison
+
+All settings remain those documented in `refined_dfg.md`: U_lattice=.01, physical mean
+U=1 m/s, D=.1 m, shear ν=.001 m²/s, fixed physical domain/body/wall planes, patch
+(40,24,320,140), shifted FP32 double buffers, no LES/forcing/Jets, finest dx=.00125 m,
+finest dt=.0000125 s, unchanged half-cell frame conversion and axis-specific phase.
+Force interval accumulation, physical leaf mass, pressure interpolation, flux sampling,
+safety thresholds, t*=5 pilot and t*=300 cap are unchanged. Existing accuracy/periodicity
+criteria are imported unchanged. Both MRT cases are fresh runs; old MRT or SRT results
+are not substituted as controls.
+
+Both cases use the **existing MRT pressure-NEQ branch**: only unknown directions q2/6/8
+are extrapolated. This differs from SRT's all-nine .5-damped branch; the retained
+`srt_neq_damp=.5` configuration value is inactive under MRT. Backflow clamp remains off.
+The outlet is still noncanonical relative to exact DFG do-nothing conditions.
+
+The reused MRT operator keeps fixed kinetic rates `s_e=s_eps=s_q=1.2` on all levels;
+only the shear rate follows existing per-level viscosity. Physical bulk relaxation
+therefore differs between coarse and fine levels. Transfer operators and subcycling
+are unchanged. This is the candidate's declared behavior, not a new fully mode-scaled
+MRT transfer scheme. Do not tune rates, the patch or gates after seeing outcomes.
+Capability smoke/parity gates do not establish DFG accuracy or speedup.
+
+## CPU and admission commands
+
+From this checkout:
+
+```sh
+python tests/unit/test_refined_dfg.py
+python tests/validation/run_refined_dfg.py --collision MRT --mode plan --output /tmp/new-mrt-dfg-plan
+```
+
+After root verifies the candidate MRT root/parity/refinement capability results and admits
+an idle device, the paired pilot may use the existing bounds and external watchdog:
+
+```sh
+CUDA_VISIBLE_DEVICES=<admitted-device> timeout --signal=TERM --kill-after=30s <hard-budget>s python tests/validation/run_refined_dfg.py --collision MRT --mode pilot --gpu-gates-approved --max-wall-seconds <soft-budget> --output <new-mrt-pilot-directory>
+```
+
+The flag is caller attestation, not evidence. An unhealthy first case prevents the second.
+Existing strict pilot receipt validation and calibrated wall-time controls remain; full
+mode is not being admitted as part of this preparation. Root must review both healthy
+pilot outcomes and separately decide any full-run admission.

docs/validation/refined_dfg_timing.md

@@ -0,0 +1,58 @@
+# Refined DFG MRT timing harness (measurement only)
+
+`tests/validation/time_refined_dfg.py` times the matched MRT refined/uniform DFG Re100 setup of
+`run_refined_dfg.py` (same configuration, geometry and initial state functions). No numerical,
+physics-gate or production change; no force or accuracy result is produced.
+
+| Variant | Loop | Force reads |
+|---|---|---|
+| A | refined `step(1)` + `read_raw_exchange` per base step | runner-like force-read workload (values discarded) |
+| B | refined `step(1)` only, no host reads (`observe_raw_exchange` stays on) | timing-only |
+| C | refined single `step(N)` | timing-only; ledger keeps the final base interval only |
+| E | uniform `sim.run(<=100)` + `read_force` | runner-like force-read workload (values discarded) |
+| E' | uniform single `sim.run(2N)` | timing-only; at most a whole-window average, different summation order |
+
+- No variant retains a valid force history; A/E omit the runner's interval aggregation and
+  validation, and no variant includes macro/probe/CSV costs. Results are not full benchmark
+  throughput.
+- Status lines: `CONSTRUCTED` (construction/JIT complete) and `WARM_READY` (first warm-up and
+  health check passed) per case.
+- Equal simulated duration: N refined base steps vs 2N uniform D80 steps.
+- Every repeat starts from the same fresh state (refined `initialize`, uniform `snapshot/restore`)
+  and the same warm-up; construction/JIT, warm-up, downloads and parity checks are untimed.
+- The actual stepping stream is synchronized explicitly before and after each timed loop; the
+  wall timer (`time.perf_counter`) is the only source. No CUDA-event span is called kernel time.
+- `--max-wall-seconds` is required for GPU runs (outer timeout, checked between variants and
+  every 100 steps in per-step loops).
+- Final populations and NaN masks must be bitwise identical to the first A repeat (A/B/C) or the
+  first E repeat (E/E'); refined and uniform grids are never compared. Leaf/fluid populations
+  must be finite. Physical health (finite leaf/fluid macros, rho in [0.5, 1.5], max Mach <= 0.2,
+  the existing runner limits) is checked after warm-up and after each timed window. Matching NaN
+  is accepted only as a storage mask; any infinity or level-count mismatch fails parity. All
+  parity/health checks are untimed.
+- Cases run sequentially (uniform, then refined) in one process with startup serialized;
+  variants are interleaved within each repeat. A-B and B-C are measured comparative
+  differences under GPU variability, not exact causal decompositions.
+
+Usage (root selects one verified idle GPU):
+
+    PYTHONPATH=src python -m unittest tests/unit/test_time_refined_dfg.py
+    python tests/validation/time_refined_dfg.py --output <new> --plan
+    CUDA_VISIBLE_DEVICES=<gpu> python tests/validation/time_refined_dfg.py --output <new> --max-wall-seconds <s>
+
+Defaults: 2000 base steps, 200 warm-up base steps, 3 repeats. Profiler capture and the
+`observe_raw_exchange=False` variant are deferred until these measurements exist.
+
+## A versus observation-window ledger (`tests/validation/time_window_dfg.py`)
+
+Measurement-only comparison on the same shifted-MRT refined setup: A (per-base `step(1)` +
+`read_raw_exchange`, summed per 80-base-step interval), W (`start_exchange_window`, `step(80)`,
+one window read) and W1 (window with 80 x `step(1)`, isolating the read pattern from call
+batching). Same fresh state, warm-up and simulated duration; populations must be bitwise equal
+to the first A repeat and per-interval total impulses within the documented engineering
+tolerance, all outside timing. The candidate source is recorded (checkout revision, dirty flag,
+file hashes); frozen refined-DFG manifests/results are not used or changed.
+
+    PYTHONPATH=src python -m unittest tests/unit/test_time_window_dfg.py
+    python tests/validation/time_window_dfg.py --output <new> --plan
+    CUDA_VISIBLE_DEVICES=<gpu> python tests/validation/time_window_dfg.py --output <new> --max-wall-seconds <s>

docs/validation/taylor_green.md

@@ -0,0 +1,161 @@
+# Periodic 2D Taylor–Green AMR campaign (prepared, not GPU-qualified)
+
+This validation harness complements DFG Re100 and Jet/body tests. It measures
+smooth 2D periodic decay and errors crossing fixed refinement interfaces at
+Re_L=100 and 1000. It does **not** certify 3D turbulence, LES, Jet performance,
+wall conditions, dynamic refinement or general high-Re flows.
+
+## Reference and conventions
+
+Primary implementation-author documentation, checked 2026-10-18:
+- TensorMesh, exact velocity and pressure with a consistent phase:
+  https://docs.tensor-mesh.com/example_gallery/fluid/taylor_green.html
+- AeroSim Nassu, 2D decay, energy and lattice-length Re definition:
+  https://docs.aerosim.io/nassu/validation/analytical/04_taylor_green_vortex_2d/README.html
+
+We use the TensorMesh phase, rescaled to length L=N root cells:
+
+    k = 2*pi/L; A(t) = U*exp(-2*nu*k*k*t)
+    ux = -A*cos(k*x)*sin(k*y)
+    uy =  A*sin(k*x)*cos(k*y)
+    p/rho0 = -A*A/4 * (cos(2*k*x)+cos(2*k*y))
+    E = <|u|^2>/2 = U*U/4 * exp(-4*nu*k*k*t)
+
+Pressure is kinematic, with rho0=1. A global velocity sign reversal leaves
+pressure unchanged; exchanging sine/cosine without the corresponding pressure
+phase can reverse its sign. CPU tests independently check momentum residuals.
+The nonlinear acceleration is a pressure gradient, **not zero**; the AeroSim
+page's prose calling advection zero is not adopted. Do not mix phase variants
+from different references. The formulas above are verified by substitution,
+not by trusting a copied pressure expression.
+
+Re_L=U*L/nu, not Re_k=U/(nu*k)=Re_L/(2*pi). U=.03 in lattice units,
+Ma=U/cs=.0519615, cs^2=1/3, nu_root=.03*N/Re_L. Acoustic refinement uses
+h_l=dt_l=2^-l: lattice velocity stays U, nu_l=2^l*nu_root and
+ tau_l=.5+3*nu_l. The worst root tau is .50288 (N32/Re1000), omega=1.98855,
+inside the existing config guard. No viscosity clamping or guard override.
+
+All initial arrays are physical FP64 equilibrium populations, supplied to the
+public shifted-FP32 postcollision-G initializer. rho=1+3p at each true cell
+center. These are valid equilibrium postcollision states, but do not include
+Chapman–Enskog nonequilibrium stress; startup acoustic/kinetic error remains in
+the measurements and is shared with uniform baselines. The incompressible
+reference is not an exact finite-Mach LBM solution. This is why resolution alone
+need not remove the full error floor.
+
+## Geometry and workload
+
+All patches use level-wide union form, so origins are zero, including fine
+arrays with holes. Coordinates are (index+.5)*2^-level. Metrics integrate only
+ACTIVE leaves, each weighted by 4^-level. Covered cells and halo cells never
+enter norms or mass. Geometry tests prove exact single coverage at root-cell
+resolution. All blocks are 8x8, and both root dimensions are periodic.
+
+At each root N=32,64:
+- 2 levels: L1 rectangle (N/8,N/8,7N/8,3N/4), in L0 indices.
+- 3 levels: same L1 plus L2 rectangle (3N/4,N/2,3N/2,5N/4), in L1 indices.
+
+These off-axis boundaries cross nonzero flow and gradients rather than sitting
+only on symmetry or stagnation lines. Whole geometry scales with N.
+
+16 cases: at each Re, uniform N=32,64,128,256 plus AMR (root32/root64)x(depth2/depth3).
+Each AMR case therefore has both a same-root uniform baseline and a uniform
+baseline at its finest spacing. Uniform N is its own lattice length, so nu and
+step count are scaled to represent the same nondimensional solution.
+Samples t*=tU/L = 0,.12,.24,.36,.48. End time is 16*N base steps, a meaningful
+.48 domain turnover or 3.016 vortex-scale turnover times (scale 1/k).
+At Re1000 energy decays about 7.3%; at Re100 about 53.1%. No plots or full fields;
+each case writes five diagnostic rows and provenance, incrementally.
+
+## Running (GPU owner must schedule launches)
+
+CPU checks, from the repository root:
+
+    PYTHONPATH=src python -m unittest discover -s tests/validation/taylor_green -p 'test_*.py' -v
+
+One GPU case, on the scheduled worker after checking occupancy:
+
+    CUDA_VISIBLE_DEVICES=<physical-id> python tests/validation/taylor_green/run_case.py --repo /path/to/reviewed/checkout --re 100 --n 32 --depth 1 --output output/taylor_green/re100_n32_d1.json
+
+Depth 1 is uniform. Run uniform cases before AMR for each Re; inspect baseline
+screening failures before proceeding. For a complete matrix, repeat the uniform
+call with N=32,64,128,256 and AMR calls with N=32,64 / depth=2,3, at each Re.
+Use serial process startup on the parent's assigned Horsea GPUs. This harness
+does not schedule or acquire GPUs. No GPU was used during preparation.
+
+    python tests/validation/taylor_green/assess.py output/taylor_green/re*_n*_d*.json --output output/taylor_green/assessment.json
+
+The local API/geometry review used clean staging commit
+cd3a657f9d3833e812c3cca173525f165b48710a. This is **not** a claim of results on d79
+or any historical commit. Each runtime output records actual Git HEAD/dirty
+state, an aggregate source-tree hash, plus SHA256 of all harness Python files. Assessment refuses dirty checkouts and mixed source
+or harness identities. Keep reports outside the checkout or .gitignore them so
+new outputs do not change its dirty flag between cases. The source-tree hash excludes bytecode caches; production qualification requires
+a clean checkout.
+The runner uses the existing PeriodicBlockStepper / FixedPatchStepper periodic
+paths only. It does not depend on the new channel-root implementation.
+
+## Preregistered assessment, before any GPU results
+
+All sampled times must satisfy finite screening ceilings:
+- Relative composite velocity L2 <= 5%
+- Relative globally gauge-adjusted pressure L2 <= 15%
+- Relative kinetic-energy error <= 5%
+- Composite mass drift from downloaded t0 <= 1e-5
+
+These are screening budgets, not literature-certified tolerances. Acceptance
+also requires all 16 cases, common source/harness, correct nondimensional times
+and viscosity, and baseline-relative checks on worst post-startup-sample errors:
+- AMR error <= max(2 * same-root uniform error, declared numerical floor).
+- Root64 AMR error <= max(.9 * root32 AMR error, floor), for each depth/Re.
+- Uniform256 error <= max(.9 * uniform32 error, floor).
+
+Floors: velocity .002, pressure .02, energy .004. They are declared low-Mach /
+equilibrium-startup budgets, not fitted to observed results. Ma^2=.0027 motivates
+their scale; it is an asymptotic expectation, not a proof of those coefficients.
+The coarse baseline sets the AMR budget because much of the domain remains
+coarse. The finest uniform error and AMR/fine ratio are always reported, with no
+claim that partial refinement should match a fully fine domain. A pass below a
+floor is bounded accuracy, not a measured convergence order. Report errors and
+ratios even for a pass. Failure requires diagnosis or a separately preregistered
+follow-up, never retroactive tolerance tuning. Higher Re is deliberately absent;
+Re100/1000 must first be stable, resolved and reviewed.
+
+Diagnostics use one global pressure gauge, direct cell-center analytic values,
+kinetic energy per unit reference mass, and actual leaf cell areas. Composite
+midpoint quadrature introduces a small initial energy integration error, retained
+rather than subtracted. Mass drift is relative to downloaded t0, so density
+initialization and mixed-grid quadrature do not masquerade as time drift.
+
+## Pilot scheduling and transfer-consistency scope
+
+Start with one uniform case and one matched AMR case; the 16-case matrix is a
+qualification plan, not authorization to launch every case automatically. The
+GPU owner reviews the pilot before expanding it. Each process prints a flushed
+READY only after initialization, a real base step exercising every level's
+kernel/halo/transfer path, and a finite positive-density download check. That
+step belongs to the measured trajectory; it is not discarded warmup. PROGRESS
+messages contain each of the five diagnostic samples. READY alone is not an
+accuracy pass. Reports include per-level active cells, hops and collision updates.
+
+The complete matrix costs 768,540,672 active collision updates and 21,504 root
+intervals, excluding ghost streaming, halos, transfers and diagnostic work.
+Uniform N256 is the largest case (268,435,456 updates / 4096 intervals per Re).
+AMR N64/depth3 needs 50,987,008 updates / 1024 intervals. These are work counts,
+not GPU wall-time estimates. Baseline comparisons remain available on failed
+screens and partial pilots; absent/nonfinite operands are null, never passing.
+
+Acoustic viscosity scaling and unchanged lattice velocity do not by themselves
+prove nonequilibrium-stress consistency across refinement interfaces. A smooth
+collocated Chapman–Enskog postcollision nonequilibrium has a (tau-1)*h factor;
+the existing implementation instead uses volumetric piecewise-constant explosion
+and stream-only ghosts. Whether that actual transfer construction achieves the
+needed stress accuracy is an unresolved question for this pilot to measure.
+It is not a proven defect, and a collocated-state scaling factor must not be
+inserted into this different transfer operator without a separate derivation
+and validation. This campaign measures the current operator and makes no
+theoretical stress-consistency certification or production-kernel change.
+
+## Repository integration
+
+The four Python harness files remain together under tests/validation/taylor_green, so runtime provenance hashes precisely this harness directory. Formatting changes produce a new harness identity; preparation and integrated receipts must not be mixed. All analytic definitions, numerical floors, assessment gates, sample times, and source-import assertions are unchanged. Production source is unchanged from the cd3 snapshot; these periodic tests do not waive the separate channel GPU gate.

docs/validation/uniform_les_flow.md

@@ -0,0 +1,90 @@
+# Bounded uniform-grid LES flow validation: prepared, GPU NOT RUN
+
+## Status and scope
+
+Prepared for independent review, then an explicitly authorized Horsea run only. No production files were edited, no normalization patch was applied to a full solver checkout, and no GPU was used. This is a small end-to-end **wall-bounded, x-periodic** shear-decay test, not a fully periodic 2D test. It exercises the actual production Simulation → Stepper → OneStep → Smagorinsky → SRT dispatch and double-buffer pull streaming, plus existing halfway bounce-back walls.
+
+The original fully periodic-y plan is unsupported at the pinned revision: both `one_step_double.cu:90` and `one_step_esopull.cu:133` replace wall-adjacent incoming populations for every fluid node on y=1 and NY-2 regardless of its neighboring flags. `tests/validation/run_poiseuille.py:1–10` already documents this. Merely changing all flags to FLUID does not remove this behavior. The present test deliberately matches those real walls; it does not modify the kernels or claim a periodic-y public API.
+
+Only unforced SRT, D2Q9, FP32 shifted storage and double-buffer streaming are included. No MRT/TRT, force correction, EsoPull, AMR, Jet, 3D turbulence or physical LES coefficient calibration is certified. In continuum this unidirectional incompressible shear is trace-free; small finite-Mach/discretization stresses in the LBM are not asserted to vanish exactly. The underlying scalar-helper audit remains the cleaner normalization isolation.
+
+## Exact source contract
+
+Baseline is explicitly repinned to the production snapshot in `cd3a657f9d3833e812c3cca173525f165b48710a`, whose reviewed AMR-channel source origin is `47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5`. This new experiment does not inherit historical d79 GPU/flow acceptance. The original preparation remains associated with d79; its results are not rewritten.
+
+`tests/validation/data/les_flow/source_pins.json` records all 101 tracked production source files from that snapshot. The runner checks every hash before importing CelerisLab or touching a GPU. It accepts only:
+
+- `--variant baseline`: all baseline hashes, including helper SHA256 `ff7f000297ee0817e255ad9fcb6a13b647e3ddd5b206518d35f38d2895a9f501`.
+- `--variant corrected`: the same source bytes except helper SHA256 `a3db46c2e7d27530b2d1436e65cdbed0b72cc7e047bfbeef2a640877271b51aa`, from the independently reviewed `tests/validation/les/data/smagorinsky_normalization.patch` artifact.
+
+The runner records the actual checkout commit and exact imported package. It never applies a patch itself. Run it against separate isolated full checkouts. The patch candidate's minimal header tree is not a runnable solver checkout. A source mismatch is a blocker, not permission to regenerate pins from an arbitrary new checkout.
+
+Relevant pinned implementation evidence:
+
+- `tests/integration/test_numerical_gpu.py:24–85`: equilibrium construction, installation of spatial postcollision state, count=1 adapter, raw population observation and x-periodic/walled flags.
+- `tests/validation/run_poiseuille.py:29–78`: production Simulation use with periodic x, halfway y walls and downloaded populations.
+- `src/CelerisLab/lbm/kernels/core/layout.cuh:87–104`: actual modulo neighbors.
+- `src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh:18–38`: incoming populations copied from opposite local input directions.
+- `src/CelerisLab/lbm/kernels/operators/helpers.cuh:53–61`: LES effective omega feeds actual SRT collision.
+- `src/CelerisLab/cuda/compiler_v2.py:56–95`: compiler snapshots into private temporary build directories; package sources are not rewritten.
+
+## Geometry, units and initial state
+
+8×34 nodes: 8 periodic-x columns, 32 fluid rows, 2 ghost-wall rows. Relative to the bottom physical wall, fluid centers are y=0.5,...,31.5; physical halfway walls are y=0 and 32. Every x edge fluid node is explicitly FLUID; all top/bottom ghost nodes are SOLID|BC_WALL. Thus unused inlet/outlet configuration has no inlet/outlet flags to act on. There are no bodies, Jets or external forces.
+
+Lattice units: dx_LBM=dt_LBM=1, rho=1, nu=0.02, tau0=0.56, U=0.04, H=32. Initial u_x(y)=U sin(pi y/H), u_y=0. The maximum Mach number is 0.0693. Height-based Re=U H/nu=64; a wavelength-based Re would differ and is not used. No dimensional physical length/time/viscosity is asserted without a user-supplied unit mapping.
+
+Initialize equilibrium populations, encode shifted FP32, upload both DB buffers and set completed-count=1 following the existing integration-test adapter. That count is storage metadata: physical elapsed time starts at 0 after installation. Observe density and momentum directly from current postcollision populations, adding the descriptor weights in float64; do not mix in cached macro fields. This uses a clearly labeled validation-only raw state/flag adapter, not a promised public `set_periodic` or `set_velocity_profile` API.
+
+Cases: LES enabled/closed-form with Cs=0 (molecular control), Cs=.16 (ordinary-coefficient sanity), then Cs=1 (amplified normalization diagnostic). Cs=1 is intentionally not a physically recommended SGS coefficient. Initial peak modeled nu_t/nu = Cs² Δ² U(pi/H)/nu = 0.19635; estimated effective tau remains near 0.56–0.5718, far from omega guards. Standard default Cs=.16 yields a much weaker signal at these safe parameters and is not used to make a decisive short-run normalization claim. No defaults are changed.
+
+Observe at elapsed steps 519, 1038, 2075, approximately 0.1, 0.2, 0.4 molecular decay times. Three cases total 6225 measured advancing launches plus 3 warm-up launches per checkout (6228 total), three snapshots each. Population buffers contain only 272×9 FP32 entries each (about 9.6 KiB), apart from ordinary driver/compiler/context overhead. No frames, trajectories or heavy output. Full result JSON is a few tens of KiB. Runtime depends on compilation/device; no performance claim is made.
+
+## Independent continuum reference
+
+For u=(u(y),0), |S|=sqrt(2 S:S)=|du/dy|. Therefore
+
+    du/dt = d/dy [(nu + (Cs Delta)^2 |du/dy|) du/dy],
+    u(0,t)=u(H,t)=0.
+
+`tests/validation/les_flow_oracle.py` integrates this equation directly with centered finite-volume face fluxes and SSPRK2, at 128 and 256 cells on the same physical H=32 domain. Stability step size uses the nonlinear flux derivative nu+2(Cs Delta)^2|gradient|. Boundary face gradients use the zero Dirichlet value across a half cell. It does not import or reproduce the LBM quadratic/helper. A separate fixed-n=256 test halves the CFL factor from .4 to .2 for Cs=0,.16,1 at every checkpoint, requiring relative changes below 1e-8; measured maximum is 1.8141e-10. This isolates temporal error from the spatial-refinement test.
+
+Crucially the SGS filter Delta=1 stays fixed as the CPU oracle mesh changes from dy=.25 to .125. Confusing oracle spacing with filter width would change the PDE and invalidate convergence. Future LBM grid refinement also needs an explicit physical-unit/filter-width scaling contract; this artifact does not silently reinterpret Cs or Delta when resizing H.
+
+Cs=0 is checked directly against U sin(pi y/H) exp[-nu(pi/H)^2 t]. The nonlinear reference includes evolving higher harmonics, rather than assuming it remains a sine. At the final time, host oracle predicts 5.1706% SGS reduction in the fundamental amplitude and 1.5409% full-profile separation between standard and old coefficient conventions. The old convention reference is computed by changing continuum Cs to Cs/2^(1/4), purely as a diagnostic; it does not replace standard acceptance.
+
+## Predeclared criteria (do not tune after seeing GPU results)
+
+- Initial decoded profile relative L2 <= 1e-5; otherwise no advancing test.
+- Molecular Cs=0 profile error against analytic decay <= 0.003 at every checkpoint.
+- Cs=1 full-profile error against standard nonlinear oracle <= 0.005 at every checkpoint.
+- Oracle 128/256-grid interpolated relative difference <= 8e-5 at every checkpoint.
+- Relative error in SGS fundamental-amplitude decrement <= 0.15 at the last two checkpoints. Decrement is measured from actual Cs=0 amplitude minus actual Cs=1 amplitude; reference uses analytic Cs=0 minus the nonlinear oracle. The first checkpoint remains a transient diagnostic for this increment gate.
+- Fluid mean-density drift <= 5e-5, max |u_y| <= 1e-5 and x-column velocity spread <= 1e-6 at every checkpoint; finite populations, strictly positive finite fluid density and finite decoded velocity/profile/diagnostics required throughout observations, including startup and warm-up.
+
+Reasons: molecular 0.3% leaves room for finite-lattice dispersion, equilibrium-start kinetic transient and FP32 error at H=32; nonlinear 0.5% is substantially below the predicted 1.54% old-vs-standard separation. The 15% increment tolerance is below the old convention's approximately 29% viscosity-coefficient deficit, while allowing continuum/lattice differences. Directly subtracting the two measured amplitudes reduces common lattice error. Oracle convergence is over 60 times tighter than the nonlinear profile bound. These are declared engineering test gates, not GPU-validated uncertainty estimates; a failure requires diagnosis/refinement and an explicitly revised experiment, never a post-hoc tolerance edit to obtain a pass.
+
+Both baseline and candidate are judged against the same standard convention. Baseline is expected to fail SGS criteria if the flow reproduces the scalar finding; an exit 1 with valid evidence is useful. The runner reports old-convention profile error but never masks baseline failure. Candidate passing alone is insufficient if baseline does not demonstrate meaningful discrimination. Review both results and source pins together.
+
+## Host verification and later GPU commands
+
+From the repository root, using existing Python/NumPy and the existing pytest environment:
+
+    PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=src python -m pytest -q tests/unit/test_les_flow_oracle.py tests/unit/test_les_flow_runner.py
+
+Nine tests cover analytic decay, independent-oracle mesh convergence, symmetry/positivity/energy decay, filter invariance, fixed-grid timestep halving, signal discrimination, invalid density/velocity rejection, source pins/candidate/tamper rejection and the no-GPU-without-opt-in CLI guard. The integrated tests preserve those checks. No PyCUDA import or GPU run is required for them. Cs=.16 uses the same 0.5% full-profile safety/sanity bound; its SGS increment is not a discriminatory normalization gate.
+
+The source-guard tests materialize temporary candidate headers with the existing reviewed LES preparation tool; the GPU entrypoint uses `tests/validation/run_uniform_les_flow.py`, `tests/validation/les_flow_oracle.py`, and `tests/validation/data/les_flow/source_pins.json` in that repository layout.
+
+After review, make isolated full baseline and corrected checkouts of the exact base. Apply only `tests/validation/les/data/smagorinsky_normalization.patch` to the corrected checkout; do not edit the benchmark/production tree. On authorized Horsea, with existing NumPy, PyCUDA, nvcc and driver, run sequentially and select the authorized physical GPU explicitly:
+
+    CUDA_VISIBLE_DEVICES=<physical-GPU> PYTHONDONTWRITEBYTECODE=1 python tests/validation/run_uniform_les_flow.py --run-gpu --device 0 --root <isolated-baseline> --variant baseline --output <new-evidence-dir>/baseline.json
+    CUDA_VISIBLE_DEVICES=<physical-GPU> PYTHONDONTWRITEBYTECODE=1 python tests/validation/run_uniform_les_flow.py --run-gpu --device 0 --root <isolated-corrected> --variant corrected --output <new-evidence-dir>/corrected.json
+
+`--device 0` is the logical ordinal within CUDA_VISIBLE_DEVICES. Do not use a shell `&&` between these: baseline's expected standard-criterion failure must not suppress the corrected run. Existing output files are refused. A JSON `READY` line is printed separately for each case only after source verification, compilation, initialization, raw-state installation, one actual advancing OneStep launch, post-warm-up macro validation, complete restoration of both DB buffers/flags/count to the initial state, profile readback and CUDA synchronization. It gives source identity, device, Cs, warm-up/reset confirmation, installed-profile error and maximum measured steps. This is the first trustworthy startup checkpoint, not a statement that the physics passed. `DONE` is printed after all three cases and all criteria; failures/compile errors before that remain incomplete runs.
+
+Deliver both JSON files and stdout/stderr with artifact hashes. Review before any production integration. AMR LES guards remain untouched.
+
+## Isolated candidate preparation
+
+Use the current validation commit for two separate full checkouts, preserving the baseline. In the corrected checkout only, apply `tests/validation/les/data/smagorinsky_normalization.patch` with `git apply`; do not commit that experimental production change or update the pins. The runner requires exactly the baseline bytes plus the reviewed helper hash. The header-only scalar diagnostic tree is insufficient for this end-to-end flow. Baseline and corrected commands must run as separate Python processes; package import caching must never combine their implementations. Both reject an imported CelerisLab package outside the explicitly supplied checkout, including an installed DynamisLab copy. Retain the actual checkout commit, variant, exact header hash, and experimental source label with each result.

src/CelerisLab/body/jets.py

@@ -204,13 +204,14 @@ def edge_to_node_frame(points):
 
 
 def surface_link_velocity(objects, cfg, flags, fluid_idx, dir, body_id, rx, ry,
-                          jets=(), hits=None):
+                          jets=(), hits=None, *, _shifted_blocked=False):
     """Per-link prescribed surface velocity (n, 2) FP32 and aperture mask.
 
     Host-only and storage-neutral: rows follow the given physical cut-link
     order. Returns (None, None) when neither an object carries surface
     actuation nor a jet is given. Jets need true FP64 ``hits`` (n, 2) and may
     not share a link with each other or with a legacy aperture.
+    Shifted storage is privately enabled only by the all-blocked stepper.
     """
     enabled = [obj for obj in objects if obj.surface_actuation is not None]
     if not enabled and not jets:
@@ -234,7 +235,9 @@ def surface_link_velocity(objects, cfg, flags, fluid_idx, dir, body_id, rx, ry,
         if hits is None or len(hits) != len(dir):
             raise ValueError("Jets require true hit points row-aligned with the cut links")
         if (cfg is None or not cfg.is_d2q9 or cfg.store_precision != "FP32"
-                or cfg.streaming != "double_buffer" or cfg.ddf_shifting):
+                or cfg.streaming != "double_buffer"
+                or (cfg.ddf_shifting and not (_shifted_blocked is True
+                    and cfg.collision == "SRT" and not cfg.les_enabled))):
             raise ValueError("Jets require D2Q9/FP32/unshifted/double_buffer")
         shape, jet_id = jet_link_shapes(jets, hits)
         driven = jet_id >= 0

src/CelerisLab/lbm/block_stepper.py

@@ -21,6 +21,7 @@ import pycuda.driver as cuda
 from ..cuda import compiler_v2 as compiler
 from .blocks import NQ, BlockLayout2D, blocked_link_tables
 from .descriptors import D2Q9_EX, D2Q9_EY, FLUID, SOLID
+from .population_storage import _StoredPopulations, _encode_physical, _decode_stored
 
 MAX_GRID_DIM = 65535  # CUDA gridDim.y/z limit (rows, blocks)
 
@@ -85,10 +86,10 @@ class PeriodicBlockStepper:
         checked.validate()
         if (not checked.is_d2q9 or checked.collision != "SRT"
                 or checked.streaming != "double_buffer" or checked.store_precision != "FP32"
-                or checked.ddf_shifting or checked.les_enabled
+                or checked.les_enabled
                 or (checked.ny, checked.nx) != layout.shape):
             raise ValueError("PeriodicBlockStepper requires matching D2Q9/SRT/FP32/"
-                             "unshifted/double_buffer config without LES")
+                             "double_buffer config without LES")
         ly, lx = layout.local_shape
         pairs = len(layout.halo_dst)
         tpb = checked.threads_per_block
@@ -99,6 +100,7 @@ class PeriodicBlockStepper:
         links = 0 if self.boundary is None else len(self.boundary["cell"])
         if links >= 2**31:
             raise ValueError("cut links exceed launch limits")
+        self._ddf_shifting = bool(checked.ddf_shifting)
         self.layout = layout
         self.ctx = context
         self._links = links
@@ -177,14 +179,20 @@ class PeriodicBlockStepper:
         """Upload dense G_0; the spare buffer starts NaN so stale halos cannot hide."""
         with self._lock, self.ctx.activate():
             self._require("new", "ready")
-            dense = np.array(g, dtype=np.float32, order="C", copy=True)
+            dense = np.array(g, dtype=np.float64 if self._ddf_shifting else np.float32,
+                             order="C", copy=True)
             if dense.shape != (NQ, *self.layout.shape):
                 raise ValueError("Expected a Q-major (9, ny, nx) array")
             fluid = np.ones(self.layout.shape, bool) if self.boundary is None else ~self.boundary["solid"]
             if not np.isfinite(dense[:, fluid]).all() or (dense[:, fluid].sum(axis=0) <= 0).any():
                 raise ValueError("Fluid populations must be finite with positive density")
             dense[:, ~fluid] = np.nan  # Solid cores are deliberately invalid.
-            blocked = self.layout.to_blocked(dense)
+            encoded = _encode_physical(dense, self._ddf_shifting)
+            values = encoded[:, fluid]
+            if (not np.isfinite(values).all()
+                    or (values.sum(axis=0, dtype=np.float64) + int(self._ddf_shifting) <= 0).any()):
+                raise ValueError("Stored fluid populations must be finite with positive density")
+            blocked = self.layout.to_blocked(encoded)
             spare = np.full_like(blocked, np.nan)
             try:
                 cuda.memcpy_htod(self._current, blocked)
@@ -225,7 +233,12 @@ class PeriodicBlockStepper:
                 raise
 
     def download(self):
-        """Return dense (9, ny, nx) FP32 core populations; no host cache."""
+        """Return physical core populations (FP64 when shifted); no host cache."""
+        snapshot = self._download_raw()
+        return _decode_stored(snapshot.values[0], snapshot.shifted)
+
+    def _download_raw(self):
+        """Return one tagged dense FP32 population array without conversion."""
         with self._lock, self.ctx.activate():
             self._require("ready")
             blocked = np.empty(self.layout.blocked_shape, np.float32)
@@ -234,7 +247,7 @@ class PeriodicBlockStepper:
             except BaseException:
                 self._status = "failed"
                 raise
-            return self.layout.to_dense(blocked)
+            return _StoredPopulations((self.layout.to_dense(blocked),), self._ddf_shifting)
 
     def close(self):
         """Release owned resources; retain them for retry if context activation fails."""

src/CelerisLab/lbm/channel.py

@@ -0,0 +1,54 @@
+"""Host contract for the blocked root's uniform SRT/MRT channel boundary.
+
+Boundary ghosts are evolved state, but are not physical fluid leaves. Halo
+ownership remains periodic to reproduce the uniform kernel's wrapped pulls.
+No boundary closure is applied on a fine level.
+"""
+import numpy as np
+
+
+def channel_masks(shape):
+    """Physical fluid and bulk-transfer masks in GLOBAL root coordinates."""
+    ny, nx = shape
+    fluid = np.zeros(shape, bool)
+    fluid[1:ny - 1, 1:nx - 1] = True
+    bulk = fluid.copy()
+    bulk[[1, ny - 2]] = False
+    return fluid, bulk
+
+
+def validate_channel_config(cfg, blocks):
+    if any(block is None for block in blocks):
+        raise ValueError("channel requires all-blocked levels")
+    if (cfg.inlet_scheme != "zou_he_local" or cfg.outlet_mode != "neq_extrap"
+            or cfg.y_wall_bc != "bounce_back" or cfg.top_wall_bc != "velocity"
+            or cfg.wall_slot_width != 0 or cfg.wall_normal_velocity != 0):
+        raise ValueError("channel requires Zou-He inlet, pressure NEQ outlet and stationary bounce-back y walls")
+
+
+def validate_channel_transfers(pair, layout, targets):
+    """Reject actual exchange/closure overlap, not a padded geometric envelope.
+
+    Explosion donors and restriction receivers must use ordinary bulk collision
+    semantics. Restriction writes receiver-local source slots; these may not
+    overwrite the channel's extra immutable-input reads (pressure donor and
+    same-node halfway-wall corrections). The slot check deliberately respects
+    block seams, rather than treating different halo images as the same slot.
+    """
+    ny, nx = pair.coarse_shape
+    _, bulk = channel_masks(pair.coarse_shape)
+    legal = bulk.ravel() & ~pair.covered.ravel()
+    receivers = pair.coalesce_coarse % (ny * nx)
+    if not legal[pair.ghost_coarse].all() or not legal[receivers].all():
+        raise ValueError("channel transfer donors and receivers must be uncovered bulk fluid")
+    owner = layout.owner
+    reads = []
+    # Outlet closes from the full previous-step state at its left neighbor.
+    donor = owner[1:-1, -1].astype(np.int64) - 1
+    for q in range(9):
+        reads.append(layout.slot(donor, q))
+    for y, directions in ((1, (4, 6, 7)), (ny - 2, (3, 5, 8))):
+        for q in directions:
+            reads.append(layout.slot(owner[y, 1:-1], q))
+    if np.intersect1d(targets, np.concatenate(reads)).size:
+        raise ValueError("channel restriction writes overlap immutable boundary closure reads")

src/CelerisLab/lbm/kernels/blocks.cu

@@ -14,12 +14,15 @@
 #include "core/obs.cuh"
 #include "operators/macro.cuh"
 #include "operators/collision_srt.cuh"
+#if COLLISION_MODEL == 2
+#include "operators/collision_mrt.cuh"
+#endif
 #include "streaming/pull_double_buffer.cuh"
 #include "boundary/curved_boundary.cuh"
 #include "boundary/planar_velocity.cuh"
 
-#if DIM != 2 || NQ != 9 || COLLISION_MODEL != 0 || STORE_PRECISION != 0 || USE_DDF_SHIFTING || USE_LES || STREAMING_MODEL != 0
-#error "Blocks require unshifted FP32 D2Q9 SRT double buffers without LES"
+#if DIM != 2 || NQ != 9 || (COLLISION_MODEL != 0 && COLLISION_MODEL != 2) || STORE_PRECISION != 0 || USE_LES || STREAMING_MODEL != 0
+#error "Blocks require FP32 D2Q9 SRT or MRT double buffers without LES"
 #endif
 
 // Fixed-patch cell kinds, matching refinement.py ACTIVE/INNER_GHOST/OUTER_GHOST.
@@ -141,11 +144,30 @@ extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int hal
                         gx, d_block_wall.nx, d_block_wall.slot_width, d_block_wall.amplitude));
                 else if (gy == d_block_wall.ny - 2) reconstruct_top_reservoir(f, RHO);
             }
+#if COLLISION_MODEL == 2
+            // Uniform MRT sequence (collide_dispatch, zero forcing); SRT path below unchanged.
+            float rho, ux, uy;
+            compute_rho_u(f, rho, ux, uy);
+            collide_mrt_no_force(f, rho, ux, uy,
+                fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega)));
+#else
             float rho, ux, uy, feq[NQ];
             compute_rho_u(f, rho, ux, uy);
+#if USE_DDF_SHIFTING
+            // Preserve the density perturbation before adding the unit reference.
+            const float drho = f[0]+f[1]+f[2]+f[3]+f[4]+f[5]+f[6]+f[7]+f[8];
+            const float jx = f[1]-f[2]+f[5]-f[6]+f[7]-f[8];
+            const float jy = f[3]-f[4]+f[5]-f[6]-f[7]+f[8];
+            compute_feq_from_momentum(drho, jx, jy, ux, uy, feq);
+#else
             compute_feq(rho, ux, uy, feq);
+#endif
             collide_srt_no_force(f, feq, d_params.omega);
+#endif
         }
     }
     stream_pull_store(k, f, out + base);
 }
+
+// Root channel closure; fine and periodic levels keep BlockStep.
+#include "step/block_channel.cu"

src/CelerisLab/lbm/kernels/boundary/inlet/common.cuh

@@ -27,18 +27,22 @@
 #define INLET_REGULARIZED_NEQ_DAMP 0.50f
 #endif
 
-__device__ __forceinline__ float inlet_target_u(float y_coord) {
+__device__ __forceinline__ float inlet_target_u(float y_coord, unsigned int ny, float velocity) {
 #if INLET_PROFILE == 0
-    return U0;
+    return velocity;
 #else
-    const float y_clamped = fminf((float)(NY - 2), fmaxf(1.0f, y_coord));
-    const float H = fmaxf((float)(NY - 2), 1.0f);
+    const float y_clamped = fminf((float)(ny - 2), fmaxf(1.0f, y_coord));
+    const float H = fmaxf((float)(ny - 2), 1.0f);
     const float eta = (y_clamped - 0.5f) / H;
     const float shape = fmaxf(0.0f, 4.0f * eta * (1.0f - eta));
-    return U0 * 1.5f * shape;
+    return velocity * 1.5f * shape;
 #endif
 }
 
+__device__ __forceinline__ float inlet_target_u(float y_coord) {
+    return inlet_target_u(y_coord, NY, U0);
+}
+
 #if LATTICE_MODEL == LATTICE_D2Q9
 __device__ __forceinline__ float west_velocity_rho_closure_d2q9(
     const float* __restrict__ f,

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

@@ -1,6 +1,6 @@
 // On-site D2Q9 Zou-He velocity closure, paired lattice ordering.
 // Input is poststream f at the arrival time. Only unknown populations change;
-// the caller owns collision/store. No force, shifting or off-lattice walls.
+// the caller owns collision/store. No force or off-lattice walls. Shifted storage uses unit reference density.
 #ifndef CELERIS_BOUNDARY_PLANAR_VELOCITY_CUH
 #define CELERIS_BOUNDARY_PLANAR_VELOCITY_CUH
 #include <math.h>
@@ -28,7 +28,11 @@ __device__ __forceinline__ float bottom_wall_normal_velocity(unsigned int x)
 
 __device__ __forceinline__ void reconstruct_bottom_velocity(float* f, float u, float v)
 {
-    const float rho = (f[0] + f[1] + f[2] + 2.0f*(f[4] + f[6] + f[7])) / (1.0f-v);
+    const float rho = (
+#if USE_DDF_SHIFTING
+        1.0f +
+#endif
+        f[0] + f[1] + f[2] + 2.0f*(f[4] + f[6] + f[7])) / (1.0f-v);
     f[3] = f[4] + (2.0f/3.0f)*rho*v;
     f[5] = f[6] + 0.5f*(f[2]-f[1]) + 0.5f*rho*u + (1.0f/6.0f)*rho*v;
     f[8] = f[7] + 0.5f*(f[1]-f[2]) - 0.5f*rho*u + (1.0f/6.0f)*rho*v;
@@ -36,7 +40,11 @@ __device__ __forceinline__ void reconstruct_bottom_velocity(float* f, float u, f
 
 __device__ __forceinline__ void reconstruct_top_velocity(float* f, float u, float v)
 {
-    const float rho = (f[0] + f[1] + f[2] + 2.0f*(f[3] + f[5] + f[8])) / (1.0f+v);
+    const float rho = (
+#if USE_DDF_SHIFTING
+        1.0f +
+#endif
+        f[0] + f[1] + f[2] + 2.0f*(f[3] + f[5] + f[8])) / (1.0f+v);
     f[4] = f[3] - (2.0f/3.0f)*rho*v;
     f[6] = f[5] + 0.5f*(f[1]-f[2]) - 0.5f*rho*u - (1.0f/6.0f)*rho*v;
     f[7] = f[8] + 0.5f*(f[2]-f[1]) + 0.5f*rho*u - (1.0f/6.0f)*rho*v;
@@ -46,7 +54,12 @@ __device__ __forceinline__ void reconstruct_top_velocity(float* f, float u, floa
 // J_y=rho*v follows from known poststream populations (ordinary D2Q9).
 __device__ __forceinline__ void reconstruct_top_reservoir(float* f, float rho)
 {
-    const float jy = f[0]+f[1]+f[2]+2.0f*(f[3]+f[5]+f[8])-rho;
+#if USE_DDF_SHIFTING
+    const float density_offset = rho - 1.0f;
+#else
+    const float density_offset = rho;
+#endif
+    const float jy = f[0]+f[1]+f[2]+2.0f*(f[3]+f[5]+f[8])-density_offset;
     f[4] = f[3] - (2.0f/3.0f)*jy;
     f[6] = f[5] + 0.5f*(f[1]-f[2]) - (1.0f/6.0f)*jy;
     f[7] = f[8] + 0.5f*(f[2]-f[1]) - (1.0f/6.0f)*jy;

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

@@ -64,11 +64,12 @@ __device__ __forceinline__ void compute_rho_u(const float* __restrict__ f,
 #if LATTICE_MODEL == LATTICE_D2Q9
 
 __device__ __forceinline__ void compute_feq(float rho, float ux, float uy,
-                                            float* __restrict__ feq)
-{
+                                            float* __restrict__ feq
 #if USE_DDF_SHIFTING
-    const float rhom1 = rho - 1.0f;          // perturbation density
+                                            , float rhom1
 #endif
+                                            )
+{
     const float ux3 = 3.0f * ux;
     const float uy3 = 3.0f * uy;
     const float c3  = -1.5f * (ux * ux + uy * uy);  // = -|u|²/(2cs²)
@@ -112,6 +113,39 @@ __device__ __forceinline__ void compute_feq(float rho, float ux, float uy,
     #undef FEQ_CORE
 }
 
+#if USE_DDF_SHIFTING
+// Shifted D2Q9 equilibrium using input momentum in the linear term.
+// Algebraically identical to the usual second-order equilibrium; explicit
+// rational divisions avoid 3*w rounding and the rho*(j/rho) round trip.
+__device__ __forceinline__ void compute_feq_from_momentum(
+    float delta, float jx, float jy, float ux, float uy,
+    float* __restrict__ feq)
+{
+    const float common = fmaf(-1.5f, fmaf(jx, ux, jy * uy), delta);
+    feq[0] = W0_VAL * common;
+    const float jp[4] = {jx, jy, jx + jy, jx - jy};
+    const float up[4] = {ux, uy, ux + uy, ux - uy};
+    #pragma unroll
+    for (int p = 0; p < 4; ++p) {
+        const float weight = p < 2 ? WS_VAL : WE_VAL;
+        const float odd = __fdiv_rn(jp[p], p < 2 ? 3.0f : 12.0f);
+        const float even = weight * fmaf(4.5f * jp[p], up[p], common);
+        feq[1 + 2*p] = even + odd;
+        feq[2 + 2*p] = even - odd;
+    }
+}
+#endif
+
+#if USE_DDF_SHIFTING
+// Existing callers recover the offset from rho; blocked shifted collision
+// supplies sum(g) directly to the overload above, retaining small perturbations.
+__device__ __forceinline__ void compute_feq(float rho, float ux, float uy,
+                                            float* __restrict__ feq)
+{
+    compute_feq(rho, ux, uy, feq, rho - 1.0f);
+}
+#endif
+
 #elif LATTICE_MODEL == LATTICE_D3Q19
 
 __device__ __forceinline__ void compute_feq(float rho, float ux, float uy, float uz,

src/CelerisLab/lbm/kernels/refinement.cu

@@ -11,6 +11,9 @@
 #include "core/obs.cuh"
 #include "operators/macro.cuh"
 #include "operators/collision_srt.cuh"
+#if COLLISION_MODEL == 2
+#include "operators/collision_mrt.cuh"
+#endif
 #include "streaming/pull_double_buffer.cuh"
 #include "boundary/planar_velocity.cuh"
 #include "boundary/curved_boundary.cuh"
@@ -19,8 +22,8 @@ extern "C" {
 #include "step/curved_double.cu"
 }
 
-#if DIM != 2 || NQ != 9 || COLLISION_MODEL != 0 || STORE_PRECISION != 0 || USE_DDF_SHIFTING || USE_LES || STREAMING_MODEL != 0
-#error "Refinement requires unshifted FP32 D2Q9 SRT double buffers without LES"
+#if DIM != 2 || NQ != 9 || (COLLISION_MODEL != 0 && COLLISION_MODEL != 2) || STORE_PRECISION != 0 || USE_DDF_SHIFTING || USE_LES || STREAMING_MODEL != 0
+#error "Refinement requires unshifted FP32 D2Q9 SRT or MRT double buffers without LES"
 #endif
 
 // Enabled only in the validation build; no diagnostic arrays in normal builds.
@@ -33,10 +36,17 @@ __device__ unsigned int incoming_counts[NQ * TOTAL_CELLS];
 #endif
 
 __device__ __forceinline__ void patch_collide(float* f, unsigned long k) {
+#if COLLISION_MODEL == 2
+    float rho, ux, uy;
+    compute_rho_u(f, rho, ux, uy);
+    collide_mrt_no_force(f, rho, ux, uy,
+        fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega)));
+#else
     float rho, ux, uy, feq[NQ];
     compute_rho_u(f, rho, ux, uy);
     compute_feq(rho, ux, uy, feq);
     collide_srt_no_force(f, feq, d_params.omega);
+#endif
 #ifdef CELERIS_REFINEMENT_AUDIT
     ++collision_counts[k];
 #endif

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

@@ -0,0 +1,74 @@
+// Root-only blocked counterpart of OneStep's D2Q9/SRT channel closure.
+// NX/NY are LOCAL strides; all physical boundary decisions use global coordinates.
+#include "../boundary/inlet/common.cuh"
+#include "../boundary/inlet/zou_he_local.cuh"
+#include "../boundary/outlet/pressure_neq.cuh"
+
+struct BlockChannel { unsigned int nx, ny; float velocity; };
+extern "C" { __constant__ BlockChannel d_block_channel; }
+
+// Same signature as BlockStep so the AMR hop/transfer schedule is unchanged.
+extern "C" __global__ void BlockChannelStep(
+    const fpxx* in, fpxx* out, unsigned int halo,
+    const uint16_t* flags, const unsigned char* kind)
+{
+    unsigned int x, y; unsigned long k;
+    index_from_thread(x, y, k);
+    if (x < halo || y < halo || x >= NX - halo || y >= NY - halo) return;
+    const unsigned int bx = NX - 2 * halo, by = NY - 2 * halo;
+    const unsigned int nbx = d_block_channel.nx / bx;
+    const unsigned int gx = (blockIdx.z % nbx) * bx + x - halo;
+    const unsigned int gy = (blockIdx.z / nbx) * by + y - halo;
+    const unsigned long base = (unsigned long)blockIdx.z * NQ * TOTAL_CELLS;
+    const unsigned long cell = (unsigned long)blockIdx.z * TOTAL_CELLS + k;
+    float f[NQ];
+    if (kind != nullptr && kind[cell] == KIND_OUTER) {
+        for (int q = 0; q < NQ; ++q) f[q] = NAN;
+    } else {
+        unsigned long j[NQ];
+        compute_neighbors(x, y, j);
+        stream_pull_load(k, f, in + base, j);
+        const bool wall = gy == 0 || gy == d_block_channel.ny - 1;
+        const bool outlet = gx == d_block_channel.nx - 1;
+        if (wall) {
+            // Wall owns all four corners. Ghosts pull and swap, never collide.
+            for (int q = 1; q < NQ; q += 2) {
+                const float t = f[q]; f[q] = f[q + 1]; f[q + 1] = t;
+            }
+        } else if (gx == 0) {
+            apply_zou_he_left_velocity_inlet_d2q9(
+                f, inlet_target_u((float)gy, d_block_channel.ny, d_block_channel.velocity), 0.0f);
+        } else if (outlet) {
+            // Immutable input donor, receiver-local (also valid across a block seam).
+            float donor[NQ];
+            for (int q = 0; q < NQ; ++q)
+                donor[q] = load_ddf(in + base, index_f(k - 1, q));
+            apply_pressure_outlet_d2q9(f, donor);
+        } else if (gy == 1) {
+            f[3] = load_ddf(in + base, index_f(k, 4));
+            f[5] = load_ddf(in + base, index_f(k, 6));
+            f[8] = load_ddf(in + base, index_f(k, 7));
+        } else if (gy == d_block_channel.ny - 2) {
+            f[4] = load_ddf(in + base, index_f(k, 3));
+            f[6] = load_ddf(in + base, index_f(k, 5));
+            f[7] = load_ddf(in + base, index_f(k, 8));
+        }
+        if (!wall && !outlet) {
+#if COLLISION_MODEL == 2
+            // Match uniform MRT arithmetic (collide_dispatch, zero forcing).
+            float rho, ux, uy;
+            compute_rho_u(f, rho, ux, uy);
+            collide_mrt_no_force(f, rho, ux, uy,
+                fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega)));
+#else
+            // Match uniform SRT arithmetic, including its shifted equilibrium path.
+            float rho, ux, uy, feq[NQ], forcing[NQ] = {};
+            compute_rho_u(f, rho, ux, uy);
+            compute_feq(rho, ux, uy, feq);
+            const float omega = fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega));
+            collide_srt(f, feq, forcing, omega);
+#endif
+        }
+    }
+    stream_pull_store(k, f, out + base);
+}

src/CelerisLab/lbm/population_storage.py

@@ -0,0 +1,24 @@
+"""Private D2Q9 FP32 storage boundary for the opt-in blocked steppers."""
+from dataclasses import dataclass
+import numpy as np
+
+# Unit-density reference weights in FP64; sum(w)=1 matches rho=1+sum(g).
+_WEIGHTS = np.array([4/9, 1/9, 1/9, 1/9, 1/9, 1/36, 1/36, 1/36, 1/36],
+                    dtype=np.float64)[:, None, None]
+
+
+@dataclass(frozen=True, eq=False)
+class _StoredPopulations:
+    values: tuple
+    shifted: bool
+
+
+def _encode_physical(values, shifted):
+    if not shifted:
+        return np.array(values, dtype=np.float32, order="C", copy=True)
+    return np.ascontiguousarray(np.asarray(values, dtype=np.float64) - _WEIGHTS,
+                                dtype=np.float32)
+
+
+def _decode_stored(values, shifted):
+    return values.astype(np.float64) + _WEIGHTS if shifted else values

src/CelerisLab/lbm/raw_exchange.py

@@ -7,16 +7,26 @@ import numpy as np
 
 
 def summarize_exchange(rows, body_ids, aperture, *, body_id, level, h,
-                       time_end, window, reference_center):
+                       time_end, window, reference_center, duration=None):
     """Reduce FP64 interval integrals or convert FP32 last-substep raw rows.
 
     dt_l=h_l under acoustic 2:1 refinement. Local momentum j -> base force
     h*j -> impulse h*dt*j; torque adds one h; mass is h**2*dm.
+    window="window" reduces an FP64 observation-window ledger spanning
+    ``duration`` (integer >= 1) base intervals; rates divide by that duration.
     """
-    if window not in ("last_base_step", "last_substep"):
+    if window not in ("last_base_step", "last_substep", "window"):
         raise ValueError("unknown exchange window")
+    if window == "window":
+        if isinstance(duration, bool) or not isinstance(duration, (int, np.integer)) or duration < 1:
+            raise ValueError("window duration must be an integer number of base intervals >= 1")
+    elif duration is not None:
+        raise ValueError("duration applies only to window='window'")
     values = np.array(rows, dtype=np.float64, copy=True)
-    duration = 1.0 if window == "last_base_step" else h
+    if window == "window":
+        duration = float(duration)
+    else:
+        duration = 1.0 if window == "last_base_step" else h
     if window == "last_substep":
         values *= np.array([h*h, h*h, h*h*h, h*h])
     selected = np.asarray(body_ids) == body_id
@@ -31,6 +41,7 @@ def summarize_exchange(rows, body_ids, aperture, *, body_id, level, h,
     return dict(wall=reduce(selected & ~aperture), aperture=reduce(selected & aperture),
                 total=reduce(selected), body_id=int(body_id), level=int(level),
                 time_start=time_end-duration, time_end=time_end, dt=h,
-                sample_count=round(1/h) if window == "last_base_step" else 1,
+                sample_count=(round(duration/h) if window == "window"
+                              else round(1/h) if window == "last_base_step" else 1),
                 window=window, units="base_lattice_per_unit_depth",
                 definition="raw_boundary_exchange", reference_center=tuple(reference_center), valid=True)

src/CelerisLab/lbm/refinement.py

@@ -52,17 +52,17 @@ class FixedPatch2D:
     def from_config(cls, cfg: LBMConfig, bounds) -> "FixedPatch2D":
         """Prepare maps without changing cfg or allocating coarse-sized tables.
 
-        The initial backend contract is D2Q9/SRT/FP32/unshifted/double_buffer,
+        The backend contract is D2Q9/SRT or MRT/FP32/double_buffer,
         without LES. Two coarse cells of margin keep all halo donors in bounds;
         this does NOT certify that those donors are fluid or periodic.
         """
         checked = replace(cfg)
         checked.validate()
-        if (not checked.is_d2q9 or checked.collision != "SRT"
+        if (not checked.is_d2q9 or checked.collision not in ("SRT", "MRT")
                 or checked.streaming != "double_buffer"
                 or checked.store_precision != "FP32"
-                or checked.ddf_shifting or checked.les_enabled):
-            raise ValueError("FixedPatch2D requires D2Q9/SRT/FP32/unshifted/double_buffer without LES")
+                or checked.les_enabled):
+            raise ValueError("FixedPatch2D requires D2Q9/SRT or MRT/FP32/double_buffer without LES")
         bounds = tuple(bounds)
         if len(bounds) != 4 or any(isinstance(v, (bool, np.bool_))
                                    or not isinstance(v, Integral) for v in bounds):
@@ -151,20 +151,17 @@ class FixedPatch2D:
             raise ValueError("levels are not synchronized: fine_steps must equal 2*coarse_steps")
 
 
-MIN_PATCH_GAP = 4  # coarse cells: the full 4-fine-cell supports stay disjoint
+MIN_PATCH_GAP = 4  # Legacy exported constant; union transfers impose no gap policy.
 
 
 @dataclass(frozen=True, eq=False)
 class FixedPatchSet2D:
-    """Several separated FixedPatch2D rectangles sharing one 2:1 fine level.
+    """Rectangles describing one 2:1 fine union, with unique transfer ownership.
 
-    The fine level is the whole (2*ny, 2*nx) grid in level coordinates; only
-    the patch supports (active region + HALO ghost band) are stored, by a
-    sparse BlockLayout2D from fine_layout(). Tables have FixedPatch2D meaning
-    and row/child order, re-addressed to level cells: ghost_fine are level
-    fine cells, coalesce_fine are q*N+cell level slots; coarse ids unchanged.
-    Patches are kept in canonical (y0, x0) order, so input order is irrelevant.
-    Active regions must be >= MIN_PATCH_GAP coarse cells apart (Chebyshev).
+    Input patches remain canonical metadata. Numerical kinds and tables are
+    built once from their union; adjacent, overlapping and concave unions are
+    accepted when their actual transfer dependencies close. Sparse storage owns
+    every whole tile intersecting the support, independent of rectangle edges.
     """
     coarse_shape: tuple[int, int]
     patches: tuple
@@ -182,42 +179,43 @@ class FixedPatchSet2D:
         if not rectangles:
             raise ValueError("at least one patch rectangle is required")
         patches = sorted((FixedPatch2D.from_config(cfg, r) for r in rectangles),
-                         key=lambda p: (p.bounds[1], p.bounds[0]))
-        for i, a in enumerate(patches):
-            for b in patches[i + 1:]:
-                (ax0, ay0, ax1, ay1), (bx0, by0, bx1, by1) = a.bounds, b.bounds
-                if max(ay0 - by1, by0 - ay1, ax0 - bx1, bx0 - ax1) < MIN_PATCH_GAP:
-                    raise ValueError(f"patches {a.bounds} and {b.bounds} are closer "
-                                     f"than {MIN_PATCH_GAP} coarse cells")
+                         key=lambda p: (p.bounds[1], p.bounds[0], p.bounds[3], p.bounds[2]))
         ny, nx = patches[0].coarse_shape
-        fy, fx = 2 * ny, 2 * nx
-        if fy * fx * 9 > np.iinfo(np.uint64).max // 2:
-            raise ValueError("fine level exceeds 64-bit indexing")
-        kind = np.full((fy, fx), OUTER_GHOST, np.uint8)
-        ghost_fine, coalesce_fine = [], []
-        for p in patches:
-            y0, x0 = p.fine_level_origin
-            ly, lx = p.fine_shape
-            kind[y0:y0 + ly, x0:x0 + lx] = p.fine_kind
-            gy, gx = np.divmod(p.ghost_fine.astype(np.int64), lx)
-            ghost_fine.append((gy + y0) * fx + gx + x0)
-            q, cell = np.divmod(p.coalesce_fine.astype(np.int64), ly * lx)
-            cy, cx = np.divmod(cell, lx)
-            coalesce_fine.append(q * fy * fx + (cy + y0) * fx + cx + x0)
-        ghost_coarse = np.concatenate([p.ghost_coarse for p in patches])
-        coalesce_coarse = np.concatenate([p.coalesce_coarse for p in patches])
         covered = np.zeros((ny, nx), bool)
         for p in patches:
             covered |= p.covered
-        cells = ny * nx
-        if (covered.ravel()[ghost_coarse].any()
-                or covered.ravel()[coalesce_coarse % np.uint64(cells)].any()
-                or len(np.unique(coalesce_coarse)) != len(coalesce_coarse)):
-            raise ValueError("patch transfers overlap another patch")
-        return cls((ny, nx), tuple(patches), patches[0].tau_coarse, patches[0].tau_fine,
-                   _readonly(kind, np.uint8), _readonly(np.concatenate(ghost_fine), np.uint64),
-                   _readonly(ghost_coarse, np.uint32), _readonly(coalesce_coarse, np.uint64),
-                   _readonly(np.concatenate(coalesce_fine), np.uint64).reshape(-1, 4))
+        fy, fx = 2 * ny, 2 * nx
+        if fy * fx * 9 > np.iinfo(np.uint64).max // 2:
+            raise ValueError("fine level exceeds 64-bit indexing")
+        active = np.kron(covered, np.ones((2, 2), bool))
+        inner = _dilate(active, 2)
+        support = _dilate(active, HALO)
+        kind = np.full((fy, fx), OUTER_GHOST, np.uint8)
+        kind[inner] = INNER_GHOST
+        kind[active] = ACTIVE
+        gy, gx = np.nonzero(support & ~active)
+        ghost_fine = (gy * fx + gx).astype(np.uint64)
+        ghost_coarse = ((gy >> 1) * nx + (gx >> 1)).astype(np.uint32)
+        targets, children = [], []
+        ry, rx = np.nonzero(_dilate(covered, 1) & ~covered)
+        for y, x in zip(ry.tolist(), rx.tolist()):
+            for q in range(1, 9):
+                if _closed_union_contains(covered, 2 * x + 1 - D2Q9_EX[q], 2 * y + 1 - D2Q9_EY[q]):
+                    targets.append(q * ny * nx + y * nx + x)
+                    children.append([q * fy * fx + (2 * y + dy) * fx + 2 * x + dx
+                                     for dy in (0, 1) for dx in (0, 1)])
+        targets = np.asarray(targets, np.uint64)
+        children = np.asarray(children, np.uint64).reshape(-1, 4)
+        cells = np.asarray(children % np.uint64(fy * fx), np.intp)
+        if not (kind.ravel()[cells] == INNER_GHOST).all():
+            raise ValueError("coalesce children must be INNER ghost cells")
+        pair = cls((ny, nx), tuple(patches), patches[0].tau_coarse, patches[0].tau_fine,
+                   _readonly(kind, np.uint8), _readonly(ghost_fine, np.uint64),
+                   _readonly(ghost_coarse, np.uint32), _readonly(targets, np.uint64),
+                   _readonly(children, np.uint64))
+        _transfer_footprint(pair)
+        _check_parent_pull_closure(pair)
+        return pair
 
     @property
     def fine_shape(self) -> tuple[int, int]:
@@ -240,22 +238,61 @@ class FixedPatchSet2D:
         return mask
 
     def fine_layout(self, block):
-        """Sparse non-periodic layout storing exactly the patch supports."""
+        """Sparse non-periodic tiles covering union support; domain divisibility remains."""
         from .blocks import BlockLayout2D
         block = tuple(block)
         if len(block) != 2 or any(isinstance(v, (bool, np.bool_)) or not isinstance(v, Integral)
                                   or v <= 0 for v in block):
             raise ValueError("fine block must be two positive integers (by, bx)")
         by, bx = map(int, block)
-        origins = set()
-        for p in self.patches:
-            y0, x0 = p.fine_level_origin
-            ly, lx = p.fine_shape
-            if y0 % by or x0 % bx or ly % by or lx % bx:
-                raise ValueError(f"patch {p.bounds} support is not aligned to fine blocks {block}")
-            origins.update((y, x) for y in range(y0, y0 + ly, by) for x in range(x0, x0 + lx, bx))
+        support = _dilate(self.fine_kind == ACTIVE, HALO)
+        ty, tx = np.nonzero(support)
+        origins = sorted({(int(y) // by * by, int(x) // bx * bx) for y, x in zip(ty, tx)})
         return BlockLayout2D.build(*self.fine_shape, by, bx, periodic=(False, False),
-                                   origins=sorted(origins))
+                                   origins=origins)
+
+
+def _dilate(mask, steps):
+    """Chebyshev dilation with a zero (non-periodic) exterior."""
+    out = np.asarray(mask, bool)
+    for _ in range(steps):
+        grown = out.copy()
+        grown[1:] |= out[:-1]
+        grown[:-1] |= out[1:]
+        wide = grown.copy()
+        wide[:, 1:] |= grown[:, :-1]
+        wide[:, :-1] |= grown[:, 1:]
+        out = wide
+    return out
+
+
+def _closed_union_contains(covered, mx, my):
+    """Does doubled point (mx, my) lie in a closed covered square [2X, 2X+2] x [2Y, 2Y+2]?"""
+    xs = (mx // 2 - 1, mx // 2) if mx % 2 == 0 else ((mx - 1) // 2,)
+    ys = (my // 2 - 1, my // 2) if my % 2 == 0 else ((my - 1) // 2,)
+    ny, nx = covered.shape
+    return any(0 <= y < ny and 0 <= x < nx and covered[y, x] for y in ys for x in xs)
+
+
+def _check_parent_pull_closure(pair):
+    """Every uncovered parent cell pulling q from a covered cell needs its (R, q) row; no row
+    may stage into an uncovered parent cell unless its link passes the closed union."""
+    ny, nx = pair.coarse_shape
+    rows = set(map(int, pair.coalesce_coarse))
+    if len(rows) != len(pair.coalesce_coarse):
+        raise ValueError("duplicate coalesce writer (receiver, q)")
+    covered = pair.covered
+    for y, x in zip(*np.nonzero(~covered)):
+        for q in range(1, 9):
+            sy, sx = y - D2Q9_EY[q], x - D2Q9_EX[q]
+            if 0 <= sy < ny and 0 <= sx < nx and covered[sy, sx]:
+                if q * ny * nx + y * nx + x not in rows:
+                    raise ValueError("a surviving parent ACTIVE pull lacks its coalesce row")
+    for slot in rows:
+        q, cell = divmod(slot, ny * nx)
+        y, x = divmod(cell, nx)
+        if covered[y, x] or not _closed_union_contains(covered, 2*x + 1 - D2Q9_EX[q], 2*y + 1 - D2Q9_EY[q]):
+            raise ValueError("coalesce row outside the union boundary")
 
 
 @dataclass(frozen=True, eq=False)
@@ -323,14 +360,14 @@ def nest_patch(cfg, parent, child_bounds):
 
 
 def nest_patch_set(cfg, parent, child_rectangles_in_parent):
-    """Separated child union in the parent's allocated L1 array coordinates.
+    """Child rectangle union in the parent's allocated L1 array coordinates.
 
     The root edge coordinate of L1 index (x, y) is
     ``parent.fine_origin + (x, y)/2``. For a set parent its origin is zero;
     a scalar parent's window includes its ghost band and has a shifted origin.
     The child set allocates a dense public extent twice that L1 array, while
-    its blocked storage omits holes. Both adjacent pairs retain gap >= 4 in
-    their respective parent-cell units. Validate support over the whole union.
+    its blocked storage omits holes. Validate the actual adjacent-pair support over the whole union;
+    no rectangle-gap policy is imposed.
     """
     child = FixedPatchSet2D.from_config(
         _middle_config(cfg, parent), child_rectangles_in_parent)
@@ -373,8 +410,19 @@ def _nested_kind(parent, child):
     return _readonly(composed, np.uint8)
 
 
+def _circle_in_coverage(covered, cx, cy, radius):
+    """Exact disk versus uncovered parent squares; no containing-rectangle rule."""
+    ny, nx = covered.shape
+    if cx-radius < 0 or cy-radius < 0 or cx+radius > nx or cy+radius > ny:
+        return False
+    y, x = np.nonzero(~covered)
+    dx = np.maximum(np.maximum(x-cx, cx-(x+1)), 0.)
+    dy = np.maximum(np.maximum(y-cy, cy-(y+1)), 0.)
+    return not np.any(dx*dx + dy*dy < radius*radius)
+
+
 def prepare_fine_boundary(cfg, patch, objects, *, wall_aware=False, jets=(),
-                          parent_origin=(0., 0.), parent_spacing=1.):
+                          parent_origin=(0., 0.), parent_spacing=1., _shifted_blocked=False):
     """Project fixed circles and reject boundary/transfer overlap before CUDA.
 
     Input centers/radii and Jets are in the root edge frame. ``parent_origin``
@@ -386,7 +434,7 @@ def prepare_fine_boundary(cfg, patch, objects, *, wall_aware=False, jets=(),
     Original objects are not retained or mutated. No runtime geometry updates.
     ``patch`` may be a FixedPatchSet2D: objects are then projected once onto
     the level-wide fine grid (fine_origin (0, 0)); each circle must lie inside
-    one patch, body ids are the positions in ``objects`` (never patch ids), and
+    the union, body ids are the positions in ``objects`` (never patch ids), and
     jets and the (receiver, dir) writer checks run once over the whole level.
     ``jets`` use the same coarse cell-edge frame as the objects; their
     selection is checked against true fine hits here. Jet amplitudes need no
@@ -403,8 +451,8 @@ def prepare_fine_boundary(cfg, patch, objects, *, wall_aware=False, jets=(),
     checked = replace(cfg)
     checked.validate()
     if (patch.coarse_shape != (cfg.ny, cfg.nx) or not checked.is_d2q9
-            or checked.collision != "SRT" or checked.streaming != "double_buffer"
-            or checked.store_precision != "FP32" or checked.ddf_shifting or checked.les_enabled
+            or checked.collision not in ("SRT", "MRT") or checked.streaming != "double_buffer"
+            or checked.store_precision != "FP32" or checked.les_enabled
             or patch.tau_coarse != 3*checked.viscosity+.5):
         raise ValueError("Fine boundary requires the matching fixed-patch configuration")
     from ..body.geometry.circle import CircleGeometry
@@ -416,7 +464,6 @@ def prepare_fine_boundary(cfg, patch, objects, *, wall_aware=False, jets=(),
     flags = np.full(nx * ny, FLUID, np.uint16)
     active = patch.fine_kind.ravel() == ACTIVE
     local, links = [], []
-    rects = patch.bounds if isinstance(patch, FixedPatchSet2D) else (patch.bounds,)
     covered = patch.covered
     dx = dt = .5 * parent_spacing
     fine_origin = origin + parent_spacing * np.asarray(patch.fine_origin)
@@ -430,9 +477,8 @@ def prepare_fine_boundary(cfg, patch, objects, *, wall_aware=False, jets=(),
         px, py = (np.asarray((cx, cy)) - origin) / parent_spacing
         pr = radius / parent_spacing
         if (not np.isfinite([cx, cy, radius]).all() or radius <= 0
-                or not any(x0 <= px-pr and px+pr <= x1 and y0 <= py-pr
-                           and py+pr <= y1 for x0, y0, x1, y1 in rects)):
-            raise ValueError("Entire fine circle must lie inside one physical patch")
+                or not _circle_in_coverage(covered, px, py, pr)):
+            raise ValueError("Entire fine circle must lie inside the physical patch union")
         center = tuple((v-o)/dx-.5 for v, o in zip((cx, cy), fine_origin))
         spec = obj.surface_actuation
         if spec is not None:
@@ -509,7 +555,7 @@ def prepare_fine_boundary(cfg, patch, objects, *, wall_aware=False, jets=(),
         jet_cfg.validate()
         fluid_idx, dirs, _, body_id, rx, ry, _, _ = packed
         surface_link_velocity(local, jet_cfg, flags, fluid_idx, dirs, body_id, rx, ry,
-                              jets=jets, hits=hits)
+                              jets=jets, hits=hits, _shifted_blocked=_shifted_blocked)
     return FineBoundary2D(tuple(local), _readonly(flags, np.uint16),
                           tuple(_readonly(a, a.dtype) for a in packed),
                           _readonly(sorted(boundary), np.uint32),

src/CelerisLab/lbm/refinement_stepper.py

@@ -16,6 +16,7 @@ import pycuda.driver as cuda
 
 from ..cuda import compiler_v2 as compiler
 from .blocks import BlockLayout2D
+from .channel import channel_masks, validate_channel_config, validate_channel_transfers
 from ..body.jets import jet_link_velocity
 from .block_stepper import MAX_GRID_DIM
 from .hierarchy import build_hierarchy, is_rectangle
@@ -25,6 +26,7 @@ from .refinement_tables import (blocked_base_tables, blocked_fine_tables,  # noq
                                 BLOCKED_BOUNDARY_DEVICE_COLUMNS)
 from .descriptors import SOLID
 from .raw_exchange import summarize_exchange
+from .population_storage import _StoredPopulations, _encode_physical, _decode_stored
 
 
 @dataclass
@@ -134,6 +136,11 @@ class FixedPatchStepper:
 
     ``coarse_boundary`` is ``periodic`` (default) or ``slot_reservoir``:
     a bottom on-site slot and top fixed-density reservoir from ``cfg``.
+    ``channel`` requires all-blocked levels and matches uniform SRT Zou-He
+    inlet, pressure-NEQ outlet and stationary bounce-back y walls. Boundary
+    ghosts evolve internally but public downloads exclude them; initialization
+    must supply their finite positive-density state. Transfers must avoid the
+    actual channel closure footprint. Channel remapping is not supported.
     Periodic mode does not apply the wall/slot boundary fields in ``cfg``.
     Outer walls remain coarse; fine halo donors must be ordinary fluid.
     ``fine_objects`` accepts fixed circles in coarse physical coordinates;
@@ -165,18 +172,21 @@ class FixedPatchStepper:
     base interval or last substep in base-lattice units per unit depth. This
     is boundary momentum reaction, not a validated nozzle/stress load. Dense
     body layouts are not supported by this new diagnostic in this candidate.
+    ``observe_window_exchange=True`` (requires observe_raw_exchange) adds a second
+    per-link FP64 ledger accumulated by the same kernel between
+    start_exchange_window() and read_exchange_window(); disabled by default.
     ``jets`` (body.jets.Jet) use the same edge frame as ``fine_objects`` and
     drive fine cut links selected at their true hits; ``set_jet_amplitude``
     re-uploads the per-link velocity column only.
 
     Several patches: pass ``bounds`` as a sequence of (x0, y0, x1, y1)
-    rectangles (FixedPatchSet2D, >= 4 coarse cells apart). This requires
+    rectangles (FixedPatchSet2D, dependency-validated union). This requires
     fine_block and base_block; fine_objects/jets are projected once onto the
-    level-wide fine grid (each circle inside one patch, body ids = object
+    level-wide fine grid (each circle inside the union, body ids = object
     positions, jet and writer checks level-wide), and
     uses one level-wide fine grid (2*ny, 2*nx): public fine arrays have that
     shape, NaN outside the patches' ACTIVE cells. The step sequence is the
-    blocked one above, with every patch's transfers in one launch each.
+    blocked one above, with unique union transfers in one launch each.
 
     Three-level separated unions: ``child_rectangles_in_parent=[...]`` is
     mutually exclusive with scalar ``child`` and accepts either parent form.
@@ -209,11 +219,11 @@ class FixedPatchStepper:
                  fine_block=None, base_block=None, jets=(), child=None, child_block=None,
                  child_objects=(), child_jets=(), child_rectangles_in_parent=None,
                  refinements=None, level_blocks=None, level_objects=None, level_jets=None,
-                 observe_raw_exchange=False):
+                 observe_raw_exchange=False, observe_window_exchange=False):
         if HALO != 4:
             raise RuntimeError("refinement.cu requires a four-cell fine halo")
-        if coarse_boundary not in ("periodic", "slot_reservoir"):
-            raise ValueError("coarse_boundary must be periodic or slot_reservoir")
+        if coarse_boundary not in ("periodic", "slot_reservoir", "channel"):
+            raise ValueError("coarse_boundary must be periodic, slot_reservoir or channel")
         if refinements is None:
             if bounds is None:
                 raise ValueError("bounds or refinements is required")
@@ -238,6 +248,12 @@ class FixedPatchStepper:
         level_blocks = [None if b is None else tuple(b) for b in level_blocks]
         _check_levels(refinements, level_blocks, level_objects, level_jets)
         depth = len(refinements) + 1
+        self._ddf_shifting = bool(cfg.ddf_shifting)
+        if self._ddf_shifting and any(b is None for b in level_blocks):
+            raise ValueError("shifted populations require all-blocked levels; dense/mixed is unsupported")
+        self._channel = coarse_boundary == "channel"
+        if self._channel:
+            validate_channel_config(cfg, level_blocks)
         self._wall_aware = coarse_boundary == "slot_reservoir"
         if self._wall_aware and (cfg.y_wall_bc != "on_site_velocity" or cfg.top_wall_bc != "reservoir"
                                  or cfg.velocity != 0 or not 2 <= cfg.wall_slot_width < cfg.nx):
@@ -264,6 +280,9 @@ class FixedPatchStepper:
                 if layout.nblocks > MAX_GRID_DIM or layout.local_shape[0] > MAX_GRID_DIM:
                     raise ValueError(f"{_level_name(level)} blocks exceed launch limits")
             self._layouts.append(layout)
+        if self._channel:
+            base = blocked_base_tables(pairs[0], self._layouts[0], np.argsort(pairs[0].coalesce_coarse))
+            validate_channel_transfers(pairs[0], self._layouts[0], base[1])
         # Kernel configurations per level (geometry configs live in self.hierarchy).
         fine_cfg = replace(cfg, nx=pairs[0].fine_shape[1], ny=pairs[0].fine_shape[0],
                            viscosity=2 * cfg.viscosity)
@@ -283,11 +302,12 @@ class FixedPatchStepper:
             pair = pairs[level - 1]
             if level == 1:
                 host = prepare_fine_boundary(cfg, pair, level_objects[1],
-                                             wall_aware=self._wall_aware, jets=level_jets[1])
+                                             wall_aware=self._wall_aware, jets=level_jets[1],
+                                             _shifted_blocked=self._ddf_shifting)
             else:
                 host = prepare_fine_boundary(self._level_cfgs[level - 1], pair, level_objects[level],
                                              jets=level_jets[level], parent_origin=origin,
-                                             parent_spacing=spacing)
+                                             parent_spacing=spacing, _shifted_blocked=self._ddf_shifting)
             self._boundary_hosts.append(host)
             origin = tuple(o + spacing * f for o, f in zip(origin, pair.fine_origin))
             spacing /= 2
@@ -298,6 +318,13 @@ class FixedPatchStepper:
                                       for b, layout in zip(self._boundary_hosts, self._layouts)):
             raise ValueError("raw exchange currently requires blocked body levels")
         self._observe_raw_exchange = observe_raw_exchange
+        if not isinstance(observe_window_exchange, bool):
+            raise ValueError("observe_window_exchange must be bool")
+        if observe_window_exchange and not observe_raw_exchange:
+            raise ValueError("observe_window_exchange requires observe_raw_exchange")
+        self._observe_window = observe_window_exchange
+        self._window_valid = False
+        self._window_start = None
         self._exchange_valid = False
         self._exchange_spacing = tuple(spacings)
         self._exchange_centers = tuple(tuple(tuple(o.geometry.center) for o in objects)
@@ -305,13 +332,17 @@ class FixedPatchStepper:
         self._level_jets = level_jets
         self._level_boundaries = [None] + [
             _prepare_blocked_boundary(self._boundary_hosts[level], self._layouts[level],
-                                      self._level_cfgs[level], level_jets[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]
         for level in range(1, depth):
             host = self._boundary_hosts[level]
             if host is not None:
                 self._leaf[level] &= (host.flags.reshape(pairs[level - 1].fine_shape) & SOLID) == 0
+        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 the bounded manual L2 continuation.
         self._remap_cfg = replace(cfg)
         self._level_blocks = level_blocks
@@ -344,6 +375,8 @@ class FixedPatchStepper:
                             count = len(b.tables["cell"])
                             b.last_raw = self._allocate(count * 4 * 4)
                             b.interval_integral = self._allocate(count * 4 * 8)
+                            if self._observe_window:
+                                b.window_integral = self._allocate(count * 4 * 8)
                             b.accumulate = level.module.get_function("BlockAccumulateExchange")
                         b.kernel = level.module.get_function("BlockCurvedBoundary")
                 if self._level_boundaries[1] is not None:
@@ -378,7 +411,7 @@ class FixedPatchStepper:
             config = replace(config, nx=layout.local_shape[1], ny=layout.local_shape[0],
                              y_wall_bc="bounce_back", top_wall_bc="velocity",
                              wall_slot_width=0, wall_normal_velocity=0.0, velocity=0.0)
-            name, source = "BlockStep", "blocks.cu"
+            name, source = ("BlockChannelStep" if index == 0 and self._channel else "BlockStep"), "blocks.cu"
         config.validate()
         if not np.isclose(config.omega, 1 / tau, rtol=1e-14, atol=0):
             raise RuntimeError("level omega/tau mismatch")
@@ -406,6 +439,12 @@ class FixedPatchStepper:
             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
+                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
                 ptr, size = module.get_global("d_block_wall")
@@ -469,24 +508,59 @@ class FixedPatchStepper:
         for array, leaf in zip(arrays, self._leaf):
             array[:, ~leaf] = np.nan
 
-    def initialize(self, *level_g, coarse_g=None, fine_g=None, child_g=None):
-        """Upload an explicit G_0, one Q-major array per level; non-leaves are discarded.
+    def _mask_nonstate(self, *arrays):
+        for array, state in zip(arrays, self._state):
+            array[:, ~state] = np.nan
 
-        Pass the level arrays positionally (L0 first), or use the original
-        two/three-level names coarse_g, fine_g[, child_g]; mixing is rejected.
+    def _check_storage(self):
+        if any(bool(c.ddf_shifting) != self._ddf_shifting for c in self._level_cfgs):
+            raise ValueError("mixed population representations are unsupported")
+        if self._ddf_shifting and any(layout is None for layout in self._layouts):
+            raise ValueError("shifted populations require all-blocked levels")
+
+    def initialize(self, *level_g, coarse_g=None, fine_g=None, child_g=None):
+        """Upload physical postcollision G, encoding before any FP32 rounding.
+
+        Shifted input should be physical FP64 to retain small perturbations.
+        Public initialization resets the step count; non-leaves are discarded.
         """
         given = _initial_levels(level_g, coarse_g, fine_g, child_g)
         with self._lock, self.ctx.activate():
             self._require("new", "ready")
+            self._check_storage()
             if len(given) != len(self._levels):
                 raise ValueError("Expected one Q-major array per level")
-            arrays = [np.array(a, dtype=np.float32, order="C", copy=True) for a in given]
+            dtype = np.float64 if self._ddf_shifting else np.float32
+            arrays = [np.array(a, dtype=dtype, order="C", copy=True) for a in given]
             if any(a.shape != level.shape for a, level in zip(arrays, self._levels)):
                 raise ValueError("Expected Q-major (9,ny,nx) arrays for every level")
-            for array, mask in zip(arrays, self._leaf):
+            for array, mask in zip(arrays, self._state):
                 if not np.isfinite(array[:, mask]).all() or (array[:, mask].sum(axis=0) <= 0).any():
                     raise ValueError("Leaf populations must be finite with positive density")
-            self._mask_nonleaves(*arrays)
+            encoded = tuple(_encode_physical(a, self._ddf_shifting) for a in arrays)
+            self._initialize_raw(_StoredPopulations(encoded, self._ddf_shifting))
+
+    def _initialize_raw(self, snapshot):
+        """Upload an internally tagged same-representation state without a codec."""
+        with self._lock, self.ctx.activate():
+            self._require("new", "ready")
+            self._check_storage()
+            if (not isinstance(snapshot, _StoredPopulations)
+                    or type(snapshot.shifted) is not bool or snapshot.shifted != self._ddf_shifting):
+                raise ValueError("stored population representation mismatch")
+            if len(snapshot.values) != len(self._levels):
+                raise ValueError("Expected one stored array per level")
+            if any(not isinstance(a, np.ndarray) or a.dtype != np.float32 for a in snapshot.values):
+                raise ValueError("stored populations must be FP32")
+            arrays = [np.array(a, order="C", copy=True) for a in snapshot.values]
+            if any(a.shape != level.shape for a, level in zip(arrays, self._levels)):
+                raise ValueError("Expected Q-major (9,ny,nx) arrays for every level")
+            for array, mask in zip(arrays, self._state):
+                values = array[:, mask]
+                if (not np.isfinite(values).all()
+                        or (values.sum(axis=0, dtype=np.float64) + int(self._ddf_shifting) <= 0).any()):
+                    raise ValueError("Leaf populations must be finite with positive density")
+            self._mask_nonstate(*arrays)
             arrays = [a if level.layout is None else level.layout.to_blocked(a)
                       for a, level in zip(arrays, self._levels)]
             try:
@@ -497,6 +571,7 @@ class FixedPatchStepper:
                     cuda.memset_d8(self._fine_curved.link_exchange_gpu, 0, self._fine_curved.count * 12)
                 self._step_count = 0
                 self._exchange_valid = False
+                self._window_valid = False
                 self._status = "ready"
             except BaseException:
                 self._status = "failed"
@@ -507,8 +582,8 @@ class FixedPatchStepper:
 
         Periodic three-level set children only; L0/L1 and block sizes stay fixed.
         A body run requires its original root-edge-frame ``child_objects``:
-        stationary circles pinned in unchanged complete L2 rectangles.
-        Separated fluid-only rectangles may be added, moved or removed.
+        stationary circles with identical physical boundary and BFL support.
+        Other fluid coverage may change, even in the same connected union.
         Current L2 Jets (including updated amplitudes)
         are preserved. Local projected boundary objects are not valid inputs.
         The lock and stream synchronization certify
@@ -518,13 +593,14 @@ class FixedPatchStepper:
         construction/upload errors never invalidate the old instance; failure
         to synchronize/read its device state does, as with step/download.
         """
-        from .synchronized_remap import (RemapTopology, remap_synchronized,
+        from .synchronized_remap import (RemapTopology, remap_synchronized, _remap_stored,
                                          same_physical_boundary, boundary_velocity)
 
         with self._lock:
             self._require("ready")
+            self._check_storage()
             pairs = self.hierarchy.pairs
-            if (self._wall_aware or self.hierarchy.depth != 3
+            if (self._wall_aware or self._channel or self.hierarchy.depth != 3
                     or not isinstance(pairs[1], FixedPatchSet2D)
                     or self._boundary_hosts[1] is not None
                     or any(b is not None for b in self._level_boundaries[:2])):
@@ -534,23 +610,25 @@ class FixedPatchStepper:
             actual = self._level_boundaries[2]
             jets = () if actual is None else actual.jets
             old = RemapTopology.from_config(cfg, pairs[0], pairs[1].bounds,
-                                           child_objects=child_objects, child_jets=jets)
+                                           child_objects=child_objects, child_jets=jets,
+                                           _shifted_blocked=self._ddf_shifting)
             if (not same_physical_boundary(old.boundary, self._boundary_hosts[2])
                     or (self._boundary_hosts[2] is not None and actual is None)):
                 raise ValueError("child_objects must reproduce the original root-edge-frame boundary")
             if actual is not None:
                 velocity = actual.tables.get("surface_velocity")
                 if velocity is None:
-                    velocity = boundary_velocity(cfg, pairs[1], actual.host, jets)
+                    velocity = boundary_velocity(cfg, pairs[1], actual.host, jets,
+                                                 _shifted_blocked=self._ddf_shifting)
                 if (not same_physical_boundary(old.boundary, actual.host)
                         or not np.array_equal(old.surface_velocity, velocity)):
                     raise ValueError("child_objects changed boundary ownership or surface velocity")
             new = RemapTopology.from_config(cfg, pairs[0], rectangles,
-                                           child_objects=child_objects, child_jets=jets)
-            if (old.pinned_rectangles != new.pinned_rectangles
-                    or not same_physical_boundary(old.boundary, new.boundary)
+                                           child_objects=child_objects, child_jets=jets,
+                                           _shifted_blocked=self._ddf_shifting)
+            if (not same_physical_boundary(old.boundary, new.boundary)
                     or not np.array_equal(old.surface_velocity, new.surface_velocity)):
-                raise ValueError("remap requires the identical pinned body rectangles and boundary")
+                raise ValueError("remap requires the identical pinned physical boundary and BFL support")
             base_block, fine_block, child_block = self._level_blocks
             # Check sparse block alignment/launch size before any device read/allocation.
             layout = new.child.fine_layout(child_block)
@@ -559,22 +637,28 @@ class FixedPatchStepper:
             with self.ctx.activate():
                 try:
                     self._stream.synchronize()
-                    source = self.download()
+                    source = self._download_raw() if self._ddf_shifting else self.download()
                 except BaseException:
                     self._status = "failed"
                     raise
             count = self._step_count
-            transferred = remap_synchronized(
+            transfer = _remap_stored if self._ddf_shifting else remap_synchronized
+            options = dict(shifted=self._ddf_shifting) if self._ddf_shifting else {}
+            transferred = transfer(
                 old, new, source, base_count=count, level_steps=(count, 2*count, 4*count),
-                phase="postcollision_g")
+                phase="postcollision_g", **options)
             candidate = type(self)(
                 cfg, pairs[0].bounds, context=self.ctx, base_block=base_block,
                 fine_block=fine_block, child_block=child_block,
                 child_rectangles_in_parent=new.child.bounds,
                 child_objects=child_objects, child_jets=jets,
-                observe_raw_exchange=self._observe_raw_exchange)
+                observe_raw_exchange=self._observe_raw_exchange,
+                observe_window_exchange=self._observe_window)
             try:
-                candidate.initialize(*transferred.g)
+                if self._ddf_shifting:
+                    candidate._initialize_raw(_StoredPopulations(transferred.g, transferred.shifted))
+                else:
+                    candidate.initialize(*transferred.g)
                 candidate._step_count = count
             except BaseException:
                 try:
@@ -701,6 +785,10 @@ class FixedPatchStepper:
             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)
+            if self._window_valid and b.window_integral is not None:
+                b.accumulate(b.last_raw, b.window_integral, np.uint32(count),
+                             np.float64(h), np.float64(h),
+                             grid=((count + self._block[0] - 1) // self._block[0], 1), **kw)
         flags = np.uintp(0) if level.boundary is None else level.boundary.flags
         level.kernel(src, dst, np.uint32(level.layout.halo), flags, level.kind, grid=level.grid, **kw)
 
@@ -734,9 +822,57 @@ class FixedPatchStepper:
                     self._exchange_valid = self._observe_raw_exchange
                 self._status = "ready"
             except BaseException:
+                self._window_valid = False
                 self._status = "failed"
                 raise
 
+    def start_exchange_window(self):
+        """Zero the window ledgers and open a window at the current completed step count."""
+        with self._lock, self.ctx.activate():
+            self._require("ready")
+            if not self._observe_window:
+                raise RuntimeError("window exchange observation is disabled")
+            self._window_valid = False
+            try:
+                for b in self._level_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:
+                self._status = "failed"
+                raise
+            self._window_start = self._step_count
+            self._window_valid = True
+
+    def read_exchange_window(self, body_id, *, level=1):
+        """Non-destructive read of [window start, current step count] in base-lattice units."""
+        with self._lock, self.ctx.activate():
+            self._require("ready")
+            if not self._observe_window:
+                raise RuntimeError("window exchange observation is disabled")
+            if not self._window_valid:
+                raise RuntimeError("no valid exchange window; call start_exchange_window()")
+            duration = self._step_count - self._window_start
+            if duration < 1:
+                raise RuntimeError("exchange window requires at least one completed base interval")
+            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]
+            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")
+            rows = np.empty((len(b.tables["cell"]), 4), np.float64)
+            try:
+                cuda.memcpy_dtoh(rows, b.window_integral)
+            except BaseException:
+                self._status = "failed"
+                raise
+            return summarize_exchange(rows, b.tables["body_id"], b.tables["aperture"],
+                                      body_id=body_id, level=level, h=self._exchange_spacing[level],
+                                      time_end=self._step_count, window="window",
+                                      reference_center=self._exchange_centers[level][body_id],
+                                      duration=duration)
+
     def read_raw_exchange(self, body_id, *, level=1, window="last_base_step"):
         """Read raw boundary reaction about its geometric center in the root edge frame.
 
@@ -773,9 +909,17 @@ class FixedPatchStepper:
                                       reference_center=self._exchange_centers[level][body_id])
 
     def download(self):
-        """Return synchronized G arrays (one per level), NaN on all non-leaves; no host cache."""
+        """Return physical G (FP64 for shifted storage), NaN on non-leaves."""
+        snapshot = self._download_raw()
+        arrays = tuple(_decode_stored(a, snapshot.shifted) for a in snapshot.values)
+        self._mask_nonleaves(*arrays)
+        return arrays
+
+    def _download_raw(self):
+        """Current owned populations with a representation tag; no add/subtract w."""
         with self._lock, self.ctx.activate():
             self._require("ready")
+            self._check_storage()
             arrays = [np.empty(level.shape if level.layout is None else level.layout.blocked_shape,
                                np.float32) for level in self._levels]
             try:
@@ -786,8 +930,8 @@ class FixedPatchStepper:
                 raise
             arrays = [a if level.layout is None else level.layout.to_dense(a)
                       for a, level in zip(arrays, self._levels)]
-            self._mask_nonleaves(*arrays)
-            return tuple(arrays)
+            self._mask_nonstate(*arrays)
+            return _StoredPopulations(tuple(arrays), self._ddf_shifting)
 
     def close(self):
         """Release owned resources; retain them for retry if context activation fails."""
@@ -807,6 +951,7 @@ class FixedPatchStepper:
                     if b is not None:
                         b.flags = b.device = b.kernel = None
                         b.last_raw = b.interval_integral = b.accumulate = None
+                        b.window_integral = None
                 self._levels.clear()
                 self._pairs = []
                 if getattr(self, "_fine_curved", None) is not None:

src/CelerisLab/lbm/refinement_tables.py

@@ -85,9 +85,10 @@ class _BlockedBoundary:
     last_raw: object = None
     interval_integral: object = None
     accumulate: object = None
+    window_integral: object = None  # optional observation-window FP64 ledger
 
 
-def _prepare_blocked_boundary(boundary, layout, cfg, jets):
+def _prepare_blocked_boundary(boundary, layout, cfg, jets, *, _shifted_blocked=False):
     if boundary is None or layout is None:
         return None
     fluid_idx, dirs, q, body_id, rx, ry, _, fallback = boundary.curved_host
@@ -102,6 +103,6 @@ def _prepare_blocked_boundary(boundary, layout, cfg, jets):
         cfg.validate()
         tables["surface_velocity"], tables["aperture"] = surface_link_velocity(
             boundary.objects, cfg, boundary.flags, fluid_idx, dirs, body_id, rx, ry,
-            jets=jets, hits=boundary.curved_hits)
+            jets=jets, hits=boundary.curved_hits, _shifted_blocked=_shifted_blocked)
     return _BlockedBoundary(boundary, tables, jets, actuated,
                             jet_link_shapes(jets, boundary.curved_hits) if jets else None)

src/CelerisLab/lbm/synchronized_remap.py

@@ -1,7 +1,7 @@
 """CPU-only whole-volume G transfer, not a runtime adaptive solver.
 
 Periodic three-level grids with fixed L0/L1 topology are supported, including
-stationary circles pinned inside unchanged L2 rectangles. The caller
+stationary circles with unchanged physical boundary/BFL support. The caller
 certifies that current owned leaves are quiescent,
 unshifted D2Q9 postcollision G at the supplied completed base interval.
 No device buffers, device boundary state, ghosts, collision or time advance
@@ -16,17 +16,19 @@ from .refinement import (ACTIVE, FixedPatch2D, FixedPatchSet2D, FineBoundary2D,
                          nest_patch_set, prepare_fine_boundary)
 from .descriptors import SOLID
 from ..body.jets import surface_link_velocity
+from .population_storage import _StoredPopulations
 
 
 @dataclass(frozen=True, eq=False)
 class RemapTopology:
-    """Validated fixed parent and separated L2 set; use ``from_config``.
+    """Validated fixed parent and L2 union; use ``from_config``.
 
     Child rectangles are in the parent's allocated L1 coordinates, including
     its window offset/halo when the parent is a scalar patch. Dense masks are
     rebuilt from geometry, never inferred from finite values in input arrays.
-    ``pinned_rectangles`` canonically lists each body-containing rectangle once;
-    fluid-only layouts have no pins.
+    ``pinned_rectangles`` is legacy metadata for single-rectangle containment.
+    Eligibility pins physical boundary/stencil data, not entire rectangles or
+    connected fluid regions; each topology independently validates transfer closure.
     """
     parent: object
     child: FixedPatchSet2D
@@ -37,7 +39,7 @@ class RemapTopology:
 
     @classmethod
     def from_config(cls, cfg, parent, child_rectangles, *, coarse_boundary="periodic", objects=(),
-                    child_objects=(), child_jets=()):
+                    child_objects=(), child_jets=(), _shifted_blocked=False):
         if coarse_boundary != "periodic" or tuple(objects) or cfg.nz != 1:
             raise ValueError("remap requires periodic 2D state; use child_objects for a pinned L2 body")
         if type(parent) not in (FixedPatch2D, FixedPatchSet2D):
@@ -49,7 +51,7 @@ class RemapTopology:
                            viscosity=2 * cfg.viscosity)
         boundary = prepare_fine_boundary(
             fine_cfg, child, child_objects, jets=child_jets,
-            parent_origin=parent.fine_origin, parent_spacing=.5)
+            parent_origin=parent.fine_origin, parent_spacing=.5, _shifted_blocked=_shifted_blocked)
         pinned, velocity = (), None
         child_leaf = child.fine_kind == ACTIVE
         if boundary is not None:
@@ -59,11 +61,14 @@ class RemapTopology:
             for obj in child_objects:
                 cx, cy = (np.asarray(obj.geometry.center) - parent.fine_origin) * 2
                 radius = obj.geometry._radius * 2
-                containing.add(next(r for r in child.bounds if
-                                    r[0] <= cx-radius and cx+radius <= r[2] and
-                                    r[1] <= cy-radius and cy+radius <= r[3]))
+                match = next((r for r in child.bounds if
+                              r[0] <= cx-radius and cx+radius <= r[2] and
+                              r[1] <= cy-radius and cy+radius <= r[3]), None)
+                if match is not None:
+                    containing.add(match)
             pinned = tuple(sorted(containing))
-            velocity = boundary_velocity(cfg, child, boundary, child_jets)
+            velocity = boundary_velocity(cfg, child, boundary, child_jets,
+                                         _shifted_blocked=_shifted_blocked)
             if velocity is not None:
                 velocity.setflags(write=False)
         masks = (~parent.covered, middle_kind == ACTIVE, child_leaf)
@@ -72,7 +77,7 @@ class RemapTopology:
         return cls(parent, child, masks, boundary, pinned, velocity)
 
 
-def boundary_velocity(cfg, child, boundary, jets=()):
+def boundary_velocity(cfg, child, boundary, jets=(), *, _shifted_blocked=False):
     """Physical cut-link order; never blocked addresses or stale Jet snapshots."""
     local_cfg = replace(cfg, nx=child.fine_shape[1], ny=child.fine_shape[0],
                         viscosity=4 * cfg.viscosity)
@@ -80,7 +85,7 @@ def boundary_velocity(cfg, child, boundary, jets=()):
     fluid, dirs, _, ids, rx, ry, _, _ = boundary.curved_host
     velocity, _ = surface_link_velocity(
         boundary.objects, local_cfg, boundary.flags, fluid, dirs, ids, rx, ry,
-        jets=jets, hits=boundary.curved_hits)
+        jets=jets, hits=boundary.curved_hits, _shifted_blocked=_shifted_blocked)
     return velocity
 
 
@@ -115,6 +120,7 @@ class RemapCandidate:
     phase: str
     origin: tuple
     spacing: float
+    shifted: bool = False
 
 
 def _same_parent(a, b):
@@ -123,7 +129,7 @@ def _same_parent(a, b):
             and a.tau_coarse == b.tau_coarse and a.tau_fine == b.tau_fine)
 
 
-def _read_leaves(topology, arrays, layouts):
+def _read_leaves(topology, arrays, layouts, *, shifted=False):
     """Unpack only owned cores; nonleaf values are not evidence of validity."""
     if len(arrays) != 3 or len(layouts) != 3:
         raise ValueError("expected three G arrays and three layouts")
@@ -137,8 +143,10 @@ def _read_leaves(topology, arrays, layouts):
             a = layout.to_dense(a)
         if a.shape != (9, *leaf.shape):
             raise ValueError("G must have Q-major (9, ny, nx) shape")
+        if shifted and a.dtype != np.float32:
+            raise ValueError("shifted stored populations must be FP32")
         values = a[:, leaf]
-        if not np.isfinite(values).all() or (values.sum(axis=0, dtype=np.float64) <= 0).any():
+        if not np.isfinite(values).all() or (values.sum(axis=0, dtype=np.float64) + int(shifted) <= 0).any():
             raise ValueError("required leaf donors must be finite with positive density")
         result.append(a)
     if len({a.dtype for a in result}) != 1:
@@ -163,14 +171,27 @@ def remap_synchronized(old, new, old_g, *, base_count, level_steps, phase,
     the source dtype. No equilibrium reconstruction or nonequilibrium scaling.
     All destination nonleaves are NaN. Inputs and their write flags are intact.
     """
+    return _remap(old, new, old_g, base_count=base_count, level_steps=level_steps,
+                  phase=phase, origin=origin, spacing=spacing, layouts=layouts)
+
+
+def _remap_stored(old, new, snapshot, *, shifted, **kwargs):
+    """Private same-representation transfer; public remap remains physical."""
+    if (not isinstance(snapshot, _StoredPopulations) or type(shifted) is not bool
+            or type(snapshot.shifted) is not bool or snapshot.shifted != shifted):
+        raise ValueError("stored population representation mismatch")
+    return _remap(old, new, snapshot.values, shifted=shifted, **kwargs)
+
+
+def _remap(old, new, old_g, *, base_count, level_steps, phase,
+           origin=(0., 0.), spacing=1., layouts=None, shifted=False):
     if (not isinstance(old, RemapTopology) or not isinstance(new, RemapTopology)
             or not _same_parent(old.parent, new.parent)
             or old.child.tau_fine != new.child.tau_fine):
         raise ValueError("remap requires unchanged L0/L1 geometry and relaxation")
-    if (old.pinned_rectangles != new.pinned_rectangles
-            or not same_physical_boundary(old.boundary, new.boundary)
+    if (not same_physical_boundary(old.boundary, new.boundary)
             or not np.array_equal(old.surface_velocity, new.surface_velocity)):
-        raise ValueError("remap requires the identical pinned body rectangles and physical boundary")
+        raise ValueError("remap requires the identical pinned physical boundary and BFL support")
     if not isinstance(level_steps, (tuple, list)):
         raise ValueError("level_steps must be three completed integer counts")
     counts = (base_count, *level_steps)
@@ -187,7 +208,7 @@ def remap_synchronized(old, new, old_g, *, base_count, level_steps, phase,
     if (frame.shape != (2,) or not np.isfinite(frame).all()
             or not np.isscalar(spacing) or not np.isfinite(spacing) or spacing <= 0):
         raise ValueError("frame needs a finite (x,y) origin and positive spacing")
-    source = _read_leaves(old, old_g, (None,)*3 if layouts is None else layouts)
+    source = _read_leaves(old, old_g, (None,)*3 if layouts is None else layouts, shifted=shifted)
     candidate = [np.full((9, *mask.shape), np.nan, dtype=source[0].dtype)
                  for mask in new.leaf_masks]
     for level, (before, after) in enumerate(zip(old.leaf_masks, new.leaf_masks)):
@@ -208,6 +229,6 @@ def remap_synchronized(old, new, old_g, *, base_count, level_steps, phase,
     # Scale first to avoid an overflowing intermediate sum for finite FP64 G.
     mean = sum(source[2][:, cy, cx].astype(np.float64) * .25 for cy, cx in children)
     candidate[1][:, y, x] = mean
-    _read_leaves(new, candidate, (None,)*3)
+    _read_leaves(new, candidate, (None,)*3, shifted=shifted)
     return RemapCandidate(new, tuple(candidate), int(base_count), tuple(map(int, level_steps)),
-                          phase, tuple(map(float, frame)), float(spacing))
+                          phase, tuple(map(float, frame)), float(spacing), shifted)

tests/unit/cpp/channel_root.cpp

@@ -0,0 +1,84 @@
+// Actual production step kernels under a CPU launch shim. No CUDA claim.
+#include <cmath>
+#include <cstdio>
+#include <cstdint>
+#include <vector>
+#define __device__
+#define __constant__
+#define __global__
+#define __forceinline__ inline
+#define __launch_bounds__(...)
+#define LATTICE_D2Q9 0
+#define LATTICE_D3Q19 1
+#define LATTICE_MODEL 0
+#define NQ 9
+#define DIM 2
+#define NZ 1
+#define NT 1
+#define COLLISION_MODEL 0
+#define USE_LES 0
+#define RHO 1.f
+#define U0 .025f
+#define OMEGA_COLLISION_MIN .01f
+#define OMEGA_COLLISION_MAX 1.99f
+#define Y_WALL_BC 0
+#define OUTLET_MODE 0
+struct Dim {unsigned int x=0,y=0,z=0;} threadIdx,blockIdx,blockDim;
+inline float __fdiv_rn(float a,float b){return a/b;}
+#include "core/precision.cuh"
+#include "core/descriptors.cuh"
+#include "core/layout.cuh"
+#include "core/params.cuh"
+#include "core/flags.cuh"
+#include "operators/macro.cuh"
+#include "operators/collision_srt.cuh"
+#include "streaming/pull_double_buffer.cuh"
+#if UNIFORM
+#include "operators/forcing_guo.cuh"
+#include "operators/helpers.cuh"
+#include "boundary/bounce_back.cuh"
+#include "boundary/inlet_outlet.cuh"
+#include "step/one_step_double.cu"
+#else
+constexpr unsigned char KIND_OUTER=2;
+#include "step/block_channel.cu"
+#endif
+int main(){
+    constexpr int gx=24,gy=20,n=gx*gy;
+    std::vector<float> state(9*n),next(9*n);
+    for(int q=0;q<9;q++)for(int y=0;y<gy;y++)for(int x=0;x<gx;x++)
+        state[q*n+y*gx+x]=d_w[q]*(1.f+.001f*std::sin(float(x+3*y+q)))-USE_DDF_SHIFTING*d_w[q];
+    d_params={1.25f,0,0,0,0};
+#if UNIFORM
+    std::vector<uint16_t> flags(n,FLAG_FLUID);
+    for(int y=0;y<gy;y++)for(int x=0;x<gx;x++){
+        if(y==0||y==gy-1) flags[y*gx+x]=FLAG_SOLID|FLAG_BC_WALL;
+        else if(x==0) flags[y*gx+x]=FLAG_SOLID|FLAG_BC_INLET;
+        else if(x==gx-1) flags[y*gx+x]=FLAG_SOLID|FLAG_BC_OUTLET;
+    }
+#else
+    constexpr int bx=NX-2,by=NY-2,nbx=gx/bx,nby=gy/by;
+    std::vector<float> input(nbx*nby*9*TOTAL_CELLS),output(input.size());
+    d_block_channel={gx,gy,U0};
+#endif
+    for(int step=0;step<5;step++){
+#if UNIFORM
+        for(int y=0;y<gy;y++)for(int x=0;x<gx;x++){
+            blockIdx={unsigned(x),unsigned(y),0};OneStep(flags.data(),state.data(),next.data());
+        }
+#else
+        // Populate cores and every wrapped halo from the same dense source.
+        for(int b=0;b<nbx*nby;b++)for(int y=0;y<NY;y++)for(int x=0;x<NX;x++)for(int q=0;q<9;q++){
+            int xx=((b%nbx)*bx+x-1+gx)%gx, yy=((b/nbx)*by+y-1+gy)%gy;
+            input[b*9*TOTAL_CELLS+q*TOTAL_CELLS+y*NX+x]=state[q*n+yy*gx+xx];
+        }
+        for(int b=0;b<nbx*nby;b++)for(int y=1;y<NY-1;y++)for(int x=1;x<NX-1;x++){
+            blockIdx={unsigned(x),unsigned(y),unsigned(b)};
+            BlockChannelStep(input.data(),output.data(),1,nullptr,nullptr);
+            for(int q=0;q<9;q++)next[q*n+((b/nbx)*by+y-1)*gx+(b%nbx)*bx+x-1]=output[b*9*TOTAL_CELLS+q*TOTAL_CELLS+y*NX+x];
+        }
+#endif
+        state.swap(next);
+    }
+    for(float f:state)printf("%.9g\n",f);
+}

tests/unit/cpp/collision_regression.cpp

@@ -13,6 +13,8 @@
 #define OMEGA_COLLISION_MAX 1.99f
 struct Params {float omega, fx, fy, fz;} d_params;
 #include "core/descriptors.cuh"
+// Host round-to-nearest shim; CUDA intrinsic behavior is tested on-device.
+inline float __fdiv_rn(float a, float b) { return a / b; }
 #include "operators/macro.cuh"
 #include "operators/collision_srt.cuh"
 #include "operators/collision_trt.cuh"

tests/unit/cpp/semantic_regression.cpp

@@ -34,6 +34,8 @@ inline int obs_torque_index(int,int){return 0;}
 #include "core/descriptors.cuh"
 #include "core/precision.cuh"
 #include "core/layout.cuh"
+// Host round-to-nearest shim; CUDA intrinsic behavior is tested on-device.
+inline float __fdiv_rn(float a, float b) { return a / b; }
 #include "operators/macro.cuh"
 #include "operators/collision_srt.cuh"
 #include "operators/collision_trt.cuh"

tests/unit/cpp/shifted_block_operators.cpp

@@ -0,0 +1,45 @@
+#include <cmath>
+#include <cstdio>
+#include <initializer_list>
+#define __device__
+#define __constant__
+#define __forceinline__ inline
+#define LATTICE_D2Q9 0
+#define LATTICE_MODEL 0
+#define NQ 9
+#define DIM 2
+#define USE_DDF_SHIFTING 1
+#define WALL_NORMAL_VELOCITY 0.f
+#define WALL_SLOT_WIDTH 0
+#define NX 16
+#include "core/descriptors.cuh"
+// Host round-to-nearest shim; CUDA intrinsic behavior is tested on-device.
+inline float __fdiv_rn(float a, float b) { return a / b; }
+#include "operators/macro.cuh"
+#include "operators/collision_srt.cuh"
+#include "boundary/planar_velocity.cuh"
+void emit(const float* f) {for(int q=0;q<9;q++) printf("%.9g ",f[q]);puts("");}
+int main() {
+    for(float delta : {0.f, 0x1p-30f, -0x1p-30f, .0002f})
+    for(float u : {0.f,.005f}) {
+        float eq[9]; compute_feq(1.f+delta,u,-.2f*u,eq,delta);
+        printf("E %.9g %.9g %.9g\n",delta,u,-.2f*u);emit(eq);
+    }
+    for(float delta : {0.f, 0x1p-30f, -0x1p-30f, .04f, -.01f})
+    for(float jx : {0.f,.005f,-.02f}) {
+        const float jy = -.2f*jx, rho = 1.f+delta;
+        const float ux = jx/rho, uy = jy/rho;
+        float eq[9];compute_feq_from_momentum(delta,jx,jy,ux,uy,eq);
+        printf("J %.9g %.9g %.9g %.9g %.9g\n",delta,jx,jy,ux,uy);emit(eq);
+    }
+    for(int mode : {0,1,2}) for(float rho : {1.f,1.04f}) {
+        float f[9];const float u=.007f,v=.003f;
+        compute_feq(rho,u,v,f,rho-1.f);
+        for(int q=0;q<9;q++)f[q]+=.00001f*(q-4);
+        printf("B %d %.9g %.9g %.9g\n",mode,rho,u,v);emit(f);
+        if(mode==0)reconstruct_bottom_velocity(f,u,v);
+        else if(mode==1)reconstruct_top_velocity(f,u,v);
+        else reconstruct_top_reservoir(f,rho);
+        emit(f);
+    }
+}

tests/unit/test_channel_parity_diagnostic.py

@@ -0,0 +1,47 @@
+"""CPU-only traceback capture checks; no CUDA or GPU test import."""
+import importlib.util
+from pathlib import Path
+from types import SimpleNamespace
+import numpy as np
+
+PATH = Path(__file__).resolve().parents[1] / 'validation/multigrid/channel_parity_diagnostic/discover.py'
+spec = importlib.util.spec_from_file_location('channel_discovery', PATH)
+discovery = importlib.util.module_from_spec(spec)
+spec.loader.exec_module(discovery)
+
+
+def test_captures_exact_original_frame_and_arrays():
+    def test_uniform_matched_root():
+        cfg = SimpleNamespace(velocity=.025)
+        shifted, profile, clamp, damping, block = True, 'parabolic', False, .8, (4, 1)
+        expected = np.array([.001], np.float32)
+        actual = np.nextafter(expected, np.float32(1))
+        np.testing.assert_array_equal(actual, expected)
+    try:
+        test_uniform_matched_root()
+    except AssertionError as exc:
+        config, expected, actual = discovery.capture_case(exc)
+    assert config['shifted'] == 1 and config['profile'] == 'parabolic'
+    assert config['clamp'] == 0 and config['damping'] == .8
+    assert config['block'] == [4, 1] and config['steps'] == 8
+    np.testing.assert_array_equal(actual, np.nextafter(expected, np.float32(1)))
+    assert expected.flags.owndata and actual.flags.owndata
+
+
+def test_other_assertion_does_not_guess_case():
+    def test_uniform_matched_root():
+        cfg = SimpleNamespace(velocity=.025)
+        shifted, profile, clamp, damping, block = False, 'uniform', True, .5, (20, 24)
+        expected = actual = np.array([np.nan], np.float32)
+        assert np.isfinite(expected).all()
+    try:
+        test_uniform_matched_root()
+    except AssertionError as exc:
+        assert discovery.capture_case(exc) is None
+
+
+def test_unrelated_array_assertion_does_not_guess_case():
+    try:
+        np.testing.assert_array_equal([1], [2])
+    except AssertionError as exc:
+        assert discovery.capture_case(exc) is None

tests/unit/test_channel_root.py

@@ -0,0 +1,101 @@
+"""Root-channel host contract and actual-header CPU parity (no CUDA evidence)."""
+from dataclasses import replace
+import numpy as np
+import pytest
+
+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 .test_numerical_operators import compile_run
+
+
+@pytest.mark.parametrize("shift", [0, 1])
+@pytest.mark.parametrize("profile", [0, 1])
+@pytest.mark.parametrize("clamp,damping", [(1, ".5f"), (0, ".8f")])
+def test_actual_kernel_uniform_parity(tmp_path, shift, profile, clamp, damping):
+    """Nonuniform ghosts, corners, top wall, global profile and outlet halo donor."""
+    defines = dict(USE_DDF_SHIFTING=shift, INLET_PROFILE=profile,
+                   OUTLET_BACKFLOW_CLAMP=clamp, OUTLET_SRT_NEQ_DAMP=damping)
+    uniform = np.array(compile_run(tmp_path, "channel_root.cpp", dict(
+        defines, UNIFORM=1, NX=24, NY=20)), np.float32)
+    assert np.isfinite(uniform).all()
+    for by, bx in ((20, 24), (5, 6), (4, 1)):
+        blocked = np.array(compile_run(tmp_path, "channel_root.cpp", dict(
+            defines, UNIFORM=0, NX=bx+2, NY=by+2)), np.float32)
+        np.testing.assert_array_equal(blocked, uniform)
+
+
+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]
+    return pair, layout, targets
+
+
+def test_masks_and_minimal_footprint():
+    fluid, bulk = channel_masks((20, 24))
+    assert fluid.sum() == 18*22 and bulk.sum() == 16*22
+    assert fluid[1, 1] and not bulk[1, 1]
+    assert bulk[2, 1] and not fluid[0, 0]
+    # No tile alignment or artificial large padding: closest legal x donor is x=1.
+    validate_channel_transfers(*transfer((3, 4, 19, 16)))
+    for bounds in ((2, 4, 19, 16), (3, 3, 19, 16), (3, 4, 22, 16), (3, 4, 19, 17)):
+        with pytest.raises(ValueError, match="bulk"):
+            validate_channel_transfers(*transfer(bounds))
+
+
+def test_actual_closure_read_conflicts():
+    pair, layout, targets = transfer((6, 6, 18, 14))
+    validate_channel_transfers(pair, layout, targets)
+    for target in (layout.slot(layout.owner[5, -1]-1, 2),
+                   layout.slot(layout.owner[1, 4], 4),
+                   layout.slot(layout.owner[-2, 4], 3)):
+        with pytest.raises(ValueError, match="immutable"):
+            validate_channel_transfers(pair, layout, np.append(targets, target))
+
+
+def test_channel_scope():
+    cfg = LBMConfig(nx=24, ny=20, viscosity=.1)
+    validate_channel_config(cfg, [(4, 6), (4, 4)])
+    with pytest.raises(ValueError, match="all-blocked"):
+        validate_channel_config(cfg, [None, (4, 4)])
+    for change in (dict(inlet_scheme="equilibrium"), dict(outlet_mode="zero_gradient"),
+                   dict(y_wall_bc="free_slip"), dict(wall_slot_width=2)):
+        with pytest.raises(ValueError, match="channel requires"):
+            validate_channel_config(replace(cfg, **change), [(4, 6), (4, 4)])
+
+
+def test_evolved_ghost_snapshot_and_physical_download():
+    from contextlib import nullcontext
+    from unittest.mock import Mock
+    from .test_refinement_three_level_multipatch import host_module
+    host = host_module()
+    storage = {}
+    host.compiler.load_module.return_value.get_global.side_effect = lambda name: (name, 12 if name == "d_block_channel" else 20)
+    host.cuda.mem_alloc.side_effect = lambda size: Mock()
+    host.cuda.memcpy_htod.side_effect = lambda ptr, a: storage.__setitem__(id(ptr), np.array(a, copy=True))
+    host.cuda.memcpy_dtoh.side_effect = lambda a, ptr: np.copyto(a, storage[id(ptr)])
+    context = Mock(activate=nullcontext, sm_arch="sm_70")
+    cfg = LBMConfig(nx=24, ny=20, viscosity=.1, velocity=.025)
+    backend = host.FixedPatchStepper(cfg, (6, 6, 18, 14), context=context,
+                                    coarse_boundary="channel", base_block=(4, 6), fine_block=(4, 4))
+    try:
+        weights = np.array([4/9]+[1/9]*4+[1/36]*4)[:, None, None]
+        arrays = [np.broadcast_to(weights, level.shape).copy() for level in backend._levels]
+        arrays[0][:, 0, 0] *= 1.001
+        backend.initialize(*arrays)
+        snapshot = backend._download_raw()
+        assert np.isfinite(snapshot.values[0][:, 0, :]).all()
+        assert np.isnan(backend.download()[0][:, 0, :]).all()
+        backend._initialize_raw(snapshot)
+        for a, b in zip(snapshot.values, backend._download_raw().values):
+            np.testing.assert_array_equal(a, b)
+        arrays[0][0, 0, 0] = np.nan
+        with pytest.raises(ValueError, match="finite"):
+            backend.initialize(*arrays)
+        assert backend._levels[0].kernel == host.compiler.load_module.return_value.get_function("BlockChannelStep")
+    finally:
+        backend.close()

tests/unit/test_hierarchy.py

@@ -50,11 +50,11 @@ class FourLevels(unittest.TestCase):
         cases = [([L1, [(10, 20, 30, 44)]], "inherited ACTIVE"),
                  ([L1, L2[:1] + [(54, 54, 58, 58)]], "inherited ACTIVE"),
                  ([L1, L2, [(40, 40, 60, 60)]], "donors/receivers"),
-                 ([L1, L2, [(41, 41, 60, 60)]], "donors/receivers"),
-                 ([L1, L2, [(56, 56, 64, 64), (66, 56, 72, 64)]], "closer than")]
+                 ([L1, L2, [(41, 41, 60, 60)]], "donors/receivers")]
         for levels, reason in cases:
             with self.subTest(levels=levels), self.assertRaisesRegex(ValueError, reason):
                 build_hierarchy(CFG, levels)
+        build_hierarchy(CFG, [L1, L2, [(56, 56, 64, 64), (66, 56, 72, 64)]])  # gap 2 closes
         build_hierarchy(CFG, [L1, L2, [(42, 42, 60, 60)]])  # minimal legal margin
 
 

tests/unit/test_jet_circle_cv.py

@@ -0,0 +1,139 @@
+"""CPU-only science, geometry, units, and source tests; no PyCUDA import."""
+
+import sys
+from pathlib import Path
+import unittest
+import numpy as np
+
+ROOT = Path(__file__).resolve().parents[2]
+REPO = ROOT
+sys.path.insert(0, str(ROOT / "tests/validation"))
+sys.path.insert(0, str(REPO / "src"))
+from run_jet_circle_cv import (
+    CX,
+    CY,
+    W,
+    case,
+    state,
+    cv,
+    aperture_probe,
+    measured_aperture_flux,
+    verify,
+)
+
+
+def eq(rho, u, v):
+    cu = CX[:, None, None] * u + CY[:, None, None] * v
+    return W[:, None, None] * rho * (1 + 3 * cu + 4.5 * cu**2 - 1.5 * (u * u + v * v))
+
+
+class Host(unittest.TestCase):
+    def test_descriptor_and_uniform_closed_flux(self):
+        from CelerisLab.lbm.descriptors import D2Q9_EX, D2Q9_EY
+
+        np.testing.assert_array_equal(CX, D2Q9_EX)
+        np.testing.assert_array_equal(CY, D2Q9_EY)
+        rho = np.full((48, 64), 1.01)
+        g = eq(rho, 0.023, -0.017)
+        inventory, flux = cv(g, np.ones((48, 64), bool), (0, 0), (5, 4, 25, 20), 0.0072)
+        np.testing.assert_allclose(
+            inventory, 320 * 1.01 * np.array([1, 0.023, -0.017]), atol=1e-12
+        )
+        np.testing.assert_allclose(flux, 0, atol=1e-14)
+
+    def test_affine_velocity_independent_divergence(self):
+        y, x = np.mgrid[:48, :64]
+        x = (x + 0.5) * 0.5
+        y = (y + 0.5) * 0.5
+        a = 0.001
+        g = eq(np.ones_like(x), a * x, -a * y)
+        _, flux = cv(g, np.ones(x.shape, bool), (0, 0), (5, 4, 25, 20), 0.0072)
+        np.testing.assert_allclose(
+            flux, [0, a * a * 320 * 15, a * a * 320 * 12], atol=1e-13
+        )
+
+    def test_quadratic_viscous_flux_sign_and_units(self):
+        y, x = np.mgrid[:48, :64]
+        y = (y + 0.5) * 0.5
+        a = 0.0001
+        nu = 0.0072
+        g = eq(np.ones_like(y), a * y * y, np.zeros_like(y))
+        _, flux = cv(g, np.ones(y.shape, bool), (0, 0), (5, 4, 25, 20), nu)
+        np.testing.assert_allclose(flux, [0, -2 * nu * a * 320, 0], atol=2e-14)
+
+    def test_geometry_adapter_match_and_apertures(self):
+        from CelerisLab.config import LBMConfig
+        from CelerisLab.body.geometry.circle import CircleGeometry
+        from CelerisLab.body.objects import SimObject
+        from CelerisLab.body.jets import Jet, Sector, Radial, jet_link_shapes
+        from CelerisLab.lbm.refinement import FixedPatch2D, prepare_fine_boundary
+        from CelerisLab.lbm.descriptors import SOLID
+
+        c = case()
+        center = np.array(c["center"])
+        r = c["diameter"] / 2
+        cfg = LBMConfig(nx=c["nx"], ny=c["ny"], viscosity=c["nu"])
+        cfg.validate()
+        patch = FixedPatch2D.from_config(cfg, c["bounds"])
+        obj = SimObject(0, CircleGeometry(*center, r), center=tuple(center), radius=r)
+        jets = [
+            Jet(
+                Sector(tuple(center), a, np.pi / 12, r + 0.1),
+                Radial(tuple(center), inward=False),
+                "uniform",
+                s * c["jet_speed"],
+            )
+            for a, s in [(np.pi / 2, 1), (3 * np.pi / 2, -1)]
+        ]
+        host = prepare_fine_boundary(cfg, patch, [obj], jets=jets)
+        shape, ids = jet_link_shapes(jets, host.curved_hits)
+        self.assertTrue((ids == 0).any() and (ids == 1).any() and (ids < 0).any())
+        self.assertGreater(host.curved_host[2].min(), 0.010001)
+        self.assertFalse(np.any(host.curved_host[-1]))
+        links = CircleGeometry(*(2 * center - 0.5), 2 * r).build_cut_links(
+            2 * c["nx"], 2 * c["ny"]
+        )
+        np.testing.assert_array_equal([v.dir for v in links], host.curved_host[1])
+        np.testing.assert_allclose(
+            [v.q for v in links], host.curved_host[2], atol=2e-7, rtol=0
+        )
+        np.testing.assert_allclose(
+            np.array([[v.hit_x, v.hit_y] for v in links]) * 0.5 + 0.25,
+            host.curved_hits,
+            atol=2e-12,
+        )
+        fluid = (host.flags.reshape(patch.fine_shape) & SOLID) == 0
+        g = state((9, *patch.fine_shape), c["u"])
+        g[:, ~fluid] = np.nan
+        for contour in c["contours"]:
+            cv(g, fluid, patch.fine_origin, contour, c["nu"])
+        probe = aperture_probe(fluid, patch.fine_origin, center, r)
+        measured = measured_aperture_flux(g, probe)
+        np.testing.assert_allclose(
+            measured, 0, atol=1e-14
+        )  # x-directed flow has zero flux on symmetric top/bottom arcs
+        print(
+            "GEOMETRY",
+            dict(
+                links=len(links),
+                aperture_counts=[int((ids == i).sum()) for i in [0, 1]],
+                q_min=float(host.curved_host[2].min()),
+                q_max=float(host.curved_host[2].max()),
+                fallbacks=int(np.count_nonzero(host.curved_host[-1])),
+                sample_distances=[p[-1] for p in probe],
+            ),
+        )
+
+    def test_diffusive_scaling(self):
+        a, b = case(1), case(2)
+        self.assertEqual(a["re_d"], b["re_d"])
+        self.assertEqual(a["nu"], b["nu"])
+        self.assertEqual(a["final_t_star"], b["final_t_star"])
+        self.assertEqual(a["steps"] * 4, b["steps"])
+
+    def test_source_pin(self):
+        verify(REPO)
+
+
+if __name__ == "__main__":
+    unittest.main()

tests/unit/test_jet_cv_measurements.py

@@ -0,0 +1,182 @@
+"""Manufactured references for continuation instrumentation, no CUDA."""
+
+import sys
+from pathlib import Path
+import tempfile
+import unittest
+import numpy as np
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "validation"))
+from jet_cv_measurements import (
+    sample_steps,
+    factors,
+    ray_geometry,
+    reconstruct_aperture,
+    contour_terms,
+    LINEAR,
+    QUADRATIC,
+    OFFSETS,
+    MeasurementRecorder,
+)
+from run_jet_circle_cv import (
+    case,
+    CX,
+    CY,
+    W,
+    cv,
+    aperture_probe,
+    measured_aperture_flux,
+)
+from analyze_jet_cv_continuation import integrate, analyze
+
+
+class Measurements(unittest.TestCase):
+    def test_common_physical_schedule(self):
+        a = sample_steps(case(1, steps=3200))
+        b = sample_steps(case(2, steps=3200))
+        np.testing.assert_array_equal(a, b / 4)
+        self.assertEqual(len(a), 876)
+        self.assertEqual(a[-1], 3200)
+        self.assertTrue(all(t in a for t in range(0, 3201, 100)))
+        self.assertEqual(len(a[a <= 100]), 101)
+
+    def test_reference_factors(self):
+        f = factors(2)
+        self.assertEqual(f["mass"], 0.25)
+        self.assertEqual(f["impulse"], 0.5)
+        self.assertEqual(f["force"], 2)
+        self.assertEqual(f["mass_rate"], 1)
+
+    def test_extrapolation_polynomials(self):
+        self.assertAlmostEqual(LINEAR @ (2 + 3 * OFFSETS), 2)
+        self.assertAlmostEqual(QUADRATIC @ (2 + 3 * OFFSETS + 4 * OFFSETS**2), 2)
+
+    def test_independent_linear_field_volume_and_mass(self):
+        h = 0.5
+        y, x = np.mgrid[:64, :64]
+        x = (x + 0.5) * h
+        y = (y + 0.5) * h
+        rho = 1 + 0.001 * x
+        vel = np.array([0.01 + 0.0001 * x, -0.02 + 0.0002 * y])
+        center = np.array([16.0, 16.0])
+        r = 5
+        fluid = (x - center[0]) ** 2 + (y - center[1]) ** 2 > r * r
+        geo = ray_geometry(fluid, (0.0, 0.0), center, r)
+        ray, est = reconstruct_aperture(rho, vel, geo)
+        points = center + r * geo["normal"]
+        rr = 1 + 0.001 * points[..., 0]
+        uv = np.stack(
+            (0.01 + 0.0001 * points[..., 0], -0.02 + 0.0002 * points[..., 1]), axis=-1
+        )
+        un = np.sum(uv * geo["normal"], axis=-1)
+        exact = np.stack(
+            (np.sum(un * geo["ds"], axis=1), np.sum(rr * un * geo["ds"], axis=1)),
+            axis=-1,
+        )
+        np.testing.assert_allclose(est, np.broadcast_to(exact, est.shape), atol=1e-15)
+        self.assertGreater(np.max(abs(est[:, :, 0] - est[:, :, 1])), 1e-4)
+
+    def test_split_contour_stress(self):
+        y, x = np.mgrid[:48, :64]
+        y = (y + 0.5) * 0.5
+        rho = np.ones_like(y)
+        v = np.array([1e-4 * y * y, np.zeros_like(y)])
+        terms = contour_terms(rho, v, (0, 0), (5, 4, 25, 20), 0.0072)
+        np.testing.assert_allclose(terms[:2], 0, atol=1e-15)
+        np.testing.assert_allclose(terms[2], [-2 * 0.0072 * 1e-4 * 320, 0], atol=1e-15)
+
+    def test_recorder_snapshot_and_reconstructible_arrays(self):
+        h = 0.5
+        y, x = np.mgrid[:64, :64]
+        x = (x + 0.5) * h
+        y = (y + 0.5) * h
+        fluid = (x - 16) ** 2 + (y - 16) ** 2 > 25
+        fine_leaf = fluid.copy()
+        fine_leaf[:, :4] = False
+        root_leaf = np.zeros((32, 32), bool)
+        root_leaf[:, :2] = True
+
+        def equilibrium(ny, nx):
+            cu = CX[:, None, None] * 0.01
+            return np.broadcast_to(
+                W[:, None, None] * (1 + 3 * cu + 4.5 * cu * cu - 1.5 * 0.01**2),
+                (9, ny, nx),
+            ).copy()
+
+        ug = equilibrium(64, 64)
+        levels = [equilibrium(32, 32), ug.copy()]
+        c = dict(
+            scale=1,
+            center=[16, 16],
+            diameter=10,
+            u=0.01,
+            nu=0.0072,
+            contours=[[9, 9, 23, 23], [7, 7, 25, 25]],
+        )
+        data = {
+            name: [cv(ug, mask, (0, 0), r, c["nu"]) for r in c["contours"]]
+            for name, mask in [("uniform_adapter", fluid), ("amr_api", fine_leaf)]
+        }
+        with tempfile.TemporaryDirectory() as tmp:
+            rec = MeasurementRecorder(
+                c, tmp, fluid, [root_leaf, fine_leaf], (0, 0), CX, CY
+            )
+            rec.capture(0, ug, levels, data, {k: np.zeros(2) for k in data}, 0.0)
+            rec.records[-1]["raw_cumulative"] = np.zeros((2, 5, 4))
+            rec.flush()
+            a = np.load(Path(tmp) / "samples_s1.npz")
+            self.assertEqual(a["arc_ray_values"].shape, (1, 2, 3, 2, 32, 3))
+            np.testing.assert_allclose(
+                a["global_mass_reference"][0, 0], a["global_mass_reference"][0, 1]
+            )
+            snap = np.load(Path(tmp) / "fields_s1_tref0.npz")
+            np.testing.assert_allclose(
+                snap["uniform_velocity_reference"][0][fluid], 0.01, rtol=1e-7
+            )
+
+    def test_offline_mass_force_normalization(self):
+        import json
+
+        c = case(2, steps=3200)
+        s = 2
+        t = np.arange(0, 3201, 100)
+        n = len(t)
+        inv = np.zeros((n, 2, 2, 3))
+        inv[:, :, :, 0] = 100 * s * s
+        inv[:, :, :, 1] = -t[:, None, None] * 0.002 * s
+        raw = np.zeros((n, 2, 5, 4))
+        raw[:, :, 2, 0] = t[:, None] * 0.002 * s
+        with tempfile.TemporaryDirectory() as tmp:
+            path = Path(tmp)
+            (path / "samples_s2.json").write_text(
+                json.dumps(dict(case=c, sample_file="samples_s2.npz"))
+            )
+            np.savez(
+                path / "samples_s2.npz",
+                t_reference=t,
+                cv_inventory=inv,
+                cv_flux=np.zeros_like(inv),
+                raw_cumulative=raw,
+                arc_estimates=np.zeros((n, 2, 2, 2, 2)),
+                global_mass_reference=np.ones((n, 2)) * 100,
+                level_stats=np.broadcast_to([1, 1, 0, 0.03, 100], (n, 3, 5)),
+                velocity_error=np.zeros(n),
+            )
+            result = analyze(path)
+            self.assertEqual(len(result["windows"]), 32)
+            np.testing.assert_allclose(
+                result["windows"][0]["normalized_cv_residual"], 0, atol=1e-14
+            )
+            f0 = 0.5 * 0.03**2 * 24.084
+            np.testing.assert_allclose(
+                result["windows"][0]["raw_reaction_coefficient"], [[0.002 / f0, 0]] * 2
+            )
+
+    def test_nonuniform_trapezoidal_time_units(self):
+        t = np.array([0.0, 1.0, 2.0, 4.0, 8.0])
+        y = np.stack((3 * t + 2, 2 * t - 7), axis=-1)
+        np.testing.assert_allclose(integrate(t, y), [112, 8], atol=1e-15)
+
+
+if __name__ == "__main__":
+    unittest.main()

tests/unit/test_les_flow_oracle.py

@@ -0,0 +1,70 @@
+import sys
+from pathlib import Path
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "validation"))
+import json
+import unittest
+import numpy as np
+from les_flow_oracle import H, U, NU, TIMES, solve, relative, amplitude
+
+
+class OracleTests(unittest.TestCase):
+    def test_linear_exact(self):
+        y, states = solve(0, n=256)
+        for t, u in states.items():
+            exact = U * np.sin(np.pi * y / H) * np.exp(-NU * (np.pi / H) ** 2 * t)
+            self.assertLess(relative(u, exact), 1e-5)
+
+    def test_nonlinear_refinement_symmetry_and_dissipation(self):
+        y, states = solve(1, n=256)
+        yc, coarse = solve(1, n=128)
+        energy = U**2 / 2
+        for t, u in states.items():
+            self.assertTrue(np.all(np.isfinite(u)))
+            self.assertGreaterEqual(float(u.min()), 0)
+            self.assertLessEqual(float(u.max()), U)
+            np.testing.assert_allclose(u, u[::-1], atol=1e-14)
+            self.assertLess(float(np.mean(u * u)), energy)
+            energy = float(np.mean(u * u))
+            self.assertLess(relative(np.interp(yc, y, u), coarse[t]), 8e-5)
+
+    def test_filter_is_independent_of_oracle_grid(self):
+        y, a = solve(1, n=128, times=(TIMES[-1],))
+        _, b = solve(0.5, n=128, times=(TIMES[-1],), delta=2)
+        np.testing.assert_array_equal(a[TIMES[-1]], b[TIMES[-1]])
+
+    def test_fixed_grid_timestep_halving(self):
+        errors = []
+        for cs in (0.0, 0.16, 1.0):
+            _, full = solve(cs, n=256, cfl=0.4)
+            _, half = solve(cs, n=256, cfl=0.2)
+            errors.extend(relative(full[t], half[t]) for t in TIMES)
+        self.assertLess(max(errors), 1e-8)
+        print(
+            json.dumps({"fixed_grid_timestep_half_max_relative": max(errors)}),
+            flush=True,
+        )
+
+    def test_signal_and_wrong_convention_separate(self):
+        y, std = solve(1, n=256, times=(TIMES[-1],))
+        _, wrong = solve(2 ** (-0.25), n=256, times=(TIMES[-1],))
+        t = TIMES[-1]
+        base = U * np.exp(-NU * (np.pi / H) ** 2 * t)
+        signal = (base - amplitude(y, std[t])) / base
+        separation = relative(wrong[t], std[t])
+        print(
+            json.dumps(
+                {
+                    "final_time": t,
+                    "standard_sgs_amplitude_signal": signal,
+                    "wrong_convention_profile_separation": separation,
+                }
+            ),
+            flush=True,
+        )
+        self.assertGreater(signal, 0.04)
+        self.assertGreater(separation, 0.01)
+
+
+if __name__ == "__main__":
+    unittest.main()

tests/unit/test_les_flow_runner.py

@@ -0,0 +1,107 @@
+"""Host-only entrypoint/source-guard checks; never imports PyCUDA."""
+
+import hashlib
+import json
+from pathlib import Path
+import shutil
+import subprocess
+import sys
+from tempfile import TemporaryDirectory
+import unittest
+import numpy as np
+from unittest.mock import patch
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "validation"))
+from run_uniform_les_flow import source_guard, decode_macros, HERE, HEADER
+
+ROOT = HERE.parents[1]
+
+
+class GuardTests(unittest.TestCase):
+    def test_decoded_density_and_velocity_guards(self):
+        w = np.array([4 / 9] + [1 / 9] * 4 + [1 / 36] * 4)
+        cx = np.array([0, 1, -1, 0, 0, 1, -1, -1, 1])
+        cy = np.array([0, 0, 0, 1, -1, 1, -1, 1, -1])
+        data = np.zeros((9, 4, 2))
+        profile, mass, vy, spread = decode_macros(data, w, cx, cy)
+        np.testing.assert_allclose(profile, 0, atol=1e-16)
+        self.assertAlmostEqual(mass, 1)
+        for scale in (1, 2):
+            invalid = np.broadcast_to(-scale * w[:, None, None], data.shape).copy()
+            with self.assertRaisesRegex(RuntimeError, "density"):
+                decode_macros(invalid, w, cx, cy)
+        invalid = data.copy()
+        invalid[0, 1, 0] = np.nan
+        with self.assertRaisesRegex(RuntimeError, "populations"):
+            decode_macros(invalid, w, cx, cy)
+        # Finite density but overflowing moment must not produce an accepted NaN/Inf norm.
+        invalid = data.copy()
+        invalid[1, :, :] = 1e308
+        invalid[2, :, :] = -1e308
+        with self.assertRaisesRegex(RuntimeError, "velocity/profile"):
+            decode_macros(invalid, w, cx, cy)
+
+    def test_baseline_pins(self):
+        self.assertEqual(
+            source_guard(ROOT, "baseline")["base_commit"],
+            "cd3a657f9d3833e812c3cca173525f165b48710a",
+        )
+
+    def test_corrected_and_tamper_guard(self):
+        pins = json.loads((HERE / "data/les_flow/source_pins.json").read_text())
+        with TemporaryDirectory() as directory:
+            root = Path(directory)
+            for name in pins["sources"]:
+                dest = root / name
+                dest.parent.mkdir(parents=True, exist_ok=True)
+                shutil.copyfile(ROOT / name, dest)
+            # Rebuild the reviewed candidate from actual pinned headers, not fixtures.
+            candidate = root / "header_candidate"
+            subprocess.run(
+                [
+                    sys.executable,
+                    str(HERE / "les/prepare.py"),
+                    "--repo",
+                    str(ROOT),
+                    "--output-dir",
+                    str(candidate),
+                ],
+                check=True,
+            )
+            shutil.copyfile(candidate / "patched" / HEADER, root / HEADER)
+            with patch(
+                "run_uniform_les_flow.subprocess.check_output",
+                return_value="isolated-fixture\n",
+            ):
+                source_guard(root, "corrected")
+                with self.assertRaisesRegex(RuntimeError, "Source pin mismatch"):
+                    source_guard(root, "baseline")
+                target = root / "src/CelerisLab/lbm/kernels/step/one_step_double.cu"
+                target.write_bytes(target.read_bytes() + b"\n// tamper\n")
+                with self.assertRaisesRegex(RuntimeError, "Source pin mismatch"):
+                    source_guard(root, "corrected")
+
+    def test_gpu_requires_opt_in_before_import(self):
+        with TemporaryDirectory() as directory:
+            result = subprocess.run(
+                [
+                    sys.executable,
+                    str(HERE / "run_uniform_les_flow.py"),
+                    "--root",
+                    str(ROOT),
+                    "--variant",
+                    "baseline",
+                    "--output",
+                    str(Path(directory) / "out.json"),
+                ],
+                capture_output=True,
+                text=True,
+            )
+            self.assertEqual(result.returncode, 2)
+            self.assertIn("Explicit --run-gpu is required", result.stderr)
+            self.assertNotIn("pycuda", result.stderr)
+            self.assertFalse((Path(directory) / "out.json").exists())
+
+
+if __name__ == "__main__":
+    unittest.main()

tests/unit/test_multiple_pinned_remap.py

@@ -81,14 +81,9 @@ class MultiplePinnedRemap(unittest.TestCase):
         self.assertEqual(old.pinned_rectangles, PINS)
         self.check_transfer(old, new, fields(old), 2e-13)
 
-    def test_direct_helper_rejects_each_changed_pin_and_body_order(self):
+    def test_direct_helper_rejects_changed_body_order(self):
         old = mesh(PINS+FLUID, objects())
         candidates = [mesh(PINS+FLUID, objects()[::-1])]
-        for index in range(2):
-            changed = list(PINS+FLUID)
-            x0, y0, x1, y1 = changed[index]
-            changed[index] = (x0-4, y0, x1, y1)
-            candidates.append(mesh(changed, objects()))
         for new in candidates:
             # Invalid donors ensure boundary validation precedes array reads.
             with self.assertRaisesRegex(ValueError, 'pinned'):
@@ -138,6 +133,7 @@ class MultiplePinnedRemap(unittest.TestCase):
             changed = list(PINS+FLUID)
             x0, y0, x1, y1 = changed[index]
             changed[index] = (x0-4, y0, x1, y1)
+            changed[index] = (x0+8, y0, x1, y1)  # removes actual body support
             cases.append((changed, bodies))
             cases.append((tuple(r for i, r in enumerate(PINS+FLUID) if i != index), bodies))
         start = len(allocations)

tests/unit/test_pinned_body_remap.py

@@ -83,7 +83,7 @@ class PinnedBodyRemap(unittest.TestCase):
         run,module,allocations,storage,objects=ready()
         start=len(allocations)
         cases=[(B,()),(B,run.child_boundary.objects),(B,(body(center=(12.1,14.)),)),
-               (((20,16,36,40),B[1]),objects),(((16,20,32,40),B[1]),objects),
+               (((24,16,36,40),B[1]),objects),(((16,28,32,40),B[1]),objects),
                ((B[1],),objects)]
         module.cuda.memcpy_dtoh.reset_mock()
         for rects,objs in cases:
@@ -124,7 +124,9 @@ class PinnedBodyRemap(unittest.TestCase):
         for a in allocations[len(saved):]: a.free.assert_called_once()
         run.close()
 
-    def test_direct_helper_rejects_changed_pinned_rectangle(self):
+    def test_direct_helper_allows_changed_rectangle_with_same_boundary(self):
         old=mesh(A,(body(),))
         new=mesh(((12,16,32,40),A[1]),(body(),))
-        with self.assertRaisesRegex(ValueError,'pinned'): remap(old,new,fields(old))
+        out = remap(old,new,fields(old))
+        self.assertTrue(same_physical_boundary(old.boundary,new.boundary))
+        self.assertIs(out.topology,new)

tests/unit/test_raw_exchange_window.py

@@ -0,0 +1,45 @@
+"""Host reduction contract for the optional observation-window exchange ledger."""
+import unittest
+
+import numpy as np
+
+from CelerisLab.lbm.raw_exchange import summarize_exchange
+
+
+def reduce(rows, window, duration=None, h=.5, time_end=10):
+    return summarize_exchange(rows, np.array([0, 0, 1]), np.array([False, True, False]),
+                              body_id=0, level=1, h=h, time_end=time_end, window=window,
+                              reference_center=(0., 0.), duration=duration)
+
+
+class ExchangeWindowReduction(unittest.TestCase):
+    def test_window_duration_scaling(self):
+        rows = np.array([[1., 2., 3., 4.], [.5, -1., 0., 1.], [9., 9., 9., 9.]])
+        out = reduce(rows, "window", duration=4)
+        np.testing.assert_array_equal(out["total"]["impulse"], [1.5, 1.])
+        np.testing.assert_array_equal(out["total"]["force"], [1.5 / 4, 1. / 4])
+        self.assertEqual(out["total"]["torque"], 3. / 4)
+        self.assertEqual(out["total"]["mass_rate"], 5. / 4)
+        np.testing.assert_array_equal(out["aperture"]["impulse"], [.5, -1.])
+        self.assertEqual((out["time_start"], out["time_end"]), (6., 10))
+        self.assertEqual(out["sample_count"], 8)  # duration / h fine substeps
+        self.assertEqual(out["window"], "window")
+
+    def test_existing_windows_unchanged(self):
+        rows = np.array([[1., 2., 3., 4.], [0., 0., 0., 0.], [0., 0., 0., 0.]])
+        base = reduce(rows, "last_base_step")
+        self.assertEqual((base["time_start"], base["sample_count"]), (9., 2))
+        np.testing.assert_array_equal(base["total"]["force"], [1., 2.])
+        sub = reduce(rows, "last_substep")
+        self.assertEqual(sub["sample_count"], 1)
+
+    def test_duration_validation(self):
+        rows = np.zeros((3, 4))
+        for bad in (None, 0, -1, 1.5, True):
+            with self.subTest(duration=bad), self.assertRaises(ValueError):
+                reduce(rows, "window", duration=bad)
+        with self.assertRaises(ValueError):
+            reduce(rows, "last_base_step", duration=2)
+
+
+if __name__ == "__main__": unittest.main()

tests/unit/test_re100.py

@@ -0,0 +1,103 @@
+import sys, json, hashlib
+from pathlib import Path
+import numpy as np
+import pytest
+
+ROOT = Path(__file__).resolve().parents[2]
+sys.path.insert(0, str(ROOT / "tests/validation"))
+import re100_helpers as h
+import periodic as p
+
+
+@pytest.mark.parametrize("n", [20, 40, 80])
+def test_units(n):
+    x = h.parameters(n, 0.01)
+    assert x["viscosity"] * x["dx"] ** 2 / x["dt"] == pytest.approx(0.001)
+    assert x["mean_speed"] * x["dx"] / x["dt"] == pytest.approx(1.0)
+    assert x["mean_speed"] * n / x["viscosity"] == pytest.approx(100)
+    assert x["peak_mach"] == pytest.approx(0.015 * np.sqrt(3))
+    assert (x["top"] - x["bottom"]) * x["dx"] == pytest.approx(0.41)
+    assert (x["nx"] - 1) * x["dx"] == pytest.approx(2.2)
+    assert (np.array(x["center"]) - [0, 0.5]) * x["dx"] == pytest.approx([0.2, 0.2])
+    cfg, _, _ = h.make_config(n, 0.01, "MRT", "neq_extrap", "zou_he_local")
+    assert cfg["method"]["inlet"]["scheme"] == "zou_he_local"
+
+
+def test_pressure_units():
+    x = h.parameters(40, 0.01)
+    ny, nx = x["ny"], x["nx"]
+    xx = np.arange(nx)[None, :]
+    rho = np.broadcast_to(1 + xx * 1e-6, (ny, nx))
+    fluid = np.ones_like(rho, dtype=bool)
+    assert h.pressure_probes(rho, fluid, x)["dp"] == pytest.approx(
+        -40e-6 / 0.01**2 / 3, abs=1e-10
+    )
+
+
+def rows(amplitude=1.0, drift=0.0):
+    out = []
+    for t in np.arange(240, 300.001, 0.02):
+        out.append(
+            dict(
+                tstar=t,
+                tstar_force=t,
+                cl=amplitude * np.sin(2 * np.pi * t / 3.3),
+                cd=3.2 + 0.03 * np.cos(4 * np.pi * t / 3.3) + drift * t,
+                dp=2.46 + 0.04 * np.sin(4 * np.pi * t / 3.3),
+                mass=1.0,
+                flux_in=1.0,
+                flux_out=1.0,
+            )
+        )
+    return out
+
+
+def test_periodic_synthetic():
+    a = p.assess(rows())
+    assert a["periodic"]
+    assert len(a["cycles"]) > 5
+    assert a["cycles"][-1]["st"] == pytest.approx(1 / 3.3, rel=1e-5)
+    assert not p.assess(rows(amplitude=0.01))["periodic"]
+    assert not p.assess(rows(drift=0.01))["periodic"]
+    assert not p.assess(rows()[:100])["periodic"]
+
+
+def test_production_unmodified():
+    m = json.loads(
+        (ROOT / "tests/validation/data/dfg_re100/source_manifest.json").read_text()
+    )
+    assert m["commit"] == h.PIN
+    for name, digest in m["production_sha256"].items():
+        assert hashlib.sha256((ROOT / name).read_bytes()).hexdigest() == digest, name
+
+
+def test_reference():
+    r = json.loads(
+        (ROOT / "tests/validation/data/dfg_re100/reference_cycle.json").read_text()
+    )
+    assert 0.301 < r["st"] < 0.303
+    assert 3.22 < r["cd_max"] < 3.23
+    assert 2.5 < r["dp_max"] < 2.51
+    assert p.compare(r, r)["reference_agreement"]
+
+
+def test_stale_cycles_rejected():
+    r = rows()
+    level = next(x["cl"] for x in r if x["tstar"] >= 284.5)
+    for x in r:
+        if x["tstar"] >= 284.5:
+            x["cl"] = level
+    assert not p.assess(r)["periodic"]
+
+
+def test_min_time_is_used(monkeypatch):
+    monkeypatch.setitem(p.GATES, "min_tstar", 290.0)
+    assert not p.assess(rows())["periodic"]
+
+
+def test_nonconsecutive_cycles_rejected(monkeypatch):
+    cycles = p.extract(rows())
+    bad = cycles[-6:]
+    bad.pop(2)
+    monkeypatch.setattr(p, "extract", lambda *a, **k: bad)
+    assert not p.assess(rows())["periodic"]

tests/unit/test_refined_dfg.py

@@ -0,0 +1,345 @@
+"""CPU-only contract tests. These are not CUDA or DFG acceptance evidence."""
+
+import json
+import sys
+from pathlib import Path
+import numpy as np
+import unittest
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "validation"))
+import run_refined_dfg as r
+
+h, periodic = r.helpers(r.DEFAULT_SOURCE)
+
+
+def test_source_pin():
+    historical = json.loads((r.HERE / "data/refined_dfg/source_manifest.json").read_text())
+    assert historical["commit"].startswith("cd3a657")
+    assert r.verify_source(r.DEFAULT_SOURCE, "MRT")["commit"].startswith("3ea131e")
+    with unittest.TestCase().assertRaisesRegex(RuntimeError, "Source mismatch"):
+        r.verify_source(r.DEFAULT_SOURCE, "SRT")
+
+
+def test_matched_units_and_planes():
+    p = r.plan(h)
+    c = p["root"]
+    f = p["uniform"]
+    assert c["nx"] == 881 and c["ny"] == 166
+    assert f["nx"] == 1761 and f["ny"] == 330
+    assert c["dx"] / 2 == f["dx"] and c["dt"] / 2 == f["dt"]
+    assert np.isclose(c["viscosity"] * c["dx"] ** 2 / c["dt"], 0.001)
+    assert np.isclose(f["viscosity"] * f["dx"] ** 2 / f["dt"], 0.001)
+    assert np.isclose((p["body_root_edge"][0] - 0.5) * c["dx"], 0.2)
+    assert np.isclose((p["body_root_edge"][1] - 1) * c["dx"], 0.2)
+    assert np.isclose((165 - 1) * c["dx"], 0.41)
+    assert np.isclose((880.5 - 0.5) * c["dx"], 2.2)
+    assert np.isclose(2 * p["radius_root"] * c["dx"], 0.1)
+
+
+def test_same_numerics():
+    c, _, _ = r.config(h, 40)
+    f, _, _ = r.config(h, 80)
+    assert c["method"] == f["method"]
+    m = c["method"]
+    assert m["collision"] == "SRT" and m["streaming"] == "double_buffer"
+    assert m["store_precision"] == "FP32" and m["ddf_shifting"]
+    assert not m["les"]["enabled"]
+    assert m["inlet"]["scheme"] == "zou_he_local"
+    assert m["outlet"]["mode"] == "neq_extrap" and m["outlet"]["srt_neq_damp"] == 0.5
+    assert not m["outlet"]["backflow_clamp"]
+
+
+def test_actual_geometry_guards():
+    cfg, pair, body, host = r.host_geometry(r.DEFAULT_SOURCE, h)
+    assert pair.fine_shape == (240, 568)
+    assert np.allclose(
+        host.objects[0].geometry.center,
+        [
+            (80.5 - pair.fine_origin[0]) / 0.5 - 0.5,
+            (81 - pair.fine_origin[1]) / 0.5 - 0.5,
+        ],
+    )
+    assert cfg.nx % r.BASE_BLOCK[1] == 0 and cfg.ny % r.BASE_BLOCK[0] == 0
+    assert len(host.boundary_cells) > 0
+
+
+def test_physical_leaf_mass_excludes_covered_and_halo():
+    c = np.full((9, 4, 4), 1 / 9)
+    f = np.full((9, 8, 8), 1 / 9)
+    cm = np.ones((4, 4), bool)
+    cm[1:3, 1:3] = False
+    fm = np.zeros((8, 8), bool)
+    fm[2:6, 2:6] = True
+    c[:, ~cm] = np.nan
+    f[:, ~fm] = np.nan
+    assert np.isclose(r.leaf_mass([c, f], 2, [cm, fm]), 64)
+    f[:, 3, 3] = np.nan
+    with unittest.TestCase().assertRaisesRegex(ValueError, "Nonfinite"):
+        r.leaf_mass([c, f], 2, [cm, fm])
+
+
+def test_force_impulse_interval_scaling():
+    from CelerisLab.lbm.raw_exchange import summarize_exchange
+
+    # Two fine substeps each j=(2,4), h=.5 => root impulse=(1,2).
+    result = summarize_exchange(
+        [[1.0, 2.0, 0.0, 0.0]],
+        [0],
+        [False],
+        body_id=0,
+        level=1,
+        h=0.5,
+        time_end=8,
+        window="last_base_step",
+        reference_center=(80.5, 81.0),
+    )
+    assert result["sample_count"] == 2 and result["time_start"] == 7
+    np.testing.assert_allclose(result["total"]["force"], [1, 2])
+    assert result["units"] == "base_lattice_per_unit_depth"
+    # Force unit scales dx^3/dt^2, thus uniform fine force is twice root.
+    root_cd = 1 / (0.5 * r.SPEED**2 * 40)
+    fine_cd = 2 / (0.5 * r.SPEED**2 * 80)
+    assert root_cd == fine_cd
+
+
+def test_pressure_local_frame_and_units():
+    _, pair, _, _ = r.host_geometry(r.DEFAULT_SOURCE, h)
+    center = [
+        (80.5 - pair.fine_origin[0]) / 0.5 - 0.5,
+        (81 - pair.fine_origin[1]) / 0.5 - 0.5,
+    ]
+    p = {"center": center, "radius": 40.0, "mean_speed": r.SPEED}
+    yy, xx = np.indices(pair.fine_shape)
+    rho = 1 + 1e-7 * xx
+    probes = h.pressure_probes(rho, np.ones(pair.fine_shape, bool), p)
+    assert np.isclose(probes["dp"], -80e-7 / (3 * r.SPEED**2))
+
+
+def test_initializer_coordinates_and_rest_ghosts():
+    g = r.initial_state((166, 881), (0, 0), 1, root=True)
+    rho, ux, uy = r.macros(g)
+    assert np.max(abs(rho - 1)) < 1e-14
+    assert np.max(abs(uy)) == 0
+    assert np.max(abs(ux[[0, -1]])) == 0
+    assert np.isclose(ux[81, 50], 6 * r.SPEED * (80.5 / 164) * (1 - 80.5 / 164))
+
+
+def test_axis_specific_sampling_phase():
+    _, pair, _, _ = r.host_geometry(r.DEFAULT_SOURCE, h)
+    assert pair.fine_origin == (38.0, 22.0)
+    j = np.arange(pair.fine_shape[1])
+    k = np.arange(pair.fine_shape[0])
+    fine_x = (pair.fine_origin[0] + (j + 0.5) / 2 - 0.5) * 0.0025
+    fine_y = (pair.fine_origin[1] + (k + 0.5) / 2 - 1) * 0.0025
+    np.testing.assert_allclose(fine_x / 0.00125, j + 75.5, atol=1e-12)
+    np.testing.assert_allclose(fine_y / 0.00125, (k + 43) - 0.5, atol=1e-12)
+    np.testing.assert_allclose(fine_x - (j + 75) * 0.00125, 0.5 * 0.00125, atol=1e-15)
+    np.testing.assert_allclose(fine_y - ((k + 43) - 0.5) * 0.00125, 0, atol=1e-15)
+
+
+def test_full_requires_completed_pilot():
+    with unittest.TestCase().assertRaisesRegex(ValueError, "pilot-results"):
+        r.verify_pilot(None, r.verify_source(r.DEFAULT_SOURCE, "MRT"))
+
+
+def pilot_fixture(path):
+    """Synthetic contract fixture only; never a physical result or run receipt."""
+    manifest = r.verify_source(r.DEFAULT_SOURCE, "MRT")
+    (path / "plan.json").write_text(
+        json.dumps(
+            dict(
+                source_manifest=manifest,
+                collision="MRT",
+                bounds_root_edge=list(r.BOUNDS),
+                runner_sha256=r.runner_hash(),
+            )
+        )
+    )
+    for case, diameter in (("uniform", 80), ("refined", 40)):
+        dest = path / case
+        dest.mkdir()
+        interval = round(0.02 * diameter / r.SPEED)
+        rows = [
+            dict(
+                step=(i + 1) * interval,
+                tstar=(i + 1) * 0.02,
+                tstar_force=(i + 0.5) * 0.02,
+                cd=3.2,
+                cl=0.1,
+                dp=2.0,
+                mass=0.8,
+                rho_min=0.99,
+                rho_max=1.01,
+                max_mach=0.03,
+                flux_in=0.4,
+                flux_out=0.4,
+                rho_front=1.0,
+                rho_back=1.0,
+            )
+            for i in range(250)
+        ]
+        (dest / "summary.json").write_text(
+            json.dumps(
+                dict(
+                    case=case,
+                    collision="MRT",
+                    mode="pilot",
+                    status="pilot_complete",
+                    healthy=True,
+                    steps=250 * interval,
+                    history=rows,
+                    startup_seconds=1.0,
+                    wall_seconds=11.0,
+                    step_seconds=9.0,
+                )
+            )
+        )
+        (dest / "ready.json").write_text(
+            json.dumps(
+                dict(
+                    case=case,
+                    collision="MRT",
+                    mode="pilot",
+                    step=interval,
+                    runner_sha256=r.runner_hash(),
+                    all_runtime_and_observation_paths_exercised=True,
+                )
+            )
+        )
+    return manifest
+
+
+def test_pilot_validation_and_calibration():
+    from tempfile import TemporaryDirectory
+
+    with TemporaryDirectory() as tmp:
+        path = Path(tmp)
+        manifest = pilot_fixture(path)
+        assert r.verify_pilot(path, manifest, "MRT") == 3606.0
+
+
+def test_reject_malformed_pilot_receipts():
+    from tempfile import TemporaryDirectory
+    from copy import deepcopy
+
+    with TemporaryDirectory() as tmp:
+        path = Path(tmp)
+        manifest = pilot_fixture(path)
+        receipt = path / "uniform/summary.json"
+        valid = json.loads(receipt.read_text())
+        mutations = [
+            lambda s: s.update(case="refined"),
+            lambda s: s.update(mode="full"),
+            lambda s: s.update(healthy="true"),
+            lambda s: s.update(steps=1),
+            lambda s: s["history"].pop(),
+            lambda s: s["history"][1].update(step=160),
+            lambda s: s["history"][-1].update(tstar=float("nan")),
+            lambda s: s["history"][0].update(cd=float("inf")),
+            lambda s: s["history"][0].update(max_mach=0.3),
+            lambda s: s["history"][0].update(tstar_force=0.02),
+            lambda s: s.update(wall_seconds=0),
+        ]
+        for mutate in mutations:
+            bad = deepcopy(valid)
+            mutate(bad)
+            receipt.write_text(json.dumps(bad))
+            with unittest.TestCase().assertRaises(ValueError):
+                r.verify_pilot(path, manifest, "MRT")
+        # Exact reviewer regression: previously accepted NaN endpoint.
+        receipt.write_text(
+            json.dumps(
+                dict(
+                    status="pilot_complete",
+                    healthy=True,
+                    history=[{"tstar": float("nan")}],
+                )
+            )
+        )
+        with unittest.TestCase().assertRaises(ValueError):
+            r.verify_pilot(path, manifest, "MRT")
+        receipt.write_text(json.dumps(valid))
+        (path / "uniform/ready.json").write_text("{}")
+        with unittest.TestCase().assertRaises(ValueError):
+            r.verify_pilot(path, manifest, "MRT")
+
+
+def test_strict_finite_json_and_deadline():
+    from tempfile import TemporaryDirectory
+
+    with TemporaryDirectory() as tmp:
+        p = Path(tmp) / "bad.json"
+        for data in ('{"x": NaN}', '{"x": Infinity}', '{"x": 1e999}'):
+            p.write_text(data)
+            with unittest.TestCase().assertRaises(ValueError):
+                r.read_finite_json(p)
+    with unittest.TestCase().assertRaises(TimeoutError):
+        r.check_deadline(0)
+    r.check_deadline(float("inf"))
+
+
+def test_pressure_probes_actual_body_mask():
+    from CelerisLab.lbm.descriptors import SOLID
+    from CelerisLab.lbm.refinement import ACTIVE
+
+    _, pair, _, host = r.host_geometry(r.DEFAULT_SOURCE, h)
+    fluid = (pair.fine_kind == ACTIVE) & (
+        (host.flags.reshape(pair.fine_shape) & SOLID) == 0
+    )
+    p = dict(center=host.objects[0].geometry.center, radius=40.0, mean_speed=r.SPEED)
+    assert h.pressure_probes(np.ones(pair.fine_shape), fluid, p)["dp"] == 0.0
+
+
+def test_matched_mrt_selection():
+    for diameter in (40, 80):
+        srt, sbody, sp = r.config(h, diameter)
+        mrt, mbody, mp = r.config(h, diameter, "MRT")
+        assert mrt["method"]["collision"] == "MRT"
+        mrt["method"]["collision"] = "SRT"
+        assert mrt == srt and mbody == sbody and mp == sp
+    srt_plan, mrt_plan = r.plan(h), r.plan(h, "MRT")
+    for key in ("root", "uniform", "bounds_root_edge", "finest_dx_m", "finest_dt_s"):
+        assert srt_plan[key] == mrt_plan[key]
+    assert "q2/6/8" in mrt_plan["outlet"] and "all nine" in srt_plan["outlet"]
+    cfg, pair, _, host = r.host_geometry(r.DEFAULT_SOURCE, h, "MRT")
+    assert cfg.collision == "MRT" and pair.fine_shape == (240, 568)
+    assert len(host.boundary_cells) > 0
+
+
+def test_reject_collision_receipts():
+    from tempfile import TemporaryDirectory
+
+    for rel in ("plan.json", "uniform/summary.json", "refined/summary.json",
+                "uniform/ready.json", "refined/ready.json"):
+        for invalid in ("SRT", None):
+            with TemporaryDirectory() as tmp:
+                path = Path(tmp)
+                manifest = pilot_fixture(path)
+                receipt = path / rel
+                record = json.loads(receipt.read_text())
+                if invalid is None:
+                    record.pop("collision")
+                else:
+                    record["collision"] = invalid
+                receipt.write_text(json.dumps(record))
+                with unittest.TestCase().assertRaises(ValueError):
+                    r.verify_pilot(path, manifest, "MRT")
+
+
+def test_record_actual_checkout_separately():
+    from source_provenance import checkout_provenance
+
+    actual = checkout_provenance(r.DEFAULT_SOURCE)
+    assert len(actual["checkout_revision"]) == 40
+    assert isinstance(actual["checkout_dirty"], bool)
+    assert r.verify_source(r.DEFAULT_SOURCE, "MRT")["production_origin_commit"].startswith("3ea131e")
+
+
+if __name__ == "__main__":
+    suite = unittest.TestSuite(
+        unittest.FunctionTestCase(fn)
+        for name, fn in sorted(globals().items())
+        if name.startswith("test_")
+    )
+    result = unittest.TextTestRunner(verbosity=2).run(suite)
+    raise SystemExit(not result.wasSuccessful())

tests/unit/test_refinement.py

@@ -142,8 +142,8 @@ class FixedPatchTests(unittest.TestCase):
                 FixedPatch2D.from_config(LBMConfig(nx=nx, ny=ny, viscosity=.1), bounds)
 
     def test_unsupported_backend_and_fine_relaxation_rejected(self):
-        for kwargs in ({"collision": "TRT"}, {"collision": "MRT"}, {"streaming": "esopull"},
-                       {"ddf_shifting": True}, {"les_enabled": True},
+        for kwargs in ({"collision": "TRT"}, {"streaming": "esopull"},
+                       {"les_enabled": True},
                        {"store_precision": "FP16S"}, {"lattice_model": "D3Q19"},
                        {"data_type": "FP64"}, {"omega_min": 1.0}):
             cfg = config(**kwargs)
@@ -181,14 +181,14 @@ class FixedPatchSetTests(unittest.TestCase):
             ly, lx = p.fine_shape
             np.testing.assert_array_equal(s.fine_kind[y0:y0 + ly, x0:x0 + lx], p.fine_kind)
             gy, gx = np.divmod(p.ghost_fine.astype(np.int64), lx)
-            np.testing.assert_array_equal(s.ghost_fine[cells:cells + len(gy)],
-                                          (gy + y0) * fx + gx + x0)
-            np.testing.assert_array_equal(s.ghost_coarse[cells:cells + len(gy)], p.ghost_coarse)
+            ghosts = dict(zip(map(int, s.ghost_fine), map(int, s.ghost_coarse)))
+            for f, c in zip((gy + y0) * fx + gx + x0, p.ghost_coarse):
+                self.assertEqual(ghosts[int(f)], int(c))
             q, c = np.divmod(p.coalesce_fine.astype(np.int64), ly * lx)
             cy, cx = np.divmod(c, lx)
-            np.testing.assert_array_equal(s.coalesce_fine[rows:rows + len(q)],
-                                          q * fy * fx + (cy + y0) * fx + cx + x0)
-            np.testing.assert_array_equal(s.coalesce_coarse[rows:rows + len(q)], p.coalesce_coarse)
+            rows_by_target = dict(zip(map(int, s.coalesce_coarse), s.coalesce_fine))
+            for target, children in zip(p.coalesce_coarse, q * fy * fx + (cy + y0) * fx + cx + x0):
+                np.testing.assert_array_equal(rows_by_target[int(target)], children)
             rows, cells = rows + len(q), cells + len(gy)
         self.assertEqual((rows, cells), (len(s.coalesce_coarse), len(s.ghost_fine)))
         self.assertEqual(int((s.fine_kind == ACTIVE).sum()), 4 * int(s.covered.sum()))
@@ -207,10 +207,12 @@ class FixedPatchSetTests(unittest.TestCase):
         np.testing.assert_array_equal(s.coalesce_coarse, p.coalesce_coarse)
         np.testing.assert_array_equal(s.ghost_coarse, p.ghost_coarse)
 
-    def test_gap_and_overlap_rejected(self):
-        for gap in (0, 2, 3):
-            with self.subTest(gap=gap), self.assertRaisesRegex(ValueError, "closer"):
-                self.build(((4, 4, 8, 12), (8 + gap, 4, 12 + gap, 12)))
+    def test_gap_and_overlap_use_dependency_checks(self):
+        for gap in (-2, 0, 1, 2, 3):
+            with self.subTest(gap=gap):
+                pair = self.build(((4, 4, 8, 12), (8 + gap, 4, 12 + gap, 12)))
+                refinement._transfer_footprint(pair)
+                refinement._check_parent_pull_closure(pair)
         with self.assertRaises(ValueError):
             self.build(())
         with self.assertRaises(ValueError):  # FixedPatch2D margin still applies
@@ -226,7 +228,7 @@ class FixedPatchSetTests(unittest.TestCase):
                 support[y0:y0 + p.fine_shape[0], x0:x0 + p.fine_shape[1]] = True
             np.testing.assert_array_equal(layout.owned, support)
             self.assertTrue(layout.owned.ravel()[s.ghost_fine.astype(np.intp)].all())
-        with self.assertRaisesRegex(ValueError, "aligned"):
+        with self.assertRaisesRegex(ValueError, "divide"):
             self.build(self.FACE).fine_layout((5, 4))
         for block in ((0, 4), (4, 0), (-4, 4), (4.0, 4), (True, 4), (4,), (4, 4, 4), (np.bool_(1), 4)):
             with self.subTest(block=block), self.assertRaisesRegex(ValueError, "positive integers"):

tests/unit/test_refinement_mrt.py

@@ -0,0 +1,44 @@
+"""Host guards and source contract for MRT reuse in refined/blocked stepping (no CUDA)."""
+from pathlib import Path
+import unittest
+
+from CelerisLab.config import LBMConfig
+from CelerisLab.lbm.refinement import FixedPatch2D
+
+KERNELS = Path(__file__).resolve().parents[2] / "src/CelerisLab/lbm/kernels"
+BOUNDS = (8, 6, 16, 14)
+
+
+class RefinementMRTHost(unittest.TestCase):
+    def test_patch_accepts_implemented_models(self):
+        for collision in ("SRT", "MRT"):
+            with self.subTest(collision=collision):
+                FixedPatch2D.from_config(LBMConfig(nx=24, ny=20, viscosity=.1, collision=collision), BOUNDS)
+
+    def test_patch_rejects_unimplemented(self):
+        for options in (dict(collision="TRT"), dict(collision="MRT", les_enabled=True)):
+            with self.subTest(options=options), self.assertRaises(ValueError):
+                FixedPatch2D.from_config(LBMConfig(nx=24, ny=20, viscosity=.1, **options), BOUNDS)
+
+    def test_srt_statements_and_guards_preserved(self):
+        """MRT is an additive compile-time branch; SRT statements and other guards remain."""
+        blocks = (KERNELS / "blocks.cu").read_text()
+        channel = (KERNELS / "step/block_channel.cu").read_text()
+        refinement = (KERNELS / "refinement.cu").read_text()
+        for text in ("compute_feq_from_momentum(drho, jx, jy, ux, uy, feq);",
+                     "compute_feq(rho, ux, uy, feq);",
+                     "collide_srt_no_force(f, feq, d_params.omega);"):
+            self.assertIn(text, blocks)
+        self.assertIn("collide_srt(f, feq, forcing, omega);", channel)
+        self.assertIn("collide_srt_no_force(f, feq, d_params.omega);", refinement)
+        for source in (blocks, refinement):
+            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)
+        self.assertIn("USE_DDF_SHIFTING ||", refinement)  # dense/mixed shifted stays rejected
+        for source in (blocks, channel, refinement):
+            self.assertEqual(source.count("collide_mrt_no_force("), 1)
+
+
+if __name__ == "__main__": unittest.main()

tests/unit/test_refinement_three_level_multipatch.py

@@ -217,7 +217,7 @@ class ConstructorAndSchedule(unittest.TestCase):
         cases=[dict(child=(8,8,16,16)),dict(child_rectangles_in_parent=[]),
                dict(child_rectangles_in_parent=[(16,16,32)]),
                dict(child_rectangles_in_parent=[(16.5,16,32,40)]),
-               dict(child_rectangles_in_parent=[(16,16,32,24),(16,27,32,40)]),
+               dict(child_rectangles_in_parent=[(16,16,32,24),(16,27,32,55)]),
                dict(child_rectangles_in_parent=[(40,16,48,40)]),
                dict(child_rectangles_in_parent=[(8,16,32,40)]),
                dict(child_block=None),dict(fine_block=None),dict(base_block=None),
@@ -226,8 +226,8 @@ class ConstructorAndSchedule(unittest.TestCase):
         for changes in cases:
             with self.subTest(changes=changes),self.assertRaises((ValueError,TypeError)):
                 host.FixedPatchStepper(CFG,PARENTS,context=context,**dict(KW,**changes))
-        with self.assertRaisesRegex(ValueError,'closer'):
-            host.FixedPatchStepper(CFG,[(4,4,20,28),(23,4,44,28)],context=context,**KW)
+        with self.assertRaisesRegex(ValueError,'margin'):
+            host.FixedPatchStepper(CFG,[(1,4,20,28),(23,4,44,28)],context=context,**KW)
         objects,jets=bodies_and_jets()
         with self.assertRaisesRegex(ValueError,'only on L2'):
             host.FixedPatchStepper(CFG,PARENTS,context=context,fine_objects=objects,**KW)

tests/unit/test_runtime_child_remap.py

@@ -67,7 +67,7 @@ class RuntimeChildRemap(unittest.TestCase):
             run._status = status
             with self.assertRaises(RuntimeError): run.remap_child(B)
         run._status = 'ready'
-        for bad in ([], [(40,16,48,40)], [(16,16,32,24),(16,27,32,40)]):
+        for bad in ([], [(40,16,48,40)], [(16,16,32,24),(16,27,32,55)]):
             with self.assertRaises(ValueError): run.remap_child(bad)
         hosts, pairs = run._boundary_hosts, run.hierarchy.pairs
         for attribute, value in (('_wall_aware', True),

tests/unit/test_shifted_block_operators.py

@@ -0,0 +1,41 @@
+"""Actual-header CPU algebra, not CUDA arithmetic/performance evidence."""
+import numpy as np
+from .test_numerical_operators import compile_run
+
+
+def test_retained_density_and_shifted_planar_headers(tmp_path):
+    rows=compile_run(tmp_path,'shifted_block_operators.cpp',{})
+    cx=np.array([0,1,-1,0,0,1,-1,1,-1]);cy=np.array([0,0,0,1,-1,1,-1,-1,1])
+    w=np.array([4/9]+[1/9]*4+[1/36]*4)
+    it=iter(rows)
+    for header in it:
+        kind,*numbers=header.split()
+        a=np.fromstring(next(it),sep=' ',dtype=np.float32).astype(float)
+        if kind=='E':
+            delta,u,v=np.array(numbers,dtype=np.float32).astype(float)
+            rho=float(np.float32(1+delta));cu=3*(cx*u+cy*v)
+            expected=w*(delta+rho*(cu+.5*cu*cu-1.5*(u*u+v*v)))
+            np.testing.assert_allclose(a,expected,rtol=3e-6,atol=1e-10)
+            if u==0:
+                np.testing.assert_allclose(a.sum(),delta,rtol=5e-8,atol=0)
+                if delta: assert np.any(a!=0)
+        elif kind=='J':
+            delta,jx,jy,ux,uy=np.array(numbers,dtype=np.float32).astype(float)
+            jd=cx*jx+cy*jy;ud=cx*ux+cy*uy
+            expected=w*(delta+3*jd+4.5*jd*ud-1.5*(jx*ux+jy*uy))
+            np.testing.assert_allclose(a,expected,rtol=3e-6,atol=2e-9)
+            actual=np.array([a.sum(),a@cx,a@cy,a@(cx*cx),a@(cx*cy),a@(cy*cy)])
+            target=np.array([delta,jx,jy,delta/3+jx*ux,jx*uy,delta/3+jy*uy])
+            np.testing.assert_allclose(actual,target,atol=4e-9,rtol=3e-6)
+            if jx==0 and delta: assert np.any(a!=0)
+        else:
+            mode=int(numbers[0]);rho,u,v=np.array(numbers[1:],dtype=np.float32).astype(float)
+            b=np.fromstring(next(it),sep=' ',dtype=np.float32).astype(float)
+            known=cy!=(1 if mode==0 else -1)
+            np.testing.assert_array_equal(a[known],b[known])
+            physical_rho=1+b.sum()
+            if mode==2:
+                assert abs(physical_rho-rho)<2e-8
+                assert abs(b@cx)<2e-9
+            else:
+                np.testing.assert_allclose([b@cx,b@cy],physical_rho*np.array([u,v]),atol=3e-9,rtol=0)

tests/unit/test_shifted_jets.py

@@ -0,0 +1,105 @@
+"""Shifted velocity Jets: bounded opt-in and host lifecycle, no CUDA numerics."""
+from dataclasses import replace
+import unittest
+from unittest.mock import patch
+import numpy as np
+from CelerisLab.body.jets import Jet, Tube, Parallel
+from CelerisLab.body.objects import SurfaceActuation
+from CelerisLab.lbm.synchronized_remap import RemapTopology
+from CelerisLab.lbm.refinement import FixedPatchSet2D
+from CelerisLab.lbm.population_storage import _StoredPopulations
+from tests.unit import test_refinement_three_level_multipatch as backend
+from tests.unit.test_pinned_body_remap import A, B, body, jet, mesh
+from tests.unit.test_shifted_storage import shifted_fields
+
+CFG=replace(backend.CFG,ddf_shifting=True)
+
+
+def topology(objects, jets, *, enabled=False, cfg=CFG):
+    return RemapTopology.from_config(cfg,FixedPatchSet2D.from_config(cfg,backend.PARENTS),A,
+        child_objects=objects,child_jets=jets,_shifted_blocked=enabled)
+
+
+def ready():
+    objects,jets=(body(),),(jet(),)
+    with patch.object(backend,'CFG',CFG),patch.object(backend,'bodies_and_jets',return_value=(objects,jets)):
+        run,module,allocations=backend.mock_backend(children=A,bodies=True)
+    storage={}
+    module.cuda.memcpy_htod.side_effect=lambda p,a:storage.__setitem__(id(p),np.array(a,copy=True))
+    module.cuda.memcpy_dtoh.side_effect=lambda a,p:np.copyto(a,storage[id(p)])
+    run._initialize_raw(_StoredPopulations(shifted_fields(mesh(A,objects,jets)),True))
+    run._step_count=7
+    return run,module,allocations,storage,objects
+
+
+class ShiftedJets(unittest.TestCase):
+    def test_default_guard_and_storage_neutral_host_rows(self):
+        with self.assertRaisesRegex(ValueError,'unshifted'):
+            topology((body(),),(jet(),))
+        shifted=topology((body(),),(jet(),),enabled=True)
+        unshifted=topology((body(),),(jet(),),cfg=backend.CFG)
+        np.testing.assert_array_equal(shifted.surface_velocity,unshifted.surface_velocity)
+        self.assertEqual(shifted.surface_velocity.dtype,np.float32)
+        legacy=body(SurfaceActuation(.7,.85,.002))
+        with self.assertRaisesRegex(ValueError,'unshifted'):
+            topology((legacy,),(),enabled=True)
+
+    def test_rejections_before_device_allocation(self):
+        module=backend.host_module()
+        moving=body();moving.state.omega=.01
+        cases=[((moving,),(jet(),)),((body(),),(jet(),jet())),
+               ((body(),),(Jet(Tube((100.,100.),0.,1.,.1),Parallel(0.)),)),
+               ((body(center=(8.,14.)),),(jet(),)),
+               ((body(),body(center=(34.,14.))),
+                (Jet(Tube((10.,14.),0.,26.,.3),Parallel(0.)),))]
+        for objects,jets in cases:
+            with self.subTest(objects=len(objects),jets=len(jets)),self.assertRaises(ValueError):
+                module.FixedPatchStepper(CFG,backend.PARENTS,context=None,
+                    base_block=(4,4),fine_block=(4,4),child_block=(4,4),
+                    child_rectangles_in_parent=A,child_objects=objects,child_jets=jets)
+        for blocks in ((None,None,None),((4,4),None,(4,4))):
+            with self.assertRaises(ValueError):
+                module.FixedPatchStepper(CFG,backend.PARENTS,context=None,
+                    base_block=blocks[0],fine_block=blocks[1],child_block=blocks[2],
+                    child_rectangles_in_parent=A,child_objects=(body(),),child_jets=(jet(),))
+        module.cuda.mem_alloc.assert_not_called()
+        module.compiler.compile_kernel.assert_not_called()
+
+    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]
+            before=run._download_raw()
+            original=boundary.tables['surface_velocity'].copy()
+            module.cuda.memcpy_htod.reset_mock()
+            run.set_jet_amplitude(0,-jet().amplitude,level=2)
+            module.cuda.memcpy_htod.assert_called_once()
+            self.assertIs(module.cuda.memcpy_htod.call_args.args[0],boundary.device['surface_velocity'])
+            np.testing.assert_array_equal(boundary.tables['surface_velocity'],-original)
+            self.assertEqual(run.step_count,7)
+            for a,b in zip(before.values,run._download_raw().values):self.assertEqual(a.tobytes(),b.tobytes())
+            clone=run.remap_child(B,child_objects=objects)
+            self.assertEqual(clone._level_jets[2][0].amplitude,-jet().amplitude)
+            np.testing.assert_array_equal(clone._level_boundaries[2].tables['surface_velocity'],
+                                          boundary.tables['surface_velocity'])
+            self.assertEqual(clone.step_count,7)
+            saved=run._download_raw()
+            with self.assertRaises(ValueError):run.remap_child(B,child_objects=(body(center=(12.1,14.)),))
+            for a,b in zip(saved.values,run._download_raw().values):self.assertEqual(a.tobytes(),b.tobytes())
+        finally:
+            if clone is not None:clone.close()
+            run.close()
+
+    def test_invalid_amplitude_does_not_upload_or_commit(self):
+        run,module,_,_,_=ready()
+        try:
+            before=run._level_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'])
+        finally:run.close()
+
+
+if __name__=='__main__':unittest.main()

tests/unit/test_shifted_storage.py

@@ -0,0 +1,184 @@
+"""CPU host contracts; mocked CUDA executes real lifecycle, not numerical kernels."""
+import ast
+import types
+from pathlib import Path
+from contextlib import nullcontext
+from dataclasses import replace
+import unittest
+from unittest.mock import Mock, patch
+import numpy as np
+from CelerisLab.lbm.population_storage import (_StoredPopulations, _WEIGHTS,
+                                               _encode_physical, _decode_stored)
+from CelerisLab.lbm.synchronized_remap import _remap_stored, remap_synchronized
+from tests.unit.test_synchronized_remap import topology, A, B
+from tests.unit import test_refinement_three_level_multipatch as backend
+
+
+def shifted_fields(mesh):
+    rng = np.random.default_rng(129)
+    arrays = []
+    for leaf in mesh.leaf_masks:
+        a = rng.normal(0, 1e-7, (9, *leaf.shape)).astype(np.float32)
+        a[:, ~leaf] = np.nan
+        y, x = np.nonzero(leaf)
+        a[0, y[:6], x[:6]] = np.array([0., -0., 2**-80, -2**-80, 2**-140, -2**-140], np.float32)
+        arrays.append(a)
+    return tuple(arrays)
+
+
+def remap(old, new, arrays, shifted=True):
+    return _remap_stored(old, new, _StoredPopulations(arrays, shifted), shifted=True,
+                         base_count=7, level_steps=(7, 14, 28), phase='postcollision_g')
+
+
+def ready():
+    with patch.object(backend, 'CFG', replace(backend.CFG, ddf_shifting=True)):
+        run, module, allocations = backend.mock_backend(children=A)
+    storage = {}
+    module.cuda.memcpy_htod.side_effect = lambda p, a: storage.__setitem__(id(p), np.array(a, copy=True))
+    module.cuda.memcpy_dtoh.side_effect = lambda a, p: np.copyto(a, storage[id(p)])
+    run._initialize_raw(_StoredPopulations(shifted_fields(topology(A)), True))
+    run._step_count = 7
+    return run, module, allocations, storage
+
+
+class ShiftedStorage(unittest.TestCase):
+    def test_codec_preserves_sub_fp32_physical_perturbation(self):
+        physical = np.broadcast_to(_WEIGHTS, (9, 2, 2)).copy()
+        physical[1] += 2**-35
+        encoded = _encode_physical(physical, True)
+        self.assertTrue((encoded[1] == np.float32(2**-35)).all())
+        np.testing.assert_array_equal(_decode_stored(encoded, True), physical)
+        self.assertEqual(_decode_stored(encoded, True).dtype, np.float64)
+        np.testing.assert_array_equal(_encode_physical(_WEIGHTS, True), np.zeros((9,1,1), np.float32))
+        self.assertEqual(_encode_physical(physical, False).dtype, np.float32)
+
+    def test_actual_fill_loop_bitwise_copy_and_single_rounding_restriction(self):
+        old, new = topology(A), topology(B)
+        arrays = shifted_fields(old)
+        saved = [a.tobytes() for a in arrays]
+        out = remap(old, new, arrays)
+        self.assertTrue(out.shifted)
+        for a, b, om, nm, before in zip(arrays, out.g, old.leaf_masks, new.leaf_masks, saved):
+            self.assertEqual(a[:, om & nm].tobytes(), b[:, om & nm].tobytes())
+            self.assertTrue(np.isnan(b[:, ~nm]).all())
+            self.assertEqual(a.tobytes(), before)
+            self.assertFalse(np.shares_memory(a,b))
+        y,x = np.nonzero(new.leaf_masks[2] & ~old.leaf_masks[2])
+        self.assertEqual(out.g[2][:,y,x].tobytes(), arrays[1][:,y//2,x//2].tobytes())
+        y,x = np.nonzero(new.leaf_masks[1] & ~old.leaf_masks[1])
+        mean = sum(arrays[2][:,2*y+dy,2*x+dx].astype(np.float64)*.25
+                   for dy in (0,1) for dx in (0,1))
+        self.assertEqual(out.g[1][:,y,x].tobytes(), mean.astype(np.float32).tobytes())
+        def integral(mesh, values):
+            return sum(a[:,m].sum(axis=1,dtype=np.float64)/4**l
+                       for l,(a,m) in enumerate(zip(values,mesh.leaf_masks)))
+        cast_error = abs(mean.astype(np.float32).astype(np.float64)-mean).sum(axis=1)/4
+        self.assertTrue(np.all(abs(integral(old, arrays)-integral(new,out.g)) <= cast_error+2e-18))
+
+    def test_shifted_density_and_tag_validation(self):
+        old,new = topology(A),topology(B)
+        a = tuple(np.where(m[None], np.float32(-.001), np.nan).repeat(9,axis=0) for m in old.leaf_masks)
+        remap(old,new,a)
+        with self.assertRaisesRegex(ValueError,'representation'): remap(old,new,a,False)
+        with self.assertRaisesRegex(ValueError,'FP32'): remap(old,new,tuple(x.astype(np.float64) for x in a))
+        a[0][:,old.leaf_masks[0]] = -1
+        with self.assertRaisesRegex(ValueError,'positive density'): remap(old,new,a)
+        rest = tuple(np.where(m[None],np.float32(0),np.nan).repeat(9,axis=0) for m in old.leaf_masks)
+        remap(old,new,rest)
+        with self.assertRaisesRegex(ValueError,'positive density'):
+            remap_synchronized(old,new,rest,base_count=7,level_steps=(7,14,28),phase='postcollision_g')
+
+    def test_runtime_remap_bypasses_public_codec_and_preserves_both_buffers(self):
+        run,module,allocations,storage = ready()
+        clone = None
+        try:
+            before = run._download_raw()
+            expected = remap(topology(A),topology(B),before.values)
+            with patch.object(run,'download',side_effect=AssertionError('public decode used')):
+                clone = run.remap_child(B)
+            self.assertEqual((run.step_count,clone.step_count),(7,7))
+            self.assertFalse(clone._exchange_valid)
+            for level, want, got in zip(clone._levels,expected.g,clone._download_raw().values):
+                self.assertEqual(want.tobytes(),got.tobytes())
+                for pointer in (level.current,level.spare):
+                    np.testing.assert_array_equal(storage[id(pointer)],level.layout.to_blocked(want))
+            for a,b in zip(before.values,run._download_raw().values): self.assertEqual(a.tobytes(),b.tobytes())
+            for a in clone.download(): self.assertEqual(a.dtype,np.float64)
+        finally:
+            if clone is not None: clone.close()
+            run.close()
+
+    def test_runtime_public_encode_and_raw_mismatch_fail_before_upload(self):
+        run,module,_,_ = ready()
+        try:
+            physical = tuple(np.broadcast_to(_WEIGHTS, level.shape).copy() for level in run._levels)
+            for a in physical: a[1] += 2**-35
+            run.initialize(*physical)
+            for a,mask in zip(run._download_raw().values,run._leaf):
+                self.assertTrue((a[1,mask] == np.float32(2**-35)).all())
+            self.assertEqual(run.step_count,0)
+            module.cuda.memcpy_htod.reset_mock()
+            with self.assertRaisesRegex(ValueError,'representation'):
+                run._initialize_raw(_StoredPopulations(run._download_raw().values,False))
+            module.cuda.memcpy_htod.assert_not_called()
+            run._level_cfgs[0] = replace(run._level_cfgs[0],ddf_shifting=False)
+            with self.assertRaisesRegex(ValueError,'mixed'): run._download_raw()
+        finally: run.close()
+
+    def test_dense_or_mixed_shifted_rejected_before_compilation(self):
+        module=backend.host_module()
+        cfg=replace(backend.CFG,ddf_shifting=True)
+        for kwargs in ({},dict(base_block=(4,4)),dict(fine_block=(4,4))):
+            with self.assertRaisesRegex(ValueError,'all-blocked|base_block requires fine_block'):
+                module.FixedPatchStepper(cfg,(4,4,20,28),context=None,**kwargs)
+        module.compiler.compile_kernel.assert_not_called()
+
+    def test_uniform_block_codec_and_raw_snapshot(self):
+        path = Path(backend.r.__file__).with_name('block_stepper.py')
+        tree = ast.parse(path.read_text())
+        tree.body = [n for n in tree.body if not (
+            isinstance(n, ast.Import) and any(a.name == 'pycuda.driver' for a in n.names)
+            or isinstance(n, ast.ImportFrom) and n.module == 'cuda')]
+        module = types.ModuleType('CelerisLab.lbm._test_block_host')
+        module.__package__ = 'CelerisLab.lbm'
+        cuda, compiler = Mock(), Mock()
+        module.__dict__.update(cuda=cuda,compiler=compiler)
+        exec(compile(tree,str(path),'exec'),module.__dict__)
+        cuda.mem_get_info.return_value=(10**9,10**9)
+        cuda.mem_alloc.side_effect=lambda size: Mock()
+        gpu=Mock();gpu.get_global.return_value=(Mock(),20)
+        compiler.load_module.return_value=gpu
+        context=Mock(activate=nullcontext,sm_arch='sm_70')
+        layout=module.BlockLayout2D.build(16,16,4,4)
+        cfg=replace(backend.CFG,nx=16,ny=16,ddf_shifting=True)
+        run=module.PeriodicBlockStepper(cfg,layout,context=context)
+        storage={}
+        cuda.memcpy_htod.side_effect=lambda p,a: storage.__setitem__(id(p),np.array(a,copy=True))
+        cuda.memcpy_dtoh.side_effect=lambda a,p: np.copyto(a,storage[id(p)])
+        try:
+            physical=np.broadcast_to(_WEIGHTS,(9,16,16)).copy();physical[1]+=2**-35
+            run.initialize(physical)
+            raw=run._download_raw()
+            self.assertTrue(raw.shifted)
+            self.assertEqual(len(raw.values),1)
+            self.assertTrue((raw.values[0][1]==np.float32(2**-35)).all())
+            np.testing.assert_array_equal(run.download(),physical)
+        finally: run.close()
+
+    def test_passive_shifted_body_raw_remap_preserves_support(self):
+        from tests.unit.test_pinned_body_remap import mesh, body, A as rects, B as moved
+        old,new=mesh(rects,(body(),)),mesh(moved,(body(),))
+        result=remap(old,new,shifted_fields(old))
+        self.assertTrue(result.shifted)
+        self.assertTrue(np.isnan(result.g[2][:,~new.leaf_masks[2]]).all())
+        changed=mesh(moved,(body(center=(12.1,14.)),))
+        with self.assertRaisesRegex(ValueError,'boundary'):
+            remap(old,changed,shifted_fields(old))
+
+    def test_fp64_physical_roundtrip_cannot_replace_raw(self):
+        tiny=np.full((9,1,1),np.float32(2**-80))
+        self.assertFalse(np.array_equal(tiny,_encode_physical(_decode_stored(tiny,True),True)))
+
+
+if __name__=='__main__': unittest.main()

tests/unit/test_synchronized_remap.py

@@ -175,12 +175,11 @@ class SynchronizedRemapTests(unittest.TestCase):
         self.assertFalse(np.array_equal(self.g[2][:, lost], returned.g[2][:, lost]))
 
     def test_invalid_geometry_counts_representation_frame(self):
-        for rectangles in ([], [(2, 2, 10, 10)], [(16, 16, 32, 40)] * 2,
-                           [(16, 16, 33.5, 40)], [(16, 16, 32, 40), (33, 16, 36, 40)]):
+        for rectangles in ([], [(2, 2, 10, 10)], [(16, 16, 33.5, 40)]):
             with self.assertRaises(ValueError):
                 topology(rectangles)
         parent = self.old.parent
-        for changes in (dict(lattice_model='D3Q19'), dict(ddf_shifting=True),
+        for changes in (dict(lattice_model='D3Q19'),
                         dict(collision='TRT'), dict(les_enabled=True), dict(nz=2)):
             with self.assertRaises(ValueError):
                 RemapTopology.from_config(replace(CFG, **changes), parent, A)

tests/unit/test_time_refined_dfg.py

@@ -0,0 +1,86 @@
+"""CPU contract of the measurement-only refined DFG timing harness (no CUDA)."""
+from pathlib import Path
+import sys
+import unittest
+
+import numpy as np
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "validation"))
+import time_refined_dfg as t  # noqa: E402
+
+
+class TimingHarnessContract(unittest.TestCase):
+    def test_equal_simulated_duration(self):
+        s = t.schedule(2000, 200)
+        self.assertEqual(s["refined"], (200, 2000))
+        self.assertEqual(s["uniform"], (400, 4000))  # one refined base step = two D80 steps
+        self.assertAlmostEqual(s["tstar_window"], 2000 * .01 / 40)
+        for bad in ((0, 1), (1, -1)):
+            with self.assertRaises(ValueError):
+                t.schedule(*bad)
+
+    def test_timing_only_labels(self):
+        for name in ("B", "C", "E'"):
+            self.assertFalse(t.VARIANTS[name]["force_history_valid"])
+            self.assertTrue(t.VARIANTS[name]["timing_only"])
+        for name in ("A", "E"):
+            self.assertFalse(t.VARIANTS[name]["force_history_valid"])  # reads discarded
+            self.assertTrue(t.VARIANTS[name]["force_reads"])
+            self.assertIn("runner-like", t.VARIANTS[name]["force_note"])
+        self.assertEqual({v["case"] for v in t.VARIANTS.values()}, {"refined", "uniform"})
+        self.assertIn("final base interval", t.VARIANTS["C"]["force_note"])
+        self.assertIn("whole-window", t.VARIANTS["E'"]["force_note"])
+
+    def test_deadline(self):
+        import time
+        with self.assertRaises(TimeoutError):
+            t.check_deadline(time.monotonic() - 1)
+        t.check_deadline(time.monotonic() + 60)
+
+    def test_parity_detects_bits_and_nan_masks(self):
+        a = np.array([[1., np.nan], [2., 3.]], np.float32)
+        self.assertIsNone(t.same_state([a], [a.copy()]))
+        b = a.copy(); b[1, 1] = np.nextafter(np.float32(3), np.float32(4))
+        self.assertIn("bits", t.same_state([a], [b]))
+        c = a.copy(); c[0, 0] = np.nan
+        self.assertIn("NaN", t.same_state([a], [c]))
+        self.assertIn("level count", t.same_state([a], [a, a]))
+        inf = a.copy(); inf[1, 0] = np.inf
+        self.assertIn("infinite", t.same_state([inf], [inf.copy()]))  # matching inf still fails
+        self.assertIn("infinite", t.same_state([a], [inf]))
+
+    def test_flat_uniform_ddf(self):
+        fluid = np.array([[True, False, True], [True, True, False]])
+        flat = np.ones(9 * fluid.size, np.float32)  # Simulation.get_ddf() layout
+        self.assertTrue(t.uniform_populations_finite(flat, fluid))
+        flat[1 * fluid.size + 1] = np.nan  # q=1, solid cell 1: ignored
+        self.assertTrue(t.uniform_populations_finite(flat, fluid))
+        flat[4 * fluid.size + 0] = np.nan  # q=4, fluid cell 0: detected
+        self.assertFalse(t.uniform_populations_finite(flat, fluid))
+        self.assertTrue(t.uniform_populations_finite(np.ones((9,) + fluid.shape), fluid))
+        with self.assertRaises(RuntimeError):
+            t.uniform_populations_finite(np.ones(9 * fluid.size + 1), fluid)
+
+    def test_physical_health_limits(self):
+        ok = (np.array([1., 1.01]), np.array([.01, .02]), np.array([0., .01]))
+        health = t.physical_health([ok])
+        self.assertAlmostEqual(health["rho_max"], 1.01)
+        for bad in ((np.array([1., np.nan]),) + ok[1:],
+                    (np.array([.4, 1.]),) + ok[1:],
+                    (ok[0], np.array([.2, 0.]), ok[2]),
+                    (np.array([]), np.array([]), np.array([]))):
+            with self.assertRaises(RuntimeError):
+                t.physical_health([bad])
+
+    def test_summary_reports_differences_without_claims(self):
+        rows = [dict(variant=v, seconds=s) for v, s in
+                (("A", 3.), ("A", 5.), ("B", 2.), ("C", 1.), ("E", 2.), ("E'", 1.5))]
+        s = t.summarize(rows, .5)
+        self.assertEqual(s["A"]["median_seconds"], 4.)
+        self.assertEqual(s["A_minus_B_seconds"], 2.)
+        self.assertEqual(s["B_minus_C_seconds"], 1.)
+        self.assertEqual(s["A_over_E"], 2.)
+        self.assertIn("not exact causal", s["interpretation"])
+
+
+if __name__ == "__main__": unittest.main()

tests/unit/test_time_window_dfg.py

@@ -0,0 +1,38 @@
+"""CPU contract of the A-versus-window timing harness (no CUDA)."""
+from pathlib import Path
+import sys
+import unittest
+
+import numpy as np
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "validation"))
+import time_window_dfg as w  # noqa: E402
+
+
+class WindowTimingContract(unittest.TestCase):
+    def test_interval_and_variants(self):
+        self.assertEqual(w.INTERVAL, 80)
+        self.assertEqual(w.VARIANTS, ("A", "W", "W1"))
+
+    def test_components(self):
+        obs = dict(total=dict(impulse=np.array([1., 2.]), angular_impulse=3., mass_into_fluid=4.))
+        np.testing.assert_array_equal(w.components(obs), [1., 2., 3., 4.])
+
+    def test_interval_comparison(self):
+        ref = np.array([[1., -1., 0., 0.]])
+        scale = np.array([[10., 10., 1., 1.]])
+        self.assertEqual(w.compare_intervals(ref, scale, ref.copy()), 0.)
+        near = ref + np.array([[5e-9, 0., 0., 0.]])
+        self.assertLessEqual(w.compare_intervals(ref, scale, near), 1.)
+        far = ref + np.array([[1e-7, 0., 0., 0.]])
+        self.assertGreater(w.compare_intervals(ref, scale, far), 1.)
+        zero = np.zeros((1, 4))
+        self.assertGreater(w.compare_intervals(zero, zero, zero + 1e-12), 1.)  # floor, no div-by-0
+
+    def test_candidate_identity_records_files(self):
+        identity = w.candidate_identity(Path(__file__).resolve().parents[2])
+        self.assertIn("checkout_revision", identity)
+        self.assertEqual(set(identity["sha256"]), set(w.CANDIDATE_FILES))
+
+
+if __name__ == "__main__": unittest.main()

tests/unit/test_union_body_support.py

@@ -0,0 +1,44 @@
+"""Union seam bodies and actual support pinning; CPU host-only."""
+import unittest
+import numpy as np
+from CelerisLab.config import LBMConfig
+from CelerisLab.lbm.refinement import FixedPatchSet2D,prepare_fine_boundary,_circle_in_coverage
+from CelerisLab.lbm.synchronized_remap import same_physical_boundary
+from tests.unit import test_refinement as refinement_tests
+from tests.unit.test_pinned_body_remap import mesh,body,ready,A
+from tests.unit.test_synchronized_remap import fields,remap,integrals
+
+class UnionBodySupport(unittest.TestCase):
+    def test_seam_circle_matches_single_rectangle(self):
+        cfg=LBMConfig(nx=32,ny=16,viscosity=.1)
+        obj=refinement_tests.PatchSetBoundaryTests.body(16.,8.)
+        def b(rects):return prepare_fine_boundary(cfg,FixedPatchSet2D.from_config(cfg,rects),[obj])
+        self.assertTrue(same_physical_boundary(b([(4,2,28,14)]),b([(4,2,16,14),(16,2,28,14)])))
+        with self.assertRaisesRegex(ValueError,'union'):
+            b([(4,2,16,14),(17,2,28,14)])
+
+    def test_continuous_disk_not_only_sampled_solids(self):
+        mask=np.ones((8,8),bool);mask[3,3]=False
+        self.assertFalse(_circle_in_coverage(mask,3.99,3.99,.02))
+        self.assertTrue(_circle_in_coverage(mask,4.1,4.1,.1))
+
+    def test_repartition_and_connected_fluid_changes_do_not_pin_rectangles(self):
+        old=mesh(A,(body(),));initial=fields(old)
+        for rects in (((16,16,24,40),(24,16,32,40),A[1]),((12,16,32,40),A[1]),((20,16,36,40),A[1])):
+            new=mesh(rects,(body(),));self.assertTrue(same_physical_boundary(old.boundary,new.boundary))
+            out=remap(old,new,initial)
+            np.testing.assert_allclose(integrals(old,initial),integrals(new,out.g),rtol=2e-13,atol=2e-12)
+            for a,b,om,nm in zip(initial,out.g,old.leaf_masks,new.leaf_masks):
+                np.testing.assert_array_equal(a[:,om&nm],b[:,om&nm])
+
+    def test_runtime_seam_repartition_preserves_boundary_and_live_jets(self):
+        run,module,allocations,storage,objects=ready()
+        try:
+            run.set_jet_amplitude(0,.017,level=2)
+            clone=run.remap_child(((16,16,24,40),(24,16,32,40),A[1]),child_objects=objects)
+            try:
+                self.assertEqual(clone._boundaries[2].jets[0].amplitude,.017)
+                self.assertTrue(same_physical_boundary(run.child_boundary,clone.child_boundary))
+                for a,b in zip(run.download(),clone.download()):np.testing.assert_array_equal(a,b)
+            finally:clone.close()
+        finally:run.close()

tests/unit/test_union_refinement.py

@@ -0,0 +1,161 @@
+"""CPU tests: union tables vs current FixedPatchSet2D; numpy + CelerisLab on PYTHONPATH.
+Optional numerics: ADAPTER_DIR=<adapter bundle dir> [ADAPTER_FILE=adapter.py or
+multilevel_numeric_adapter.py (default: whichever exists)]."""
+from dataclasses import replace
+import importlib.util
+import os
+import unittest
+
+import numpy as np
+from CelerisLab.config import LBMConfig
+from CelerisLab.lbm.refinement import (ACTIVE, OUTER_GHOST, FixedPatchSet2D, _middle_config,
+                                       _nested_kind, _transfer_footprint)
+from CelerisLab.lbm.refinement_tables import blocked_kind
+
+from CelerisLab.lbm.refinement import (FixedPatchSet2D as UnionPatchSet2D,
+    _check_parent_pull_closure as check_parent_pull_closure, _dilate as dilate)
+
+def table_sets(pair):
+    return (pair.fine_kind.tobytes(), pair.covered.tobytes(),
+            frozenset(zip(map(int,pair.ghost_fine),map(int,pair.ghost_coarse))),
+            frozenset((int(t),tuple(map(int,c))) for t,c in zip(pair.coalesce_coarse,pair.coalesce_fine)))
+
+CFG = LBMConfig(nx=30, ny=30, viscosity=.1)
+A, B, U = (8, 8, 16, 20), (16, 8, 24, 20), (8, 8, 24, 20)
+
+
+def union(*rects, cfg=CFG):
+    return UnionPatchSet2D.from_config(cfg, rects)
+
+
+class Equivalence(unittest.TestCase):
+    def test_abutting_boxes_equal_their_union_rectangle(self):
+        reference = table_sets(FixedPatchSet2D.from_config(CFG, [U]))
+        for rects in ((A, B), (B, A), (U,)):
+            with self.subTest(rects=rects):
+                self.assertEqual(table_sets(union(*rects)), reference)
+
+    def test_vertical_abutting(self):
+        self.assertEqual(table_sets(union((8, 8, 20, 14), (8, 14, 20, 22))),
+                         table_sets(FixedPatchSet2D.from_config(CFG, [(8, 8, 20, 22)])))
+
+    def test_separated_rectangles_reproduce_current_tables(self):
+        rects = [(4, 4, 10, 10), (16, 16, 26, 26)]  # gap 6 >= MIN_PATCH_GAP
+        self.assertEqual(table_sets(union(*rects)),
+                         table_sets(FixedPatchSet2D.from_config(CFG, rects)))
+
+    def test_layout_tiles_are_free_of_rectangle_alignment(self):
+        pair = union(A, B)
+        support = dilate(pair.fine_kind == ACTIVE, 4)
+        reference = FixedPatchSet2D.from_config(CFG, [U]).fine_layout((4, 4))
+        np.testing.assert_array_equal(pair.fine_layout((4, 4)).owned, reference.owned)
+        layout = pair.fine_layout((6, 10))  # not aligned to the rectangles
+        self.assertTrue(layout.owned[support].all())
+        kind = blocked_kind(layout, pair.fine_kind)
+        owned = layout.owned
+        np.testing.assert_array_equal(kind[layout.owner[owned]], pair.fine_kind[owned])
+
+
+class NewlyLegal(unittest.TestCase):
+    def check(self, pair):
+        _transfer_footprint(pair)
+        check_parent_pull_closure(pair)
+        ghosts = pair.ghost_fine.astype(np.intp)
+        self.assertFalse((pair.fine_kind.ravel()[ghosts] == ACTIVE).any())
+        self.assertFalse(pair.covered.ravel()[pair.ghost_coarse.astype(np.intp)].any())
+
+    def test_small_gaps_and_l_shape(self):
+        for rects in (((8, 8, 16, 20), (17, 8, 24, 20)),   # gap 1
+                      ((8, 8, 16, 20), (18, 8, 24, 20)),   # gap 2
+                      ((8, 8, 20, 14), (8, 14, 14, 22))):  # L-shape
+            with self.subTest(rects=rects):
+                self.check(union(*rects))
+
+    def test_corner_contact_only_by_dependency_checks(self):
+        # Policy: no geometric ban; accepted iff the direction-specific checks pass.
+        try:
+            pair = union((8, 8, 16, 16), (16, 16, 24, 24))
+        except ValueError as error:
+            self.assertRegex(str(error), "outer ghost|allocation|INNER|surviving|coalesce")
+        else:
+            self.check(pair)
+
+
+class Negatives(unittest.TestCase):
+    def test_inputs_are_validated_before_conversion(self):
+        for bad in ((8.5, 8, 16, 20), ("8", 8, 16, 20), (True, 8, 16, 20), "abcd", (8, 8, 16)):
+            with self.subTest(bad=bad), self.assertRaises(ValueError):
+                union(bad, B)
+        self.assertEqual(union(tuple(np.int64(v) for v in A), B).bounds, (A, B))
+        pair = union(A, B)
+        for block in ((4.0, 4), (True, 4), (0, 4), (4,), ("4", 4)):
+            with self.subTest(block=block), self.assertRaises(ValueError):
+                pair.fine_layout(block)
+
+    def test_missing_donor_at_domain_margin(self):
+        with self.assertRaisesRegex(ValueError, "margin"):
+            union((1, 8, 10, 20), (10, 8, 20, 20))
+
+    def test_missing_or_duplicate_coalesce_row(self):
+        pair = union(A, B)
+        with self.assertRaisesRegex(ValueError, "lacks its coalesce row"):
+            check_parent_pull_closure(replace(pair, coalesce_coarse=pair.coalesce_coarse[1:],
+                                              coalesce_fine=pair.coalesce_fine[1:]))
+        doubled = replace(pair, coalesce_coarse=np.r_[pair.coalesce_coarse, pair.coalesce_coarse[:1]],
+                          coalesce_fine=np.r_[pair.coalesce_fine, pair.coalesce_fine[:1]])
+        with self.assertRaisesRegex(ValueError, "duplicate"):
+            check_parent_pull_closure(doubled)
+
+    def test_direction_specific_dependency_on_outer_cell(self):
+        pair = union(A, B)
+        active = pair.fine_kind == ACTIVE
+        kind = pair.fine_kind.copy()
+        kind[dilate(active, 2) & ~dilate(active, 1)] = OUTER_GHOST  # thin the INNER band
+        with self.assertRaisesRegex(ValueError, "outer ghost"):
+            _transfer_footprint(replace(pair, fine_kind=kind))
+
+    def test_nested_union_child_donors_must_be_parent_active(self):
+        parent = FixedPatchSet2D.from_config(CFG, [(6, 6, 24, 24)])  # L1 ACTIVE 12..48
+        middle = _middle_config(CFG, parent)
+        with self.assertRaisesRegex(ValueError, "donors/receivers"):
+            _nested_kind(parent, UnionPatchSet2D.from_config(middle, [(12, 20, 30, 40), (30, 20, 40, 40)]))
+        child = UnionPatchSet2D.from_config(middle, [(16, 20, 30, 40), (30, 20, 40, 40)])
+        reference = FixedPatchSet2D.from_config(middle, [(16, 20, 40, 40)])
+        self.assertEqual(table_sets(child), table_sets(reference))
+        np.testing.assert_array_equal(_nested_kind(parent, child), _nested_kind(parent, reference))
+
+
+@unittest.skipUnless(os.environ.get("ADAPTER_DIR"), "set ADAPTER_DIR for the CPU D2Q9 adapter run")
+class AdapterNumerics(unittest.TestCase):
+    def test_eight_steps_bitwise(self):
+        folder = os.environ["ADAPTER_DIR"]
+        name = os.environ.get("ADAPTER_FILE") or next(
+            n for n in ("multilevel_numeric_adapter.py", "adapter.py")
+            if os.path.exists(os.path.join(folder, n)))
+        path = os.path.join(folder, name)
+        spec = importlib.util.spec_from_file_location("multilevel_numeric_adapter", path)
+        adapter = importlib.util.module_from_spec(spec)
+        spec.loader.exec_module(adapter)
+        results = []
+        for pair in (union(A, B), FixedPatchSet2D.from_config(CFG, [U])):
+            kinds = [np.where(pair.covered, OUTER_GHOST, ACTIVE).astype(np.uint8), pair.fine_kind]
+            taus = [pair.tau_coarse, pair.tau_fine]
+            initial = []
+            for level, kind in enumerate(kinds):
+                ny, nx = kind.shape
+                y = (np.arange(ny) + .5) / 2 ** level
+                ux = 1e-3 * np.sin(2 * np.pi * y / 30)[:, None] * np.ones((1, nx))
+                cx = np.array([0, 1, -1, 0, 0, 1, -1, 1, -1], float)[:, None, None]
+                w = np.array([4/9] + [1/9]*4 + [1/36]*4)[:, None, None]
+                g = w * (1 + 3*cx*ux + 4.5*(cx*ux)**2 - 1.5*ux*ux)
+                g[:, kind != ACTIVE] = np.nan
+                initial.append(g)
+            run = adapter.make_scheduler(kinds, taus, [pair], initial)
+            run.step(8)
+            results.append([level.current.copy() for level in run.levels])
+        for a, b in zip(*results):
+            np.testing.assert_array_equal(a, b)
+
+
+if __name__ == "__main__":
+    unittest.main()

tests/unit/test_validation_source_provenance.py

@@ -0,0 +1,74 @@
+"""Explicit experimental repin and historical-source provenance contracts."""
+
+import hashlib
+import json
+from pathlib import Path
+import subprocess
+import sys
+
+ROOT = Path(__file__).resolve().parents[2]
+sys.path.insert(0, str(ROOT / "tests/validation"))
+from source_provenance import checkout_provenance
+import re100_helpers
+import run_jet_circle_cv
+
+ORIGIN = "47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5"
+
+
+def test_manifests_cover_current_source_and_preserve_origin():
+    names = subprocess.check_output(
+        ["git", "-C", str(ROOT), "ls-files", "src"], text=True
+    ).splitlines()
+    for group, hashes_key, revision_key in (
+        ("dfg_re100", "production_sha256", "commit"),
+        ("jet_circle_cv", "files", "revision"),
+    ):
+        manifest = json.loads(
+            (
+                ROOT / "tests/validation/data" / group / "source_manifest.json"
+            ).read_text()
+        )
+        assert manifest[revision_key] == ORIGIN
+        assert set(manifest[hashes_key]) == set(names)
+        assert manifest["source_origin_commits"][-1] == ORIGIN
+        assert (
+            manifest["historical_checkpoints"]["production"]
+            == "d79e3f6f0a424f5bcc3a1a6d610313e49daf97a2"
+        )
+        assert "not inherited from d79" in manifest["experiment"]
+        for name, expected in manifest[hashes_key].items():
+            assert (
+                hashlib.sha256((ROOT / name).read_bytes()).hexdigest() == expected
+            ), name
+    assert re100_helpers.PIN == run_jet_circle_cv.PIN == ORIGIN
+
+
+def test_runtime_checkout_identity_is_separate_from_source_origin(tmp_path):
+    subprocess.run(["git", "init", "-q", str(tmp_path)], check=True)
+    (tmp_path / "test.txt").write_text("tracked\n")
+    subprocess.run(["git", "-C", str(tmp_path), "add", "test.txt"], check=True)
+    subprocess.run(
+        [
+            "git",
+            "-C",
+            str(tmp_path),
+            "-c",
+            "user.name=Test",
+            "-c",
+            "user.email=test@example.invalid",
+            "commit",
+            "-qm",
+            "fixture",
+        ],
+        check=True,
+    )
+    expected = subprocess.check_output(
+        ["git", "-C", str(tmp_path), "rev-parse", "HEAD"], text=True
+    ).strip()
+    assert checkout_provenance(tmp_path) == dict(
+        checkout_revision=expected, checkout_dirty=False
+    )
+    (tmp_path / "test.txt").write_text("changed\n")
+    assert checkout_provenance(tmp_path) == dict(
+        checkout_revision=expected, checkout_dirty=True
+    )

tests/validation/analyze_jet_cv_continuation.py

@@ -0,0 +1,168 @@
+#!/usr/bin/env python3
+"""Offline common-unit/sampling sensitivity comparison; never grants physics acceptance."""
+
+import argparse
+import hashlib
+import json
+from pathlib import Path
+import numpy as np
+
+
+def integrate(t, y):
+    return np.sum(
+        np.diff(t).reshape((-1,) + (1,) * (y.ndim - 1)) * (y[1:] + y[:-1]) * 0.5, axis=0
+    )
+
+
+def analyze(directory):
+    meta_path = next(Path(directory).glob("samples_s*.json"))
+    meta = json.loads(meta_path.read_text())
+    c = meta["case"]
+    s = c["scale"]
+    sample_path = meta_path.parent / meta["sample_file"]
+    checksum_verified = False
+    if "sample_sha256" in meta:
+        if (
+            hashlib.sha256(sample_path.read_bytes()).hexdigest()
+            != meta["sample_sha256"]
+        ):
+            raise ValueError("Sample file hash mismatch")
+        checksum_verified = True
+    a = np.load(sample_path)
+    t = a["t_reference"]
+    if not len(t) or t[0] != 0 or np.any(np.diff(t) <= 0):
+        raise ValueError("Sample reference times must increase strictly from zero")
+    from jet_cv_measurements import sample_steps
+
+    complete_schedule = np.array_equal(t, sample_steps(c) / s**2)
+    records = []
+    # Input retains root units; convert before all budgets/comparisons.
+    inv = a["cv_inventory"] * np.array([1 / s**2, 1 / s, 1 / s])
+    flux = a["cv_flux"] * np.array([1.0, s, s])
+    raw = a["raw_cumulative"] * np.array([1 / s, 1 / s, 1.0, 1 / s**2])
+    raw[:, :, :, 2] /= (
+        s**2
+    )  # angular impulse reference = root/s^2 (not used in force tests)
+    q0 = 2 * np.pi / 12 * (c["diameter"] / 2) * c["jet_speed"]
+    f0 = 0.5 * (c["u"] * s) ** 2 * (c["diameter"] / s)
+    for lo in range(0, 3200, 100):
+        hi = lo + 100
+        ids = np.flatnonzero((t >= lo) & (t <= hi))
+        if len(ids) < 2 or t[ids[0]] != lo or t[ids[-1]] != hi:
+            continue
+        first, last = ids[0], ids[-1]
+        delta = inv[last] - inv[first]
+        ledger = raw[last] - raw[first]
+        cv_variants = []
+        for stride in (1, 2, 4):
+            selected = np.unique(np.r_[ids[::stride], ids[-1]])
+            total = integrate(t[selected], flux[selected])
+            m = delta[:, :, 0] + total[:, :, 0]
+            j = -delta[:, :, 1:] - total[:, :, 1:]
+            cv_variants.append(
+                np.concatenate(
+                    (
+                        (m - ledger[:, 2, 3, None])[:, :, None] / (100 * q0),
+                        (j - ledger[:, 2, None, :2]) / (100 * f0),
+                    ),
+                    axis=-1,
+                )
+            )
+        arc = (
+            integrate(t[ids], a["arc_estimates"][ids]) / 100
+        )  # Q/M rate conversion is unity
+        # Each estimator, aperture, Q/M; compare against signed prescribed Q (rho_ref=1 for M only).
+        arc_ratio = arc / np.array([q0, -q0])[None, None, :, None]
+        global_delta = (
+            a["global_mass_reference"][last] - a["global_mass_reference"][first]
+        )
+        records.append(
+            dict(
+                t_reference=[lo, hi],
+                startup=hi <= 800,
+                normalized_cv_residual=np.array(cv_variants).tolist(),
+                temporal_sensitivity_max=float(
+                    np.max(abs(np.array(cv_variants)[1:] - cv_variants[0]))
+                ),
+                arc_ratio_to_command_reference_density=arc_ratio.tolist(),
+                global_mass_residual_over_nominal=(
+                    (global_delta - ledger[:, 2, 3]) / (100 * q0)
+                ).tolist(),
+                raw_reaction_coefficient=(ledger[:, 2, :2] / (100 * f0)).tolist(),
+                raw_gross_mass_over_nominal=(
+                    (abs(ledger[:, 3, 3]) + abs(ledger[:, 4, 3])) / (100 * q0)
+                ).tolist(),
+                raw_wall_mass_over_nominal=(ledger[:, 0, 3] / (100 * q0)).tolist(),
+            )
+        )
+    return dict(
+        scale=s,
+        complete_schedule=bool(complete_schedule),
+        sample_checksum_verified=checksum_verified,
+        windows=records,
+        sampled_density_range=[
+            float(a["level_stats"][..., 0].min()),
+            float(a["level_stats"][..., 1].max()),
+        ],
+        sampled_max_speed_reference=float(a["level_stats"][..., 3].max()),
+        global_mass_initial_final=a["global_mass_reference"][[0, -1]].tolist(),
+        max_sampled_fine_velocity_error=float(a["velocity_error"].max()),
+        physics_validation_pass=False,
+    )
+
+
+def main():
+    p = argparse.ArgumentParser(description=__doc__)
+    p.add_argument("directories", nargs="+", type=Path)
+    p.add_argument("--output", type=Path, required=True)
+    a = p.parse_args()
+    results = [analyze(d) for d in a.directories]
+    payload = dict(
+        status="TRANSIENT_DIAGNOSTIC_COMPARISON",
+        results=results,
+        physics_validation_pass=False,
+    )
+    if sorted(x["scale"] for x in results) == [1, 2] and all(
+        x["complete_schedule"] for x in results
+    ):
+        ordered = sorted(results, key=lambda x: x["scale"])
+        summary = []
+        for startup in (True, False):
+            r = [
+                np.array(
+                    [
+                        w["normalized_cv_residual"][0]
+                        for w in item["windows"]
+                        if w["startup"] == startup
+                    ]
+                )
+                for item in ordered
+            ]
+            e = [
+                np.array(
+                    [
+                        w["temporal_sensitivity_max"]
+                        for w in item["windows"]
+                        if w["startup"] == startup
+                    ]
+                )
+                for item in ordered
+            ]
+            rms = [np.sqrt(np.mean(x * x, axis=(0, 1, 2))) for x in r]
+            unc = max(float(x.max()) for x in e)
+            summary.append(
+                dict(
+                    startup=startup,
+                    rms_normalized_residual_mass_Cx_Cy=[v.tolist() for v in rms],
+                    max_temporal_sensitivity=unc,
+                    reduction_exceeds_twice_temporal_proxy=(
+                        rms[0] - rms[1] > 2 * unc
+                    ).tolist(),
+                )
+            )
+        payload["refinement_trend"] = summary
+    a.output.write_text(json.dumps(payload, indent=2) + "\n")
+
+
+if __name__ == "__main__":
+    main()

tests/validation/data/dfg_re100/reference_cycle.json

@@ -0,0 +1,818 @@
+{
+  "start": 96.60907453162486,
+  "end": 99.92207699254307,
+  "period": 3.313002460918213,
+  "st": 0.301841007302737,
+  "cd_min": 3.164265915836695,
+  "cd_max": 3.2273924509469625,
+  "cd_midrange": 3.1958291833918286,
+  "cd_amplitude": 0.06312653511026767,
+  "cd_time_mean": 3.1958873825129652,
+  "cd_waveform": [
+    3.2192714249444543,
+    3.2202803099943376,
+    3.2212299416471293,
+    3.222118601683268,
+    3.2229440924252426,
+    3.2237043728228487,
+    3.224397563502949,
+    3.225021951759953,
+    3.2255759960759045,
+    3.226058329556411,
+    3.2264677639287034,
+    3.226803291773589,
+    3.2270640893003435,
+    3.227249518358305,
+    3.227359127231524,
+    3.2273924509469625,
+    3.2273485621938347,
+    3.22722863562996,
+    3.227032979281687,
+    3.2267620869419815,
+    3.2264166368706584,
+    3.2259974896059336,
+    3.225505684975984,
+    3.2249424394952255,
+    3.224309142817072,
+    3.2236073535652885,
+    3.222838795285557,
+    3.222005351549716,
+    3.2211090610907003,
+    3.220152110736951,
+    3.21913591342886,
+    3.2180639425025355,
+    3.2169388049876204,
+    3.2157632320581224,
+    3.2145400725256823,
+    3.213272285099224,
+    3.2119629314217737,
+    3.210615168117056,
+    3.209232239485073,
+    3.2078174691606125,
+    3.2063742523121657,
+    3.204906047188526,
+    3.203416367201774,
+    3.201908772307293,
+    3.20038671156229,
+    3.1988540415270714,
+    3.197314448828512,
+    3.1957715956753288,
+    3.194229147910405,
+    3.1926907667041973,
+    3.19116010031245,
+    3.1896407755313154,
+    3.188136389507111,
+    3.1866505018003437,
+    3.18518662592634,
+    3.1837482219382216,
+    3.1823386881615674,
+    3.1809613539268415,
+    3.1796199070610616,
+    3.178317365693067,
+    3.1770566177141566,
+    3.1758406383588778,
+    3.1746722981220814,
+    3.1735543560753845,
+    3.1724894540381894,
+    3.171480110448416,
+    3.1705287147383436,
+    3.169637521817997,
+    3.168808647409722,
+    3.168044062789466,
+    3.16734559035941,
+    3.166714899466848,
+    3.1661541138486857,
+    3.165664557595598,
+    3.1652468672004876,
+    3.1649020564911114,
+    3.164630969966524,
+    3.164434280442736,
+    3.1643124874918267,
+    3.164265915836695,
+    3.164294714155108,
+    3.164398854572734,
+    3.164578132207916,
+    3.164832165115698,
+    3.1651603948311298,
+    3.165562087003554,
+    3.166036763487302,
+    3.166583632347454,
+    3.1672006737836313,
+    3.1678864585242907,
+    3.1686393915089255,
+    3.1694577152711365,
+    3.17033951339357,
+    3.171282714587482,
+    3.1722850969178604,
+    3.1733442928081175,
+    3.17445779387909,
+    3.175622956745662,
+    3.176837008657004,
+    3.178097054026256,
+    3.1794002164290207,
+    3.1807436978583974,
+    3.182123709232136,
+    3.1835369281843153,
+    3.1849799460334114,
+    3.1864492760653347,
+    3.1879413618349814,
+    3.1894525859099345,
+    3.19097927879319,
+    3.1925177278192263,
+    3.1940641867753663,
+    3.195614884812199,
+    3.197166036575152,
+    3.198713851202913,
+    3.2002545337987547,
+    3.2017840862384728,
+    3.2032988764817527,
+    3.2047951919540822,
+    3.2062693641104913,
+    3.207717777711432,
+    3.2091368805426974,
+    3.2105231925781985,
+    3.2118733148571565,
+    3.2131839390335397,
+    3.2144518552407257,
+    3.2156739612615417,
+    3.216847269859638,
+    3.217968917317378,
+    3.219036170327386,
+    3.220045562724451,
+    3.220995334928343,
+    3.2218832986850447,
+    3.2227073170066918,
+    3.2234654153623845,
+    3.224155785778181,
+    3.2247767917974066,
+    3.2253269721252744,
+    3.2258050437055585,
+    3.2262099052701587,
+    3.2265406390922298,
+    3.226796513255351,
+    3.2269769825521033,
+    3.227081689768787,
+    3.22710955220279,
+    3.227061168669443,
+    3.2269370924239436,
+    3.2267377237295465,
+    3.2264636466881744,
+    3.2261156276975886,
+    3.2256946121982577,
+    3.2252017217835,
+    3.224638250833764,
+    3.2240056622744246,
+    3.223305583413101,
+    3.222539801329067,
+    3.2217102575995633,
+    3.2208190428742154,
+    3.219867854353387,
+    3.218859259026148,
+    3.2177961858992568,
+    3.2166812675225214,
+    3.2155172535532928,
+    3.2143070040629325,
+    3.2130534817570826,
+    3.211759744984257,
+    3.210428939916766,
+    3.2090642928856465,
+    3.207669102594001,
+    3.2062467320151002,
+    3.2048006003874154,
+    3.2033341751083237,
+    3.2018508739555798,
+    3.2003542779095873,
+    3.1988481127850616,
+    3.1973359573002664,
+    3.1958213973977445,
+    3.194308017670146,
+    3.19279939346063,
+    3.1912990827323733,
+    3.189810618404239,
+    3.188337500433207,
+    3.1868831881313007,
+    3.185451092560893,
+    3.1840445691293837,
+    3.1826669104651053,
+    3.1813214859380303,
+    3.1800117901097846,
+    3.178740483661364,
+    3.177510533708781,
+    3.1763248101985377,
+    3.1751860802710605,
+    3.1740970016898973,
+    3.173060117741568,
+    3.172077851413086,
+    3.171152500485684,
+    3.1702862322665037,
+    3.169481079259431,
+    3.1687389344676062,
+    3.1680615476209555,
+    3.1674506253282964,
+    3.1669085513260424,
+    3.1664356220663086,
+    3.1660329701689696,
+    3.165701568133256,
+    3.165442226438076,
+    3.165255590934966,
+    3.16514214162142,
+    3.1651021912290123,
+    3.1651358837399672,
+    3.1652431944171915,
+    3.165423928768434,
+    3.165677722868927,
+    3.166004043509568,
+    3.1664021891744154,
+    3.1668723687065503,
+    3.16741250732336,
+    3.168021321234566,
+    3.1686974411192694,
+    3.1694393380619967,
+    3.170245326209458,
+    3.171113566351732,
+    3.172042068956519,
+    3.173028698759894,
+    3.1740711784468334,
+    3.1751670940119223,
+    3.176313898883396,
+    3.177508920042188,
+    3.17874936360477,
+    3.1800328259079818,
+    3.181355841861699,
+    3.1827151303301267,
+    3.1841074729782854,
+    3.1855295642750754,
+    3.186978019262098,
+    3.1884493811374868,
+    3.1899401298427543,
+    3.191446689954596,
+    3.192965439405816,
+    3.1944927184101823,
+    3.1960248382013705,
+    3.1975580901006118,
+    3.199088754756803,
+    3.200613009960425,
+    3.20212708603014,
+    3.2036273295775386,
+    3.205110074007536,
+    3.2065716893821072,
+    3.208008592293136,
+    3.209417254433106,
+    3.210794212466389,
+    3.212136076499309,
+    3.213439539354214,
+    3.214701385159765,
+    3.215918497504247,
+    3.2170878683639637,
+    3.2182066051468428,
+    3.2192715160336944
+  ],
+  "cl_min": -1.021285623433455,
+  "cl_max": 0.9865706078511706,
+  "cl_midrange": -0.017357507791142257,
+  "cl_amplitude": 2.0078562312846255,
+  "cl_time_mean": -0.018055778744564696,
+  "cl_waveform": [
+    -1.0212846006341447,
+    -1.0209838590222233,
+    -1.0200727383007768,
+    -1.0185590327265965,
+    -1.0164441167317841,
+    -1.01372972461554,
+    -1.010417946688192,
+    -1.0065112246383783,
+    -1.0020123475697247,
+    -0.9969244470282416,
+    -0.9912509924392147,
+    -0.984995786253974,
+    -0.9781629589771562,
+    -0.9707569640900402,
+    -0.9627825730042762,
+    -0.9542433948273606,
+    -0.9451386764809715,
+    -0.9354819238811002,
+    -0.925279142068636,
+    -0.9145366279059466,
+    -0.9032609650239106,
+    -0.8914590187491076,
+    -0.8791379311317419,
+    -0.8663051160407256,
+    -0.8529682543310683,
+    -0.8391352890960927,
+    -0.8248144210074706,
+    -0.8100141037255332,
+    -0.7947430394336983,
+    -0.7790101644090639,
+    -0.7628173731216689,
+    -0.7461816969118068,
+    -0.7291127993693902,
+    -0.7116205711659391,
+    -0.6937151260047051,
+    -0.6754067966482753,
+    -0.6567061309580705,
+    -0.637623888008147,
+    -0.6181710341714578,
+    -0.5983587392102484,
+    -0.5781983723659854,
+    -0.5577014984612264,
+    -0.5368798738456273,
+    -0.5157454424421091,
+    -0.4943063234695201,
+    -0.4725784485709182,
+    -0.4505750883115271,
+    -0.4283089024165666,
+    -0.405792716370432,
+    -0.38303951669758823,
+    -0.36006244613658056,
+    -0.3368747985998418,
+    -0.3134900140010978,
+    -0.28992167285852,
+    -0.2661834907138685,
+    -0.24228931233131965,
+    -0.21825310567974987,
+    -0.19408895564887746,
+    -0.16980993582635662,
+    -0.14543133426087812,
+    -0.12096814315114233,
+    -0.09643485175457375,
+    -0.0718460310133605,
+    -0.04721632579948139,
+    -0.022560446976959612,
+    0.0021068368194982773,
+    0.02677070764406316,
+    0.05141630692334945,
+    0.07602874438036183,
+    0.10059310719281325,
+    0.1250944693656024,
+    0.1495179012836364,
+    0.17384754771042515,
+    0.19806822588832526,
+    0.22216578848678734,
+    0.2461253934945549,
+    0.2699322466441472,
+    0.2935716120906063,
+    0.3170288231832102,
+    0.34028929328125385,
+    0.36333852677889195,
+    0.3861621301373943,
+    0.40874582307035934,
+    0.43107544954464694,
+    0.4531369891575652,
+    0.47491656822324385,
+    0.4963986071823978,
+    0.5175678949127962,
+    0.5384141735804333,
+    0.5589242775445442,
+    0.5790852440767277,
+    0.5988843242125788,
+    0.6183089934932646,
+    0.6373469625245218,
+    0.6559861873413382,
+    0.674214879706298,
+    0.6920215170336557,
+    0.7093948521976959,
+    0.7263239230232448,
+    0.742798061573706,
+    0.7588052844335239,
+    0.7743307471351503,
+    0.7893708273667098,
+    0.803916135973999,
+    0.8179576199407215,
+    0.8314865695778171,
+    0.8444946252982337,
+    0.8569737840555796,
+    0.8689164052636061,
+    0.8803152163390694,
+    0.8911633177499939,
+    0.9014541876864335,
+    0.9111816861332408,
+    0.9203400586319765,
+    0.9289234860093117,
+    0.936918266184497,
+    0.9443288660701289,
+    0.9511511188545412,
+    0.9573812545471582,
+    0.9630159006719111,
+    0.9680520825648866,
+    0.972487223197885,
+    0.9763191426551175,
+    0.9795460569396854,
+    0.9821665765670635,
+    0.9841797046166384,
+    0.9855848344669915,
+    0.986381747045911,
+    0.9865706078511706,
+    0.9861430751758404,
+    0.9851078843149434,
+    0.9834674674373464,
+    0.9812235079135473,
+    0.978378053362095,
+    0.974933510845897,
+    0.9708926418403944,
+    0.9662585570602749,
+    0.9610347110820011,
+    0.9552248968839854,
+    0.9488332402249186,
+    0.941864193965974,
+    0.9343225322984421,
+    0.9262133449652685,
+    0.9175353367316478,
+    0.9082985884427632,
+    0.8985116225127334,
+    0.888180745849068,
+    0.8773125515396398,
+    0.8659139131624679,
+    0.8539919791562397,
+    0.841554167318266,
+    0.8286081593376501,
+    0.8151618954310706,
+    0.8012235691354387,
+    0.7868016221216644,
+    0.7719047392049175,
+    0.7565418433711996,
+    0.7407178806800201,
+    0.7244435174512682,
+    0.707731607286953,
+    0.6905919887305398,
+    0.6730347151255153,
+    0.6550700503410322,
+    0.6367084643889164,
+    0.617960629252382,
+    0.5988374147135249,
+    0.5793498841758384,
+    0.5595092905790806,
+    0.5393270722633181,
+    0.5188148488816408,
+    0.4979844172354143,
+    0.4768457513845936,
+    0.45541040875158884,
+    0.43369366814580723,
+    0.41170800001396685,
+    0.38946603027749815,
+    0.3669805357724072,
+    0.3442644395643103,
+    0.3213308061849796,
+    0.2981928366787376,
+    0.2748638635506233,
+    0.2513573454944852,
+    0.22768686201881655,
+    0.20386610779915632,
+    0.17990888695762292,
+    0.15582865541317795,
+    0.13163881464575702,
+    0.10735496845190624,
+    0.08299130047733835,
+    0.05856207148298293,
+    0.034081612308419704,
+    0.009564316642990896,
+    -0.014975366427826247,
+    -0.039522940080061256,
+    -0.06406386762101547,
+    -0.08858358056301793,
+    -0.11306748690571801,
+    -0.1375009796129617,
+    -0.16186944527282868,
+    -0.186158127286089,
+    -0.21035090571304071,
+    -0.23443440343427951,
+    -0.2583940736125971,
+    -0.28221540953832974,
+    -0.30588395428926807,
+    -0.3293853105136673,
+    -0.35270515033803884,
+    -0.37582922531739726,
+    -0.39874337658360604,
+    -0.4214335449640186,
+    -0.44388578123531364,
+    -0.4660862563718654,
+    -0.4880212718764244,
+    -0.509677270114379,
+    -0.531036136112504,
+    -0.5520885997128909,
+    -0.5728219120147833,
+    -0.5932231908464959,
+    -0.6132797468099019,
+    -0.6329790933854313,
+    -0.652308956917678,
+    -0.6712572864574565,
+    -0.6898122635200815,
+    -0.7079623116377989,
+    -0.7256961057871241,
+    -0.7430025816000796,
+    -0.7598709443367393,
+    -0.7762906777176882,
+    -0.7922455571624138,
+    -0.8077296357950401,
+    -0.8227349965022713,
+    -0.8372523342423263,
+    -0.8512726627071948,
+    -0.8647873213967461,
+    -0.8777879823088283,
+    -0.8902666562070795,
+    -0.902215698680155,
+    -0.9136278156586564,
+    -0.9244960686140175,
+    -0.9348138793173668,
+    -0.9445750342635408,
+    -0.9537736885083689,
+    -0.9623985419578303,
+    -0.9704465786953729,
+    -0.977916689291491,
+    -0.9848045517658248,
+    -0.99110622613045,
+    -0.9968181555444723,
+    -1.0019371672305317,
+    -1.0064604728414646,
+    -1.0103856687137327,
+    -1.0137107350135544,
+    -1.0164340351998324,
+    -1.01855431505123,
+    -1.0200707006187215,
+    -1.0209826964966897,
+    -1.021285623433455
+  ],
+  "dp_min": 2.4203712005726827,
+  "dp_max": 2.507012902468643,
+  "dp_midrange": 2.463692051520663,
+  "dp_amplitude": 0.08664170189596021,
+  "dp_time_mean": 2.4639514695178963,
+  "dp_waveform": [
+    2.485558537878651,
+    2.487354668330821,
+    2.489094038859844,
+    2.4907730268211923,
+    2.4923875927671357,
+    2.493933872161179,
+    2.4954081841697953,
+    2.4968070398569093,
+    2.498127149841111,
+    2.4993654310604727,
+    2.5005190130830766,
+    2.5015852434313093,
+    2.502561692495173,
+    2.50344615769302,
+    2.5042366667973424,
+    2.504931223854115,
+    2.5055272123906085,
+    2.5060247031878675,
+    2.5064226784653965,
+    2.5067203585292273,
+    2.506917201491131,
+    2.507012902468643,
+    2.507007392300157,
+    2.5069008355614204,
+    2.506693628251957,
+    2.5063863949686396,
+    2.5059799860275485,
+    2.505475473290949,
+    2.504874146643661,
+    2.5041775076417783,
+    2.503385760083982,
+    2.5025023981520067,
+    2.501529545753141,
+    2.500469519417466,
+    2.4993248229208986,
+    2.49809814052029,
+    2.4967923311909144,
+    2.495410422113925,
+    2.493955601552369,
+    2.492431212574658,
+    2.490840745393346,
+    2.489187830493204,
+    2.4874762310352,
+    2.4857098352761477,
+    2.483891999686845,
+    2.4820274911095415,
+    2.480120662304518,
+    2.4781758269143728,
+    2.476197383057849,
+    2.474189804509867,
+    2.47215763257422,
+    2.47010546680718,
+    2.4680379564006794,
+    2.465959791179922,
+    2.4638756921667406,
+    2.46179040254991,
+    2.4597086782017046,
+    2.457635278124787,
+    2.455575155334104,
+    2.4535330651884157,
+    2.451513647441801,
+    2.4495215742730494,
+    2.447561471628255,
+    2.4456379085376456,
+    2.443755387861672,
+    2.4419183356738783,
+    2.4401310914614514,
+    2.438397897721916,
+    2.436722890392385,
+    2.4351100884293455,
+    2.4335633840618365,
+    2.432086533067437,
+    2.4306838103185973,
+    2.4293586878640068,
+    2.428113869269998,
+    2.4269524606913686,
+    2.4258773845443202,
+    2.424891371376568,
+    2.423996951340259,
+    2.4231964462177444,
+    2.4224919620804686,
+    2.4218853823843163,
+    2.421378360808599,
+    2.4209723155762863,
+    2.420668423809868,
+    2.4204676161372016,
+    2.4203712005726827,
+    2.420380068704988,
+    2.420493289903554,
+    2.420710759315692,
+    2.4210321100509016,
+    2.421456711924972,
+    2.4219836715842935,
+    2.4226118326498916,
+    2.423339777335187,
+    2.424165828339519,
+    2.4250880518614637,
+    2.4261042616582924,
+    2.4272120230151923,
+    2.4284086587121188,
+    2.4296915368281606,
+    2.4310582510675482,
+    2.432504302921731,
+    2.4340261048245524,
+    2.435619866098322,
+    2.4372816029434006,
+    2.4390071492364305,
+    2.440792167926676,
+    2.442632162995385,
+    2.444522492162555,
+    2.4464583798729533,
+    2.448434930859836,
+    2.450447144296273,
+    2.4524899276311465,
+    2.4545581275894874,
+    2.4566467249172272,
+    2.4587500444816213,
+    2.460862775988026,
+    2.462979593476627,
+    2.4650951703383326,
+    2.4672041949249874,
+    2.469301385189656,
+    2.4713815036357163,
+    2.4734393717724394,
+    2.475469884333746,
+    2.477468023387656,
+    2.4794288713654966,
+    2.4813476246082624,
+    2.483219605457596,
+    2.485039458918412,
+    2.4868033804993956,
+    2.4885072430697144,
+    2.4901469959998734,
+    2.4917187690881257,
+    2.4932188815730543,
+    2.4946438498431105,
+    2.49599039508062,
+    2.4972554500082174,
+    2.498436164504503,
+    2.4995299110592493,
+    2.5005342891380167,
+    2.501447129364351,
+    2.502266495991754,
+    2.5029895274516014,
+    2.5036154604052916,
+    2.5041435094498383,
+    2.5045727041834747,
+    2.504902313989728,
+    2.5051318472519433,
+    2.505261050325141,
+    2.5052899056087403,
+    2.5052186294050998,
+    2.5050476691328045,
+    2.5047776999609392,
+    2.504409621696276,
+    2.5039445544003622,
+    2.503383834013878,
+    2.5027281389852822,
+    2.5019794867872425,
+    2.501140507540692,
+    2.500213360279151,
+    2.4992003931281483,
+    2.498104137409639,
+    2.4969273011438373,
+    2.495672762685901,
+    2.494343563852958,
+    2.492942903234649,
+    2.491474128676318,
+    2.4899407305330525,
+    2.488346334061092,
+    2.486694691617208,
+    2.484989340356667,
+    2.4832342121680133,
+    2.4814339161814924,
+    2.479592639951934,
+    2.4777146583088454,
+    2.4758043245876524,
+    2.473866062448869,
+    2.471904357327426,
+    2.4699237475913485,
+    2.4679288158413883,
+    2.46592417988613,
+    2.4639144839871556,
+    2.461904389322368,
+    2.459898565346157,
+    2.4579017367770137,
+    2.455918764942065,
+    2.453954173148493,
+    2.4520125671214457,
+    2.4500985125535784,
+    2.448216525842054,
+    2.4463710641219887,
+    2.444566515822548,
+    2.4428071908357323,
+    2.4410973107465668,
+    2.439440999504861,
+    2.4378422733498937,
+    2.4363050317847414,
+    2.434833047645694,
+    2.43343006971473,
+    2.432100619437783,
+    2.430846995094142,
+    2.429672363555455,
+    2.428579719706815,
+    2.427571878080977,
+    2.426651464455842,
+    2.4258209081907864,
+    2.425082434921927,
+    2.4244380590290993,
+    2.423889577520045,
+    2.4234385632168207,
+    2.423086359460975,
+    2.4228340742719734,
+    2.4226825762746023,
+    2.4226340587576405,
+    2.4226874437485133,
+    2.422842736063271,
+    2.4230997975159894,
+    2.4234582347462745,
+    2.4239173992440146,
+    2.424476386499562,
+    2.4251340367353267,
+    2.425888936234961,
+    2.4267394193027876,
+    2.427683571153293,
+    2.428719231244257,
+    2.429843998149402,
+    2.431055234079425,
+    2.432351101512926,
+    2.4337276833933763,
+    2.4351813534956155,
+    2.43670858740822,
+    2.4383056624813797,
+    2.4399686671928897,
+    2.4416935114716893,
+    2.4434759376701103,
+    2.4453115320216026,
+    2.447195736733967,
+    2.4491238624275873,
+    2.451091101018784,
+    2.4530925394872574,
+    2.4551231733170757,
+    2.4571781139809588,
+    2.4592520095527344,
+    2.4613395134290164,
+    2.4634353978388575,
+    2.46553441673564,
+    2.467631320555607,
+    2.469720870921841,
+    2.4717978552593056,
+    2.4738571012373916,
+    2.4758934909645545,
+    2.477901974977891,
+    2.479877585675147,
+    2.481815450926397,
+    2.4837108064957207,
+    2.485558618881201
+  ],
+  "mass_min": 1.0,
+  "mass_max": 1.0,
+  "mass_midrange": 1.0,
+  "mass_amplitude": 0.0,
+  "mass_time_mean": 1.0,
+  "flux_in_min": 1.0,
+  "flux_in_max": 1.0,
+  "flux_in_midrange": 1.0,
+  "flux_in_amplitude": 0.0,
+  "flux_in_time_mean": 1.0,
+  "flux_out_min": 1.0,
+  "flux_out_max": 1.0,
+  "flux_out_midrange": 1.0,
+  "flux_out_amplitude": 0.0,
+  "flux_out_time_mean": 1.0,
+  "flux_relative": 0.0,
+  "source": "https://wwwold.mathematik.tu-dortmund.de/~featflow/media/dfg_bench2new_2d/20141124_dfg_bench2_official.zip",
+  "archive_sha256": "08db59605b2f60efc512d2bbc79628ea5d1cc720c0986938a16920884a0215be",
+  "selection": "lv6 dt4 last complete lift-minimum-to-minimum cycle; force columns 2,4,5; pressure columns 2,7-12; physical time multiplied by 10",
+  "note": "Newer official high-order line-integration data; do not mix older website table with this reference."
+}

tests/validation/data/dfg_re100/source_manifest.json

@@ -0,0 +1,116 @@
+{
+  "commit": "47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5",
+  "production_sha256": {
+    "src/CelerisLab/__init__.py": "21cc469bd6e970a1f9ff5970c8eea4602065e41fc881854759ee565ee55c10c6",
+    "src/CelerisLab/body/__init__.py": "e4e0e8f0769a94428d5b02fac790443997556cfc5515a5f1ed1595d2c694d69d",
+    "src/CelerisLab/body/action_smoother.py": "b1a23ee6bbd4850b172cb8738c5c3546ee465a54189e6af9d8044f023d4687dc",
+    "src/CelerisLab/body/coupling/__init__.py": "c931f7eb318d48408c3e2febcc6917c88c8dcc1697cbf4329b1f340d94c35763",
+    "src/CelerisLab/body/coupling/soa_packer.py": "8ec26fd55ef74801bbd5ab6bd3d0a7dabc750cbd498aaac2ca4c61a0108ad3c9",
+    "src/CelerisLab/body/ddf_patch.py": "60fcfb7441e73bf290bcff41ca3a83da3e7ddad765d7247464e8f02e3043b4e2",
+    "src/CelerisLab/body/geometry/__init__.py": "ada88980ab672959aca9aa3a4c4aade3219524b2ab68110cd2a2f41e3e0ca90a",
+    "src/CelerisLab/body/geometry/base.py": "cceca14361ca2f34d34a3e43730e1a79b8ca81be00f0dd62e884db32b479419f",
+    "src/CelerisLab/body/geometry/circle.py": "7605ca7c21103f1d801ac73ab8dd18dffbdac425b2dee7b8858b283628a9e436",
+    "src/CelerisLab/body/jets.py": "4104c24af4651f4bb4f8d8fd36fb6d88f41c6e4469822f36e22e74c9ede31dac",
+    "src/CelerisLab/body/manager.py": "01b653d25bd9b4333379c8bfd12383596b686d2029f2e91b1c91b618540ec6b3",
+    "src/CelerisLab/body/objects.py": "8e63cee8faa9cc6243bcff428be628dfad73a02a9a3b8ece7d78c0a63406a08d",
+    "src/CelerisLab/body/preprocess/__init__.py": "dfd9d58bc07f626ab23a70053e31b88a6d9165d2e4c76a4e9f48abe674d3d2ec",
+    "src/CelerisLab/body/preprocess/flag_overlay.py": "831fc9e9736b758847759a085ff133eb49ec91318ace4fda76f85960138ee300",
+    "src/CelerisLab/body/registry.py": "fbb7327780adbdc6da15ebdb799980daf32df22ee974af0893d7409f177bbb68",
+    "src/CelerisLab/body/sync_plan.py": "3a15bfe9ebeb91482dbfc68850761ef39c872cc12d2efc795babd0c01e392323",
+    "src/CelerisLab/common/__init__.py": "d4dbde78bc97fccc9eabac9abae8f0f15c112c10ff8e9b45bd9401d4723df1ef",
+    "src/CelerisLab/common/_types.py": "aad2ff1a9033d4fd68783ad7123afb47201ddc58785cb6dc0246947037072a5f",
+    "src/CelerisLab/common/checkpoint.py": "b2f960aecd5ac27a3a7130b809ffad4709c44d522edef0359ab952654625410b",
+    "src/CelerisLab/common/pathline.py": "a1b6efcc0a076d0f1dc8d49a595779426d7d5626b00072a59bac710ff02ad046",
+    "src/CelerisLab/common/preprocess.py": "7fc78a08857cf8e98519500b094a1d896af189b6b15658ff9b388ed51bed4515",
+    "src/CelerisLab/common/render.py": "0e055476257eb2bcd11e733766e52195de02f8fad228f570a3511a7c831bdc29",
+    "src/CelerisLab/common/streakline/__init__.py": "13ff1cfa416fdd721a08a659e76b5a5a99b8d3201add10a83b7b0bc654bc1eec",
+    "src/CelerisLab/common/streakline/_config.py": "a9185ef68e9c7b5d5a1dc0b04a28341e2ad91d13c145a7c106636a878cf0a605",
+    "src/CelerisLab/common/streakline/_integrate.py": "3fbd4368b2c8d0e7fda40dc65aac5df02b085008b4c68acd83cb7ae45d33a13b",
+    "src/CelerisLab/common/streakline/_render.py": "236cd637414ee318c4d66f9aed32110d8dc42a100d65fde06172f7cd0ba62c65",
+    "src/CelerisLab/common/streakline/_streakline.py": "1a3e08daa706b5f83cdeb337e9a945a36ddf7583c40af1d0bb9d825ac954df37",
+    "src/CelerisLab/config.py": "2f1bba7515f5bc08cb3a8d140bcafe7fa0689e01b3b06494697bb2956a71530e",
+    "src/CelerisLab/configs/CONFIG.md": "41cafdbe525366c08aa9d0416ea026ccf0026ba7123b8a9b65f968a80d928e3a",
+    "src/CelerisLab/configs/config_body.json": "055f261f0828ca1d4f2bb694bd15b72bbcabe31d552326208f28e37f761adab1",
+    "src/CelerisLab/configs/config_lbm.json": "17e4311f1f250c6ece59177a8b6ff856b69a016c7c10267e774b030a0a390865",
+    "src/CelerisLab/configs/config_lbm_flume_mrt_zouhe.json": "8f593321db05a563a42931603a228d61cbcba06dfbe113262d5cd0ad93c0fe79",
+    "src/CelerisLab/configs/config_lbm_flume_trt_regularized.json": "f4bba8e6a4339fc3c4eb14f6e2f57c4b22993c73703c773d853edfbe34736e56",
+    "src/CelerisLab/configs/config_lbm_three_cylinder_triangle.json": "eeb8e7149dbe4aa6b3a9562802735aa950750dfbe2d972a0727e788573afa817",
+    "src/CelerisLab/cuda/__init__.py": "0d01dff9f0399c8baba822aab5ca0122a135f011008ed8a499ebd9f7e5b06aea",
+    "src/CelerisLab/cuda/compiler_v2.py": "c4cc3ec4445c8d75e6a253dd87963877117a722bc73b9d4811806fc6045d08c1",
+    "src/CelerisLab/cuda/context.py": "013dd4ec6d01d2fbd11afdfb103eac6a273bd9f798d663d761b851eeb27885e9",
+    "src/CelerisLab/cuda/obs_layout.py": "5ffb2c3aeeaaaaeb64f2cd8febfbcf4ad9919e5414020ad753cfade0683b21c2",
+    "src/CelerisLab/lbm/__init__.py": "7f60fb3aaae81501b510d08c2fedff44fcb05657260eb10b6f5cc91e7204306b",
+    "src/CelerisLab/lbm/block_stepper.py": "00b313daddc6e63d8ab8423c9190ca7e92ad291de20a2f6b80eb53e3f2d85ec7",
+    "src/CelerisLab/lbm/blocks.py": "1e644c5719993f418185ad4a10b6950a30636fc912befd9eab41ff48b4cf85cd",
+    "src/CelerisLab/lbm/channel.py": "2d5b5ff43414cb35633ac4ce90ce2c3a3353950b25c168e1c2ab58d498dd8b0f",
+    "src/CelerisLab/lbm/curved_links.py": "b6957fa13d13f864dba4f7364e93a39905d062dbb3ee9b3f5191add81fff0732",
+    "src/CelerisLab/lbm/descriptors.py": "c9f025dfc14da43b989047b84cc22bd05d5cc92c9ec695c1da68fdfb788e3f2b",
+    "src/CelerisLab/lbm/equilibrium.py": "2fb2f2bb7ff326d23abce80fc2a4c0907b17740ea3e29fe6ae079c5044783826",
+    "src/CelerisLab/lbm/field.py": "0e727eacc3706add4d8d15c418e9b631050871f831698d3ac689dd5f2c701b9b",
+    "src/CelerisLab/lbm/hierarchy.py": "f234975074f6b6d289d573ff6ea6832feae0e27cb04f80b37fe81792b13a12e7",
+    "src/CelerisLab/lbm/initializers.py": "dacf9586963e862d6d0a98ef1a066036be6f91e9d423720661d8ef2c1fb2df6a",
+    "src/CelerisLab/lbm/kernels/blocks.cu": "e3b9606f861bfea8c4b0f34e8356bb230081f899c584f67be89178d889099b17",
+    "src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh": "aba3748c564ed91b161d268ef8b737b7c7be96a823063d201f3da684dd845233",
+    "src/CelerisLab/lbm/kernels/boundary/curved_boundary.cuh": "aeb16bfd483514fb03e15a2fed437d673c3b8b2fe68ab83edf4244ef593878ea",
+    "src/CelerisLab/lbm/kernels/boundary/esopull_curved_helpers.cuh": "b9b474bc4edaff05681c5443235472d04072d0d775fc0eb9c5967faef52b0258",
+    "src/CelerisLab/lbm/kernels/boundary/ibm_kernels.cuh": "e6909b547dc8d64cdd8a76a90660243a8f5fee535caa161b52b4dd8b7c29934b",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/channel_stabilized.cuh": "efe077afcbb6b731256a987c2fd79a1dbc470c902f87070d60421ea6cb13d4e2",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/common.cuh": "59b38a89439878e00de7de16b9c7b03feff443fcc5a421593f7d2536dcbb7ab8",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/equilibrium.cuh": "f0964c811b2c41a5aaa1c6d99c490b06e850c634eb8749f1f05ef61f93c2de1d",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/regularized.cuh": "943d4a571818bacb89c70bccfaac202691121e3e385803ce6309123a0a03ddaf",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/zou_he_local.cuh": "e73e66b853b1e4e0e30d8f4185d832aa414ce670a1813ce2b5459f2887f71f89",
+    "src/CelerisLab/lbm/kernels/boundary/inlet_outlet.cuh": "de8b213a6da42c2e5e9f8966199fc66ea4541b85144a90f3f5dbad3c6ab63b97",
+    "src/CelerisLab/lbm/kernels/boundary/outlet/pressure_neq.cuh": "c8b2adc66d12fe602076ddc9693c5c208711fcaddb58aa9a31401a269f51c4c4",
+    "src/CelerisLab/lbm/kernels/boundary/planar_velocity.cuh": "501b2759d0e2e4280cb6fdd127d66a50fa89f6de48fc2cae9a7eafb11ae90bed",
+    "src/CelerisLab/lbm/kernels/config.h": "7d1124125e0b354558297934460e64c760fb68cdd7bf6d01a1238bb2eaab8c37",
+    "src/CelerisLab/lbm/kernels/config/config_grid.h": "1db3e030ebdc116cd7f02362d1fd077daac55e8f0b4329f6285948c224d3ca5d",
+    "src/CelerisLab/lbm/kernels/config/config_method.h": "4374e04a130f90f97d17cc20c40d2f9ea0ed0dd8603817d07f215d02266e247c",
+    "src/CelerisLab/lbm/kernels/config/config_objects.h": "db793af052c756561a190a3ff2cb93ee5fe28462ee4f97c6af80882b4dc8b072",
+    "src/CelerisLab/lbm/kernels/config/config_obs.h": "c11e2249359145b5c4cbaaf0487506fa77c1f4b82c7c169ddc574aa0babcd727",
+    "src/CelerisLab/lbm/kernels/config/config_physics.h": "8ac0990b168e1058588676cc05c80df8634cfd671dbc92a89246c3c372d83a4e",
+    "src/CelerisLab/lbm/kernels/core/descriptors.cuh": "8440bc14857da9007982650f828a23608eec9f7cc4bf772e8c8dc1239cb45356",
+    "src/CelerisLab/lbm/kernels/core/flags.cuh": "0025705f24710330d9d9195f713879bfb650a85a3449baef0e196b831af104ab",
+    "src/CelerisLab/lbm/kernels/core/layout.cuh": "1fbd820021c89ac3efcde24145996cb0147769e607e498e61f308ed3e0b13706",
+    "src/CelerisLab/lbm/kernels/core/obs.cuh": "b742b77c9f61e1f020fe9922c3838ae9d855e17b381154958b58f10a22938922",
+    "src/CelerisLab/lbm/kernels/core/params.cuh": "6e18bf730d936d940e3d85452f0714a2eaae31bf41002f9fd520ec5b372878dd",
+    "src/CelerisLab/lbm/kernels/core/precision.cuh": "849df14b4cce50af52399e4b8f3e997196f3ae5f1163f68e28df3e6b2507a8c4",
+    "src/CelerisLab/lbm/kernels/kernel_v2.cu": "c7bd3fb705b77261ba2441c52a24a7e6a7865c165cba088a1b75201c4d29c408",
+    "src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh": "aee444eb60c47a3fd518dca975d1af05425a0818a0ec6c3fb3f6150f60c2cb44",
+    "src/CelerisLab/lbm/kernels/operators/collision_srt.cuh": "6b7ff842f392c8f853d2d1a97cc69bfe50f21283f73e659404da07cf6c95f5e1",
+    "src/CelerisLab/lbm/kernels/operators/collision_trt.cuh": "9725b9840b9ad8b4ec7be61ff5d06d3151e86d1a87585c13f2fbf934d0d2c5f6",
+    "src/CelerisLab/lbm/kernels/operators/forcing_guo.cuh": "969a430c1af03f32e3bd0df95d805e64f56334717c44dd4ba208ba503062a24b",
+    "src/CelerisLab/lbm/kernels/operators/helpers.cuh": "0b3a4d95d7d19f5fec09e5423b8f4cc99f8a226b541967425967330da5a051ec",
+    "src/CelerisLab/lbm/kernels/operators/macro.cuh": "2be20eb2b7d7f5522daaad7f4b57b1ff8cdf883efe4a67846a96783b2fe846a3",
+    "src/CelerisLab/lbm/kernels/operators/turbulence_smag.cuh": "ff7f000297ee0817e255ad9fcb6a13b647e3ddd5b206518d35f38d2895a9f501",
+    "src/CelerisLab/lbm/kernels/refinement.cu": "129f3d91da93b4324f88a0da73a4da1e5390f72c14cce574ee63306712da9895",
+    "src/CelerisLab/lbm/kernels/step/aux_kernels.cu": "578a5b20a157241ee678547059abb1372dd5a589ad06bb2c15b987eb01ca1013",
+    "src/CelerisLab/lbm/kernels/step/block_channel.cu": "befbf6d757e532a59c166d41f920d16d71557cd3a574478e7aa83da737c48cbd",
+    "src/CelerisLab/lbm/kernels/step/curved_double.cu": "bfa7a1df1e939020314d49755045f7a85fe4a19cc1c68dbe62653c66c5e7c2b3",
+    "src/CelerisLab/lbm/kernels/step/esopull_macro.cu": "c73116020ca217ede1e4de27c8269c31f9d91e33433ac8252a21e9b14187c50f",
+    "src/CelerisLab/lbm/kernels/step/esopull_sync.cu": "c0ff3e72900c3ab89a05dd03684e96d3b740cf1672d565727fc2af165933d087",
+    "src/CelerisLab/lbm/kernels/step/init_flow.cu": "305be3bb12fc0d039acec1701baabcf0e8a35e683c49c34f2af8ca9867d2b29c",
+    "src/CelerisLab/lbm/kernels/step/one_step_double.cu": "f9f8242b38e8e0cf8425b75e2b0215bc15363c62eaa0fb43c761b8e0931bfdc9",
+    "src/CelerisLab/lbm/kernels/step/one_step_esopull.cu": "aa223b2744ae6c9125673df311a96ce588032c1ba58bebb11fd33ecb990f8ed4",
+    "src/CelerisLab/lbm/kernels/streaming/esopull_semantic_helpers.cuh": "dd5a9091628b079efd97e53fd49e52d09dcee88529235db77a5648e02937af6e",
+    "src/CelerisLab/lbm/kernels/streaming/esopull_single_buffer.cuh": "77ed4ab9a0ebc5ff62a91a7c39066ef395391ae36e9040a69fb2bd8b1626dc79",
+    "src/CelerisLab/lbm/kernels/streaming/pull_double_buffer.cuh": "b65d741ed78b7b62f99531ba0c0ce1936663d8dc9640e7d7d16c707f1baf6549",
+    "src/CelerisLab/lbm/population_storage.py": "a6a86c29c52c892edfe4802fb773c387212e3f4cc86708d80e154986ddf88044",
+    "src/CelerisLab/lbm/raw_exchange.py": "8694202466eba848563e298656353b0f5db94e21a01fa05f51aad6dc2b9fd311",
+    "src/CelerisLab/lbm/refinement.py": "08c2a0557276a98cbcaa459d9bb7a9735d8bbbf64a6114c57af2fda61481a687",
+    "src/CelerisLab/lbm/refinement_stepper.py": "f2bb34fa90805ad21c3e7cf9f9a67d88d3ae3e249b4517d3898567b4bf9e364d",
+    "src/CelerisLab/lbm/refinement_tables.py": "fdd344024b89df3a83174dc810ca5e32db5e19eb8b1609f8f73b0821f56546b6",
+    "src/CelerisLab/lbm/stepper.py": "0dcb7cd7ae87e33e3dc67a38b6aa43042fbe673e37d9ad165342ee1d695933ce",
+    "src/CelerisLab/lbm/synchronized_remap.py": "7400fc4beaffd4649d6a9c524d38ce5fd6fa7eeb32eb5332e4030a4b6648f197",
+    "src/CelerisLab/simulation.py": "0a03904dc6de88667d7af699005a289adf8b09f0976cb0d099b041f4830a8f4c"
+  },
+  "experiment": "AMR-channel development source; GPU/flow validation not inherited from d79",
+  "source_origin_commits": [
+    "5338864748ac67fe2544487c059ba149ebf916da",
+    "47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5"
+  ],
+  "historical_checkpoints": {
+    "production": "d79e3f6f0a424f5bcc3a1a6d610313e49daf97a2",
+    "re100_jet": "69e9b78de9483a5ff21e2d0c74311d240bfaf88c",
+    "les_diagnostic": "20ae6770924df15f31305aad56e44a54a3dd4d88"
+  }
+}

tests/validation/data/jet_circle_cv/source_manifest.json

@@ -0,0 +1,116 @@
+{
+  "revision": "47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5",
+  "files": {
+    "src/CelerisLab/__init__.py": "21cc469bd6e970a1f9ff5970c8eea4602065e41fc881854759ee565ee55c10c6",
+    "src/CelerisLab/body/__init__.py": "e4e0e8f0769a94428d5b02fac790443997556cfc5515a5f1ed1595d2c694d69d",
+    "src/CelerisLab/body/action_smoother.py": "b1a23ee6bbd4850b172cb8738c5c3546ee465a54189e6af9d8044f023d4687dc",
+    "src/CelerisLab/body/coupling/__init__.py": "c931f7eb318d48408c3e2febcc6917c88c8dcc1697cbf4329b1f340d94c35763",
+    "src/CelerisLab/body/coupling/soa_packer.py": "8ec26fd55ef74801bbd5ab6bd3d0a7dabc750cbd498aaac2ca4c61a0108ad3c9",
+    "src/CelerisLab/body/ddf_patch.py": "60fcfb7441e73bf290bcff41ca3a83da3e7ddad765d7247464e8f02e3043b4e2",
+    "src/CelerisLab/body/geometry/__init__.py": "ada88980ab672959aca9aa3a4c4aade3219524b2ab68110cd2a2f41e3e0ca90a",
+    "src/CelerisLab/body/geometry/base.py": "cceca14361ca2f34d34a3e43730e1a79b8ca81be00f0dd62e884db32b479419f",
+    "src/CelerisLab/body/geometry/circle.py": "7605ca7c21103f1d801ac73ab8dd18dffbdac425b2dee7b8858b283628a9e436",
+    "src/CelerisLab/body/jets.py": "4104c24af4651f4bb4f8d8fd36fb6d88f41c6e4469822f36e22e74c9ede31dac",
+    "src/CelerisLab/body/manager.py": "01b653d25bd9b4333379c8bfd12383596b686d2029f2e91b1c91b618540ec6b3",
+    "src/CelerisLab/body/objects.py": "8e63cee8faa9cc6243bcff428be628dfad73a02a9a3b8ece7d78c0a63406a08d",
+    "src/CelerisLab/body/preprocess/__init__.py": "dfd9d58bc07f626ab23a70053e31b88a6d9165d2e4c76a4e9f48abe674d3d2ec",
+    "src/CelerisLab/body/preprocess/flag_overlay.py": "831fc9e9736b758847759a085ff133eb49ec91318ace4fda76f85960138ee300",
+    "src/CelerisLab/body/registry.py": "fbb7327780adbdc6da15ebdb799980daf32df22ee974af0893d7409f177bbb68",
+    "src/CelerisLab/body/sync_plan.py": "3a15bfe9ebeb91482dbfc68850761ef39c872cc12d2efc795babd0c01e392323",
+    "src/CelerisLab/common/__init__.py": "d4dbde78bc97fccc9eabac9abae8f0f15c112c10ff8e9b45bd9401d4723df1ef",
+    "src/CelerisLab/common/_types.py": "aad2ff1a9033d4fd68783ad7123afb47201ddc58785cb6dc0246947037072a5f",
+    "src/CelerisLab/common/checkpoint.py": "b2f960aecd5ac27a3a7130b809ffad4709c44d522edef0359ab952654625410b",
+    "src/CelerisLab/common/pathline.py": "a1b6efcc0a076d0f1dc8d49a595779426d7d5626b00072a59bac710ff02ad046",
+    "src/CelerisLab/common/preprocess.py": "7fc78a08857cf8e98519500b094a1d896af189b6b15658ff9b388ed51bed4515",
+    "src/CelerisLab/common/render.py": "0e055476257eb2bcd11e733766e52195de02f8fad228f570a3511a7c831bdc29",
+    "src/CelerisLab/common/streakline/__init__.py": "13ff1cfa416fdd721a08a659e76b5a5a99b8d3201add10a83b7b0bc654bc1eec",
+    "src/CelerisLab/common/streakline/_config.py": "a9185ef68e9c7b5d5a1dc0b04a28341e2ad91d13c145a7c106636a878cf0a605",
+    "src/CelerisLab/common/streakline/_integrate.py": "3fbd4368b2c8d0e7fda40dc65aac5df02b085008b4c68acd83cb7ae45d33a13b",
+    "src/CelerisLab/common/streakline/_render.py": "236cd637414ee318c4d66f9aed32110d8dc42a100d65fde06172f7cd0ba62c65",
+    "src/CelerisLab/common/streakline/_streakline.py": "1a3e08daa706b5f83cdeb337e9a945a36ddf7583c40af1d0bb9d825ac954df37",
+    "src/CelerisLab/config.py": "2f1bba7515f5bc08cb3a8d140bcafe7fa0689e01b3b06494697bb2956a71530e",
+    "src/CelerisLab/configs/CONFIG.md": "41cafdbe525366c08aa9d0416ea026ccf0026ba7123b8a9b65f968a80d928e3a",
+    "src/CelerisLab/configs/config_body.json": "055f261f0828ca1d4f2bb694bd15b72bbcabe31d552326208f28e37f761adab1",
+    "src/CelerisLab/configs/config_lbm.json": "17e4311f1f250c6ece59177a8b6ff856b69a016c7c10267e774b030a0a390865",
+    "src/CelerisLab/configs/config_lbm_flume_mrt_zouhe.json": "8f593321db05a563a42931603a228d61cbcba06dfbe113262d5cd0ad93c0fe79",
+    "src/CelerisLab/configs/config_lbm_flume_trt_regularized.json": "f4bba8e6a4339fc3c4eb14f6e2f57c4b22993c73703c773d853edfbe34736e56",
+    "src/CelerisLab/configs/config_lbm_three_cylinder_triangle.json": "eeb8e7149dbe4aa6b3a9562802735aa950750dfbe2d972a0727e788573afa817",
+    "src/CelerisLab/cuda/__init__.py": "0d01dff9f0399c8baba822aab5ca0122a135f011008ed8a499ebd9f7e5b06aea",
+    "src/CelerisLab/cuda/compiler_v2.py": "c4cc3ec4445c8d75e6a253dd87963877117a722bc73b9d4811806fc6045d08c1",
+    "src/CelerisLab/cuda/context.py": "013dd4ec6d01d2fbd11afdfb103eac6a273bd9f798d663d761b851eeb27885e9",
+    "src/CelerisLab/cuda/obs_layout.py": "5ffb2c3aeeaaaaeb64f2cd8febfbcf4ad9919e5414020ad753cfade0683b21c2",
+    "src/CelerisLab/lbm/__init__.py": "7f60fb3aaae81501b510d08c2fedff44fcb05657260eb10b6f5cc91e7204306b",
+    "src/CelerisLab/lbm/block_stepper.py": "00b313daddc6e63d8ab8423c9190ca7e92ad291de20a2f6b80eb53e3f2d85ec7",
+    "src/CelerisLab/lbm/blocks.py": "1e644c5719993f418185ad4a10b6950a30636fc912befd9eab41ff48b4cf85cd",
+    "src/CelerisLab/lbm/channel.py": "2d5b5ff43414cb35633ac4ce90ce2c3a3353950b25c168e1c2ab58d498dd8b0f",
+    "src/CelerisLab/lbm/curved_links.py": "b6957fa13d13f864dba4f7364e93a39905d062dbb3ee9b3f5191add81fff0732",
+    "src/CelerisLab/lbm/descriptors.py": "c9f025dfc14da43b989047b84cc22bd05d5cc92c9ec695c1da68fdfb788e3f2b",
+    "src/CelerisLab/lbm/equilibrium.py": "2fb2f2bb7ff326d23abce80fc2a4c0907b17740ea3e29fe6ae079c5044783826",
+    "src/CelerisLab/lbm/field.py": "0e727eacc3706add4d8d15c418e9b631050871f831698d3ac689dd5f2c701b9b",
+    "src/CelerisLab/lbm/hierarchy.py": "f234975074f6b6d289d573ff6ea6832feae0e27cb04f80b37fe81792b13a12e7",
+    "src/CelerisLab/lbm/initializers.py": "dacf9586963e862d6d0a98ef1a066036be6f91e9d423720661d8ef2c1fb2df6a",
+    "src/CelerisLab/lbm/kernels/blocks.cu": "e3b9606f861bfea8c4b0f34e8356bb230081f899c584f67be89178d889099b17",
+    "src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh": "aba3748c564ed91b161d268ef8b737b7c7be96a823063d201f3da684dd845233",
+    "src/CelerisLab/lbm/kernels/boundary/curved_boundary.cuh": "aeb16bfd483514fb03e15a2fed437d673c3b8b2fe68ab83edf4244ef593878ea",
+    "src/CelerisLab/lbm/kernels/boundary/esopull_curved_helpers.cuh": "b9b474bc4edaff05681c5443235472d04072d0d775fc0eb9c5967faef52b0258",
+    "src/CelerisLab/lbm/kernels/boundary/ibm_kernels.cuh": "e6909b547dc8d64cdd8a76a90660243a8f5fee535caa161b52b4dd8b7c29934b",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/channel_stabilized.cuh": "efe077afcbb6b731256a987c2fd79a1dbc470c902f87070d60421ea6cb13d4e2",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/common.cuh": "59b38a89439878e00de7de16b9c7b03feff443fcc5a421593f7d2536dcbb7ab8",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/equilibrium.cuh": "f0964c811b2c41a5aaa1c6d99c490b06e850c634eb8749f1f05ef61f93c2de1d",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/regularized.cuh": "943d4a571818bacb89c70bccfaac202691121e3e385803ce6309123a0a03ddaf",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/zou_he_local.cuh": "e73e66b853b1e4e0e30d8f4185d832aa414ce670a1813ce2b5459f2887f71f89",
+    "src/CelerisLab/lbm/kernels/boundary/inlet_outlet.cuh": "de8b213a6da42c2e5e9f8966199fc66ea4541b85144a90f3f5dbad3c6ab63b97",
+    "src/CelerisLab/lbm/kernels/boundary/outlet/pressure_neq.cuh": "c8b2adc66d12fe602076ddc9693c5c208711fcaddb58aa9a31401a269f51c4c4",
+    "src/CelerisLab/lbm/kernels/boundary/planar_velocity.cuh": "501b2759d0e2e4280cb6fdd127d66a50fa89f6de48fc2cae9a7eafb11ae90bed",
+    "src/CelerisLab/lbm/kernels/config.h": "7d1124125e0b354558297934460e64c760fb68cdd7bf6d01a1238bb2eaab8c37",
+    "src/CelerisLab/lbm/kernels/config/config_grid.h": "1db3e030ebdc116cd7f02362d1fd077daac55e8f0b4329f6285948c224d3ca5d",
+    "src/CelerisLab/lbm/kernels/config/config_method.h": "4374e04a130f90f97d17cc20c40d2f9ea0ed0dd8603817d07f215d02266e247c",
+    "src/CelerisLab/lbm/kernels/config/config_objects.h": "db793af052c756561a190a3ff2cb93ee5fe28462ee4f97c6af80882b4dc8b072",
+    "src/CelerisLab/lbm/kernels/config/config_obs.h": "c11e2249359145b5c4cbaaf0487506fa77c1f4b82c7c169ddc574aa0babcd727",
+    "src/CelerisLab/lbm/kernels/config/config_physics.h": "8ac0990b168e1058588676cc05c80df8634cfd671dbc92a89246c3c372d83a4e",
+    "src/CelerisLab/lbm/kernels/core/descriptors.cuh": "8440bc14857da9007982650f828a23608eec9f7cc4bf772e8c8dc1239cb45356",
+    "src/CelerisLab/lbm/kernels/core/flags.cuh": "0025705f24710330d9d9195f713879bfb650a85a3449baef0e196b831af104ab",
+    "src/CelerisLab/lbm/kernels/core/layout.cuh": "1fbd820021c89ac3efcde24145996cb0147769e607e498e61f308ed3e0b13706",
+    "src/CelerisLab/lbm/kernels/core/obs.cuh": "b742b77c9f61e1f020fe9922c3838ae9d855e17b381154958b58f10a22938922",
+    "src/CelerisLab/lbm/kernels/core/params.cuh": "6e18bf730d936d940e3d85452f0714a2eaae31bf41002f9fd520ec5b372878dd",
+    "src/CelerisLab/lbm/kernels/core/precision.cuh": "849df14b4cce50af52399e4b8f3e997196f3ae5f1163f68e28df3e6b2507a8c4",
+    "src/CelerisLab/lbm/kernels/kernel_v2.cu": "c7bd3fb705b77261ba2441c52a24a7e6a7865c165cba088a1b75201c4d29c408",
+    "src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh": "aee444eb60c47a3fd518dca975d1af05425a0818a0ec6c3fb3f6150f60c2cb44",
+    "src/CelerisLab/lbm/kernels/operators/collision_srt.cuh": "6b7ff842f392c8f853d2d1a97cc69bfe50f21283f73e659404da07cf6c95f5e1",
+    "src/CelerisLab/lbm/kernels/operators/collision_trt.cuh": "9725b9840b9ad8b4ec7be61ff5d06d3151e86d1a87585c13f2fbf934d0d2c5f6",
+    "src/CelerisLab/lbm/kernels/operators/forcing_guo.cuh": "969a430c1af03f32e3bd0df95d805e64f56334717c44dd4ba208ba503062a24b",
+    "src/CelerisLab/lbm/kernels/operators/helpers.cuh": "0b3a4d95d7d19f5fec09e5423b8f4cc99f8a226b541967425967330da5a051ec",
+    "src/CelerisLab/lbm/kernels/operators/macro.cuh": "2be20eb2b7d7f5522daaad7f4b57b1ff8cdf883efe4a67846a96783b2fe846a3",
+    "src/CelerisLab/lbm/kernels/operators/turbulence_smag.cuh": "ff7f000297ee0817e255ad9fcb6a13b647e3ddd5b206518d35f38d2895a9f501",
+    "src/CelerisLab/lbm/kernels/refinement.cu": "129f3d91da93b4324f88a0da73a4da1e5390f72c14cce574ee63306712da9895",
+    "src/CelerisLab/lbm/kernels/step/aux_kernels.cu": "578a5b20a157241ee678547059abb1372dd5a589ad06bb2c15b987eb01ca1013",
+    "src/CelerisLab/lbm/kernels/step/block_channel.cu": "befbf6d757e532a59c166d41f920d16d71557cd3a574478e7aa83da737c48cbd",
+    "src/CelerisLab/lbm/kernels/step/curved_double.cu": "bfa7a1df1e939020314d49755045f7a85fe4a19cc1c68dbe62653c66c5e7c2b3",
+    "src/CelerisLab/lbm/kernels/step/esopull_macro.cu": "c73116020ca217ede1e4de27c8269c31f9d91e33433ac8252a21e9b14187c50f",
+    "src/CelerisLab/lbm/kernels/step/esopull_sync.cu": "c0ff3e72900c3ab89a05dd03684e96d3b740cf1672d565727fc2af165933d087",
+    "src/CelerisLab/lbm/kernels/step/init_flow.cu": "305be3bb12fc0d039acec1701baabcf0e8a35e683c49c34f2af8ca9867d2b29c",
+    "src/CelerisLab/lbm/kernels/step/one_step_double.cu": "f9f8242b38e8e0cf8425b75e2b0215bc15363c62eaa0fb43c761b8e0931bfdc9",
+    "src/CelerisLab/lbm/kernels/step/one_step_esopull.cu": "aa223b2744ae6c9125673df311a96ce588032c1ba58bebb11fd33ecb990f8ed4",
+    "src/CelerisLab/lbm/kernels/streaming/esopull_semantic_helpers.cuh": "dd5a9091628b079efd97e53fd49e52d09dcee88529235db77a5648e02937af6e",
+    "src/CelerisLab/lbm/kernels/streaming/esopull_single_buffer.cuh": "77ed4ab9a0ebc5ff62a91a7c39066ef395391ae36e9040a69fb2bd8b1626dc79",
+    "src/CelerisLab/lbm/kernels/streaming/pull_double_buffer.cuh": "b65d741ed78b7b62f99531ba0c0ce1936663d8dc9640e7d7d16c707f1baf6549",
+    "src/CelerisLab/lbm/population_storage.py": "a6a86c29c52c892edfe4802fb773c387212e3f4cc86708d80e154986ddf88044",
+    "src/CelerisLab/lbm/raw_exchange.py": "8694202466eba848563e298656353b0f5db94e21a01fa05f51aad6dc2b9fd311",
+    "src/CelerisLab/lbm/refinement.py": "08c2a0557276a98cbcaa459d9bb7a9735d8bbbf64a6114c57af2fda61481a687",
+    "src/CelerisLab/lbm/refinement_stepper.py": "f2bb34fa90805ad21c3e7cf9f9a67d88d3ae3e249b4517d3898567b4bf9e364d",
+    "src/CelerisLab/lbm/refinement_tables.py": "fdd344024b89df3a83174dc810ca5e32db5e19eb8b1609f8f73b0821f56546b6",
+    "src/CelerisLab/lbm/stepper.py": "0dcb7cd7ae87e33e3dc67a38b6aa43042fbe673e37d9ad165342ee1d695933ce",
+    "src/CelerisLab/lbm/synchronized_remap.py": "7400fc4beaffd4649d6a9c524d38ce5fd6fa7eeb32eb5332e4030a4b6648f197",
+    "src/CelerisLab/simulation.py": "0a03904dc6de88667d7af699005a289adf8b09f0976cb0d099b041f4830a8f4c"
+  },
+  "experiment": "AMR-channel development source; GPU/flow validation not inherited from d79",
+  "source_origin_commits": [
+    "5338864748ac67fe2544487c059ba149ebf916da",
+    "47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5"
+  ],
+  "historical_checkpoints": {
+    "production": "d79e3f6f0a424f5bcc3a1a6d610313e49daf97a2",
+    "re100_jet": "69e9b78de9483a5ff21e2d0c74311d240bfaf88c",
+    "les_diagnostic": "20ae6770924df15f31305aad56e44a54a3dd4d88"
+  }
+}

tests/validation/data/les_flow/source_pins.json

@@ -0,0 +1,110 @@
+{
+  "base_commit": "cd3a657f9d3833e812c3cca173525f165b48710a",
+  "sources": {
+    "src/CelerisLab/__init__.py": "21cc469bd6e970a1f9ff5970c8eea4602065e41fc881854759ee565ee55c10c6",
+    "src/CelerisLab/body/__init__.py": "e4e0e8f0769a94428d5b02fac790443997556cfc5515a5f1ed1595d2c694d69d",
+    "src/CelerisLab/body/action_smoother.py": "b1a23ee6bbd4850b172cb8738c5c3546ee465a54189e6af9d8044f023d4687dc",
+    "src/CelerisLab/body/coupling/__init__.py": "c931f7eb318d48408c3e2febcc6917c88c8dcc1697cbf4329b1f340d94c35763",
+    "src/CelerisLab/body/coupling/soa_packer.py": "8ec26fd55ef74801bbd5ab6bd3d0a7dabc750cbd498aaac2ca4c61a0108ad3c9",
+    "src/CelerisLab/body/ddf_patch.py": "60fcfb7441e73bf290bcff41ca3a83da3e7ddad765d7247464e8f02e3043b4e2",
+    "src/CelerisLab/body/geometry/__init__.py": "ada88980ab672959aca9aa3a4c4aade3219524b2ab68110cd2a2f41e3e0ca90a",
+    "src/CelerisLab/body/geometry/base.py": "cceca14361ca2f34d34a3e43730e1a79b8ca81be00f0dd62e884db32b479419f",
+    "src/CelerisLab/body/geometry/circle.py": "7605ca7c21103f1d801ac73ab8dd18dffbdac425b2dee7b8858b283628a9e436",
+    "src/CelerisLab/body/jets.py": "4104c24af4651f4bb4f8d8fd36fb6d88f41c6e4469822f36e22e74c9ede31dac",
+    "src/CelerisLab/body/manager.py": "01b653d25bd9b4333379c8bfd12383596b686d2029f2e91b1c91b618540ec6b3",
+    "src/CelerisLab/body/objects.py": "8e63cee8faa9cc6243bcff428be628dfad73a02a9a3b8ece7d78c0a63406a08d",
+    "src/CelerisLab/body/preprocess/__init__.py": "dfd9d58bc07f626ab23a70053e31b88a6d9165d2e4c76a4e9f48abe674d3d2ec",
+    "src/CelerisLab/body/preprocess/flag_overlay.py": "831fc9e9736b758847759a085ff133eb49ec91318ace4fda76f85960138ee300",
+    "src/CelerisLab/body/registry.py": "fbb7327780adbdc6da15ebdb799980daf32df22ee974af0893d7409f177bbb68",
+    "src/CelerisLab/body/sync_plan.py": "3a15bfe9ebeb91482dbfc68850761ef39c872cc12d2efc795babd0c01e392323",
+    "src/CelerisLab/common/__init__.py": "d4dbde78bc97fccc9eabac9abae8f0f15c112c10ff8e9b45bd9401d4723df1ef",
+    "src/CelerisLab/common/_types.py": "aad2ff1a9033d4fd68783ad7123afb47201ddc58785cb6dc0246947037072a5f",
+    "src/CelerisLab/common/checkpoint.py": "b2f960aecd5ac27a3a7130b809ffad4709c44d522edef0359ab952654625410b",
+    "src/CelerisLab/common/pathline.py": "a1b6efcc0a076d0f1dc8d49a595779426d7d5626b00072a59bac710ff02ad046",
+    "src/CelerisLab/common/preprocess.py": "7fc78a08857cf8e98519500b094a1d896af189b6b15658ff9b388ed51bed4515",
+    "src/CelerisLab/common/render.py": "0e055476257eb2bcd11e733766e52195de02f8fad228f570a3511a7c831bdc29",
+    "src/CelerisLab/common/streakline/__init__.py": "13ff1cfa416fdd721a08a659e76b5a5a99b8d3201add10a83b7b0bc654bc1eec",
+    "src/CelerisLab/common/streakline/_config.py": "a9185ef68e9c7b5d5a1dc0b04a28341e2ad91d13c145a7c106636a878cf0a605",
+    "src/CelerisLab/common/streakline/_integrate.py": "3fbd4368b2c8d0e7fda40dc65aac5df02b085008b4c68acd83cb7ae45d33a13b",
+    "src/CelerisLab/common/streakline/_render.py": "236cd637414ee318c4d66f9aed32110d8dc42a100d65fde06172f7cd0ba62c65",
+    "src/CelerisLab/common/streakline/_streakline.py": "1a3e08daa706b5f83cdeb337e9a945a36ddf7583c40af1d0bb9d825ac954df37",
+    "src/CelerisLab/config.py": "2f1bba7515f5bc08cb3a8d140bcafe7fa0689e01b3b06494697bb2956a71530e",
+    "src/CelerisLab/configs/CONFIG.md": "41cafdbe525366c08aa9d0416ea026ccf0026ba7123b8a9b65f968a80d928e3a",
+    "src/CelerisLab/configs/config_body.json": "055f261f0828ca1d4f2bb694bd15b72bbcabe31d552326208f28e37f761adab1",
+    "src/CelerisLab/configs/config_lbm.json": "17e4311f1f250c6ece59177a8b6ff856b69a016c7c10267e774b030a0a390865",
+    "src/CelerisLab/configs/config_lbm_flume_mrt_zouhe.json": "8f593321db05a563a42931603a228d61cbcba06dfbe113262d5cd0ad93c0fe79",
+    "src/CelerisLab/configs/config_lbm_flume_trt_regularized.json": "f4bba8e6a4339fc3c4eb14f6e2f57c4b22993c73703c773d853edfbe34736e56",
+    "src/CelerisLab/configs/config_lbm_three_cylinder_triangle.json": "eeb8e7149dbe4aa6b3a9562802735aa950750dfbe2d972a0727e788573afa817",
+    "src/CelerisLab/cuda/__init__.py": "0d01dff9f0399c8baba822aab5ca0122a135f011008ed8a499ebd9f7e5b06aea",
+    "src/CelerisLab/cuda/compiler_v2.py": "c4cc3ec4445c8d75e6a253dd87963877117a722bc73b9d4811806fc6045d08c1",
+    "src/CelerisLab/cuda/context.py": "013dd4ec6d01d2fbd11afdfb103eac6a273bd9f798d663d761b851eeb27885e9",
+    "src/CelerisLab/cuda/obs_layout.py": "5ffb2c3aeeaaaaeb64f2cd8febfbcf4ad9919e5414020ad753cfade0683b21c2",
+    "src/CelerisLab/lbm/__init__.py": "7f60fb3aaae81501b510d08c2fedff44fcb05657260eb10b6f5cc91e7204306b",
+    "src/CelerisLab/lbm/block_stepper.py": "00b313daddc6e63d8ab8423c9190ca7e92ad291de20a2f6b80eb53e3f2d85ec7",
+    "src/CelerisLab/lbm/blocks.py": "1e644c5719993f418185ad4a10b6950a30636fc912befd9eab41ff48b4cf85cd",
+    "src/CelerisLab/lbm/channel.py": "2d5b5ff43414cb35633ac4ce90ce2c3a3353950b25c168e1c2ab58d498dd8b0f",
+    "src/CelerisLab/lbm/curved_links.py": "b6957fa13d13f864dba4f7364e93a39905d062dbb3ee9b3f5191add81fff0732",
+    "src/CelerisLab/lbm/descriptors.py": "c9f025dfc14da43b989047b84cc22bd05d5cc92c9ec695c1da68fdfb788e3f2b",
+    "src/CelerisLab/lbm/equilibrium.py": "2fb2f2bb7ff326d23abce80fc2a4c0907b17740ea3e29fe6ae079c5044783826",
+    "src/CelerisLab/lbm/field.py": "0e727eacc3706add4d8d15c418e9b631050871f831698d3ac689dd5f2c701b9b",
+    "src/CelerisLab/lbm/hierarchy.py": "f234975074f6b6d289d573ff6ea6832feae0e27cb04f80b37fe81792b13a12e7",
+    "src/CelerisLab/lbm/initializers.py": "dacf9586963e862d6d0a98ef1a066036be6f91e9d423720661d8ef2c1fb2df6a",
+    "src/CelerisLab/lbm/kernels/blocks.cu": "e3b9606f861bfea8c4b0f34e8356bb230081f899c584f67be89178d889099b17",
+    "src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh": "aba3748c564ed91b161d268ef8b737b7c7be96a823063d201f3da684dd845233",
+    "src/CelerisLab/lbm/kernels/boundary/curved_boundary.cuh": "aeb16bfd483514fb03e15a2fed437d673c3b8b2fe68ab83edf4244ef593878ea",
+    "src/CelerisLab/lbm/kernels/boundary/esopull_curved_helpers.cuh": "b9b474bc4edaff05681c5443235472d04072d0d775fc0eb9c5967faef52b0258",
+    "src/CelerisLab/lbm/kernels/boundary/ibm_kernels.cuh": "e6909b547dc8d64cdd8a76a90660243a8f5fee535caa161b52b4dd8b7c29934b",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/channel_stabilized.cuh": "efe077afcbb6b731256a987c2fd79a1dbc470c902f87070d60421ea6cb13d4e2",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/common.cuh": "59b38a89439878e00de7de16b9c7b03feff443fcc5a421593f7d2536dcbb7ab8",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/equilibrium.cuh": "f0964c811b2c41a5aaa1c6d99c490b06e850c634eb8749f1f05ef61f93c2de1d",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/regularized.cuh": "943d4a571818bacb89c70bccfaac202691121e3e385803ce6309123a0a03ddaf",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/zou_he_local.cuh": "e73e66b853b1e4e0e30d8f4185d832aa414ce670a1813ce2b5459f2887f71f89",
+    "src/CelerisLab/lbm/kernels/boundary/inlet_outlet.cuh": "de8b213a6da42c2e5e9f8966199fc66ea4541b85144a90f3f5dbad3c6ab63b97",
+    "src/CelerisLab/lbm/kernels/boundary/outlet/pressure_neq.cuh": "c8b2adc66d12fe602076ddc9693c5c208711fcaddb58aa9a31401a269f51c4c4",
+    "src/CelerisLab/lbm/kernels/boundary/planar_velocity.cuh": "501b2759d0e2e4280cb6fdd127d66a50fa89f6de48fc2cae9a7eafb11ae90bed",
+    "src/CelerisLab/lbm/kernels/config.h": "7d1124125e0b354558297934460e64c760fb68cdd7bf6d01a1238bb2eaab8c37",
+    "src/CelerisLab/lbm/kernels/config/config_grid.h": "1db3e030ebdc116cd7f02362d1fd077daac55e8f0b4329f6285948c224d3ca5d",
+    "src/CelerisLab/lbm/kernels/config/config_method.h": "4374e04a130f90f97d17cc20c40d2f9ea0ed0dd8603817d07f215d02266e247c",
+    "src/CelerisLab/lbm/kernels/config/config_objects.h": "db793af052c756561a190a3ff2cb93ee5fe28462ee4f97c6af80882b4dc8b072",
+    "src/CelerisLab/lbm/kernels/config/config_obs.h": "c11e2249359145b5c4cbaaf0487506fa77c1f4b82c7c169ddc574aa0babcd727",
+    "src/CelerisLab/lbm/kernels/config/config_physics.h": "8ac0990b168e1058588676cc05c80df8634cfd671dbc92a89246c3c372d83a4e",
+    "src/CelerisLab/lbm/kernels/core/descriptors.cuh": "8440bc14857da9007982650f828a23608eec9f7cc4bf772e8c8dc1239cb45356",
+    "src/CelerisLab/lbm/kernels/core/flags.cuh": "0025705f24710330d9d9195f713879bfb650a85a3449baef0e196b831af104ab",
+    "src/CelerisLab/lbm/kernels/core/layout.cuh": "1fbd820021c89ac3efcde24145996cb0147769e607e498e61f308ed3e0b13706",
+    "src/CelerisLab/lbm/kernels/core/obs.cuh": "b742b77c9f61e1f020fe9922c3838ae9d855e17b381154958b58f10a22938922",
+    "src/CelerisLab/lbm/kernels/core/params.cuh": "6e18bf730d936d940e3d85452f0714a2eaae31bf41002f9fd520ec5b372878dd",
+    "src/CelerisLab/lbm/kernels/core/precision.cuh": "849df14b4cce50af52399e4b8f3e997196f3ae5f1163f68e28df3e6b2507a8c4",
+    "src/CelerisLab/lbm/kernels/kernel_v2.cu": "c7bd3fb705b77261ba2441c52a24a7e6a7865c165cba088a1b75201c4d29c408",
+    "src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh": "aee444eb60c47a3fd518dca975d1af05425a0818a0ec6c3fb3f6150f60c2cb44",
+    "src/CelerisLab/lbm/kernels/operators/collision_srt.cuh": "6b7ff842f392c8f853d2d1a97cc69bfe50f21283f73e659404da07cf6c95f5e1",
+    "src/CelerisLab/lbm/kernels/operators/collision_trt.cuh": "9725b9840b9ad8b4ec7be61ff5d06d3151e86d1a87585c13f2fbf934d0d2c5f6",
+    "src/CelerisLab/lbm/kernels/operators/forcing_guo.cuh": "969a430c1af03f32e3bd0df95d805e64f56334717c44dd4ba208ba503062a24b",
+    "src/CelerisLab/lbm/kernels/operators/helpers.cuh": "0b3a4d95d7d19f5fec09e5423b8f4cc99f8a226b541967425967330da5a051ec",
+    "src/CelerisLab/lbm/kernels/operators/macro.cuh": "2be20eb2b7d7f5522daaad7f4b57b1ff8cdf883efe4a67846a96783b2fe846a3",
+    "src/CelerisLab/lbm/kernels/operators/turbulence_smag.cuh": "ff7f000297ee0817e255ad9fcb6a13b647e3ddd5b206518d35f38d2895a9f501",
+    "src/CelerisLab/lbm/kernels/refinement.cu": "129f3d91da93b4324f88a0da73a4da1e5390f72c14cce574ee63306712da9895",
+    "src/CelerisLab/lbm/kernels/step/aux_kernels.cu": "578a5b20a157241ee678547059abb1372dd5a589ad06bb2c15b987eb01ca1013",
+    "src/CelerisLab/lbm/kernels/step/block_channel.cu": "befbf6d757e532a59c166d41f920d16d71557cd3a574478e7aa83da737c48cbd",
+    "src/CelerisLab/lbm/kernels/step/curved_double.cu": "bfa7a1df1e939020314d49755045f7a85fe4a19cc1c68dbe62653c66c5e7c2b3",
+    "src/CelerisLab/lbm/kernels/step/esopull_macro.cu": "c73116020ca217ede1e4de27c8269c31f9d91e33433ac8252a21e9b14187c50f",
+    "src/CelerisLab/lbm/kernels/step/esopull_sync.cu": "c0ff3e72900c3ab89a05dd03684e96d3b740cf1672d565727fc2af165933d087",
+    "src/CelerisLab/lbm/kernels/step/init_flow.cu": "305be3bb12fc0d039acec1701baabcf0e8a35e683c49c34f2af8ca9867d2b29c",
+    "src/CelerisLab/lbm/kernels/step/one_step_double.cu": "f9f8242b38e8e0cf8425b75e2b0215bc15363c62eaa0fb43c761b8e0931bfdc9",
+    "src/CelerisLab/lbm/kernels/step/one_step_esopull.cu": "aa223b2744ae6c9125673df311a96ce588032c1ba58bebb11fd33ecb990f8ed4",
+    "src/CelerisLab/lbm/kernels/streaming/esopull_semantic_helpers.cuh": "dd5a9091628b079efd97e53fd49e52d09dcee88529235db77a5648e02937af6e",
+    "src/CelerisLab/lbm/kernels/streaming/esopull_single_buffer.cuh": "77ed4ab9a0ebc5ff62a91a7c39066ef395391ae36e9040a69fb2bd8b1626dc79",
+    "src/CelerisLab/lbm/kernels/streaming/pull_double_buffer.cuh": "b65d741ed78b7b62f99531ba0c0ce1936663d8dc9640e7d7d16c707f1baf6549",
+    "src/CelerisLab/lbm/population_storage.py": "a6a86c29c52c892edfe4802fb773c387212e3f4cc86708d80e154986ddf88044",
+    "src/CelerisLab/lbm/raw_exchange.py": "8694202466eba848563e298656353b0f5db94e21a01fa05f51aad6dc2b9fd311",
+    "src/CelerisLab/lbm/refinement.py": "08c2a0557276a98cbcaa459d9bb7a9735d8bbbf64a6114c57af2fda61481a687",
+    "src/CelerisLab/lbm/refinement_stepper.py": "f2bb34fa90805ad21c3e7cf9f9a67d88d3ae3e249b4517d3898567b4bf9e364d",
+    "src/CelerisLab/lbm/refinement_tables.py": "fdd344024b89df3a83174dc810ca5e32db5e19eb8b1609f8f73b0821f56546b6",
+    "src/CelerisLab/lbm/stepper.py": "0dcb7cd7ae87e33e3dc67a38b6aa43042fbe673e37d9ad165342ee1d695933ce",
+    "src/CelerisLab/lbm/synchronized_remap.py": "7400fc4beaffd4649d6a9c524d38ce5fd6fa7eeb32eb5332e4030a4b6648f197",
+    "src/CelerisLab/simulation.py": "0a03904dc6de88667d7af699005a289adf8b09f0976cb0d099b041f4830a8f4c"
+  },
+  "candidate_header_sha256": "a3db46c2e7d27530b2d1436e65cdbed0b72cc7e047bfbeef2a640877271b51aa",
+  "production_origin_commit": "47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5",
+  "source_experiment": "Explicit channel47/cd3 source experiment; not inherited d79 GPU acceptance. Corrected variant permits only the reviewed Smagorinsky header patch.",
+  "historical_preparation_commit": "d79e3f6f0a424f5bcc3a1a6d610313e49daf97a2"
+}

tests/validation/data/refined_dfg/source_manifest.json

@@ -0,0 +1,110 @@
+{
+  "commit": "cd3a657f9d3833e812c3cca173525f165b48710a",
+  "production_origin_commit": "47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5",
+  "sha256": {
+    "src/CelerisLab/__init__.py": "21cc469bd6e970a1f9ff5970c8eea4602065e41fc881854759ee565ee55c10c6",
+    "src/CelerisLab/body/__init__.py": "e4e0e8f0769a94428d5b02fac790443997556cfc5515a5f1ed1595d2c694d69d",
+    "src/CelerisLab/body/action_smoother.py": "b1a23ee6bbd4850b172cb8738c5c3546ee465a54189e6af9d8044f023d4687dc",
+    "src/CelerisLab/body/coupling/__init__.py": "c931f7eb318d48408c3e2febcc6917c88c8dcc1697cbf4329b1f340d94c35763",
+    "src/CelerisLab/body/coupling/soa_packer.py": "8ec26fd55ef74801bbd5ab6bd3d0a7dabc750cbd498aaac2ca4c61a0108ad3c9",
+    "src/CelerisLab/body/ddf_patch.py": "60fcfb7441e73bf290bcff41ca3a83da3e7ddad765d7247464e8f02e3043b4e2",
+    "src/CelerisLab/body/geometry/__init__.py": "ada88980ab672959aca9aa3a4c4aade3219524b2ab68110cd2a2f41e3e0ca90a",
+    "src/CelerisLab/body/geometry/base.py": "cceca14361ca2f34d34a3e43730e1a79b8ca81be00f0dd62e884db32b479419f",
+    "src/CelerisLab/body/geometry/circle.py": "7605ca7c21103f1d801ac73ab8dd18dffbdac425b2dee7b8858b283628a9e436",
+    "src/CelerisLab/body/jets.py": "4104c24af4651f4bb4f8d8fd36fb6d88f41c6e4469822f36e22e74c9ede31dac",
+    "src/CelerisLab/body/manager.py": "01b653d25bd9b4333379c8bfd12383596b686d2029f2e91b1c91b618540ec6b3",
+    "src/CelerisLab/body/objects.py": "8e63cee8faa9cc6243bcff428be628dfad73a02a9a3b8ece7d78c0a63406a08d",
+    "src/CelerisLab/body/preprocess/__init__.py": "dfd9d58bc07f626ab23a70053e31b88a6d9165d2e4c76a4e9f48abe674d3d2ec",
+    "src/CelerisLab/body/preprocess/flag_overlay.py": "831fc9e9736b758847759a085ff133eb49ec91318ace4fda76f85960138ee300",
+    "src/CelerisLab/body/registry.py": "fbb7327780adbdc6da15ebdb799980daf32df22ee974af0893d7409f177bbb68",
+    "src/CelerisLab/body/sync_plan.py": "3a15bfe9ebeb91482dbfc68850761ef39c872cc12d2efc795babd0c01e392323",
+    "src/CelerisLab/common/__init__.py": "d4dbde78bc97fccc9eabac9abae8f0f15c112c10ff8e9b45bd9401d4723df1ef",
+    "src/CelerisLab/common/_types.py": "aad2ff1a9033d4fd68783ad7123afb47201ddc58785cb6dc0246947037072a5f",
+    "src/CelerisLab/common/checkpoint.py": "b2f960aecd5ac27a3a7130b809ffad4709c44d522edef0359ab952654625410b",
+    "src/CelerisLab/common/pathline.py": "a1b6efcc0a076d0f1dc8d49a595779426d7d5626b00072a59bac710ff02ad046",
+    "src/CelerisLab/common/preprocess.py": "7fc78a08857cf8e98519500b094a1d896af189b6b15658ff9b388ed51bed4515",
+    "src/CelerisLab/common/render.py": "0e055476257eb2bcd11e733766e52195de02f8fad228f570a3511a7c831bdc29",
+    "src/CelerisLab/common/streakline/__init__.py": "13ff1cfa416fdd721a08a659e76b5a5a99b8d3201add10a83b7b0bc654bc1eec",
+    "src/CelerisLab/common/streakline/_config.py": "a9185ef68e9c7b5d5a1dc0b04a28341e2ad91d13c145a7c106636a878cf0a605",
+    "src/CelerisLab/common/streakline/_integrate.py": "3fbd4368b2c8d0e7fda40dc65aac5df02b085008b4c68acd83cb7ae45d33a13b",
+    "src/CelerisLab/common/streakline/_render.py": "236cd637414ee318c4d66f9aed32110d8dc42a100d65fde06172f7cd0ba62c65",
+    "src/CelerisLab/common/streakline/_streakline.py": "1a3e08daa706b5f83cdeb337e9a945a36ddf7583c40af1d0bb9d825ac954df37",
+    "src/CelerisLab/config.py": "2f1bba7515f5bc08cb3a8d140bcafe7fa0689e01b3b06494697bb2956a71530e",
+    "src/CelerisLab/configs/CONFIG.md": "41cafdbe525366c08aa9d0416ea026ccf0026ba7123b8a9b65f968a80d928e3a",
+    "src/CelerisLab/configs/config_body.json": "055f261f0828ca1d4f2bb694bd15b72bbcabe31d552326208f28e37f761adab1",
+    "src/CelerisLab/configs/config_lbm.json": "17e4311f1f250c6ece59177a8b6ff856b69a016c7c10267e774b030a0a390865",
+    "src/CelerisLab/configs/config_lbm_flume_mrt_zouhe.json": "8f593321db05a563a42931603a228d61cbcba06dfbe113262d5cd0ad93c0fe79",
+    "src/CelerisLab/configs/config_lbm_flume_trt_regularized.json": "f4bba8e6a4339fc3c4eb14f6e2f57c4b22993c73703c773d853edfbe34736e56",
+    "src/CelerisLab/configs/config_lbm_three_cylinder_triangle.json": "eeb8e7149dbe4aa6b3a9562802735aa950750dfbe2d972a0727e788573afa817",
+    "src/CelerisLab/cuda/__init__.py": "0d01dff9f0399c8baba822aab5ca0122a135f011008ed8a499ebd9f7e5b06aea",
+    "src/CelerisLab/cuda/compiler_v2.py": "c4cc3ec4445c8d75e6a253dd87963877117a722bc73b9d4811806fc6045d08c1",
+    "src/CelerisLab/cuda/context.py": "013dd4ec6d01d2fbd11afdfb103eac6a273bd9f798d663d761b851eeb27885e9",
+    "src/CelerisLab/cuda/obs_layout.py": "5ffb2c3aeeaaaaeb64f2cd8febfbcf4ad9919e5414020ad753cfade0683b21c2",
+    "src/CelerisLab/lbm/__init__.py": "7f60fb3aaae81501b510d08c2fedff44fcb05657260eb10b6f5cc91e7204306b",
+    "src/CelerisLab/lbm/block_stepper.py": "00b313daddc6e63d8ab8423c9190ca7e92ad291de20a2f6b80eb53e3f2d85ec7",
+    "src/CelerisLab/lbm/blocks.py": "1e644c5719993f418185ad4a10b6950a30636fc912befd9eab41ff48b4cf85cd",
+    "src/CelerisLab/lbm/channel.py": "2d5b5ff43414cb35633ac4ce90ce2c3a3353950b25c168e1c2ab58d498dd8b0f",
+    "src/CelerisLab/lbm/curved_links.py": "b6957fa13d13f864dba4f7364e93a39905d062dbb3ee9b3f5191add81fff0732",
+    "src/CelerisLab/lbm/descriptors.py": "c9f025dfc14da43b989047b84cc22bd05d5cc92c9ec695c1da68fdfb788e3f2b",
+    "src/CelerisLab/lbm/equilibrium.py": "2fb2f2bb7ff326d23abce80fc2a4c0907b17740ea3e29fe6ae079c5044783826",
+    "src/CelerisLab/lbm/field.py": "0e727eacc3706add4d8d15c418e9b631050871f831698d3ac689dd5f2c701b9b",
+    "src/CelerisLab/lbm/hierarchy.py": "f234975074f6b6d289d573ff6ea6832feae0e27cb04f80b37fe81792b13a12e7",
+    "src/CelerisLab/lbm/initializers.py": "dacf9586963e862d6d0a98ef1a066036be6f91e9d423720661d8ef2c1fb2df6a",
+    "src/CelerisLab/lbm/kernels/blocks.cu": "e3b9606f861bfea8c4b0f34e8356bb230081f899c584f67be89178d889099b17",
+    "src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh": "aba3748c564ed91b161d268ef8b737b7c7be96a823063d201f3da684dd845233",
+    "src/CelerisLab/lbm/kernels/boundary/curved_boundary.cuh": "aeb16bfd483514fb03e15a2fed437d673c3b8b2fe68ab83edf4244ef593878ea",
+    "src/CelerisLab/lbm/kernels/boundary/esopull_curved_helpers.cuh": "b9b474bc4edaff05681c5443235472d04072d0d775fc0eb9c5967faef52b0258",
+    "src/CelerisLab/lbm/kernels/boundary/ibm_kernels.cuh": "e6909b547dc8d64cdd8a76a90660243a8f5fee535caa161b52b4dd8b7c29934b",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/channel_stabilized.cuh": "efe077afcbb6b731256a987c2fd79a1dbc470c902f87070d60421ea6cb13d4e2",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/common.cuh": "59b38a89439878e00de7de16b9c7b03feff443fcc5a421593f7d2536dcbb7ab8",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/equilibrium.cuh": "f0964c811b2c41a5aaa1c6d99c490b06e850c634eb8749f1f05ef61f93c2de1d",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/regularized.cuh": "943d4a571818bacb89c70bccfaac202691121e3e385803ce6309123a0a03ddaf",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/zou_he_local.cuh": "e73e66b853b1e4e0e30d8f4185d832aa414ce670a1813ce2b5459f2887f71f89",
+    "src/CelerisLab/lbm/kernels/boundary/inlet_outlet.cuh": "de8b213a6da42c2e5e9f8966199fc66ea4541b85144a90f3f5dbad3c6ab63b97",
+    "src/CelerisLab/lbm/kernels/boundary/outlet/pressure_neq.cuh": "c8b2adc66d12fe602076ddc9693c5c208711fcaddb58aa9a31401a269f51c4c4",
+    "src/CelerisLab/lbm/kernels/boundary/planar_velocity.cuh": "501b2759d0e2e4280cb6fdd127d66a50fa89f6de48fc2cae9a7eafb11ae90bed",
+    "src/CelerisLab/lbm/kernels/config/config_grid.h": "1db3e030ebdc116cd7f02362d1fd077daac55e8f0b4329f6285948c224d3ca5d",
+    "src/CelerisLab/lbm/kernels/config/config_method.h": "4374e04a130f90f97d17cc20c40d2f9ea0ed0dd8603817d07f215d02266e247c",
+    "src/CelerisLab/lbm/kernels/config/config_objects.h": "db793af052c756561a190a3ff2cb93ee5fe28462ee4f97c6af80882b4dc8b072",
+    "src/CelerisLab/lbm/kernels/config/config_obs.h": "c11e2249359145b5c4cbaaf0487506fa77c1f4b82c7c169ddc574aa0babcd727",
+    "src/CelerisLab/lbm/kernels/config/config_physics.h": "8ac0990b168e1058588676cc05c80df8634cfd671dbc92a89246c3c372d83a4e",
+    "src/CelerisLab/lbm/kernels/config.h": "7d1124125e0b354558297934460e64c760fb68cdd7bf6d01a1238bb2eaab8c37",
+    "src/CelerisLab/lbm/kernels/core/descriptors.cuh": "8440bc14857da9007982650f828a23608eec9f7cc4bf772e8c8dc1239cb45356",
+    "src/CelerisLab/lbm/kernels/core/flags.cuh": "0025705f24710330d9d9195f713879bfb650a85a3449baef0e196b831af104ab",
+    "src/CelerisLab/lbm/kernels/core/layout.cuh": "1fbd820021c89ac3efcde24145996cb0147769e607e498e61f308ed3e0b13706",
+    "src/CelerisLab/lbm/kernels/core/obs.cuh": "b742b77c9f61e1f020fe9922c3838ae9d855e17b381154958b58f10a22938922",
+    "src/CelerisLab/lbm/kernels/core/params.cuh": "6e18bf730d936d940e3d85452f0714a2eaae31bf41002f9fd520ec5b372878dd",
+    "src/CelerisLab/lbm/kernels/core/precision.cuh": "849df14b4cce50af52399e4b8f3e997196f3ae5f1163f68e28df3e6b2507a8c4",
+    "src/CelerisLab/lbm/kernels/kernel_v2.cu": "c7bd3fb705b77261ba2441c52a24a7e6a7865c165cba088a1b75201c4d29c408",
+    "src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh": "aee444eb60c47a3fd518dca975d1af05425a0818a0ec6c3fb3f6150f60c2cb44",
+    "src/CelerisLab/lbm/kernels/operators/collision_srt.cuh": "6b7ff842f392c8f853d2d1a97cc69bfe50f21283f73e659404da07cf6c95f5e1",
+    "src/CelerisLab/lbm/kernels/operators/collision_trt.cuh": "9725b9840b9ad8b4ec7be61ff5d06d3151e86d1a87585c13f2fbf934d0d2c5f6",
+    "src/CelerisLab/lbm/kernels/operators/forcing_guo.cuh": "969a430c1af03f32e3bd0df95d805e64f56334717c44dd4ba208ba503062a24b",
+    "src/CelerisLab/lbm/kernels/operators/helpers.cuh": "0b3a4d95d7d19f5fec09e5423b8f4cc99f8a226b541967425967330da5a051ec",
+    "src/CelerisLab/lbm/kernels/operators/macro.cuh": "2be20eb2b7d7f5522daaad7f4b57b1ff8cdf883efe4a67846a96783b2fe846a3",
+    "src/CelerisLab/lbm/kernels/operators/turbulence_smag.cuh": "ff7f000297ee0817e255ad9fcb6a13b647e3ddd5b206518d35f38d2895a9f501",
+    "src/CelerisLab/lbm/kernels/refinement.cu": "129f3d91da93b4324f88a0da73a4da1e5390f72c14cce574ee63306712da9895",
+    "src/CelerisLab/lbm/kernels/step/aux_kernels.cu": "578a5b20a157241ee678547059abb1372dd5a589ad06bb2c15b987eb01ca1013",
+    "src/CelerisLab/lbm/kernels/step/block_channel.cu": "befbf6d757e532a59c166d41f920d16d71557cd3a574478e7aa83da737c48cbd",
+    "src/CelerisLab/lbm/kernels/step/curved_double.cu": "bfa7a1df1e939020314d49755045f7a85fe4a19cc1c68dbe62653c66c5e7c2b3",
+    "src/CelerisLab/lbm/kernels/step/esopull_macro.cu": "c73116020ca217ede1e4de27c8269c31f9d91e33433ac8252a21e9b14187c50f",
+    "src/CelerisLab/lbm/kernels/step/esopull_sync.cu": "c0ff3e72900c3ab89a05dd03684e96d3b740cf1672d565727fc2af165933d087",
+    "src/CelerisLab/lbm/kernels/step/init_flow.cu": "305be3bb12fc0d039acec1701baabcf0e8a35e683c49c34f2af8ca9867d2b29c",
+    "src/CelerisLab/lbm/kernels/step/one_step_double.cu": "f9f8242b38e8e0cf8425b75e2b0215bc15363c62eaa0fb43c761b8e0931bfdc9",
+    "src/CelerisLab/lbm/kernels/step/one_step_esopull.cu": "aa223b2744ae6c9125673df311a96ce588032c1ba58bebb11fd33ecb990f8ed4",
+    "src/CelerisLab/lbm/kernels/streaming/esopull_semantic_helpers.cuh": "dd5a9091628b079efd97e53fd49e52d09dcee88529235db77a5648e02937af6e",
+    "src/CelerisLab/lbm/kernels/streaming/esopull_single_buffer.cuh": "77ed4ab9a0ebc5ff62a91a7c39066ef395391ae36e9040a69fb2bd8b1626dc79",
+    "src/CelerisLab/lbm/kernels/streaming/pull_double_buffer.cuh": "b65d741ed78b7b62f99531ba0c0ce1936663d8dc9640e7d7d16c707f1baf6549",
+    "src/CelerisLab/lbm/population_storage.py": "a6a86c29c52c892edfe4802fb773c387212e3f4cc86708d80e154986ddf88044",
+    "src/CelerisLab/lbm/raw_exchange.py": "8694202466eba848563e298656353b0f5db94e21a01fa05f51aad6dc2b9fd311",
+    "src/CelerisLab/lbm/refinement.py": "08c2a0557276a98cbcaa459d9bb7a9735d8bbbf64a6114c57af2fda61481a687",
+    "src/CelerisLab/lbm/refinement_stepper.py": "f2bb34fa90805ad21c3e7cf9f9a67d88d3ae3e249b4517d3898567b4bf9e364d",
+    "src/CelerisLab/lbm/refinement_tables.py": "fdd344024b89df3a83174dc810ca5e32db5e19eb8b1609f8f73b0821f56546b6",
+    "src/CelerisLab/lbm/stepper.py": "0dcb7cd7ae87e33e3dc67a38b6aa43042fbe673e37d9ad165342ee1d695933ce",
+    "src/CelerisLab/lbm/synchronized_remap.py": "7400fc4beaffd4649d6a9c524d38ce5fd6fa7eeb32eb5332e4030a4b6648f197",
+    "src/CelerisLab/simulation.py": "0a03904dc6de88667d7af699005a289adf8b09f0976cb0d099b041f4830a8f4c",
+    "tests/validation/re100_helpers.py": "289c6e01075ff361573f862b90cfd61e5980c9a8efcbc78aef69979b4004fd60",
+    "tests/validation/periodic.py": "a2212c0e2705126327a5019e9cafe64eeca398da1b75061ff23630c9843697a6",
+    "tests/validation/data/dfg_re100/reference_cycle.json": "073d1bda1629697efebad95667bfe80598e217d708dbffd7d5f6c6999379c55f"
+  }
+}

tests/validation/data/refined_dfg/source_manifest_mrt.json

@@ -0,0 +1,112 @@
+{
+  "commit": "3ea131e92c057843034f11a1c574301a2430dfce",
+  "production_origin_commit": "3ea131e92c057843034f11a1c574301a2430dfce",
+  "collision": "MRT",
+  "sha256": {
+    "src/CelerisLab/__init__.py": "21cc469bd6e970a1f9ff5970c8eea4602065e41fc881854759ee565ee55c10c6",
+    "src/CelerisLab/body/__init__.py": "e4e0e8f0769a94428d5b02fac790443997556cfc5515a5f1ed1595d2c694d69d",
+    "src/CelerisLab/body/action_smoother.py": "b1a23ee6bbd4850b172cb8738c5c3546ee465a54189e6af9d8044f023d4687dc",
+    "src/CelerisLab/body/coupling/__init__.py": "c931f7eb318d48408c3e2febcc6917c88c8dcc1697cbf4329b1f340d94c35763",
+    "src/CelerisLab/body/coupling/soa_packer.py": "8ec26fd55ef74801bbd5ab6bd3d0a7dabc750cbd498aaac2ca4c61a0108ad3c9",
+    "src/CelerisLab/body/ddf_patch.py": "60fcfb7441e73bf290bcff41ca3a83da3e7ddad765d7247464e8f02e3043b4e2",
+    "src/CelerisLab/body/geometry/__init__.py": "ada88980ab672959aca9aa3a4c4aade3219524b2ab68110cd2a2f41e3e0ca90a",
+    "src/CelerisLab/body/geometry/base.py": "cceca14361ca2f34d34a3e43730e1a79b8ca81be00f0dd62e884db32b479419f",
+    "src/CelerisLab/body/geometry/circle.py": "7605ca7c21103f1d801ac73ab8dd18dffbdac425b2dee7b8858b283628a9e436",
+    "src/CelerisLab/body/jets.py": "4104c24af4651f4bb4f8d8fd36fb6d88f41c6e4469822f36e22e74c9ede31dac",
+    "src/CelerisLab/body/manager.py": "01b653d25bd9b4333379c8bfd12383596b686d2029f2e91b1c91b618540ec6b3",
+    "src/CelerisLab/body/objects.py": "8e63cee8faa9cc6243bcff428be628dfad73a02a9a3b8ece7d78c0a63406a08d",
+    "src/CelerisLab/body/preprocess/__init__.py": "dfd9d58bc07f626ab23a70053e31b88a6d9165d2e4c76a4e9f48abe674d3d2ec",
+    "src/CelerisLab/body/preprocess/flag_overlay.py": "831fc9e9736b758847759a085ff133eb49ec91318ace4fda76f85960138ee300",
+    "src/CelerisLab/body/registry.py": "fbb7327780adbdc6da15ebdb799980daf32df22ee974af0893d7409f177bbb68",
+    "src/CelerisLab/body/sync_plan.py": "3a15bfe9ebeb91482dbfc68850761ef39c872cc12d2efc795babd0c01e392323",
+    "src/CelerisLab/common/__init__.py": "d4dbde78bc97fccc9eabac9abae8f0f15c112c10ff8e9b45bd9401d4723df1ef",
+    "src/CelerisLab/common/_types.py": "aad2ff1a9033d4fd68783ad7123afb47201ddc58785cb6dc0246947037072a5f",
+    "src/CelerisLab/common/checkpoint.py": "b2f960aecd5ac27a3a7130b809ffad4709c44d522edef0359ab952654625410b",
+    "src/CelerisLab/common/pathline.py": "a1b6efcc0a076d0f1dc8d49a595779426d7d5626b00072a59bac710ff02ad046",
+    "src/CelerisLab/common/preprocess.py": "7fc78a08857cf8e98519500b094a1d896af189b6b15658ff9b388ed51bed4515",
+    "src/CelerisLab/common/render.py": "0e055476257eb2bcd11e733766e52195de02f8fad228f570a3511a7c831bdc29",
+    "src/CelerisLab/common/streakline/__init__.py": "13ff1cfa416fdd721a08a659e76b5a5a99b8d3201add10a83b7b0bc654bc1eec",
+    "src/CelerisLab/common/streakline/_config.py": "a9185ef68e9c7b5d5a1dc0b04a28341e2ad91d13c145a7c106636a878cf0a605",
+    "src/CelerisLab/common/streakline/_integrate.py": "3fbd4368b2c8d0e7fda40dc65aac5df02b085008b4c68acd83cb7ae45d33a13b",
+    "src/CelerisLab/common/streakline/_render.py": "236cd637414ee318c4d66f9aed32110d8dc42a100d65fde06172f7cd0ba62c65",
+    "src/CelerisLab/common/streakline/_streakline.py": "1a3e08daa706b5f83cdeb337e9a945a36ddf7583c40af1d0bb9d825ac954df37",
+    "src/CelerisLab/config.py": "2f1bba7515f5bc08cb3a8d140bcafe7fa0689e01b3b06494697bb2956a71530e",
+    "src/CelerisLab/configs/CONFIG.md": "41cafdbe525366c08aa9d0416ea026ccf0026ba7123b8a9b65f968a80d928e3a",
+    "src/CelerisLab/configs/config_body.json": "055f261f0828ca1d4f2bb694bd15b72bbcabe31d552326208f28e37f761adab1",
+    "src/CelerisLab/configs/config_lbm.json": "17e4311f1f250c6ece59177a8b6ff856b69a016c7c10267e774b030a0a390865",
+    "src/CelerisLab/configs/config_lbm_flume_mrt_zouhe.json": "8f593321db05a563a42931603a228d61cbcba06dfbe113262d5cd0ad93c0fe79",
+    "src/CelerisLab/configs/config_lbm_flume_trt_regularized.json": "f4bba8e6a4339fc3c4eb14f6e2f57c4b22993c73703c773d853edfbe34736e56",
+    "src/CelerisLab/configs/config_lbm_three_cylinder_triangle.json": "eeb8e7149dbe4aa6b3a9562802735aa950750dfbe2d972a0727e788573afa817",
+    "src/CelerisLab/cuda/__init__.py": "0d01dff9f0399c8baba822aab5ca0122a135f011008ed8a499ebd9f7e5b06aea",
+    "src/CelerisLab/cuda/compiler_v2.py": "c4cc3ec4445c8d75e6a253dd87963877117a722bc73b9d4811806fc6045d08c1",
+    "src/CelerisLab/cuda/context.py": "013dd4ec6d01d2fbd11afdfb103eac6a273bd9f798d663d761b851eeb27885e9",
+    "src/CelerisLab/cuda/obs_layout.py": "5ffb2c3aeeaaaaeb64f2cd8febfbcf4ad9919e5414020ad753cfade0683b21c2",
+    "src/CelerisLab/lbm/__init__.py": "7f60fb3aaae81501b510d08c2fedff44fcb05657260eb10b6f5cc91e7204306b",
+    "src/CelerisLab/lbm/block_stepper.py": "00b313daddc6e63d8ab8423c9190ca7e92ad291de20a2f6b80eb53e3f2d85ec7",
+    "src/CelerisLab/lbm/blocks.py": "1e644c5719993f418185ad4a10b6950a30636fc912befd9eab41ff48b4cf85cd",
+    "src/CelerisLab/lbm/channel.py": "40d0354e5dc97bcdb7e7b5c1e7f9303a109b7134d8a3af88b8ff4cd75ab5a095",
+    "src/CelerisLab/lbm/curved_links.py": "b6957fa13d13f864dba4f7364e93a39905d062dbb3ee9b3f5191add81fff0732",
+    "src/CelerisLab/lbm/descriptors.py": "c9f025dfc14da43b989047b84cc22bd05d5cc92c9ec695c1da68fdfb788e3f2b",
+    "src/CelerisLab/lbm/equilibrium.py": "2fb2f2bb7ff326d23abce80fc2a4c0907b17740ea3e29fe6ae079c5044783826",
+    "src/CelerisLab/lbm/field.py": "0e727eacc3706add4d8d15c418e9b631050871f831698d3ac689dd5f2c701b9b",
+    "src/CelerisLab/lbm/hierarchy.py": "f234975074f6b6d289d573ff6ea6832feae0e27cb04f80b37fe81792b13a12e7",
+    "src/CelerisLab/lbm/initializers.py": "dacf9586963e862d6d0a98ef1a066036be6f91e9d423720661d8ef2c1fb2df6a",
+    "src/CelerisLab/lbm/kernels/blocks.cu": "2fa894b20a29f006911ac3e3aad0793cc4aaf9ccd7a67d01c6320a0836b1d8ef",
+    "src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh": "aba3748c564ed91b161d268ef8b737b7c7be96a823063d201f3da684dd845233",
+    "src/CelerisLab/lbm/kernels/boundary/curved_boundary.cuh": "aeb16bfd483514fb03e15a2fed437d673c3b8b2fe68ab83edf4244ef593878ea",
+    "src/CelerisLab/lbm/kernels/boundary/esopull_curved_helpers.cuh": "b9b474bc4edaff05681c5443235472d04072d0d775fc0eb9c5967faef52b0258",
+    "src/CelerisLab/lbm/kernels/boundary/ibm_kernels.cuh": "e6909b547dc8d64cdd8a76a90660243a8f5fee535caa161b52b4dd8b7c29934b",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/channel_stabilized.cuh": "efe077afcbb6b731256a987c2fd79a1dbc470c902f87070d60421ea6cb13d4e2",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/common.cuh": "59b38a89439878e00de7de16b9c7b03feff443fcc5a421593f7d2536dcbb7ab8",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/equilibrium.cuh": "f0964c811b2c41a5aaa1c6d99c490b06e850c634eb8749f1f05ef61f93c2de1d",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/regularized.cuh": "943d4a571818bacb89c70bccfaac202691121e3e385803ce6309123a0a03ddaf",
+    "src/CelerisLab/lbm/kernels/boundary/inlet/zou_he_local.cuh": "e73e66b853b1e4e0e30d8f4185d832aa414ce670a1813ce2b5459f2887f71f89",
+    "src/CelerisLab/lbm/kernels/boundary/inlet_outlet.cuh": "de8b213a6da42c2e5e9f8966199fc66ea4541b85144a90f3f5dbad3c6ab63b97",
+    "src/CelerisLab/lbm/kernels/boundary/outlet/pressure_neq.cuh": "c8b2adc66d12fe602076ddc9693c5c208711fcaddb58aa9a31401a269f51c4c4",
+    "src/CelerisLab/lbm/kernels/boundary/planar_velocity.cuh": "501b2759d0e2e4280cb6fdd127d66a50fa89f6de48fc2cae9a7eafb11ae90bed",
+    "src/CelerisLab/lbm/kernels/config/config_grid.h": "1db3e030ebdc116cd7f02362d1fd077daac55e8f0b4329f6285948c224d3ca5d",
+    "src/CelerisLab/lbm/kernels/config/config_method.h": "4374e04a130f90f97d17cc20c40d2f9ea0ed0dd8603817d07f215d02266e247c",
+    "src/CelerisLab/lbm/kernels/config/config_objects.h": "db793af052c756561a190a3ff2cb93ee5fe28462ee4f97c6af80882b4dc8b072",
+    "src/CelerisLab/lbm/kernels/config/config_obs.h": "c11e2249359145b5c4cbaaf0487506fa77c1f4b82c7c169ddc574aa0babcd727",
+    "src/CelerisLab/lbm/kernels/config/config_physics.h": "8ac0990b168e1058588676cc05c80df8634cfd671dbc92a89246c3c372d83a4e",
+    "src/CelerisLab/lbm/kernels/config.h": "7d1124125e0b354558297934460e64c760fb68cdd7bf6d01a1238bb2eaab8c37",
+    "src/CelerisLab/lbm/kernels/core/descriptors.cuh": "8440bc14857da9007982650f828a23608eec9f7cc4bf772e8c8dc1239cb45356",
+    "src/CelerisLab/lbm/kernels/core/flags.cuh": "0025705f24710330d9d9195f713879bfb650a85a3449baef0e196b831af104ab",
+    "src/CelerisLab/lbm/kernels/core/layout.cuh": "1fbd820021c89ac3efcde24145996cb0147769e607e498e61f308ed3e0b13706",
+    "src/CelerisLab/lbm/kernels/core/obs.cuh": "b742b77c9f61e1f020fe9922c3838ae9d855e17b381154958b58f10a22938922",
+    "src/CelerisLab/lbm/kernels/core/params.cuh": "6e18bf730d936d940e3d85452f0714a2eaae31bf41002f9fd520ec5b372878dd",
+    "src/CelerisLab/lbm/kernels/core/precision.cuh": "849df14b4cce50af52399e4b8f3e997196f3ae5f1163f68e28df3e6b2507a8c4",
+    "src/CelerisLab/lbm/kernels/kernel_v2.cu": "c7bd3fb705b77261ba2441c52a24a7e6a7865c165cba088a1b75201c4d29c408",
+    "src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh": "aee444eb60c47a3fd518dca975d1af05425a0818a0ec6c3fb3f6150f60c2cb44",
+    "src/CelerisLab/lbm/kernels/operators/collision_srt.cuh": "6b7ff842f392c8f853d2d1a97cc69bfe50f21283f73e659404da07cf6c95f5e1",
+    "src/CelerisLab/lbm/kernels/operators/collision_trt.cuh": "9725b9840b9ad8b4ec7be61ff5d06d3151e86d1a87585c13f2fbf934d0d2c5f6",
+    "src/CelerisLab/lbm/kernels/operators/forcing_guo.cuh": "969a430c1af03f32e3bd0df95d805e64f56334717c44dd4ba208ba503062a24b",
+    "src/CelerisLab/lbm/kernels/operators/helpers.cuh": "0b3a4d95d7d19f5fec09e5423b8f4cc99f8a226b541967425967330da5a051ec",
+    "src/CelerisLab/lbm/kernels/operators/macro.cuh": "2be20eb2b7d7f5522daaad7f4b57b1ff8cdf883efe4a67846a96783b2fe846a3",
+    "src/CelerisLab/lbm/kernels/operators/turbulence_smag.cuh": "ff7f000297ee0817e255ad9fcb6a13b647e3ddd5b206518d35f38d2895a9f501",
+    "src/CelerisLab/lbm/kernels/refinement.cu": "b45aad9c022cecb64f7e968ec3cae2d8385a3cc430ee409700c404a3f4dd8cf8",
+    "src/CelerisLab/lbm/kernels/step/aux_kernels.cu": "578a5b20a157241ee678547059abb1372dd5a589ad06bb2c15b987eb01ca1013",
+    "src/CelerisLab/lbm/kernels/step/block_channel.cu": "ee8a31398b8e5233d977c816e6b8e060c1eb2ec71320e594226dbca55593c0d0",
+    "src/CelerisLab/lbm/kernels/step/curved_double.cu": "bfa7a1df1e939020314d49755045f7a85fe4a19cc1c68dbe62653c66c5e7c2b3",
+    "src/CelerisLab/lbm/kernels/step/esopull_macro.cu": "c73116020ca217ede1e4de27c8269c31f9d91e33433ac8252a21e9b14187c50f",
+    "src/CelerisLab/lbm/kernels/step/esopull_sync.cu": "c0ff3e72900c3ab89a05dd03684e96d3b740cf1672d565727fc2af165933d087",
+    "src/CelerisLab/lbm/kernels/step/init_flow.cu": "305be3bb12fc0d039acec1701baabcf0e8a35e683c49c34f2af8ca9867d2b29c",
+    "src/CelerisLab/lbm/kernels/step/one_step_double.cu": "f9f8242b38e8e0cf8425b75e2b0215bc15363c62eaa0fb43c761b8e0931bfdc9",
+    "src/CelerisLab/lbm/kernels/step/one_step_esopull.cu": "aa223b2744ae6c9125673df311a96ce588032c1ba58bebb11fd33ecb990f8ed4",
+    "src/CelerisLab/lbm/kernels/streaming/esopull_semantic_helpers.cuh": "dd5a9091628b079efd97e53fd49e52d09dcee88529235db77a5648e02937af6e",
+    "src/CelerisLab/lbm/kernels/streaming/esopull_single_buffer.cuh": "77ed4ab9a0ebc5ff62a91a7c39066ef395391ae36e9040a69fb2bd8b1626dc79",
+    "src/CelerisLab/lbm/kernels/streaming/pull_double_buffer.cuh": "b65d741ed78b7b62f99531ba0c0ce1936663d8dc9640e7d7d16c707f1baf6549",
+    "src/CelerisLab/lbm/population_storage.py": "a6a86c29c52c892edfe4802fb773c387212e3f4cc86708d80e154986ddf88044",
+    "src/CelerisLab/lbm/raw_exchange.py": "8694202466eba848563e298656353b0f5db94e21a01fa05f51aad6dc2b9fd311",
+    "src/CelerisLab/lbm/refinement.py": "100d5e46ac894f6008b61364cdad5d37c2ca7ed949d8d71405572d6366526417",
+    "src/CelerisLab/lbm/refinement_stepper.py": "f2bb34fa90805ad21c3e7cf9f9a67d88d3ae3e249b4517d3898567b4bf9e364d",
+    "src/CelerisLab/lbm/refinement_tables.py": "fdd344024b89df3a83174dc810ca5e32db5e19eb8b1609f8f73b0821f56546b6",
+    "src/CelerisLab/lbm/stepper.py": "0dcb7cd7ae87e33e3dc67a38b6aa43042fbe673e37d9ad165342ee1d695933ce",
+    "src/CelerisLab/lbm/synchronized_remap.py": "7400fc4beaffd4649d6a9c524d38ce5fd6fa7eeb32eb5332e4030a4b6648f197",
+    "src/CelerisLab/simulation.py": "0a03904dc6de88667d7af699005a289adf8b09f0976cb0d099b041f4830a8f4c",
+    "tests/validation/re100_helpers.py": "289c6e01075ff361573f862b90cfd61e5980c9a8efcbc78aef69979b4004fd60",
+    "tests/validation/periodic.py": "a2212c0e2705126327a5019e9cafe64eeca398da1b75061ff23630c9843697a6",
+    "tests/validation/data/dfg_re100/reference_cycle.json": "073d1bda1629697efebad95667bfe80598e217d708dbffd7d5f6c6999379c55f",
+    "tests/validation/source_provenance.py": "6dfe8f2d819dba0e240ba761d31162a5033685d492d09db75eb2b8aaab3a5a4c"
+  }
+}

tests/validation/jet_cv_measurements.py

@@ -0,0 +1,285 @@
+"""Measurement-only continuation diagnostics. Never changes populations or commands."""
+
+import hashlib
+import json
+from pathlib import Path
+import numpy as np
+
+OFFSETS = np.array([1.5, 2.5, 3.5])
+LINEAR = np.array([2.5, -1.5, 0.0])
+QUADRATIC = np.array([4.375, -5.25, 1.875])
+
+
+def sample_steps(c):
+    """Same reference times at both scales; dense startup retained."""
+    s2 = c["scale"] ** 2
+    reference = sorted(set(range(0, 101)) | set(range(100, 3201, 4)))
+    return np.array(reference, int) * s2
+
+
+def factors(scale):
+    """Lref=one s=1 root length, Tref=one s=1 root time, rho_ref=1."""
+    return dict(
+        length=1 / scale,
+        time=1 / scale**2,
+        velocity=scale,
+        mass=1 / scale**2,
+        momentum=1 / scale,
+        impulse=1 / scale,
+        force=scale,
+        volume_rate=1.0,
+        mass_rate=1.0,
+        pressure=scale**2,
+    )
+
+
+def primitives(g, cx, cy):
+    f = np.asarray(g, np.float64)
+    rho = f.sum(axis=0)
+    mom = np.stack([np.einsum("q,qyx->yx", v, f) for v in (cx, cy)])
+    return rho, mom / rho
+
+
+def contour_terms(rho, vel, origin, rectangle, nu, h=0.5):
+    """Independent split of total outward momentum flux into advective,p,viscous."""
+    x0, y0, x1, y1 = np.rint((np.array(rectangle) - np.tile(origin, 2)) / h).astype(int)
+    duy, dux = np.gradient(vel[0], h)
+    dvy, dvx = np.gradient(vel[1], h)
+    adv = rho * vel[:, None] * vel[None, :]
+    pressure = np.zeros_like(adv)
+    pressure[0, 0] = pressure[1, 1] = (rho - 1) / 3
+    viscous = -rho * nu * np.array([[2 * dux, duy + dvx], [duy + dvx, 2 * dvy]])
+    out = []
+    for field in (adv, pressure, viscous):
+        flux = np.zeros(2)
+        for i, sign in ((x0, -1), (x1, 1)):
+            flux += (
+                sign
+                * 0.5
+                * (field[:, 0, y0:y1, i - 1] + field[:, 0, y0:y1, i]).sum(axis=1)
+                * h
+            )
+        for j, sign in ((y0, -1), (y1, 1)):
+            flux += (
+                sign
+                * 0.5
+                * (field[:, 1, j - 1, x0:x1] + field[:, 1, j, x0:x1]).sum(axis=1)
+                * h
+            )
+        out.append(flux)
+    return np.array(out)
+
+
+def ray_geometry(fluid, origin, center, radius, h=0.5):
+    z, w = np.polynomial.legendre.leggauss(32)
+    theta = np.array([np.pi / 2, 3 * np.pi / 2])[:, None] + np.pi / 12 * z
+    normal = np.stack((np.cos(theta), np.sin(theta)), axis=-1)
+    positions = np.asarray(center) + normal[:, :, None, :] * (
+        radius + h * OFFSETS[None, None, :, None]
+    )
+    local = (positions - origin) / h - 0.5
+    ix = np.floor(local[..., 0]).astype(int)
+    iy = np.floor(local[..., 1]).astype(int)
+    a = local[..., 0] - ix
+    b = local[..., 1] - iy
+    if not (
+        np.all(ix >= 0)
+        and np.all(iy >= 0)
+        and np.all(ix + 1 < fluid.shape[1])
+        and np.all(iy + 1 < fluid.shape[0])
+    ):
+        raise ValueError("Ray interpolation leaves allocated grid")
+    for dy, dx in ((0, 0), (0, 1), (1, 0), (1, 1)):
+        if not fluid[iy + dy, ix + dx].all():
+            raise ValueError("Ray stencil must be active fluid")
+    return dict(
+        ix=ix,
+        iy=iy,
+        a=a,
+        b=b,
+        normal=normal,
+        ds=w * radius * np.pi / 12,
+        positions=positions,
+        origin=np.asarray(origin),
+        h=h,
+    )
+
+
+def reconstruct_aperture(rho, vel, geo):
+    field = np.concatenate((rho[None], vel), axis=0)
+    x, y, a, b = geo["ix"], geo["iy"], geo["a"], geo["b"]
+    ray = (
+        field[:, y, x] * (1 - a) * (1 - b)
+        + field[:, y, x + 1] * a * (1 - b)
+        + field[:, y + 1, x] * (1 - a) * b
+        + field[:, y + 1, x + 1] * a * b
+    )
+    estimates = []
+    for coeff in (LINEAR, QUADRATIC):
+        wall = np.einsum("cajk,k->caj", ray, coeff)
+        un = np.einsum("caj,ajc->aj", wall[1:], geo["normal"])
+        q = np.sum(un * geo["ds"], axis=1)
+        mass = np.sum(wall[0] * un * geo["ds"], axis=1)
+        estimates.append(np.stack((q, mass), axis=-1))
+    return ray, np.array(estimates)
+
+
+class MeasurementRecorder:
+    """Retain reconstructible per-sample terms and sparse full primitive snapshots."""
+
+    def __init__(self, c, output, fluid_u, leaves, origin_a, cx, cy):
+        self.c = c
+        self.output = Path(output)
+        self.fluid_u = fluid_u
+        self.leaves = leaves
+        self.origin_a = np.asarray(origin_a)
+        self.cx = cx
+        self.cy = cy
+        self.records = []
+        self.rays = [
+            ray_geometry(fluid_u, (0.0, 0.0), np.array(c["center"]), c["diameter"] / 2),
+            ray_geometry(leaves[1], origin_a, np.array(c["center"]), c["diameter"] / 2),
+        ]
+        self.snapshots = []
+
+    def capture(self, step, ug, levels, cv_data, nearest_flux, velocity_error):
+        c = self.c
+        s = c["scale"]
+        record = dict(
+            step=step,
+            t_reference=step / s**2,
+            t_star=step * c["u"] / c["diameter"],
+            velocity_error=velocity_error,
+        )
+        arrays = [ug, *levels]
+        masks = [self.fluid_u, *self.leaves]
+        hs = [0.5, 1.0, 0.5]
+        primitive = [primitives(g, self.cx, self.cy) for g in arrays]
+        stats = []
+        for (rho, vel), mask, h in zip(primitive, masks, hs):
+            if (
+                not np.isfinite(rho[mask]).all()
+                or rho[mask].min() <= 0
+                or not np.isfinite(vel[:, mask]).all()
+            ):
+                raise RuntimeError(
+                    "Invalid density/velocity on an active composite leaf"
+                )
+            stats.append(
+                [
+                    rho[mask].min(),
+                    rho[mask].max(),
+                    np.max(abs(rho[mask] - 1)),
+                    np.linalg.norm(vel[:, mask], axis=0).max() * s,
+                    rho[mask].sum() * h * h / s**2,
+                ]
+            )
+        record["level_stats"] = stats  # uniform, AMR root leaf, AMR fine leaf
+        record["global_mass_reference"] = [stats[0][-1], stats[1][-1] + stats[2][-1]]
+        record["cv_inventory"] = []
+        record["cv_flux"] = []
+        record["momentum_flux_terms"] = []
+        record["arc_ray_values"] = []
+        record["arc_estimates"] = []
+        record["nearest_mass_flux"] = [
+            nearest_flux[k] for k in ("uniform_adapter", "amr_api")
+        ]
+        for idx, (name, p, origin) in enumerate(
+            [
+                ("uniform_adapter", primitive[0], (0.0, 0.0)),
+                ("amr_api", primitive[2], self.origin_a),
+            ]
+        ):
+            rho, vel = p
+            record["cv_inventory"].append([x[0] for x in cv_data[name]])
+            record["cv_flux"].append([x[1] for x in cv_data[name]])
+            terms = np.array(
+                [contour_terms(rho, vel, origin, r, c["nu"]) for r in c["contours"]]
+            )
+            np.testing.assert_allclose(
+                terms.sum(axis=1),
+                np.array(record["cv_flux"][-1])[:, 1:],
+                atol=2e-12,
+                rtol=2e-10,
+            )
+            record["momentum_flux_terms"].append(terms)
+            ray, estimates = reconstruct_aperture(rho, vel, self.rays[idx])
+            record["arc_ray_values"].append(ray)
+            record["arc_estimates"].append(estimates)
+        self.records.append(record)
+        if step // s**2 in (0, 100, 800, 1600, 3200) and step % s**2 == 0:
+            fields = {}
+            for name, (rho, vel), mask, h, origin in zip(
+                ("uniform", "amr_root", "amr_fine"),
+                primitive,
+                masks,
+                hs,
+                ((0.0, 0.0), (0.0, 0.0), self.origin_a),
+            ):
+                fields[name + "_rho"] = np.where(mask, rho, np.nan).astype(np.float32)
+                fields[name + "_velocity_reference"] = np.where(
+                    mask[None], vel * s, np.nan
+                ).astype(np.float32)
+                fields[name + "_origin_reference"] = np.asarray(origin) / s
+                fields[name + "_h_reference"] = np.array(h / s)
+            file = self.output / f"fields_s{s}_tref{step//s**2}.npz"
+            np.savez_compressed(file, **fields)
+            self.snapshots.append(
+                dict(
+                    file=file.name,
+                    sha256=hashlib.sha256(file.read_bytes()).hexdigest(),
+                    bytes=file.stat().st_size,
+                    t_reference=step / s**2,
+                )
+            )
+
+    def flush(self):
+        path = self.output / f"samples_s{self.c['scale']}.npz"
+        arrays = {k: np.array([r[k] for r in self.records]) for k in self.records[0]}
+        for i, g in enumerate(self.rays):
+            for k in ("normal", "ds", "positions"):
+                arrays[f"ray_{i}_{k}"] = g[k]
+        arrays["ray_offsets_fine_cells"] = OFFSETS
+        arrays["linear_coefficients"] = LINEAR
+        arrays["quadratic_coefficients"] = QUADRATIC
+        np.savez_compressed(path, **arrays)
+        metadata = dict(
+            schema=1,
+            case=self.c,
+            reference_conversion=factors(self.c["scale"]),
+            sample_file=path.name,
+            sample_sha256=hashlib.sha256(path.read_bytes()).hexdigest(),
+            sample_count=len(self.records),
+            snapshots=self.snapshots,
+            level_stats_columns=[
+                "rho_min",
+                "rho_max",
+                "max_abs_rho_minus_1",
+                "max_speed_reference",
+                "mass_reference",
+            ],
+            raw_partition_order=["wall", "aperture", "total", "blowing", "suction"],
+            raw_columns=["Jx", "Jy", "angular_impulse", "mass_into_fluid"],
+            backend_order=["uniform_adapter", "amr_api"],
+            level_order=["uniform", "amr_root_leaf", "amr_fine_leaf"],
+            momentum_term_order=["advection", "pressure", "negative_viscous_traction"],
+            arc_estimate_axes=[
+                "sample",
+                "backend",
+                "linear_or_quadratic",
+                "blowing_or_suction",
+                "Q_or_M",
+            ],
+            arc_ray_axes=[
+                "sample",
+                "backend",
+                "rho_ux_uy",
+                "aperture",
+                "gauss_point",
+                "normal_offset",
+            ],
+            warning="Extrapolation and nearest-node values are estimator-dependent, not exact wall values. Fields/ray arrays use root-lattice units except explicitly named reference arrays.",
+        )
+        (self.output / f"samples_s{self.c['scale']}.json").write_text(
+            json.dumps(metadata, indent=2) + "\n"
+        )

tests/validation/les/audit.py

@@ -0,0 +1,247 @@
+"""Bounded host-only audit of exact production Smagorinsky helper; stdlib only."""
+
+import argparse, ctypes, hashlib, json, math, pathlib, subprocess, tempfile
+
+HERE = pathlib.Path(__file__).resolve().parent
+DEFAULT_ROOT = HERE.parents[2]
+HEADERS = ["core/descriptors.cuh", "operators/turbulence_smag.cuh"]
+
+
+def directions(dim):
+    result = [(0,) * dim]
+    for axis in range(dim):
+        v = [0] * dim
+        v[axis] = 1
+        result += [tuple(v), tuple(-x for x in v)]
+    for sign in (1, -1):
+        for a in range(dim):
+            for b in range(a + 1, dim):
+                v = [0] * dim
+                v[a] = 1
+                v[b] = sign
+                result += [tuple(v), tuple(-x for x in v)]
+    return result
+
+
+def populations(dim, kind, p=0.02):
+    c = directions(dim)
+    eq = [
+        (
+            (4 / 9 if dim == 2 else 1 / 3)
+            if not any(v)
+            else (
+                (1 / 9 if dim == 2 else 1 / 18)
+                if sum(x * x for x in v) == 1
+                else 1 / 36
+            )
+        )
+        for v in c
+    ]
+    neq = [0.0] * len(c)
+    if kind == "shear":
+        for i, v in enumerate(c):
+            if v[0] and v[1]:
+                neq[i] = p * v[0] * v[1] / 4
+    elif kind == "isotropic":
+        for i, v in enumerate(c):
+            if sum(x * x for x in v) == 1:
+                neq[i] = p / 2
+        neq[0] = -dim * p
+    elif kind != "equilibrium":
+        raise ValueError(kind)
+    return [a + b for a, b in zip(eq, neq)], eq
+
+
+def stress(f, eq, dim):
+    c = directions(dim)
+    return [
+        [sum((a - b) * v[i] * v[j] for a, b, v in zip(f, eq, c)) for j in range(dim)]
+        for i in range(dim)
+    ]
+
+
+def norm(matrix, dev=False):
+    n = len(matrix)
+    tr = sum(matrix[i][i] for i in range(n)) / n if dev else 0
+    return math.sqrt(
+        sum(
+            (matrix[i][j] - (tr if i == j else 0)) ** 2
+            for i in range(n)
+            for j in range(n)
+        )
+    )
+
+
+def tau_closed(tau0, cs, pnorm, rho=1.0, standard=True):
+    factor = math.sqrt(2) if standard else 1.0
+    return 0.5 * (tau0 + math.sqrt(tau0 * tau0 + 18 * factor * cs * cs * pnorm / rho))
+
+
+def iterative(tau0, cs, pnorm_dev, rho=1.0, iters=3, cap=20.0, relax=0.7):
+    tau = tau0
+    nu0 = (tau0 - 0.5) / 3
+    for _ in range(iters):
+        s = 3 * pnorm_dev / (math.sqrt(2) * rho * tau)
+        nut = min(cs * cs * s, cap * nu0)
+        new = min(100.0, max(0.500001, 0.5 + 3 * (nu0 + nut)))
+        tau = relax * new + (1 - relax) * tau
+    return min(1.99, max(0.01, 1 / tau))
+
+
+def wrapper(dim, closed, cs, iters=3, gpu=False):
+    prefix = (
+        ""
+        if gpu
+        else "#define __device__\n#define __constant__\n#define __forceinline__ inline\n"
+    )
+    declarations = f"""#include <math.h>
+{prefix}
+#define LATTICE_D2Q9 0
+#define LATTICE_D3Q19 1
+#define LATTICE_MODEL {0 if dim==2 else 1}
+#define DIM {dim}
+#define NQ {9 if dim==2 else 19}
+#define OMEGA_COLLISION_MIN .01f
+#define OMEGA_COLLISION_MAX 1.99f
+#define LES_CS {cs:.9f}f
+#define LES_CLOSED_FORM {int(closed)}
+#define LES_FP_ITERS {iters}
+#include "core/descriptors.cuh"
+#include "operators/turbulence_smag.cuh"
+"""
+    if gpu:
+        return (
+            declarations
+            + """extern "C" __global__ void probe(const float* f,const float* eq,float rho,float omega,float* out) { if(threadIdx.x==0 && blockIdx.x==0) out[0]=compute_omega_smag(f,eq,rho,omega); }\n"""
+        )
+    return (
+        declarations
+        + """extern "C" float probe(const float* f,const float* eq,float rho,float omega) {return compute_omega_smag(f,eq,rho,omega);}\n"""
+    )
+
+
+def compile_host(root, directory, dim, closed, cs, iters):
+    stem = f"d{dim}_closed{closed}_cs{cs}_i{iters}"
+    src = directory / (stem + ".cpp")
+    lib = directory / (stem + ".so")
+    src.write_text(wrapper(dim, closed, cs, iters))
+    subprocess.run(
+        [
+            "g++",
+            "-std=c++11",
+            "-O2",
+            "-shared",
+            "-fPIC",
+            "-I",
+            str(root / "src/CelerisLab/lbm/kernels"),
+            str(src),
+            "-o",
+            str(lib),
+        ],
+        check=True,
+        capture_output=True,
+    )
+    dll = ctypes.CDLL(str(lib))
+    fn = dll.probe
+    fn.argtypes = [ctypes.POINTER(ctypes.c_float)] * 2 + [ctypes.c_float] * 2
+    fn.restype = ctypes.c_float
+
+    def call(f, eq, tau0):
+        arr = ctypes.c_float * len(f)
+        return float(fn(arr(*f), arr(*eq), 1.0, 1 / tau0))
+
+    return call
+
+
+def run(root):
+    rows = []
+    checks = 0
+    with tempfile.TemporaryDirectory(prefix="les-host-") as tmp:
+        for dim in (2, 3):
+            for cs in (0.16, 0.0):
+                kernels = {
+                    (closed, iters): compile_host(
+                        root, pathlib.Path(tmp), dim, closed, cs, iters
+                    )
+                    for closed, iters in ((1, 3), (0, 3), (0, 200))
+                }
+                for case, kind, p, tau0 in [
+                    ("equilibrium", "equilibrium", 0.0, 0.8),
+                    ("pure_shear", "shear", 0.02, 0.8),
+                    ("isotropic", "isotropic", 0.02, 0.8),
+                    ("cap_active", "shear", 0.2, 0.503),
+                ]:
+                    f, eq = populations(dim, kind, p)
+                    tensor = stress(f, eq, dim)
+                    pf, pd = norm(tensor), norm(tensor, True)
+                    actual = {
+                        f"closed{c}_iterations{i}": fn(f, eq, tau0)
+                        for (c, i), fn in kernels.items()
+                    }
+                    expected = {
+                        "closed1_iterations3": min(
+                            1.99,
+                            max(0.01, 1 / tau_closed(tau0, cs, pf, standard=False)),
+                        ),
+                        "closed0_iterations3": iterative(tau0, cs, pd),
+                        "closed0_iterations200": iterative(tau0, cs, pd, iters=200),
+                    }
+                    for key in expected:
+                        assert abs(actual[key] - expected[key]) < 3e-6, (
+                            dim,
+                            cs,
+                            case,
+                            key,
+                            actual,
+                            expected,
+                        )
+                        checks += 1
+                    standard = 1 / tau_closed(tau0, cs, pd)
+                    if case == "pure_shear" and cs:
+                        assert abs(actual["closed0_iterations200"] - standard) < 3e-6
+                        assert abs(actual["closed1_iterations3"] - standard) > 1e-3
+                        checks += 2
+                    if case == "isotropic" and cs:
+                        assert abs(actual["closed0_iterations3"] - 1 / tau0) < 3e-6
+                        assert actual["closed1_iterations3"] < 1 / tau0 - 1e-3
+                        checks += 2
+                    rows.append(
+                        dict(
+                            dim=dim,
+                            cs=cs,
+                            case=case,
+                            tau0=tau0,
+                            stress=tensor,
+                            actual_omega=actual,
+                            expected_omega=expected,
+                            standard_uncapped_deviatoric_omega=standard,
+                        )
+                    )
+    hashes = {
+        name: hashlib.sha256(
+            (root / "src/CelerisLab/lbm/kernels" / name).read_bytes()
+        ).hexdigest()
+        for name in HEADERS
+    }
+    return dict(
+        source_root=str(root),
+        source_headers=hashes,
+        compiler=subprocess.check_output(["g++", "--version"], text=True).splitlines()[
+            0
+        ],
+        host_checks=checks,
+        gpu_executed=False,
+        results=rows,
+    )
+
+
+if __name__ == "__main__":
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--root", type=pathlib.Path, default=DEFAULT_ROOT)
+    ap.add_argument("--output", type=pathlib.Path, required=True)
+    args = ap.parse_args()
+    report = run(args.root.resolve())
+    args.output.write_text(json.dumps(report, indent=2) + "\n")
+    print(
+        f"PASS: {report['host_checks']} host checks; exact production header, no GPU; {args.output}"
+    )

tests/validation/les/data/smagorinsky_normalization.patch

@@ -0,0 +1,37 @@
+--- a/src/CelerisLab/lbm/kernels/operators/turbulence_smag.cuh
++++ b/src/CelerisLab/lbm/kernels/operators/turbulence_smag.cuh
+@@ -5,7 +5,9 @@
+ //   nu_t = (C_s * Delta)^2 * |S|,  Delta = 1 (lattice unit)
+ //   omega_eff = 1 / (3*(nu + nu_t) + 0.5)
+ //
+-// |S| is estimated from non-equilibrium second-order moments using tau0.
++// |S| = sqrt(2 S_ij S_ij), with Delta = 1 in lattice units.
++// The stress estimate uses self-consistent tau_eff; tau0 is molecular.
++// Closed form uses full stress. Iteration uses deviatoric stress and a cap.
+ // ============================================================================
+ 
+ #ifndef CELERIS_OPERATORS_TURBULENCE_SMAG_CUH
+@@ -65,12 +67,12 @@
+ 
+ #if LES_CLOSED_FORM
+     // Closed-form Smagorinsky (no iteration):
+-    //   ω = 2 / (τ₀ + √(τ₀² + 18·Cs²·√Q / ρ))
++    //   ω = 2 / (τ₀ + √(τ₀² + 18·√2·Cs²·√Q / ρ))
+     // where Q = Πneq_ij · Πneq_ij (Frobenius norm squared)
+     const float Q = pixx*pixx + piyy*piyy + 2.0f*pixy*pixy;
+     const float sqrtQ = sqrtf(fmaxf(Q, 0.0f));
+     const float smag_const_sq = LES_CS * LES_CS;
+-    const float discriminant = tau0*tau0 + 18.0f * smag_const_sq * sqrtQ / rho_safe;
++    const float discriminant = tau0*tau0 + 18.0f * sqrtf(2.0f) * smag_const_sq * sqrtQ / rho_safe;
+     const float tau_eff = 0.5f * (tau0 + sqrtf(fmaxf(discriminant, tau0*tau0)));
+     return clamp_omega(1.0f / fmaxf(tau_eff, 0.500001f));
+ #else
+@@ -140,7 +142,7 @@
+                   + 2.0f*(pixy*pixy + pixz*pixz + piyz*piyz);
+     const float sqrtQ = sqrtf(fmaxf(Q, 0.0f));
+     const float smag_const_sq = LES_CS * LES_CS;
+-    const float discriminant = tau0*tau0 + 18.0f * smag_const_sq * sqrtQ / rho_safe;
++    const float discriminant = tau0*tau0 + 18.0f * sqrtf(2.0f) * smag_const_sq * sqrtQ / rho_safe;
+     const float tau_eff = 0.5f * (tau0 + sqrtf(fmaxf(discriminant, tau0*tau0)));
+     return clamp_omega(1.0f / fmaxf(tau_eff, 0.500001f));
+ #else

tests/validation/les/data/source_manifest.json

@@ -0,0 +1,10 @@
+{
+  "baseline_commit": "d79e3f6f0a424f5bcc3a1a6d610313e49daf97a2",
+  "files": {
+    "baseline/src/CelerisLab/lbm/kernels/operators/turbulence_smag.cuh": "ff7f000297ee0817e255ad9fcb6a13b647e3ddd5b206518d35f38d2895a9f501",
+    "baseline/src/CelerisLab/lbm/kernels/core/descriptors.cuh": "8440bc14857da9007982650f828a23608eec9f7cc4bf772e8c8dc1239cb45356",
+    "patched/src/CelerisLab/lbm/kernels/operators/turbulence_smag.cuh": "a3db46c2e7d27530b2d1436e65cdbed0b72cc7e047bfbeef2a640877271b51aa",
+    "patched/src/CelerisLab/lbm/kernels/core/descriptors.cuh": "8440bc14857da9007982650f828a23608eec9f7cc4bf772e8c8dc1239cb45356"
+  },
+  "gpu_run": false
+}

tests/validation/les/gpu_probe.py

@@ -0,0 +1,84 @@
+"""Optional exact-header GPU helper probe. Not run in this audit.
+Run ONLY through the authorized Horsea executor. Uses existing PyCUDA/NumPy.
+This is helper parity, not full collide_dispatch or turbulence validation.
+"""
+
+import argparse, json, pathlib
+from audit import DEFAULT_ROOT, populations, wrapper
+
+
+def main():
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--root", type=pathlib.Path, default=DEFAULT_ROOT)
+    ap.add_argument("--device", type=int, default=0)
+    ap.add_argument("--run-gpu", action="store_true")
+    ap.add_argument("--output", type=pathlib.Path, required=True)
+    args = ap.parse_args()
+    if not args.run_gpu:
+        ap.error("Explicit --run-gpu required; use authorized Horsea executor only")
+    import numpy as np
+    import pycuda.driver as cuda
+    from pycuda.compiler import SourceModule
+
+    cuda.init()
+    ctx = cuda.Device(args.device).make_context()
+    rows = []
+    try:
+        for dim in (2, 3):
+            for cs in (0.16, 0.0):
+                for closed, iters in ((1, 3), (0, 3), (0, 200)):
+                    module = SourceModule(
+                        wrapper(dim, closed, cs, iters, gpu=True),
+                        include_dirs=[
+                            str(args.root.resolve() / "src/CelerisLab/lbm/kernels")
+                        ],
+                        options=["-std=c++11", "--use_fast_math"],
+                        no_extern_c=True,
+                    )
+                    fn = module.get_function("probe")
+                    for case, kind, p, tau in [
+                        ("equilibrium", "equilibrium", 0.0, 0.8),
+                        ("pure_shear", "shear", 0.02, 0.8),
+                        ("isotropic", "isotropic", 0.02, 0.8),
+                        ("cap_active", "shear", 0.2, 0.503),
+                    ]:
+                        f, eq = populations(dim, kind, p)
+                        out = np.zeros(1, np.float32)
+                        fn(
+                            cuda.In(np.asarray(f, np.float32)),
+                            cuda.In(np.asarray(eq, np.float32)),
+                            np.float32(1),
+                            np.float32(1 / tau),
+                            cuda.Out(out),
+                            block=(1, 1, 1),
+                            grid=(1, 1),
+                        )
+                        ctx.synchronize()
+                        rows.append(
+                            dict(
+                                dim=dim,
+                                cs=cs,
+                                closed=closed,
+                                iters=iters,
+                                case=case,
+                                omega=float(out[0]),
+                            )
+                        )
+        args.output.write_text(
+            json.dumps(
+                dict(
+                    device=cuda.Device(args.device).name(),
+                    diagnostic_only=True,
+                    results=rows,
+                ),
+                indent=2,
+            )
+            + "\n"
+        )
+    finally:
+        ctx.pop()
+        ctx.detach()
+
+
+if __name__ == "__main__":
+    main()

tests/validation/les/prepare.py

@@ -0,0 +1,63 @@
+"""Materialize pinned baseline/candidate headers under a new explicit output path."""
+
+import argparse
+import hashlib
+import json
+from pathlib import Path
+import shutil
+import subprocess
+
+HERE = Path(__file__).resolve().parent
+REPO = HERE.parents[2]
+DATA = HERE / "data"
+PIN = "d79e3f6f0a424f5bcc3a1a6d610313e49daf97a2"
+
+
+def manifest():
+    result = json.loads((DATA / "source_manifest.json").read_text())
+    if result["baseline_commit"] != PIN:
+        raise RuntimeError("Wrong baseline pin")
+    return result
+
+
+def verify_candidate(root):
+    for name, expected in manifest()["files"].items():
+        if hashlib.sha256((root / name).read_bytes()).hexdigest() != expected:
+            raise RuntimeError("Candidate header hash mismatch: " + name)
+
+
+def materialize(repo, output):
+    """Copy only the two exact source headers, then patch the isolated copy."""
+    expected = manifest()
+    headers = [
+        Path(name).relative_to("baseline")
+        for name in expected["files"]
+        if name.startswith("baseline/")
+    ]
+    for relative in headers:
+        wanted = expected["files"][str(Path("baseline") / relative)]
+        if hashlib.sha256((repo / relative).read_bytes()).hexdigest() != wanted:
+            raise RuntimeError("Production header does not match d79: " + str(relative))
+    output.mkdir(parents=True, exist_ok=False)
+    for mode in ("baseline", "patched"):
+        for relative in headers:
+            target = output / mode / relative
+            target.parent.mkdir(parents=True, exist_ok=True)
+            shutil.copyfile(repo / relative, target)
+    patch = DATA / "smagorinsky_normalization.patch"
+    subprocess.run(
+        ["git", "apply", "--check", str(patch)], cwd=output / "patched", check=True
+    )
+    subprocess.run(["git", "apply", str(patch)], cwd=output / "patched", check=True)
+    verify_candidate(output)
+    (output / "source_manifest.json").write_text(json.dumps(expected, indent=2) + "\n")
+    return output
+
+
+if __name__ == "__main__":
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--repo", type=Path, default=REPO)
+    parser.add_argument("--output-dir", type=Path, required=True)
+    args = parser.parse_args()
+    materialize(args.repo.resolve(), args.output_dir.resolve())
+    print("Pinned baseline/candidate headers verified; production files unchanged.")

tests/validation/les/run_gpu.py

@@ -0,0 +1,44 @@
+"""Serial pinned baseline/candidate helper diagnostic; existing CUDA environment only."""
+
+import argparse
+from pathlib import Path
+import subprocess
+import sys
+from prepare import HERE, REPO, materialize, verify_candidate
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--repo", type=Path, default=REPO)
+    parser.add_argument("--device", type=int, default=0)
+    parser.add_argument("--run-gpu", action="store_true")
+    parser.add_argument("--output-dir", type=Path, required=True)
+    args = parser.parse_args()
+    if not args.run_gpu:
+        parser.error("Explicit --run-gpu required; authorized GPU executor only")
+    output = args.output_dir.resolve()
+    output.mkdir(parents=True, exist_ok=False)
+    candidate = materialize(args.repo.resolve(), output / "candidate")
+    verify_candidate(candidate)
+    for mode in ("baseline", "patched"):
+        subprocess.run(
+            [
+                sys.executable,
+                str(HERE / "gpu_probe.py"),
+                "--root",
+                str(candidate / mode),
+                "--device",
+                str(args.device),
+                "--run-gpu",
+                "--output",
+                str(output / (mode + ".json")),
+            ],
+            check=True,
+        )
+    print(
+        "Both raw helper probes completed; this is not solver or turbulence acceptance."
+    )
+
+
+if __name__ == "__main__":
+    main()

tests/validation/les/test_candidate.py

@@ -0,0 +1,216 @@
+"""Independent tensor/oracle tests against host-compiled exact copied headers."""
+
+import argparse, ctypes, hashlib, itertools, json, math, pathlib, subprocess, sys, tempfile
+
+HERE = pathlib.Path(__file__).resolve().parent
+from audit import directions, wrapper
+
+
+def tensor(dim, kind, p, a=0, b=1):
+    t = [[0.0] * dim for _ in range(dim)]
+    if kind == "shear":
+        t[a][b] = t[b][a] = p
+    elif kind == "extension":
+        t[a][a] = p
+        t[b][b] = -p
+    elif kind == "isotropic":
+        for i in range(dim):
+            t[i][i] = p
+    elif kind != "equilibrium":
+        raise ValueError(kind)
+    return t
+
+
+def state(dim, rho, t):
+    # Second-order Hermite projection: sum c_i c_i fneq_i = prescribed tensor.
+    # D2Q9/D3Q19 fourth-order isotropy makes this a distinct construction from audit.py.
+    c = directions(dim)
+    w = [
+        (
+            (4 / 9 if dim == 2 else 1 / 3)
+            if not any(v)
+            else (
+                (1 / 9 if dim == 2 else 1 / 18)
+                if sum(x * x for x in v) == 1
+                else 1 / 36
+            )
+        )
+        for v in c
+    ]
+    eq = [rho * wi for wi in w]
+    neq = [
+        4.5
+        * wi
+        * sum(
+            (v[a] * v[b] - (1 / 3 if a == b else 0)) * t[a][b]
+            for a in range(dim)
+            for b in range(dim)
+        )
+        for wi, v in zip(w, c)
+    ]
+    assert abs(sum(neq)) < 1e-12
+    for a in range(dim):
+        assert abs(sum(v[a] * f for v, f in zip(c, neq))) < 1e-12
+        for b in range(dim):
+            assert abs(sum(v[a] * v[b] * f for v, f in zip(c, neq)) - t[a][b]) < 1e-12
+    return [x + y for x, y in zip(eq, neq)], eq
+
+
+def norm(t):
+    return math.sqrt(sum(x * x for row in t for x in row))
+
+
+def standard_tau(tau0, cs, p, rho):
+    # Solve tau*(tau-tau0)=9 Cs^2 P/(sqrt(2) rho), independently via bisection.
+    k = 9 * cs * cs * p / (math.sqrt(2) * rho)
+    lo = tau0
+    hi = tau0 + max(1.0, k / tau0)
+    for _ in range(90):
+        mid = (lo + hi) / 2
+        if mid * (mid - tau0) > k:
+            hi = mid
+        else:
+            lo = mid
+    return (lo + hi) / 2
+
+
+def build(root, mode, dim, cs, closed, iters, tmp):
+    name = f"{mode}_d{dim}_cs{cs}_c{closed}_i{iters}"
+    src = tmp / (name + ".cpp")
+    lib = tmp / (name + ".so")
+    src.write_text(wrapper(dim, closed, cs, iters))
+    subprocess.run(
+        [
+            "g++",
+            "-std=c++11",
+            "-O2",
+            "-shared",
+            "-fPIC",
+            "-I",
+            str(root / mode / "src/CelerisLab/lbm/kernels"),
+            str(src),
+            "-o",
+            str(lib),
+        ],
+        check=True,
+        capture_output=True,
+    )
+    dll = ctypes.CDLL(str(lib))
+    fn = dll.probe
+    fn.argtypes = [ctypes.POINTER(ctypes.c_float)] * 2 + [ctypes.c_float] * 2
+    fn.restype = ctypes.c_float
+
+    def call(f, eq, rho, tau0):
+        ar = ctypes.c_float * len(f)
+        return float(fn(ar(*f), ar(*eq), rho, 1 / tau0))
+
+    return call
+
+
+def run(root):
+    n = 0
+    maxerr = 0.0
+    maxres = 0.0
+    examples = []
+    with tempfile.TemporaryDirectory(prefix="les-candidate-") as tmp:
+        for dim, cs in itertools.product((2, 3), (0.0, 0.08, 0.16, 0.24)):
+            kernels = {
+                (mode, c, i): build(root, mode, dim, cs, c, i, pathlib.Path(tmp))
+                for mode in ("baseline", "patched")
+                for c, i in ((1, 3), (0, 3), (0, 200))
+            }
+            for rho, tau0 in itertools.product((0.7, 1.0, 1.3), (0.6, 0.8, 1.2)):
+                # Pure shear/extension all coordinate planes and both signs; Pi/rho fixed under density scaling.
+                cases = [("equilibrium", 0.0, 0, 1), ("isotropic", 0.01 * rho, 0, 1)]
+                cases += [
+                    (kind, sign * 0.02 * rho, a, b)
+                    for kind in ("shear", "extension")
+                    for sign in (-1, 1)
+                    for a, b in itertools.combinations(range(dim), 2)
+                ]
+                for kind, p, a, b in cases:
+                    t = tensor(dim, kind, p, a, b)
+                    f, eq = state(dim, rho, t)
+                    P = norm(t)
+                    out = {key: fn(f, eq, rho, tau0) for key, fn in kernels.items()}
+                    want = 1 / standard_tau(tau0, cs, P, rho)
+                    got = out["patched", 1, 3]
+                    err = abs(got - want)
+                    maxerr = max(maxerr, err)
+                    assert err < 3e-6, (dim, cs, rho, tau0, kind, p, a, b, got, want)
+                    tau = 1 / got
+                    res = abs(
+                        tau * (tau - tau0) - 9 * cs * cs * P / (math.sqrt(2) * rho)
+                    )
+                    maxres = max(maxres, res)
+                    assert res < 2e-6, (res, dim, kind)
+                    # No edits to iterative path: exact float results must remain identical.
+                    for it in (3, 200):
+                        assert out["baseline", 0, it] == out["patched", 0, it]
+                    if kind in ("shear", "extension", "equilibrium") or cs == 0:
+                        assert abs(out["patched", 0, 200] - want) < 3e-6
+                        # Three iterations need not equal closed form; they approach the same root.
+                        assert (
+                            abs(out["patched", 0, 200] - want)
+                            <= abs(out["patched", 0, 3] - want) + 3e-7
+                        )
+                    if kind == "equilibrium" or cs == 0:
+                        assert abs(out["baseline", 1, 3] - got) < 3e-6
+                    elif kind in ("shear", "extension"):
+                        assert out["baseline", 1, 3] > got
+                    if (
+                        rho == 1.0
+                        and tau0 == 0.8
+                        and cs == 0.16
+                        and p >= 0
+                        and a == 0
+                        and b == 1
+                    ):
+                        examples.append(
+                            dict(
+                                dim=dim,
+                                kind=kind,
+                                baseline_closed=out["baseline", 1, 3],
+                                patched_closed=got,
+                                iterative3=out["patched", 0, 3],
+                                iterative200=out["patched", 0, 200],
+                                standard=want,
+                            )
+                        )
+                    n += 1
+            # Explicit fixed-P/nonunit-rho test, distinct from scaling Pi and rho together.
+            for rho in (0.7, 1.3):
+                t = tensor(dim, "shear", 0.02)
+                f, eq = state(dim, rho, t)
+                got = kernels["patched", 1, 3](f, eq, rho, 0.8)
+                assert abs(got - 1 / standard_tau(0.8, cs, norm(t), rho)) < 3e-6
+                n += 1
+            # Cap path and trace behavior unchanged; arbitrary stress is a helper-only diagnostic.
+            t = tensor(dim, "shear", 0.2)
+            f, eq = state(dim, 1.0, t)
+            for it in (3, 200):
+                assert kernels["baseline", 0, it](f, eq, 1.0, 0.503) == kernels[
+                    "patched", 0, it
+                ](f, eq, 1.0, 0.503)
+    return dict(
+        parameter_cases=n,
+        max_closed_omega_error=maxerr,
+        max_quadratic_residual=maxres,
+        examples=examples,
+        gpu_run=False,
+    )
+
+
+if __name__ == "__main__":
+    from prepare import verify_candidate
+
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--candidate-dir", type=pathlib.Path, required=True)
+    ap.add_argument("--output", type=pathlib.Path, required=True)
+    args = ap.parse_args()
+    if args.output.exists():
+        ap.error("Refusing existing evidence file")
+    verify_candidate(args.candidate_dir)
+    result = run(args.candidate_dir.resolve())
+    args.output.write_text(json.dumps(result, indent=2) + "\n")
+    print(json.dumps(result, indent=2))

tests/validation/les_flow_oracle.py

@@ -0,0 +1,56 @@
+"""Independent 1D continuum oracle; no LBM closure/header or GPU imports."""
+
+import numpy as np
+
+H, U, NU, DELTA = 32.0, 0.04, 0.02, 1.0
+TIMES = tuple(round(v * H**2 / (np.pi**2 * NU)) for v in (0.1, 0.2, 0.4))
+
+
+def solve(cs, n=256, times=TIMES, *, h=H, u0=U, nu=NU, delta=DELTA, cfl=0.4):
+    """FV centered face flux plus SSPRK2, zero Dirichlet at 0,h.
+
+    du/dt = d/dy[(nu + (cs*delta)^2*|du/dy|) du/dy].
+    The oracle refinement dx=h/n MUST NOT replace the LES filter delta.
+    """
+    if not np.isfinite(cfl) or not 0 < cfl <= 0.4:
+        raise ValueError("cfl must be finite and in (0, .4]")
+    dx = h / n
+    y = (np.arange(n) + 0.5) * dx
+    u = u0 * np.sin(np.pi * y / h)
+    a = (cs * delta) ** 2
+
+    def gradients(v):
+        g = np.empty(n + 1)
+        g[1:-1] = np.diff(v) / dx
+        g[0], g[-1] = 2 * v[0] / dx, -2 * v[-1] / dx
+        return g
+
+    def rhs(v):
+        g = gradients(v)
+        flux = (nu + a * np.abs(g)) * g
+        return np.diff(flux) / dx
+
+    t = 0.0
+    out = {}
+    for target in times:
+        if target < t:
+            raise ValueError("times must be nondecreasing and nonnegative")
+        while t < target:
+            # Nonlinear flux derivative is nu+2*a*abs(gradient).
+            dt = min(
+                target - t, cfl * dx * dx / (nu + 2 * a * np.max(np.abs(gradients(u))))
+            )
+            v = u + dt * rhs(u)
+            u = 0.5 * u + 0.5 * (v + dt * rhs(v))
+            t += dt
+        out[int(target)] = u.copy()
+    return y, out
+
+
+def amplitude(y, u, h=H):
+    s = np.sin(np.pi * y / h)
+    return float(np.dot(u, s) / np.dot(s, s))
+
+
+def relative(a, b):
+    return float(np.linalg.norm(a - b) / np.linalg.norm(b))

tests/validation/multigrid/channel_parity_diagnostic/README.md

@@ -0,0 +1,32 @@
+# Root-channel parity diagnostic
+
+This is a read-only diagnostic, not a replacement acceptance gate. It requires a
+caller-admitted GPU and the unchanged checkout containing
+`test_channel_root_gpu.py`. It does not run unrelated body/Jet tests.
+
+From the checkout, run:
+
+```
+python /path/to/channel_parity_diagnostic/discover.py --checkout "$PWD" --output /new/evidence/directory
+```
+
+Discovery invokes only the original `test_uniform_matched_root` method, with its
+original loops and exact assertion. It stops at the first assertion failure,
+extracts the configuration and arrays from that original traceback, saves the
+failure verbatim, and serially starts `replay.py` in a fresh Python process.
+The original failure remains a failure: exit 1 means the assertion was captured,
+even if the targeted replay succeeds or does not reproduce it. Exit 2 means the
+original test did not reproduce the failure; exit 3 means a different assertion
+failed and no configuration was guessed. Other errors are retained and raised.
+
+Targeted replay compiles only the uniform and one failed blocked configuration,
+serially, caching the two modules. It saves PTX and compiler metadata, all
+numeric and bit-only mismatches with coordinates/hex/ULPs at each step, first
+numeric divergence arrays, same-input localization, and comparison against the
+original multistep call. Signed-zero-only bit changes cannot trigger the first
+numeric divergence. The final replay consistency checks establish whether the
+one-step observation sequence agrees with the original call structure.
+
+Outputs retain the original source revision and gate SHA256. No source,
+assertion, tolerance, compiler flags, collision math, or kernel scheduling is
+changed. Findings must be reviewed before any production fix or gate change.

tests/validation/multigrid/channel_parity_diagnostic/discover.py

@@ -0,0 +1,104 @@
+"""Discover the first failure of the unchanged matched-root GPU test, then replay.
+
+This script never runs the body/Jet test, changes the solver, or changes its
+assertions. GPU admission remains the caller's responsibility.
+"""
+import argparse
+import hashlib
+import importlib.util
+import json
+from pathlib import Path
+import subprocess
+import sys
+import traceback
+
+
+def capture_case(exc):
+    """Extract the original test frame; no inferred profile/block/configuration."""
+    frames = []
+    tb = exc.__traceback__
+    while tb is not None:
+        frames.append(tb.tb_frame)
+        tb = tb.tb_next
+    caller = next((f for f in frames if f.f_code.co_name == 'test_uniform_matched_root'), None)
+    if caller is None or not any(f.f_code.co_name == 'assert_array_equal' for f in frames):
+        return None
+    local = caller.f_locals
+    required = ('cfg', 'shifted', 'profile', 'clamp', 'damping', 'block', 'expected', 'actual')
+    if not all(key in local for key in required):
+        return None
+    cfg = local['cfg']
+    return dict(shifted=int(local['shifted']), profile=local['profile'],
+                clamp=int(local['clamp']), damping=float(local['damping']),
+                block=[int(v) for v in local['block']], steps=8,
+                config_repr=repr(cfg)), local['expected'].copy(), local['actual'].copy()
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument('--checkout', type=Path, default=Path.cwd())
+    parser.add_argument('--output', type=Path, required=True, help='New directory; existing evidence is never overwritten')
+    args = parser.parse_args()
+    checkout, output = args.checkout.resolve(), args.output.resolve()
+    output.mkdir(parents=True, exist_ok=False)
+    import os
+    os.chdir(checkout)
+    sys.path.insert(0, str(checkout / 'src'))
+    import numpy as np
+    gate_path = checkout / 'tests/validation/multigrid/test_channel_root_gpu.py'
+    spec = importlib.util.spec_from_file_location('original_channel_gpu_gate', gate_path)
+    gate = importlib.util.module_from_spec(spec)
+    spec.loader.exec_module(gate)
+    revision = subprocess.run(['git', 'rev-parse', 'HEAD'], capture_output=True, text=True, check=True).stdout.strip()
+    manifest = dict(checkout=str(checkout), revision=revision,
+                    gate_sha256=hashlib.sha256(gate_path.read_bytes()).hexdigest(),
+                    test='ChannelRootGPU.test_uniform_matched_root')
+    (output / 'manifest.json').write_text(json.dumps(manifest, indent=2))
+    test = gate.ChannelRootGPU('test_uniform_matched_root')
+    try:
+        # The original method and np.testing.assert_array_equal are unchanged.
+        # Its existing loops naturally stop at the first failing assertion.
+        test.test_uniform_matched_root()
+    except AssertionError as exc:
+        original_trace = traceback.format_exc()
+        (output / 'original_failure.txt').write_text(original_trace)
+        print(original_trace, file=sys.stderr, flush=True)
+        case = capture_case(exc)
+        if case is None:
+            (output / 'status.json').write_text(json.dumps(dict(status='different_assertion_failure', replay_started=False), indent=2))
+            return 3
+        config, expected, actual = case
+        (output / 'failed_case.json').write_text(json.dumps(config, indent=2))
+        np.savez(output / 'original_failure_arrays.npz', expected=expected, actual=actual)
+        print('CAPTURED_CASE ' + json.dumps(config), flush=True)
+    except Exception:
+        (output / 'discovery_error.txt').write_text(traceback.format_exc())
+        (output / 'status.json').write_text(json.dumps(dict(status='discovery_error', replay_started=False), indent=2))
+        raise
+    else:
+        (output / 'status.json').write_text(json.dumps(dict(status='original_failure_not_reproduced', replay_started=False), indent=2))
+        print('The original matched-root test completed without reproducing the failure. No targeted replay was started.', flush=True)
+        return 2
+
+    # Discovery's context has closed. Run targeted diagnostics in a fresh process,
+    # serially, never concurrently with another compilation or GPU run here.
+    replay = Path(__file__).resolve().with_name('replay.py')
+    command = [sys.executable, str(replay), '--shifted', str(config['shifted']),
+               '--profile', config['profile'], '--clamp', str(config['clamp']),
+               '--damping', str(config['damping']), '--block', *map(str, config['block']),
+               '--steps', str(config['steps']), '--output', str(output / 'targeted')]
+    (output / 'replay_command.json').write_text(json.dumps(command, indent=2))
+    with (output / 'targeted_stdout.txt').open('w') as stdout, (output / 'targeted_stderr.txt').open('w') as stderr:
+        result = subprocess.run(command, cwd=checkout, stdout=stdout, stderr=stderr)
+    targeted = output / 'targeted/result.json'
+    details = json.loads(targeted.read_text()) if targeted.exists() else None
+    status = dict(status='original_assertion_failure_retained', replay_started=True,
+                  replay_returncode=result.returncode, replay_result=details)
+    (output / 'status.json').write_text(json.dumps(status, indent=2))
+    print(json.dumps(status), flush=True)
+    # Preserve the original failed gate even if replay fails or does not reproduce.
+    return 1
+
+
+if __name__ == '__main__':
+    raise SystemExit(main())

tests/validation/multigrid/channel_parity_diagnostic/replay.py

@@ -0,0 +1,102 @@
+"""Read-only production-kernel parity diagnostic; no altered math or thresholds.
+Run in checkout with PYTHONPATH=src. Supply the exact failed subcase.
+"""
+import argparse
+import importlib.util
+import json
+from pathlib import Path
+import shutil
+import sys
+import numpy as np
+
+p = argparse.ArgumentParser()
+p.add_argument('--shifted', type=int, choices=(0, 1), required=True)
+p.add_argument('--profile', choices=('uniform', 'parabolic'), required=True)
+p.add_argument('--clamp', type=int, choices=(0, 1), required=True)
+p.add_argument('--damping', type=float, required=True)
+p.add_argument('--block', type=int, nargs=2, required=True, metavar=('BY', 'BX'))
+p.add_argument('--steps', type=int, default=8)
+p.add_argument('--output', type=Path, required=True)
+a = p.parse_args()
+a.output.mkdir(parents=True, exist_ok=False)
+sys.path.insert(0, str(Path('src').resolve()))
+path = Path('tests/validation/multigrid/test_channel_root_gpu.py').resolve()
+spec = importlib.util.spec_from_file_location('channel_gpu_gate', path)
+gate = importlib.util.module_from_spec(spec); spec.loader.exec_module(gate)
+
+# Preserve the exact NVCC-generated PTX before the existing compiler cleans up.
+original_compile = gate.compiler.compile_kernel
+original_module = gate.compile_module
+cache = {}
+def retain_compile(*args, **kwargs):
+    result = original_compile(*args, **kwargs)
+    source = kwargs.get('source', 'kernel_v2.cu')
+    shutil.copyfile(result, a.output / (Path(source).stem + '.ptx'))
+    metadata = Path(result + '.json')
+    if metadata.exists(): shutil.copyfile(metadata, a.output / (Path(source).stem + '.ptx.json'))
+    return result
+def cached_module(cfg, context, source):
+    cfg.validate()
+    key = (repr(cfg), source)
+    if key not in cache: cache[key] = original_module(cfg, context, source)
+    return cache[key]
+gate.compiler.compile_kernel = retain_compile
+gate.compile_module = cached_module
+
+cfg = gate.LBMConfig(nx=24, ny=20, viscosity=.1, velocity=.025,
+                    ddf_shifting=bool(a.shifted), inlet_profile=a.profile,
+                    outlet_backflow_clamp=bool(a.clamp), outlet_srt_neq_damp=a.damping)
+config = dict(shifted=a.shifted, profile=a.profile, clamp=a.clamp,
+              damping=a.damping, block=a.block, steps=a.steps)
+(a.output / 'config.json').write_text(json.dumps(config, indent=2))
+def ordered(bits):
+    return int((~bits & 0xffffffff) if bits & 0x80000000 else (bits | 0x80000000))
+def report(expected, actual, step, label):
+    ub = expected.view(np.uint32); ab = actual.view(np.uint32)
+    indices = np.argwhere(ub != ab)
+    numeric = ~((expected == actual) | (np.isnan(expected) & np.isnan(actual)))
+    numeric_count = int(np.count_nonzero(numeric))
+    rows = []
+    for q, y, x in indices:
+        u, b = int(ub[q,y,x]), int(ab[q,y,x])
+        category = ('wall_ghost' if y in (0,cfg.ny-1) else 'inlet_ghost' if x==0
+                    else 'outlet_ghost' if x==cfg.nx-1 else 'wall_adjacent_fluid' if y in (1,cfg.ny-2) else 'bulk_fluid')
+        rows.append(dict(numeric_mismatch=bool(numeric[q,y,x]), q=int(q), y=int(y), x=int(x), category=category,
+                         uniform=float(expected[q,y,x]), blocked=float(actual[q,y,x]),
+                         uniform_hex=f'{u:08x}', blocked_hex=f'{b:08x}',
+                         ulps=abs(ordered(u)-ordered(b))))
+    record = dict(label=label, step=step, bit_count=len(rows), numeric_count=numeric_count,
+                  max_abs=float(np.max(np.abs(actual.astype(float)-expected.astype(float)))),
+                  max_ulps=max((r['ulps'] for r in rows), default=0),
+                  all_finite=bool(np.isfinite(expected).all() and np.isfinite(actual).all()),
+                  differences=rows)
+    (a.output / f'{label}_{step:03d}.json').write_text(json.dumps(record,indent=2))
+    print(json.dumps({k:v for k,v in record.items() if k!='differences'}), flush=True)
+    if rows: print(json.dumps(rows[:12]), flush=True)
+    return numeric_count > 0
+
+with gate.CudaContext(0) as context:
+    initial = gate._encode_physical(gate.state(cfg.ny,cfg.nx), bool(a.shifted))
+    uniform, blocked = initial.copy(), initial.copy()
+    found = False
+    for step in range(1,a.steps+1):
+        previous_uniform, previous_blocked = uniform, blocked
+        uniform = gate.run_root(cfg, context, previous_uniform, steps=1)
+        blocked = gate.run_root(cfg, context, previous_blocked, tuple(a.block), steps=1)
+        differs = report(uniform, blocked, step, 'propagated')
+        if differs and not found:
+            found = True
+            same_input_blocked = gate.run_root(cfg, context, previous_uniform, tuple(a.block), steps=1)
+            report(uniform, same_input_blocked, step, 'same_input')
+            np.savez(a.output/'first_divergence.npz', initial=initial,
+                     previous_uniform=previous_uniform, previous_blocked=previous_blocked,
+                     uniform=uniform, blocked=blocked, same_input_blocked=same_input_blocked)
+    # Verify repeated one-step diagnostics still reproduce the original eight-hop call.
+    direct_uniform = gate.run_root(cfg, context, initial, steps=a.steps)
+    direct_blocked = gate.run_root(cfg, context, initial, tuple(a.block), steps=a.steps)
+    reproduced = report(direct_uniform, direct_blocked, a.steps, 'original_multistep')
+    report(direct_uniform, uniform, a.steps, 'uniform_replay_consistency')
+    report(direct_blocked, blocked, a.steps, 'blocked_replay_consistency')
+
+(a.output / 'result.json').write_text(json.dumps(dict(numeric_failure_reproduced=reproduced, first_numeric_divergence_found=found), indent=2))
+print(json.dumps(dict(numeric_failure_reproduced=reproduced)), flush=True)

tests/validation/multigrid/test_channel_root_gpu.py

@@ -0,0 +1,133 @@
+"""Short root-channel GPU gates. Caller owns GPU admission (CUDA_VISIBLE_DEVICES).
+
+Run: python -m unittest discover -s tests/validation/multigrid -p test_channel_root_gpu.py
+Uniform parity covers every ghost/corner and bx=1 outlet halo donors. AMR gates
+are split invariance and finite evolution, not external DFG benchmark validation.
+"""
+from dataclasses import replace
+import struct
+import unittest
+import numpy as np
+import pycuda.driver as cuda
+
+from CelerisLab.config import LBMConfig
+from CelerisLab.cuda import compiler_v2 as compiler
+from CelerisLab.cuda.context import CudaContext
+from CelerisLab.lbm.blocks import BlockLayout2D
+from CelerisLab.lbm.descriptors import FLUID, SOLID, BC_WALL, BC_INLET, BC_OUTLET
+from CelerisLab.lbm.population_storage import _encode_physical
+from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+
+
+def state(ny, nx):
+    q, y, x = np.mgrid[:9, :ny, :nx]
+    w = np.array([4/9]+[1/9]*4+[1/36]*4)[:, None, None]
+    return w*(1+.001*np.sin(x+3*y+q))
+
+
+def compile_module(cfg, context, source):
+    cfg.validate()
+    build = compiler.generate_config(cfg)
+    try:
+        module = compiler.load_module(compiler.compile_kernel(context.sm_arch, source=source, build_dir=build))
+    finally:
+        compiler.cleanup_build_dir(build)
+    cuda.memcpy_htod(module.get_global("d_params")[0], struct.pack("=ffffI", cfg.omega, 0, 0, 0, 0))
+    return module
+
+
+def run_root(cfg, context, initial, block=None, steps=8):
+    """Launch the production uniform or blocked kernel, without a test-specific backend."""
+    allocations = []
+    def upload(a):
+        a = np.ascontiguousarray(a)
+        p = cuda.mem_alloc(a.nbytes); allocations.append(p); cuda.memcpy_htod(p, a)
+        return p
+    try:
+        layout = None if block is None else BlockLayout2D.build(cfg.ny, cfg.nx, *block)
+        local = cfg if layout is None else replace(cfg, nx=layout.local_shape[1], ny=layout.local_shape[0])
+        module = compile_module(local, context, "kernel_v2.cu" if layout is None else "blocks.cu")
+        data = initial if layout is None else layout.to_blocked(initial)
+        current, spare = upload(data), upload(np.full_like(data, np.nan))
+        tpb = cfg.threads_per_block
+        kw = dict(block=(tpb, 1, 1))
+        if layout is None:
+            flags = np.full((cfg.ny,cfg.nx), FLUID, np.uint16)
+            flags[:, 0] = SOLID|BC_INLET; flags[:, -1] = SOLID|BC_OUTLET
+            flags[[0, -1], :] = SOLID|BC_WALL
+            flag = upload(flags)
+            step = module.get_function("OneStep")
+        else:
+            dst, src = upload(layout.halo_dst), upload(layout.halo_src)
+            gather = module.get_function("BlockHaloCopy")
+            step = module.get_function("BlockChannelStep")
+            cuda.memcpy_htod(module.get_global("d_block_channel")[0], struct.pack("=IIf", cfg.nx, cfg.ny, cfg.velocity))
+        grid = ((local.nx+tpb-1)//tpb, local.ny, 1 if layout is None else layout.nblocks)
+        for _ in range(steps):
+            if layout is None:
+                step(flag, current, spare, grid=grid, **kw)
+            else:
+                gather(current, dst, src, np.uint32(len(layout.halo_dst)),
+                       grid=((len(layout.halo_dst)+tpb-1)//tpb, 1), **kw)
+                step(current, spare, np.uint32(1), np.uintp(0), np.uintp(0), grid=grid, **kw)
+            current, spare = spare, current
+        cuda.Context.synchronize()
+        result = np.empty_like(data); cuda.memcpy_dtoh(result, current)
+        return result if layout is None else layout.to_dense(result)
+    finally:
+        for pointer in allocations: pointer.free()
+
+
+class ChannelRootGPU(unittest.TestCase):
+    def test_uniform_matched_root(self):
+        with CudaContext(0) as context:
+            for shifted in (False, True):
+                for profile, clamp, damping in (("uniform", True, .5), ("parabolic", True, .5),
+                                                ("parabolic", False, .8)):
+                    cfg = LBMConfig(nx=24, ny=20, viscosity=.1, velocity=.025,
+                                    ddf_shifting=shifted, inlet_profile=profile,
+                                    outlet_backflow_clamp=clamp, outlet_srt_neq_damp=damping)
+                    initial = _encode_physical(state(cfg.ny, cfg.nx), shifted)
+                    expected = run_root(cfg, context, initial)
+                    self.assertTrue(np.isfinite(expected).all())
+                    for block in ((20,24), (5,6), (4,1)):
+                        actual = run_root(cfg, context, initial, block)
+                        np.testing.assert_array_equal(actual, expected)
+
+    def test_amr_channel_body_jet_smoke(self):
+        from CelerisLab.body.geometry.circle import CircleGeometry
+        from CelerisLab.body.objects import SimObject, SurfaceActuation
+        center = (24., 16.)
+        body = SimObject(0, geometry=CircleGeometry(*center, 2.15), center=center, radius=2.15)
+        jet = SurfaceActuation(.7, .85, .001).as_jet(center, 2.15)
+        with CudaContext(0) as context:
+            for shifted in (False, True):
+                for bodies in (False, True):
+                    cfg = LBMConfig(nx=48, ny=32, viscosity=.1, velocity=.015, ddf_shifting=shifted)
+                    reference = None
+                    for block in ((32,48), (8,8), (4,1)):
+                        with FixedPatchStepper(cfg, (8,6,40,26), context=context,
+                                coarse_boundary="channel", base_block=block, fine_block=(8,8),
+                                fine_objects=(body,) if bodies else (), jets=(jet,) if bodies else ()) as backend:
+                            arrays = [state(*level.shape[1:]) for level in backend._levels]
+                            backend.initialize(*arrays)
+                            backend.step(4)
+                            if bodies: backend.set_jet_amplitude(0, -.001)
+                            backend.step(4)
+                            raw = backend._download_raw()
+                            physical = backend.download()
+                            for a, mask in zip(raw.values, backend._state):
+                                self.assertTrue(np.isfinite(a[:,mask]).all())
+                            self.assertTrue(np.isfinite(raw.values[0][:,0,:]).all())
+                            self.assertTrue(np.isnan(physical[0][:,0,:]).all())
+                            backend._initialize_raw(raw)
+                            for a,b in zip(raw.values, backend._download_raw().values):
+                                np.testing.assert_array_equal(a,b)
+                            if reference is None: reference = raw.values
+                            else:
+                                for a,b in zip(reference, raw.values): np.testing.assert_array_equal(a,b)
+                            with self.assertRaisesRegex(ValueError, "periodic"):
+                                backend.remap_child([])
+
+
+if __name__ == "__main__": unittest.main()

tests/validation/multigrid/test_exchange_window.py

@@ -0,0 +1,260 @@
+"""Optional observation-window exchange ledger: host contract (mocked driver) and GPU gate.
+
+Run CPU and GPU classes in separate processes; GPU callers select the device.
+ExchangeWindowGPU uses the SRT/unshifted Jet fixture; ExchangeWindowShiftedMRTGPU covers a
+shifted MRT, all-blocked, stationary-body case without Jets (no shifted-MRT Jet claim).
+GPU tolerance (conservative engineering tolerance, not a universal proof): the window ledger
+sums h*dt*raw per link over fine substeps, while the reference sums per-base interval ledgers
+over links and then over base steps. Both are FP64 reorderings of the same FP32-exact substep
+contributions, so the scale is the sum of the ACTUAL absolute substep contributions (read from
+last_raw after every fine hop), not of per-base totals that could hide substep cancellation:
+|window - reference| <= 4 * m * eps * sum|substep terms| + 1e-300, m = links * substeps.
+"""
+from contextlib import ExitStack
+import unittest
+from unittest.mock import patch
+
+import numpy as np
+import pycuda.driver as cuda
+
+from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+from test_refinement_blocked_base import mocked_backend
+from test_refinement_blocked_fine import BOUNDS, circle, config
+
+EPS = np.finfo(np.float64).eps
+
+
+def build(m, window=True, raw=True):
+    return FixedPatchStepper(config("periodic"), BOUNDS, context=m["context"], base_block=(3, 4),
+                             fine_block=(16, 20), fine_objects=[circle()],
+                             observe_raw_exchange=raw, observe_window_exchange=window)
+
+
+def same_bits(test, a, b):
+    """Identical NaN masks, no infinities, identical finite FP32 bits (incl. signed zero)."""
+    a, b = np.asarray(a), np.asarray(b)
+    test.assertEqual(a.dtype, b.dtype)
+    test.assertFalse(np.isinf(a).any() or np.isinf(b).any())
+    np.testing.assert_array_equal(np.isnan(a), np.isnan(b))
+    finite = np.isfinite(a)
+    np.testing.assert_array_equal(a[finite].view(np.uint32), b[finite].view(np.uint32))
+
+
+def accumulations(m):
+    return [c for c in m["calls"] if c[0] == "BlockAccumulateExchange"]
+
+
+class ExchangeWindowCPU(unittest.TestCase):
+    def test_constructor_validation(self):
+        with mocked_backend() as m:
+            with self.assertRaisesRegex(ValueError, "requires observe_raw_exchange"):
+                build(m, window=True, raw=False)
+            with self.assertRaisesRegex(ValueError, "must be bool"):
+                build(m, window=1)
+
+    def test_state_machine_and_launches(self):
+        with mocked_backend() as m, ExitStack() as stack:
+            stack.enter_context(patch.object(cuda, "memset_d8"))
+            stack.enter_context(patch.object(cuda, "memset_d8_async"))
+            stack.enter_context(patch.object(cuda, "memcpy_dtoh",
+                                             side_effect=lambda out, src: out.fill(0)))
+            backend = build(m)
+            b = backend._level_boundaries[1]
+            self.assertIsNotNone(b.window_integral)
+            with self.assertRaises(RuntimeError):
+                backend.start_exchange_window()  # not initialized
+            arrays = [np.full(level.shape, 1 / 9, np.float32) for level in backend._levels]
+            backend.initialize(*arrays)
+            with self.assertRaisesRegex(RuntimeError, "no valid exchange window"):
+                backend.read_exchange_window(0)
+            backend.step(1)
+            self.assertEqual(len(accumulations(m)), 2)  # closed window: interval ledger only
+            backend.start_exchange_window()
+            with self.assertRaisesRegex(RuntimeError, "at least one completed"):
+                backend.read_exchange_window(0)
+            m["calls"].clear()
+            backend.step(2)
+            calls = accumulations(m)
+            self.assertEqual(len(calls), 8)  # two fine hops x two ledgers x two base steps
+            self.assertEqual({id(c[1][1]) for c in calls},
+                             {id(b.interval_integral), id(b.window_integral)})
+            out = backend.read_exchange_window(0)
+            self.assertEqual((out["time_start"], out["time_end"], out["sample_count"]), (1., 3, 4))
+            self.assertEqual(backend.read_exchange_window(0)["time_end"], 3)  # non-destructive
+            backend.initialize(*arrays)  # invalidates
+            with self.assertRaisesRegex(RuntimeError, "no valid exchange window"):
+                backend.read_exchange_window(0)
+            backend.start_exchange_window()
+            with patch.object(backend, "_hop", side_effect=RuntimeError("injected")):
+                with self.assertRaises(RuntimeError):
+                    backend.step(1)
+            self.assertFalse(backend._window_valid)
+            with self.assertRaises(RuntimeError):
+                backend.read_exchange_window(0)  # failed status
+
+    def test_disabled_default_has_no_window_cost(self):
+        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)
+            backend.initialize(*[np.full(l.shape, 1 / 9, np.float32) for l in backend._levels])
+            backend.step(1)
+            self.assertEqual(len(accumulations(m)), 2)
+            with self.assertRaisesRegex(RuntimeError, "disabled"):
+                backend.start_exchange_window()
+
+
+class ExchangeWindowGPU(unittest.TestCase):
+    def test_window_matches_summed_base_reads_and_populations_unchanged(self):
+        from test_raw_exchange_gpu import BOUNDS as RB, CFG, circle as rc, jet, state
+        from CelerisLab.cuda.context import CudaContext
+        with CudaContext(0) as ctx:
+            kw = dict(context=ctx, fine_block=(16, 20), fine_objects=[rc()], jets=[jet()],
+                      observe_raw_exchange=True)
+            with FixedPatchStepper(CFG, RB, observe_window_exchange=True, **kw) as windowed, \
+                 FixedPatchStepper(CFG, RB, **kw) as plain:
+                arrays = [state(level.shape) for level in windowed._levels]
+                windowed.initialize(*arrays)
+                plain.initialize(*arrays)
+                windowed.step(2)
+                plain.step(2)
+                windowed.start_exchange_window()
+                b = windowed._level_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)
+                substeps = 0
+                kernel = windowed._levels[1].kernel
+
+                def audited(src, dst, *args, **kw):
+                    nonlocal scale, substeps
+                    windowed._stream.synchronize()  # BFL and both accumulations are complete
+                    raw = np.empty((count, 4), np.float32)
+                    cuda.memcpy_dtoh(raw, b.last_raw)
+                    scale = scale + (np.abs(raw.astype(np.float64)) * factors).sum(axis=0)
+                    substeps += 1
+                    kernel(src, dst, *args, **kw)
+                windowed._levels[1].kernel = audited
+                keys = ("impulse", "angular_impulse", "mass_into_fluid")
+                summed = {part: [0., 0., 0., 0.] for part in ("total", "wall", "aperture")}
+                steps = 0
+                for speed in (0., .002, .002, -.001, 0., 0.):  # amplitude updates stay in-window
+                    windowed.set_jet_amplitude(0, speed)
+                    plain.set_jet_amplitude(0, speed)
+                    windowed.step(1)
+                    plain.step(1)
+                    steps += 1
+                    obs = windowed.read_raw_exchange(0)
+                    for part in summed:
+                        o = obs[part]
+                        summed[part] = [summed[part][0] + o["impulse"][0], summed[part][1] + o["impulse"][1],
+                                        summed[part][2] + o["angular_impulse"],
+                                        summed[part][3] + o["mass_into_fluid"]]
+                    for a, e in zip(windowed._download_raw().values, plain._download_raw().values):
+                        same_bits(self, a, e)
+                out = windowed.read_exchange_window(0)
+                self.assertEqual((out["time_end"] - out["time_start"], out["sample_count"]),
+                                 (steps, 2 * steps))
+                self.assertEqual(substeps, 2 * steps)
+                bound = 4 * (count * substeps) * EPS * scale + 1e-300
+                for part in summed:
+                    o = out[part]
+                    got = np.array([o["impulse"][0], o["impulse"][1], o["angular_impulse"],
+                                    o["mass_into_fluid"]])
+                    diff = np.abs(got - np.array(summed[part]))
+                    self.assertTrue((diff <= bound).all(), f"{part}: {diff} > {bound}")
+                    np.testing.assert_array_equal(o["force"], o["impulse"] / steps)
+                again = windowed.read_exchange_window(0)
+                np.testing.assert_array_equal(again["total"]["impulse"], out["total"]["impulse"])
+                windowed.initialize(*arrays)
+                with self.assertRaises(RuntimeError):
+                    windowed.read_exchange_window(0)
+
+
+def audit_substeps(stepper, b, count, record):
+    """Wrap the level-1 step kernel: after BFL and accumulation, add |substep terms| to record."""
+    factors = np.array([.25, .25, .125, .25])
+    kernel = stepper._levels[1].kernel
+
+    def audited(src, dst, *args, **kw):
+        stepper._stream.synchronize()
+        raw = np.empty((count, 4), np.float32)
+        cuda.memcpy_dtoh(raw, b.last_raw)
+        record["scale"] = record["scale"] + (np.abs(raw.astype(np.float64)) * factors).sum(axis=0)
+        record["substeps"] += 1
+        kernel(src, dst, *args, **kw)
+    stepper._levels[1].kernel = audited
+
+
+def components(obs):
+    o = obs["total"]
+    return np.array([o["impulse"][0], o["impulse"][1], o["angular_impulse"], o["mass_into_fluid"]])
+
+
+class ExchangeWindowShiftedMRTGPU(unittest.TestCase):
+    """Shifted MRT, all-blocked, stationary body, NO Jet (scoped separately from SRT/Jet)."""
+
+    def test_window_step0_reset_and_split_step(self):
+        from test_raw_exchange_gpu import BOUNDS as RB, circle as rc, state
+        from CelerisLab.config import LBMConfig
+        from CelerisLab.cuda.context import CudaContext
+        cfg = LBMConfig(nx=32, ny=24, viscosity=.1, collision="MRT", ddf_shifting=True)
+        with CudaContext(0) as ctx:
+            kw = dict(context=ctx, base_block=(12, 8), fine_block=(16, 20), fine_objects=[rc()],
+                      observe_raw_exchange=True)
+            with FixedPatchStepper(cfg, RB, observe_window_exchange=True, **kw) as windowed, \
+                 FixedPatchStepper(cfg, RB, observe_window_exchange=True, **kw) as split, \
+                 FixedPatchStepper(cfg, RB, **kw) as plain:
+                arrays = [state(level.shape) for level in windowed._levels]
+                for run in (windowed, split, plain):
+                    run.initialize(*arrays)
+                    run.step(2)
+                    if run is not plain:
+                        run.start_exchange_window()
+                b = windowed._level_boundaries[1]
+                count = len(b.tables["cell"])
+                record = dict(scale=np.zeros(4), substeps=0)
+                audit_substeps(windowed, b, count, record)
+                steps, summed = 5, np.zeros(4)
+                for _ in range(steps):
+                    windowed.step(1)
+                    plain.step(1)
+                    summed += components(windowed.read_raw_exchange(0))
+                    for a, e in zip(windowed._download_raw().values, plain._download_raw().values):
+                        same_bits(self, a, e)
+                split.step(steps)  # one call: same ledger arithmetic as five step(1) calls
+                for a, e in zip(windowed._download_raw().values, split._download_raw().values):
+                    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)
+                np.testing.assert_array_equal(rows[0].view(np.uint64), rows[1].view(np.uint64))
+
+                out = windowed.read_exchange_window(0)
+                self.assertEqual(record["substeps"], 2 * steps)
+                self.assertEqual((out["time_end"] - out["time_start"], out["sample_count"]),
+                                 (steps, 2 * steps))
+                bound = 4 * (count * record["substeps"]) * EPS * record["scale"] + 1e-300
+                diff = np.abs(components(out) - summed)
+                self.assertTrue((diff <= bound).all(), f"{diff} > {bound}")
+                np.testing.assert_array_equal(split.read_exchange_window(0)["total"]["impulse"],
+                                              out["total"]["impulse"])
+
+                windowed.step(0)  # no interval: window and its duration unchanged
+                again = windowed.read_exchange_window(0)
+                self.assertEqual(again["time_end"], out["time_end"])
+                np.testing.assert_array_equal(again["total"]["impulse"], out["total"]["impulse"])
+
+                windowed.start_exchange_window()  # reset after a window
+                with self.assertRaisesRegex(RuntimeError, "at least one completed"):
+                    windowed.read_exchange_window(0)
+                record.update(scale=np.zeros(4), substeps=0)
+                windowed.step(1)
+                single = components(windowed.read_raw_exchange(0))
+                reset = windowed.read_exchange_window(0)
+                self.assertEqual(reset["time_end"] - reset["time_start"], 1)
+                bound = 4 * (count * record["substeps"]) * EPS * record["scale"] + 1e-300
+                self.assertTrue((np.abs(components(reset) - single) <= bound).all())
+
+
+if __name__ == "__main__": unittest.main()

tests/validation/multigrid/test_refinement_mrt_gpu.py

@@ -0,0 +1,95 @@
+"""MRT reuse GPU gates. Caller owns GPU admission (CUDA_VISIBLE_DEVICES).
+
+Run: python -m unittest discover -s tests/validation/multigrid -p test_refinement_mrt_gpu.py
+Parity compares the blocked MRT channel root with uniform MRT OneStep: the existing
+pressure-NEQ outlet branch differs by collision model, so SRT is never the reference.
+AMR checks are finite-evolution/split-invariance smoke tests, not benchmark validation.
+"""
+from dataclasses import replace
+import unittest
+import numpy as np
+
+from CelerisLab.config import LBMConfig
+from CelerisLab.cuda.context import CudaContext
+from CelerisLab.lbm.population_storage import _encode_physical
+from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+from test_channel_root_gpu import run_root, state
+from test_refinement_blocked_base import raw, exact
+
+
+class RefinementMRTGPU(unittest.TestCase):
+    def test_uniform_matched_mrt_root(self):
+        with CudaContext(0) as context:
+            for shifted in (False, True):
+                for profile in ("uniform", "parabolic"):
+                    cfg = LBMConfig(nx=24, ny=20, viscosity=.1, velocity=.025, collision="MRT",
+                                    ddf_shifting=shifted, inlet_profile=profile)
+                    initial = _encode_physical(state(cfg.ny, cfg.nx), shifted)
+                    expected = run_root(cfg, context, initial)
+                    self.assertTrue(np.isfinite(expected).all())
+                    for block in ((20, 24), (5, 6), (4, 1)):
+                        np.testing.assert_array_equal(run_root(cfg, context, initial, block), expected)
+
+    def test_amr_mrt_smoke(self):
+        from CelerisLab.body.geometry.circle import CircleGeometry
+        from CelerisLab.body.objects import SimObject
+        center = (24., 16.)
+        body = SimObject(0, geometry=CircleGeometry(*center, 2.15), center=center, radius=2.15)
+        with CudaContext(0) as context:
+            for boundary in ("periodic", "channel"):
+                for shifted in (False, True):
+                    cfg = LBMConfig(nx=48, ny=32, viscosity=.1, velocity=.015 if boundary == "channel" else 0.,
+                                    collision="MRT", ddf_shifting=shifted)
+                    reference = None
+                    for block in ((32, 48), (8, 8), (4, 1)):
+                        with FixedPatchStepper(cfg, (8, 6, 40, 26), context=context, coarse_boundary=boundary,
+                                               base_block=block, fine_block=(8, 8), fine_objects=(body,)) as backend:
+                            backend.initialize(*[state(*level.shape[1:]) for level in backend._levels])
+                            backend.step(8)
+                            raw = backend._download_raw()
+                            for a, mask in zip(raw.values, backend._state):
+                                self.assertTrue(np.isfinite(a[:, mask]).all())
+                            if reference is None: reference = raw.values
+                            else:
+                                for a, b in zip(reference, raw.values): np.testing.assert_array_equal(a, b)
+
+    def test_dense_versus_blocked_mrt_two_level(self):
+        """Unshifted MRT: all-dense PatchCoarseStep/PatchFineStep vs mixed/all-blocked storage."""
+        from test_refinement_blocked_base import BASE_BLOCKS, FINE_BLOCK
+        from test_refinement_blocked_fine import NX, NY, BOUNDS
+        import test_refinement_blocked_fine as fixture
+        cfg = LBMConfig(nx=NX, ny=NY, viscosity=.1, collision="MRT")
+        with CudaContext(0) as context:
+            results = {}
+            layouts = ((None, None), *((base, FINE_BLOCK) for base in (None, *BASE_BLOCKS)))
+            for base_block, fine_block in layouts:
+                with FixedPatchStepper(cfg, BOUNDS, context=context, fine_block=fine_block,
+                                       base_block=base_block) as backend:
+                    backend.initialize(*[fixture.state(*level.shape[1:], seed) for seed, level in enumerate(backend._levels)])
+                    backend.step(4)
+                    results[base_block, fine_block] = raw(backend)
+            for layout in layouts[1:]:
+                for level, (a, b) in enumerate(zip(results[layout], results[None, None])):
+                    exact(self, a, b, f"layout={layout} level={level}")
+
+    def test_shifted_all_blocked_three_level_mrt_smoke(self):
+        """Per-level shear rate omega_l = 1/(3*2**l*nu + 1/2); finite ACTIVE leaves."""
+        cfg = LBMConfig(nx=16, ny=16, viscosity=.1, collision="MRT", ddf_shifting=True)
+        kw = dict(fine_block=(4, 4), base_block=(4, 4), child=(8, 8, 16, 16), child_block=(4, 4))
+        with CudaContext(0) as context:
+            with FixedPatchStepper(cfg, (4, 4, 12, 12), context=context, **kw) as backend:
+                self.assertEqual(len(backend._level_cfgs), 3)
+                for level, level_cfg in enumerate(backend._level_cfgs):
+                    checked = replace(level_cfg)
+                    checked.validate()  # omega is derived only by validate()
+                    self.assertEqual(checked.collision, "MRT")
+                    self.assertTrue(checked.ddf_shifting)
+                    self.assertTrue(np.isclose(checked.omega, 1 / (3 * 2**level * .1 + .5), rtol=1e-14))
+                backend.initialize(*[state(*level.shape[1:]) for level in backend._levels])
+                backend.step(4)
+                values = backend._download_raw().values
+                for a, mask in zip(values, backend._state):
+                    self.assertTrue(np.isfinite(a[:, mask]).all())
+
+
+if __name__ == "__main__": unittest.main()

tests/validation/multigrid/test_refinement_multipatch.py

@@ -50,8 +50,8 @@ class MultipatchCPU(unittest.TestCase):
                    dict(fine_block=(5, 4), base_block=(4, 4))):
             with self.subTest(kw=kw), self.assertRaises(ValueError):
                 FixedPatchStepper(cfg, ((4, 4, 8, 12), (14, 4, 18, 12)), context=context, **kw)
-        with self.assertRaisesRegex(ValueError, "closer"):
-            FixedPatchStepper(cfg, ((4, 4, 8, 12), (10, 4, 14, 12)), context=context,
+        with self.assertRaisesRegex(ValueError, "margin"):
+            FixedPatchStepper(cfg, ((1, 4, 8, 12), (10, 4, 14, 12)), context=context,
                               fine_block=(8, 4), base_block=(4, 4))
         from CelerisLab.body.jets import Jet, Parallel, Tube
         cfg32 = LBMConfig(nx=32, ny=16, viscosity=.1)

tests/validation/multigrid/test_shifted_jets_gpu.py

@@ -0,0 +1,208 @@
+"""Shifted velocity-Jet GPU gates. Run only on an authorized device.
+
+PYTHONPATH=src python tests/validation/multigrid/test_shifted_jets_gpu.py -v
+Uniform initial curved-body states rho=.99/1.01; amplitude 0,+.002,-.001,0;
+four-level substep ledger; three-level live-Jet raw remap. No pressure actuator.
+The census is an implementation-exchange ledger, not an independent physical
+traction/nozzle-thrust model. Frozen inherited caps: raw rows rtol=2e-6/atol=2e-7; fluid ledger
+rtol=2e-4/atol=8e-5. Observer on/off, interval accumulation and raw copies exact.
+Report actual residuals too: passing old loose caps alone is not an accuracy gain.
+"""
+from contextlib import ExitStack
+import unittest
+import numpy as np
+import pycuda.driver as cuda
+from CelerisLab.config import LBMConfig
+from CelerisLab.cuda.context import CudaContext
+from CelerisLab.body.geometry.circle import CircleGeometry
+from CelerisLab.body.objects import SimObject
+from CelerisLab.body.jets import Jet, Tube, Parallel
+from CelerisLab.lbm.descriptors import D2Q9_EX, D2Q9_EY, SOLID
+from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+from CelerisLab.lbm.population_storage import _StoredPopulations
+from CelerisLab.lbm.synchronized_remap import RemapTopology, _remap_stored
+
+CX,CY=np.asarray(D2Q9_EX),np.asarray(D2Q9_EY)
+OPP=np.array([np.flatnonzero((CX==-x)&(CY==-y))[0] for x,y in zip(CX,CY)])
+W=np.array([4/9,*([1/9]*4),*([1/36]*4)],np.float64)
+SPEEDS=(0.,.002,-.001,0.)
+
+
+def circle(center,radius):
+    return SimObject(17,geometry=CircleGeometry(*center,radius),center=center,radius=radius)
+
+
+def drive(center,radius):
+    return Jet(Tube(center,0.,radius*1.5,radius*.32),Parallel(0.),amplitude=0.)
+
+
+def initial(run,rho,u=.013,v=-.007):
+    cu=3*(CX*u+CY*v)
+    g=W*((rho-1)+rho*(cu+.5*cu*cu-1.5*(u*u+v*v)))
+    return _StoredPopulations(tuple(np.broadcast_to(g[:,None,None],l.shape).astype(np.float32).copy()
+                                    for l in run._levels),True)
+
+
+def download(pointer,shape,dtype=np.float32):
+    out=np.empty(shape,dtype);cuda.memcpy_dtoh(out,pointer);return out
+
+
+def raw_totals(a,leaf):
+    g=a[:,leaf].astype(np.float64)
+    # Only differences at the same fixed leaf mask are used, so mass's unit
+    # background cancels. Momentum never requires adding weights.
+    return np.array([(g*CX[:,None]).sum(),(g*CY[:,None]).sum(),g.sum()])
+
+
+def census(run,index,buffer):
+    """Independent physical link/face census of actual staged receiver windows."""
+    level=run._levels[index];layout=level.layout;b=level.boundary
+    stored=download(buffer,layout.blocked_shape);dense=layout.to_dense(stored)
+    leaf=run._leaf[index];ny,nx=leaf.shape
+    solid=(b.host.flags.reshape(ny,nx)&SOLID)!=0
+    y,x=np.nonzero(leaf);owner=layout.owner[y,x].astype(np.intp)
+    cells=int(np.prod(layout.local_shape));block,cell=np.divmod(owner,cells)
+    flat=stored.reshape(layout.nblocks,9,cells);external=np.zeros(3,np.float64)
+    for q,(cx,cy) in enumerate(zip(CX,CY)):
+        sy,sx=y-cy,x-cx;dy,dx=y+cy,x+cx
+        incoming=flat[block,q,cell-cy*layout.local_shape[1]-cx].astype(np.float64)+W[q]
+        outgoing=flat[block,q,cell].astype(np.float64)+W[q]
+        entering=~leaf[sy,sx]&~solid[sy,sx];leaving=~leaf[dy,dx]&~solid[dy,dx]
+        external+=np.array([cx,cy,1])*(incoming[entering].sum()-outgoing[leaving].sum())
+    receiver,direction,_,_,rx,ry,_,_=b.host.curved_host
+    owner=layout.owner.ravel()[receiver].astype(np.intp);block,cell=np.divmod(owner,cells)
+    outgoing=flat[block,direction,cell].astype(np.float64)
+    source=cell+CY[direction]*layout.local_shape[1]+CX[direction]
+    reflected=flat[block,OPP[direction],source].astype(np.float64)
+    physical_sum=outgoing+reflected+2*W[direction]
+    jx,jy=CX[direction]*physical_sum,CY[direction]*physical_sum
+    rows=np.column_stack((jx,jy,rx*jy-ry*jx,reflected-outgoing))
+    return raw_totals(dense,leaf),external,rows
+
+
+class Ledger:
+    def __init__(self,test,run,index):
+        self.test,self.run,self.index=test,run,index
+        self.samples=[];self.max_row=0.;self.max_balance=0.
+        level=run._levels[index];self.original=level.kernel;b=level.boundary
+        def audited(src,dst,*args,**kwargs):
+            run._stream.synchronize()
+            before,external,independent=census(run,index,src)
+            raw=download(b.last_raw,(len(b.tables['cell']),4))
+            np.testing.assert_allclose(raw,independent,rtol=2e-6,atol=2e-7)
+            self.max_row=max(self.max_row,float(abs(raw.astype(float)-independent).max()))
+            self.original(src,dst,*args,**kwargs);run._stream.synchronize()
+            after=level.layout.to_dense(download(dst,level.layout.blocked_shape))
+            change=raw_totals(after,run._leaf[index])-before
+            expected=external+np.array([-independent[:,0].sum(),-independent[:,1].sum(),independent[:,3].sum()])
+            np.testing.assert_allclose(change,expected,rtol=2e-4,atol=8e-5)
+            self.max_balance=max(self.max_balance,float(abs(change-expected).max()))
+            self.samples.append(raw.copy())
+        level.kernel=audited
+
+    def step(self):
+        run=self.run;self.samples.clear();run.step()
+        b=run._level_boundaries[self.index];h=run._exchange_spacing[self.index]
+        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])
+        for raw in self.samples:expected+=raw.astype(np.float64)*weights
+        actual=download(b.interval_integral,expected.shape,np.float64)
+        np.testing.assert_array_equal(actual,expected)
+        obs=run.read_raw_exchange(0,level=self.index)
+        for name in ('impulse','angular_impulse','mass_into_fluid'):
+            np.testing.assert_allclose(obs['total'][name],np.asarray(obs['wall'][name])+obs['aperture'][name],rtol=0,atol=2e-13)
+        self.test.assertEqual(obs['sample_count'],2**self.index)
+        for a,leaf in zip(run._download_raw().values,run._leaf):
+            self.test.assertTrue(np.isfinite(a[:,leaf]).all())
+            self.test.assertTrue((1+a[:,leaf].sum(axis=0,dtype=np.float64)>0).all())
+        return obs
+
+    def report(self,label):
+        print(f'{label}: max_abs_raw_row={self.max_row:.9g}, max_abs_fluid_ledger={self.max_balance:.9g}',flush=True)
+
+
+class ShiftedJetsGPU(unittest.TestCase):
+    def exact_raw(self,a,b):
+        self.assertTrue(a.shifted and b.shifted)
+        for x,y in zip(a.values,b.values):self.assertEqual(x.tobytes(),y.tobytes())
+
+    def test_matching_whole_circle_parallel_velocity_preserves_uniform_state(self):
+        cfg=LBMConfig(nx=32,ny=24,viscosity=.1,ddf_shifting=True)
+        center=(16.,12.);radius=2.15;obj=circle(center,radius)
+        # Rectangle contains every circle hit; prescribed surface velocity equals
+        # the uniform fluid velocity. Geometry itself remains fixed.
+        whole=Jet(Tube((center[0]-radius-1,center[1]),0.,2*radius+2,radius+1),
+                  Parallel(0.),amplitude=.013)
+        for rho in (.99,1.01):
+            with self.subTest(rho=rho),CudaContext(0) as ctx:
+                with FixedPatchStepper(cfg,(8,6,24,18),context=ctx,base_block=(4,4),fine_block=(4,4),
+                        fine_objects=(obj,),jets=(whole,),observe_raw_exchange=True) as run:
+                    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())
+                    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))
+                    self.assertLessEqual(error,1e-6)  # inherited uniform-state cap
+                    print(f'matched whole-circle rho={rho}: max_abs_raw_state={error:.9g}',flush=True)
+                    ledger.report(f'matched whole-circle rho={rho}')
+
+    def test_uniform_initial_curved_rho_both_signs_switch_and_observer(self):
+        cfg=LBMConfig(nx=32,ny=24,viscosity=.1,ddf_shifting=True)
+        obj=circle((16.,12.),2.15);j=drive((16.,12.),2.15)
+        for rho in (.99,1.01):
+            with self.subTest(rho=rho),CudaContext(0) as ctx,ExitStack() as owned:
+                settings=dict(bounds=(8,6,24,18),context=ctx,base_block=(4,4),fine_block=(4,4),fine_objects=(obj,),jets=(j,))
+                measured=owned.enter_context(FixedPatchStepper(cfg,observe_raw_exchange=True,**settings))
+                plain=owned.enter_context(FixedPatchStepper(cfg,**settings))
+                measured._initialize_raw(initial(measured,rho));plain._initialize_raw(initial(plain,rho))
+                ledger=Ledger(self,measured,1)
+                for speed in SPEEDS:
+                    before=measured._download_raw();count=measured.step_count
+                    measured.set_jet_amplitude(0,speed);plain.set_jet_amplitude(0,speed)
+                    self.exact_raw(before,measured._download_raw());self.assertEqual(measured.step_count,count)
+                    ledger.step();plain.step()
+                    self.exact_raw(measured._download_raw(),plain._download_raw())
+                ledger.report(f'L1 rho={rho}')
+
+    def test_four_level_interval_ledger(self):
+        cfg=LBMConfig(nx=48,ny=48,viscosity=.02,ddf_shifting=True)
+        obj=circle((16.,16.),.5);j=drive((16.,16.),.5)
+        with CudaContext(0) as ctx,ExitStack() as owned:
+            settings=dict(context=ctx,refinements=[[(6,6,26,26)],[(20,20,44,44)],[(56,56,72,72)]],
+                          level_blocks=[(4,4)]*4,level_objects=[(),(),(),(obj,)],level_jets=[(),(),(),(j,)])
+            measured=owned.enter_context(FixedPatchStepper(cfg,observe_raw_exchange=True,**settings))
+            plain=owned.enter_context(FixedPatchStepper(cfg,**settings))
+            measured._initialize_raw(initial(measured,1.01));plain._initialize_raw(initial(plain,1.01))
+            ledger=Ledger(self,measured,3)
+            for speed in SPEEDS:
+                measured.set_jet_amplitude(0,speed,level=3);plain.set_jet_amplitude(0,speed,level=3)
+                ledger.step();plain.step();self.exact_raw(measured._download_raw(),plain._download_raw())
+            ledger.report('L3 rho=1.01')
+
+    def test_live_jet_remap_then_both_ledgers_continue(self):
+        cfg=LBMConfig(nx=48,ny=32,viscosity=.1,ddf_shifting=True)
+        parents=((4,4,20,28),(24,4,44,28));a=((16,16,32,40),(56,16,72,40));b=(a[0],(60,16,76,40))
+        obj=circle((12.,14.),.8);j=drive((12.,14.),.8)
+        with CudaContext(0) as ctx,ExitStack() as owned:
+            old=owned.enter_context(FixedPatchStepper(cfg,parents,context=ctx,base_block=(4,4),fine_block=(4,4),
+                child_block=(4,4),child_rectangles_in_parent=a,child_objects=(obj,),child_jets=(j,),observe_raw_exchange=True))
+            old._initialize_raw(initial(old,.99));ledger=Ledger(self,old,2)
+            for speed in (.002,-.001):old.set_jet_amplitude(0,speed,level=2);ledger.step()
+            before=old._download_raw();count=old.step_count
+            topology=lambda rects:RemapTopology.from_config(cfg,old.patch,rects,child_objects=(obj,),
+                child_jets=old._level_jets[2],_shifted_blocked=True)
+            expected=_remap_stored(topology(a),topology(b),before,shifted=True,base_count=count,
+                level_steps=(count,2*count,4*count),phase='postcollision_g')
+            new=owned.enter_context(old.remap_child(b,child_objects=(obj,)))
+            self.exact_raw(before,old._download_raw())
+            self.exact_raw(_StoredPopulations(expected.g,True),new._download_raw())
+            self.assertEqual(new._level_jets[2][0].amplitude,-.001)
+            np.testing.assert_array_equal(new._level_boundaries[2].tables['surface_velocity'],old._level_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')
+
+
+if __name__=='__main__':unittest.main()

tests/validation/multigrid/test_shifted_remap_gpu.py

@@ -0,0 +1,113 @@
+"""Opt-in GPU raw shifted remap lifecycle gates; no precision/performance claim.
+
+Run only on an authorized GPU:
+  PYTHONPATH=src python tests/validation/multigrid/test_shifted_remap_gpu.py -v
+Frozen gates: bitwise raw overlap/copy and identity continuation; moved-remap
+matches the existing host fill loop exactly; finite owned leaves after stepping.
+"""
+from contextlib import ExitStack
+import unittest
+import numpy as np
+import pycuda.driver as cuda
+from CelerisLab.config import LBMConfig
+from CelerisLab.cuda.context import CudaContext
+from CelerisLab.lbm.population_storage import _StoredPopulations
+from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+from CelerisLab.lbm.synchronized_remap import RemapTopology, _remap_stored
+
+CFG=LBMConfig(nx=48,ny=32,viscosity=.1,ddf_shifting=True)
+PARENTS=((4,4,20,28),(24,4,44,28))
+A=((16,16,32,40),)
+B=((20,16,36,40),)
+
+
+def make(ctx):
+    return FixedPatchStepper(CFG,PARENTS,context=ctx,base_block=(4,4),
+                             fine_block=(4,4),child_block=(4,4),
+                             child_rectangles_in_parent=A)
+
+
+def initial(run):
+    # Non-unit density, nonzero velocity and nonequilibrium perturbations.
+    w=np.array([4/9,*([1/9]*4),*([1/36]*4)],np.float64)
+    cx=np.array([0,1,-1,0,0,1,-1,1,-1]);cy=np.array([0,0,0,1,-1,1,-1,-1,1])
+    rho,u,v=.99,.013,-.007
+    cu=3*(cx*u+cy*v)
+    eq=w*((rho-1)+rho*(cu+.5*cu*cu-1.5*(u*u+v*v)))
+    rng=np.random.default_rng(8473)
+    arrays=tuple((np.broadcast_to(eq[:,None,None],level.shape)+
+                  rng.normal(0,1e-7,level.shape)).astype(np.float32) for level in run._levels)
+    return _StoredPopulations(arrays,True)
+
+
+class ShiftedRemapGPU(unittest.TestCase):
+    def exact(self,a,b,label):
+        self.assertEqual(a.shape,b.shape,label)
+        self.assertEqual(a.dtype,b.dtype,label)
+        self.assertEqual(a.tobytes(),b.tobytes(),label)
+
+    def snapshot_equal(self,a,b,label):
+        self.assertTrue(a.shifted and b.shifted)
+        self.assertEqual(len(a.values),len(b.values))
+        for x,y in zip(a.values,b.values):self.exact(x,y,label)
+
+    def check_reset(self,run,wanted):
+        for level,want in zip(run._levels,wanted):
+            blocked=level.layout.to_blocked(want)
+            for pointer in (level.current,level.spare):
+                actual=np.empty(level.layout.blocked_shape,np.float32)
+                cuda.memcpy_dtoh(actual,pointer)
+                self.exact(actual,blocked,'fresh current and spare')
+
+    def check_finite(self,run):
+        raw=run._download_raw()
+        self.assertTrue(raw.shifted)
+        for a,leaf in zip(raw.values,run._leaf):
+            self.assertTrue(np.isfinite(a[:,leaf]).all())
+            self.assertTrue(np.isnan(a[:,~leaf]).all())
+            self.assertTrue((1+a[:,leaf].sum(axis=0,dtype=np.float64)>0).all())
+
+    def test_identity_next_step_bitwise(self):
+        with CudaContext(0) as ctx,ExitStack() as owned:
+            old=owned.enter_context(make(ctx))
+            old._initialize_raw(initial(old));old.step(2)
+            before=old._download_raw()
+            new=owned.enter_context(old.remap_child(A))
+            self.check_reset(new,before.values)
+            self.snapshot_equal(old._download_raw(),before,'old unchanged')
+            self.snapshot_equal(new._download_raw(),before,'identity raw copy')
+            old.step();new.step()
+            self.assertEqual((old.step_count,new.step_count),(3,3))
+            self.snapshot_equal(old._download_raw(),new._download_raw(),'identity next step')
+            self.check_finite(old);self.check_finite(new)
+
+    def test_moved_remap_matches_raw_host_and_old_remains_usable(self):
+        with CudaContext(0) as ctx,ExitStack() as owned:
+            old=owned.enter_context(make(ctx))
+            old._initialize_raw(initial(old));old.step(2)
+            before=old._download_raw();count=old.step_count
+            old_mesh=RemapTopology.from_config(CFG,old.patch,A)
+            new_mesh=RemapTopology.from_config(CFG,old.patch,B)
+            expected=_remap_stored(old_mesh,new_mesh,before,shifted=True,
+                base_count=count,level_steps=(count,2*count,4*count),phase='postcollision_g')
+            new=owned.enter_context(old.remap_child(B))
+            self.check_reset(new,expected.g)
+            self.snapshot_equal(new._download_raw(),_StoredPopulations(expected.g,True),'host raw fill')
+            for a,b,om,nm in zip(before.values,new._download_raw().values,old_mesh.leaf_masks,new_mesh.leaf_masks):
+                self.exact(a[:,om&nm],b[:,om&nm],'physical same-level overlap')
+            self.snapshot_equal(old._download_raw(),before,'source unchanged')
+            old.step();new.step(2)
+            self.assertEqual((old.step_count,new.step_count),(3,4))
+            self.check_finite(old);self.check_finite(new)
+
+    def test_rejected_remap_does_not_mutate_source(self):
+        with CudaContext(0) as ctx,make(ctx) as run:
+            run._initialize_raw(initial(run));run.step(2)
+            before=run._download_raw()
+            with self.assertRaises(ValueError):run.remap_child([])
+            self.snapshot_equal(run._download_raw(),before,'rejection preserves source')
+            self.assertEqual(run.step_count,2)
+            run.step();self.check_finite(run)
+
+
+if __name__=='__main__':unittest.main()

tests/validation/periodic.py

@@ -0,0 +1,130 @@
+"""CPU-only cycle extraction. All criteria frozen before any Re100 GPU result."""
+
+import numpy as np
+
+GATES = dict(
+    min_tstar=250.0,
+    cycles=5,
+    min_period=2.0,
+    max_period=5.0,
+    min_lift_amplitude=0.2,
+    period_relative_range=0.005,
+    extrema_absolute_range=0.005,
+    mass_relative_range=1e-4,
+    cycle_mean_flux_relative=0.005,
+)
+REFERENCE_GATES = dict(
+    cd_relative=0.02, cl_absolute=0.03, dp_relative=0.02, st_relative=0.02
+)
+
+
+def minima(t, y):
+    out = []
+    for i in range(1, len(y) - 1):
+        if y[i] < y[i - 1] and y[i] <= y[i + 1]:
+            # Quadratic sub-sample location; only separated physical cycles accepted.
+            coef = np.polyfit(t[i - 1 : i + 2] - t[i], y[i - 1 : i + 2], 2)
+            v = t[i] - coef[1] / (2 * coef[0]) if coef[0] > 0 else t[i]
+            if t[i - 1] <= v <= t[i + 1] and (not out or v - out[-1] > 1.0):
+                out.append(float(v))
+    return out
+
+
+def extract(rows, min_time=250.0):
+    if len(rows) < 3:
+        return []
+    t = np.array([r["tstar_force"] for r in rows])
+    cl = np.array([r["cl"] for r in rows])
+    ends = minima(t, cl)
+    cycles = []
+    for a, b in zip(ends, ends[1:]):
+        if a < min_time or not GATES["min_period"] < b - a < GATES["max_period"]:
+            continue
+        grid = np.linspace(a, b, 257)
+        c = dict(start=a, end=b, period=b - a, st=1 / (b - a))
+        for key in ("cd", "cl", "dp", "mass", "flux_in", "flux_out"):
+            times = np.array(
+                [r["tstar" if key not in ("cd", "cl") else "tstar_force"] for r in rows]
+            )
+            # No endpoint extrapolation: pressure/mass must cover the whole cycle.
+            if a < times[0] or b > times[-1]:
+                break
+            v = np.interp(grid, times, [r[key] for r in rows])
+            c[key + "_min"] = float(v.min())
+            c[key + "_max"] = float(v.max())
+            c[key + "_midrange"] = float((v.min() + v.max()) / 2)
+            c[key + "_amplitude"] = float(v.max() - v.min())
+            c[key + "_time_mean"] = float(
+                np.sum((v[1:] + v[:-1]) * np.diff(grid) / 2) / (b - a)
+            )
+            if key in ("cd", "cl", "dp"):
+                c[key + "_waveform"] = v.tolist()
+        else:
+            c["flux_relative"] = abs(
+                c["flux_in_time_mean"] - c["flux_out_time_mean"]
+            ) / max(abs(c["flux_in_time_mean"]), 1e-30)
+            cycles.append(c)
+    return cycles
+
+
+def assess(rows):
+    cycles = extract(rows, min_time=GATES["min_tstar"])
+    tail = cycles[-GATES["cycles"] :]
+    checks = {}
+    checks["enough_cycles"] = len(tail) == GATES["cycles"]
+    if checks["enough_cycles"]:
+        checks["consecutive"] = all(
+            abs(a["end"] - b["start"]) < 1e-8 for a, b in zip(tail, tail[1:])
+        )
+        checks["recent"] = (
+            0 <= rows[-1]["tstar_force"] - tail[-1]["end"] <= GATES["max_period"]
+        )
+        checks["shedding"] = (
+            min(c["cl_amplitude"] for c in tail) >= GATES["min_lift_amplitude"]
+        )
+        periods = [c["period"] for c in tail]
+        checks["period"] = (
+            np.ptp(periods) / np.mean(periods) <= GATES["period_relative_range"]
+        )
+        checks["extrema"] = all(
+            np.ptp([c[k + "_" + e] for c in tail]) <= GATES["extrema_absolute_range"]
+            for k in ("cd", "cl", "dp")
+            for e in ("min", "max")
+        )
+        mass = [c["mass_time_mean"] for c in tail]
+        checks["mass"] = (
+            np.ptp(mass) / abs(np.mean(mass)) <= GATES["mass_relative_range"]
+        )
+        checks["flux"] = (
+            max(c["flux_relative"] for c in tail) <= GATES["cycle_mean_flux_relative"]
+        )
+    return dict(
+        periodic=bool(checks["enough_cycles"] and all(checks.values())),
+        checks={k: bool(v) for k, v in checks.items()},
+        cycles=cycles,
+    )
+
+
+def compare(cycle, ref):
+    errors = {}
+    passed = True
+    for key in ("cd", "cl", "dp"):
+        for ext in ("min", "max"):
+            k = key + "_" + ext
+            e = abs(cycle[k] - ref[k])
+            e = e if key == "cl" else e / abs(ref[k])
+            errors[k] = e
+            passed &= (
+                e
+                <= REFERENCE_GATES[key + ("_absolute" if key == "cl" else "_relative")]
+            )
+        errors[key + "_waveform_rms"] = float(
+            np.sqrt(
+                np.mean(
+                    (np.array(cycle[key + "_waveform"]) - ref[key + "_waveform"]) ** 2
+                )
+            )
+        )
+    errors["st_relative"] = abs(cycle["st"] - ref["st"]) / ref["st"]
+    passed &= errors["st_relative"] <= REFERENCE_GATES["st_relative"]
+    return dict(reference_agreement=bool(passed), errors=errors)

tests/validation/re100_helpers.py

@@ -0,0 +1,128 @@
+"""DFG Re100 periodic reference comparison with declared pressure-NEQ outlet, not exact do-nothing.
+
+CPU helpers import without CUDA. GPU execution verifies the explicitly pinned
+AMR-channel experimental source snapshot. Historical d79 evidence remains tied to
+its original commits. Default pilot is too short to certify periodic convergence.
+"""
+
+import argparse
+import csv
+import hashlib
+import json
+import math
+import os
+from pathlib import Path
+import time
+from tempfile import TemporaryDirectory
+import numpy as np
+
+ROOT = Path(__file__).resolve().parents[2]
+PIN = "47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5"
+
+
+def geometry(n):
+    if isinstance(n, bool) or n < 10 or n % 10:
+        raise ValueError("Diameter must be a positive multiple of 10, at least 10")
+    return dict(
+        diameter=n,
+        nx=22 * n + 1,
+        ny=41 * n // 10 + 2,
+        center=[2 * n, 2 * n + 0.5],
+        radius=n / 2,
+        dx=0.1 / n,
+        bottom=0.5,
+        top=41 * n / 10 + 0.5,
+        inlet_source_x=0.0,
+        outlet_source_x=22.0 * n,
+    )
+
+
+def parameters(n, speed):
+    g = geometry(n)
+    if not 0 < speed <= 0.04:
+        raise ValueError("Require 0 < lattice mean speed <= .04")
+    return dict(
+        **g,
+        mean_speed=speed,
+        peak_speed=1.5 * speed,
+        viscosity=speed * n / 100,
+        dt=speed * g["dx"],
+        peak_mach=1.5 * speed * math.sqrt(3),
+        re=100.0,
+    )
+
+
+def make_config(n, speed, collision, outlet, inlet="regularized"):
+    p = parameters(n, speed)
+    cfg = json.loads((ROOT / "src/CelerisLab/configs/config_lbm.json").read_text())
+    cfg["grid"].update(nx=p["nx"], ny=p["ny"], nz=1, lattice_model="D2Q9")
+    cfg["physics"].update(velocity=speed, viscosity=p["viscosity"], rho=1.0)
+    cfg["method"].update(
+        collision=collision,
+        streaming="double_buffer",
+        store_precision="FP32",
+        ddf_shifting=True,
+        y_wall_bc="bounce_back",
+    )
+    cfg["method"]["les"]["enabled"] = False
+    cfg["method"]["inlet"].update(profile="parabolic", scheme=inlet)
+    cfg["method"]["outlet"].update(mode=outlet, backflow_clamp=False, blend_alpha=0.7)
+    body = {
+        "objects": [
+            dict(type="cylinder", center=p["center"], radius=p["radius"], omega=0.0)
+        ]
+    }
+    return cfg, body, p
+
+
+def bilinear_fluid(a, fluid, x, y):
+    ix, iy = math.floor(x), math.floor(y)
+    tx, ty = x - ix, y - iy
+    result = 0.0
+    for xx, yy, w in (
+        (ix, iy, (1 - tx) * (1 - ty)),
+        (ix + 1, iy, tx * (1 - ty)),
+        (ix, iy + 1, (1 - tx) * ty),
+        (ix + 1, iy + 1, tx * ty),
+    ):
+        if w == 0:
+            continue
+        if not (0 <= yy < a.shape[0] and 0 <= xx < a.shape[1]) or not fluid[yy, xx]:
+            raise ValueError("Pressure interpolation touched a non-fluid/outside node")
+        result += w * float(a[yy, xx])
+    return result
+
+
+def pressure_probes(rho, fluid, p):
+    # Linear extrapolation from distances h and 2h outside the cylinder to
+    # each exact wall probe, after bilinear interpolation using only fluid.
+    cx, cy = p["center"]
+    r = p["radius"]
+    values = []
+    for sign in (-1, 1):
+        wall = cx + sign * r
+        near = bilinear_fluid(rho, fluid, wall + sign, cy)
+        far = bilinear_fluid(rho, fluid, wall + 2 * sign, cy)
+        values.append(2 * near - far)
+    scale = (1.0 / p["mean_speed"]) ** 2 / 3
+    return dict(
+        rho_front=values[0], rho_back=values[1], dp=(values[0] - values[1]) * scale
+    )
+
+
+def outlet_diagnostic(rho, ux, uy, p):
+    # Extrapolate fluid values to the declared x=NX-1 source plane. Diagnostic
+    # excludes two rows beside each wall; not an exact discrete BC identity.
+    rows = slice(3, -3)
+    i = -2
+    dudx = (5 * ux[rows, i] - 8 * ux[rows, i - 1] + 3 * ux[rows, i - 2]) / 2
+    dvdx = (5 * uy[rows, i] - 8 * uy[rows, i - 1] + 3 * uy[rows, i - 2]) / 2
+    rho_plane = 2 * rho[rows, i] - rho[rows, i - 1]
+    pressure = (rho_plane - 1) / 3
+    scale = p["mean_speed"] ** 2
+    residual = np.stack((p["viscosity"] * dudx - pressure, p["viscosity"] * dvdx))
+    return dict(
+        rms_over_rho_U2=float(np.sqrt(np.mean(residual**2)) / scale),
+        max_over_rho_U2=float(np.max(abs(residual)) / scale),
+        definition="fluid-side one-sided estimate of nu*d_x(u)-p*n at declared outlet",
+    )

tests/validation/run_dfg_re100.py

@@ -0,0 +1,221 @@
+"""Uniform DFG 2D-2 Re100, existing Zou-He inlet / noncanonical pressure-NEQ outlet."""
+
+import argparse, csv, hashlib, json, math, os, time
+from pathlib import Path
+from tempfile import TemporaryDirectory
+import numpy as np
+from re100_helpers import ROOT, PIN, make_config, pressure_probes, outlet_diagnostic
+from periodic import assess, compare, GATES, REFERENCE_GATES
+from source_provenance import checkout_provenance
+
+
+def run(a):
+    import CelerisLab
+    from CelerisLab.simulation import Simulation
+    from CelerisLab.lbm.descriptors import FLUID
+
+    if Path(CelerisLab.__file__).resolve() != ROOT / "src/CelerisLab/__init__.py":
+        raise RuntimeError("Wrong source checkout")
+    manifest = json.loads(
+        (ROOT / "tests/validation/data/dfg_re100/source_manifest.json").read_text()
+    )
+    for rel, digest in manifest["production_sha256"].items():
+        if hashlib.sha256((ROOT / rel).read_bytes()).hexdigest() != digest:
+            raise RuntimeError("Production mismatch: " + rel)
+    if a.output.exists() and any(a.output.iterdir()):
+        raise ValueError("Output must be empty/new")
+    a.output.mkdir(parents=True, exist_ok=True)
+    cfg, body, p = make_config(a.diameter, a.speed, "MRT", "neq_extrap", "zou_he_local")
+    cap = 5.0 if a.mode == "pilot" else 300.0
+    max_steps = math.ceil(cap * p["diameter"] / p["mean_speed"])
+    result = dict(
+        case="DFG 2D-2 Re100, pressure-NEQ outlet deviation",
+        pin=PIN,
+        source_experiment=manifest["experiment"],
+        **checkout_provenance(ROOT),
+        production_verified=True,
+        parameters=p,
+        mode=a.mode,
+        cap_tstar=cap,
+        inlet=cfg["method"]["inlet"],
+        outlet=cfg["method"]["outlet"],
+        canonical_outlet="nu*d_n(u)-p*n=0; NOT imposed exactly",
+        initialization="unchanged production channel-flow initialization, not zero initial velocity",
+        force_phase="batch-averaged link momentum exchange, timestamp at batch midpoint",
+        pressure_phase="postcollision end-sample; interpolated onto force-cycle times for comparison",
+        gates=GATES,
+        reference_gates=REFERENCE_GATES,
+        status="running",
+        history=[],
+        suite_validated=False,
+    )
+    reference = json.loads(
+        (ROOT / "tests/validation/data/dfg_re100/reference_cycle.json").read_text()
+    )
+    started = time.monotonic()
+    step_seconds = 0.0
+    healthy = True
+    with TemporaryDirectory() as tmp:
+        cp, bp = Path(tmp) / "lbm.json", Path(tmp) / "body.json"
+        cp.write_text(json.dumps(cfg))
+        bp.write_text(json.dumps(body))
+        with Simulation(str(cp), str(bp), device_id=0) as sim:
+            sim.initialize()
+            result["startup_seconds"] = time.monotonic() - started
+            fluid = (sim.field.flag.reshape(p["ny"], p["nx"]) & FLUID) != 0
+            np.savez_compressed(
+                a.output / "geometry.npz",
+                flag=sim.field.flag.reshape(p["ny"], p["nx"]),
+                x=np.arange(p["nx"]) * p["dx"],
+                y=(np.arange(p["ny"]) - 0.5) * p["dx"],
+            )
+            with (a.output / "history.csv").open("w") as f:
+                writer = None
+                while sim.stepper.step_count < max_steps:
+                    begin = sim.stepper.step_count
+                    tstar = begin * p["mean_speed"] / p["diameter"]
+                    # Pilot exercises the expensive dense measurement path too.
+                    width = 0.02 if a.mode == "pilot" or tstar >= 240 else 0.5
+                    n = min(
+                        max(1, round(width * p["diameter"] / p["mean_speed"])),
+                        max_steps - begin,
+                    )
+                    impulse = np.zeros(2)
+                    timer = time.monotonic()
+                    done = 0
+                    while done < n:
+                        batch = min(100, n - done)
+                        sim.run(batch)
+                        impulse += sim.read_force(0).astype(np.float64) * batch
+                        done += batch
+                    step_seconds += time.monotonic() - timer
+                    macro = sim.get_macroscopic()
+                    rho, ux, uy = [
+                        np.asarray(macro[k], np.float64).reshape(p["ny"], p["nx"])
+                        for k in ("rho", "ux", "uy")
+                    ]
+                    healthy = bool(
+                        all(np.isfinite(v[fluid]).all() for v in (rho, ux, uy))
+                        and np.isfinite(impulse).all()
+                        and (rho[fluid] > 0).all()
+                    )
+                    if not healthy:
+                        break
+                    den = p["mean_speed"] ** 2 * p["diameter"] / 2
+                    row = dict(
+                        step=sim.stepper.step_count,
+                        tstar=sim.stepper.step_count * p["mean_speed"] / p["diameter"],
+                        tstar_force=(begin + n / 2) * p["mean_speed"] / p["diameter"],
+                        cd=float(impulse[0] / n / den),
+                        cl=float(impulse[1] / n / den),
+                        mass=float(rho[fluid].sum()),
+                        rho_min=float(rho[fluid].min()),
+                        rho_max=float(rho[fluid].max()),
+                        max_mach=float(
+                            np.sqrt(3 * (ux[fluid] ** 2 + uy[fluid] ** 2)).max()
+                        ),
+                        flux_in=float(np.sum(rho[1:-1, 2] * ux[1:-1, 2])),
+                        flux_out=float(np.sum(rho[1:-1, -3] * ux[1:-1, -3])),
+                        **pressure_probes(rho, fluid, p),
+                    )
+                    result["history"].append(row)
+                    if (
+                        row["max_mach"] > 0.2
+                        or row["rho_min"] < 0.5
+                        or row["rho_max"] > 1.5
+                    ):
+                        healthy = False
+                        result["safety_stop"] = "Mach > .2 or density outside [.5,1.5]"
+                    if writer is None:
+                        writer = csv.DictWriter(f, fieldnames=row)
+                        writer.writeheader()
+                    writer.writerow(row)
+                    if not healthy:
+                        f.flush()
+                        break
+                    if len(result["history"]) == 1:
+                        print(
+                            "READY "
+                            + json.dumps(
+                                dict(
+                                    pid=os.getpid(),
+                                    step=row["step"],
+                                    tstar=row["tstar"],
+                                    all_runtime_and_macro_paths_exercised=True,
+                                )
+                            ),
+                            flush=True,
+                        )
+                    if len(result["history"]) % 100 == 0:
+                        f.flush()
+                        print(json.dumps(row), flush=True)
+                result.update(
+                    steps=sim.stepper.step_count,
+                    step_seconds=step_seconds,
+                    wall_seconds=time.monotonic() - started,
+                )
+                result["measurement_seconds"] = max(
+                    0.0,
+                    result["wall_seconds"] - result["startup_seconds"] - step_seconds,
+                )
+                result["mlups"] = (
+                    p["nx"] * p["ny"] * result["steps"] / max(step_seconds, 1e-30) / 1e6
+                )
+                result.update(assess(result["history"]))
+                if not healthy:
+                    result["periodic"] = False
+                result["status"] = (
+                    "invalid_state"
+                    if not healthy
+                    else (
+                        "pilot_complete"
+                        if a.mode == "pilot"
+                        else (
+                            "periodic"
+                            if result["periodic"]
+                            else "incomplete_periodicity"
+                        )
+                    )
+                )
+                if healthy:
+                    np.savez_compressed(
+                        a.output / "final_field.npz", rho=rho, ux=ux, uy=uy, fluid=fluid
+                    )
+                    result["outlet_residual"] = outlet_diagnostic(rho, ux, uy, p)
+                if result["cycles"]:
+                    result["reference_comparison"] = compare(
+                        result["cycles"][-1], reference
+                    )
+                result["case_reference_pass"] = bool(
+                    healthy
+                    and result["periodic"]
+                    and result.get("reference_comparison", {}).get(
+                        "reference_agreement", False
+                    )
+                )
+    (a.output / "summary.json").write_text(
+        json.dumps(result, indent=2, allow_nan=False)
+    )
+    print(
+        json.dumps({k: v for k, v in result.items() if k not in ("history", "cycles")}),
+        flush=True,
+    )
+    return (
+        1 if not healthy else (2 if a.mode == "full" and not result["periodic"] else 0)
+    )
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--diameter", type=int, choices=[20, 40, 80], default=20)
+    parser.add_argument("--speed", type=float, choices=[0.01, 0.02], default=0.01)
+    parser.add_argument("--mode", choices=["pilot", "full"], default="pilot")
+    parser.add_argument("--output", type=Path, required=True)
+    a = parser.parse_args()
+    if not os.environ.get("CUDA_VISIBLE_DEVICES"):
+        raise SystemExit("Select physical GPU with CUDA_VISIBLE_DEVICES")
+    return run(a)
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())

tests/validation/run_jet_circle_cv.py

@@ -0,0 +1,625 @@
+#!/usr/bin/env python3
+"""Bounded jet-on periodic cylinder: actual AMR Jet API vs blocked validation adapter."""
+
+import argparse
+import hashlib
+import json
+import os
+from pathlib import Path
+import sys
+import time
+import numpy as np
+from source_provenance import checkout_provenance
+
+PIN = "47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5"
+HERE = Path(__file__).resolve().parent
+CX = np.array([0, 1, -1, 0, 0, 1, -1, 1, -1])
+CY = np.array([0, 0, 0, 1, -1, 1, -1, -1, 1])
+W = np.array([4 / 9] + [1 / 9] * 4 + [1 / 36] * 4)
+
+
+def case(scale=1, re=100, steps=800):
+    u = 0.03 / scale
+    d = 24.084 * scale
+    return dict(
+        scale=scale,
+        re_d=re,
+        beta=0.1,
+        nx=192 * scale,
+        ny=128 * scale,
+        diameter=d,
+        center=[96 * scale, 64 * scale],
+        u=u,
+        jet_speed=0.1 * u,
+        nu=u * d / re,
+        tau_root=0.5 + 3 * u * d / re,
+        tau_fine=0.5 + 6 * u * d / re,
+        bounds=[48 * scale, 32 * scale, 144 * scale, 96 * scale],
+        contours=[
+            [72 * scale, 44 * scale, 120 * scale, 84 * scale],
+            [64 * scale, 40 * scale, 128 * scale, 88 * scale],
+        ],
+        steps=steps * scale**2,
+        window=100 * scale**2,
+        dx_root=1 / scale,
+        dt_root=1 / scale**2,
+        fine_dx_root_units=0.5,
+        fine_dt_root_units=0.5,
+        mach=np.sqrt(3) * u,
+        final_t_star=steps * scale**2 * u / d,
+    )
+
+
+def state(shape, u):
+    cu = CX[:, None, None] * u
+    return (
+        np.broadcast_to(
+            W[:, None, None] * (1 + 3 * cu + 4.5 * cu**2 - 1.5 * u * u), shape
+        )
+        .astype(np.float32)
+        .copy()
+    )
+
+
+def moments(g):
+    f = np.asarray(g, np.float64)
+    rho = f.sum(axis=0)
+    momentum = np.stack([np.einsum("q,qyx->yx", v, f) for v in (CX, CY)])
+    return rho, momentum, momentum / rho
+
+
+def cv(g, fluid, origin, rectangle, nu, h=0.5):
+    """Cell-centred inventory and midpoint face quadrature, physical root units.
+
+    Cauchy stress from centred velocity gradients, NOT postcollision non-equilibrium.
+    Rectangle faces lie on cell edges; fluid mask excludes the staircase cylinder.
+    """
+    rho, mom, vel = moments(g)
+    x0, y0, x1, y1 = np.rint((np.array(rectangle) - np.tile(origin, 2)) / h).astype(int)
+    np.testing.assert_allclose(
+        np.array([x0, y0, x1, y1]) * h + np.tile(origin, 2), rectangle, atol=1e-12
+    )
+    ny, nx = fluid.shape
+    if not (2 <= x0 < x1 <= nx - 2 and 2 <= y0 < y1 <= ny - 2):
+        raise ValueError("Contour needs two fluid-node stencil layers")
+    # Every face and its velocity-gradient stencil must be active fluid.
+    for ys, xs in [
+        (slice(y0 - 2, y1 + 2), slice(x0 - 2, x0 + 2)),
+        (slice(y0 - 2, y1 + 2), slice(x1 - 2, x1 + 2)),
+        (slice(y0 - 2, y0 + 2), slice(x0 - 2, x1 + 2)),
+        (slice(y1 - 2, y1 + 2), slice(x0 - 2, x1 + 2)),
+    ]:
+        if not fluid[ys, xs].all():
+            raise ValueError("Contour stencil intersects solid/inactive region")
+    use = fluid[y0:y1, x0:x1]
+    inventory = (
+        np.array(
+            [
+                rho[y0:y1, x0:x1][use].sum(),
+                mom[0, y0:y1, x0:x1][use].sum(),
+                mom[1, y0:y1, x0:x1][use].sum(),
+            ]
+        )
+        * h
+        * h
+    )
+    duy, dux = np.gradient(vel[0], h)
+    dvy, dvx = np.gradient(vel[1], h)
+    mu = rho * nu
+    # Subtract constant reference pressure, whose closed-contour integral is zero.
+    p = (rho - 1) / 3
+    tensor = np.array(
+        [
+            [
+                rho * vel[0] ** 2 + p - 2 * mu * dux,
+                rho * vel[0] * vel[1] - mu * (duy + dvx),
+            ],
+            [
+                rho * vel[0] * vel[1] - mu * (duy + dvx),
+                rho * vel[1] ** 2 + p - 2 * mu * dvy,
+            ],
+        ]
+    )
+    transport = np.concatenate(
+        [mom[None], tensor], axis=0
+    )  # [mass,Px,Py] x normal axis
+    flux = np.zeros(3)
+    for i, sign in [(x0, -1), (x1, 1)]:
+        flux += (
+            sign
+            * 0.5
+            * (transport[:, 0, y0:y1, i - 1] + transport[:, 0, y0:y1, i]).sum(axis=1)
+            * h
+        )
+    for j, sign in [(y0, -1), (y1, 1)]:
+        flux += (
+            sign
+            * 0.5
+            * (transport[:, 1, j - 1, x0:x1] + transport[:, 1, j, x0:x1]).sum(axis=1)
+            * h
+        )
+    if not np.isfinite(np.r_[inventory, flux]).all():
+        raise RuntimeError("Nonfinite CV quantities")
+    return inventory, flux
+
+
+def aperture_probe(fluid, origin, center, radius, h=0.5):
+    """Independent true-arc quadrature with nearest active-fluid sampling.
+
+    This is a first-order near-wall diagnostic, not certified wall reconstruction.
+    No target velocity or raw-population ledger is used to measure the flux.
+    """
+    z, w = np.polynomial.legendre.leggauss(64)
+    result = []
+    for angle in (np.pi / 2, 3 * np.pi / 2):
+        theta = angle + np.pi / 12 * z
+        er = np.column_stack((np.cos(theta), np.sin(theta)))
+        points = np.asarray(center) + radius * er
+        selected = []
+        distance = []
+        for point in points:
+            ix, iy = np.rint((point - np.asarray(origin)) / h - 0.5).astype(int)
+            candidates = [
+                (y, x)
+                for y in range(iy - 3, iy + 4)
+                for x in range(ix - 3, ix + 4)
+                if 0 <= y < fluid.shape[0] and 0 <= x < fluid.shape[1] and fluid[y, x]
+            ]
+            if not candidates:
+                raise ValueError("No active near-wall sampling node")
+            yy, xx = np.array(candidates).T
+            dist = np.linalg.norm(
+                np.column_stack((xx, yy)) * h + np.asarray(origin) + h / 2 - point,
+                axis=1,
+            )
+            k = np.argmin(dist)
+            selected.append(candidates[k])
+            distance.append(dist[k])
+        yy, xx = np.array(selected).T
+        result.append((yy, xx, er, w * radius * np.pi / 12, float(max(distance))))
+    return result
+
+
+def measured_aperture_flux(g, probe):
+    _, momentum, _ = moments(g)
+    return np.array(
+        [
+            np.sum(np.einsum("in,ni->n", momentum[:, y, x], normal) * ds)
+            for y, x, normal, ds, _ in probe
+        ]
+    )
+
+
+def verify(repo):
+    manifest = json.loads(
+        (HERE / "data/jet_circle_cv/source_manifest.json").read_text()
+    )
+    if manifest["revision"] != PIN:
+        raise RuntimeError("Wrong pinned revision")
+    for path, sha in manifest["files"].items():
+        if hashlib.sha256((repo / path).read_bytes()).hexdigest() != sha:
+            raise RuntimeError("Source hash mismatch: " + path)
+    return manifest
+
+
+def run_case(c, ctx, output, provenance, continuation=False):
+    from CelerisLab.body.geometry.circle import CircleGeometry
+    from CelerisLab.body.objects import SimObject
+    from CelerisLab.body.jets import Jet, Sector, Radial
+    from CelerisLab.config import LBMConfig
+    from CelerisLab.lbm.blocks import BlockLayout2D
+    from CelerisLab.lbm.block_stepper import PeriodicBlockStepper
+    from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+    from CelerisLab.lbm.descriptors import D2Q9_EX, D2Q9_EY
+    import pycuda.driver as cuda
+
+    np.testing.assert_array_equal(CX, D2Q9_EX)
+    np.testing.assert_array_equal(CY, D2Q9_EY)
+    nx, ny, u, nu = c["nx"], c["ny"], c["u"], c["nu"]
+    center = np.array(c["center"], float)
+    radius = c["diameter"] / 2
+    cfg = LBMConfig(
+        nx=nx,
+        ny=ny,
+        viscosity=nu,
+        velocity=0.0,
+        collision="SRT",
+        streaming="double_buffer",
+        store_precision="FP32",
+        ddf_shifting=False,
+        les_enabled=False,
+    )
+    fine_cfg = LBMConfig(
+        nx=2 * nx,
+        ny=2 * ny,
+        viscosity=2 * nu,
+        velocity=0.0,
+        collision="SRT",
+        streaming="double_buffer",
+        store_precision="FP32",
+        ddf_shifting=False,
+        les_enabled=False,
+    )
+    obj = SimObject(
+        0, CircleGeometry(*center, radius), center=tuple(center), radius=radius
+    )
+    fc = 2 * center - 0.5
+    fobj = SimObject(
+        0, CircleGeometry(*fc, 2 * radius), center=tuple(fc), radius=2 * radius
+    )
+    jets = [
+        Jet(
+            Sector(tuple(center), angle, np.pi / 12, radius + 0.1),
+            Radial(tuple(center), inward=False),
+            "uniform",
+            sign * c["jet_speed"],
+        )
+        for angle, sign in [(np.pi / 2, 1), (3 * np.pi / 2, -1)]
+    ]
+    layout = BlockLayout2D.build(2 * ny, 2 * nx, 8, 8)
+    with (
+        PeriodicBlockStepper(fine_cfg, layout, context=ctx, objects=[fobj]) as uniform,
+        FixedPatchStepper(
+            cfg,
+            c["bounds"],
+            context=ctx,
+            base_block=(8, 8),
+            fine_block=(8, 8),
+            fine_objects=[obj],
+            jets=jets,
+            observe_raw_exchange=True,
+        ) as amr,
+    ):
+        b = amr._level_boundaries[1]
+        count = len(b.tables["cell"])
+        if np.any(b.tables["q"] <= 0.010001) or np.any(b.tables["fallback"]):
+            raise RuntimeError("Pilot requires unclamped q and no fallback")
+        links = fobj.geometry.build_cut_links(2 * nx, 2 * ny)
+        np.testing.assert_array_equal([v.dir for v in links], b.tables["dir"])
+        np.testing.assert_allclose(
+            [v.q for v in links], b.tables["q"], atol=2e-7, rtol=0
+        )
+        np.testing.assert_array_equal([v.fallback for v in links], b.tables["fallback"])
+        np.testing.assert_allclose(
+            np.array([[v.hit_x, v.hit_y] for v in links]) * 0.5 + 0.25,
+            b.host.curved_hits,
+            atol=2e-12,
+            rtol=0,
+        )
+        # Same pre-existing optional-input adapter as test_raw_exchange_gpu.py.
+        velocity = uniform._upload(b.tables["surface_velocity"])
+        rx = uniform._upload(b.tables["rx"])
+        ry = uniform._upload(b.tables["ry"])
+        raw = uniform._allocate(count * 16)
+        original = uniform._curved
+        samples = []
+
+        def boundary(*args, **kw):
+            original(*args[:7], velocity, rx, ry, raw, **kw)
+            uniform._stream.synchronize()
+            a = np.empty((count, 4), np.float32)
+            cuda.memcpy_dtoh(a, raw)
+            samples.append(a)
+
+        uniform._curved = boundary
+
+        def initialize():
+            uniform.initialize(state((9, 2 * ny, 2 * nx), u))
+            amr.initialize(*(state(l.shape, u) for l in amr._levels))
+
+        initialize()
+        # Compile/launch all stepping and diagnostic paths, then reset physical time.
+        amr.step()
+        uniform.step(2)
+        uniform.download()
+        amr.download()
+        amr.read_raw_exchange(0)
+        initialize()
+        samples.clear()
+        jet_id = b.jet_cache[1]
+        masks = {
+            "wall": jet_id < 0,
+            "aperture": jet_id >= 0,
+            "total": np.ones(count, bool),
+            "blowing": jet_id == 0,
+            "suction": jet_id == 1,
+        }
+        if not masks["blowing"].any() or not masks["suction"].any():
+            raise RuntimeError("Both apertures must resolve cut links")
+        fluid_u = ~uniform.boundary["solid"]
+        fluid_a = amr._leaf[1]
+        origin_a = np.array(amr.patch.fine_origin)
+        probes = {
+            "uniform_adapter": aperture_probe(fluid_u, (0.0, 0.0), center, radius),
+            "amr_api": aperture_probe(fluid_a, origin_a, center, radius),
+        }
+
+        recorder = None
+        if continuation:
+            from jet_cv_measurements import MeasurementRecorder, sample_steps
+
+            recorder = MeasurementRecorder(
+                c, output, fluid_u, amr._leaf, origin_a, CX, CY
+            )
+            scheduled = set(sample_steps(c).tolist())
+        else:
+            scheduled = set(range(c["steps"] + 1))
+
+        cumulative_raw = {
+            name: {k: np.zeros(4) for k in masks}
+            for name in ("uniform_adapter", "amr_api")
+        }
+
+        def capture(step):
+            ug = uniform.download()
+            levels = amr.download()
+            ag = levels[1]
+            data = {}
+            aperture = {}
+            for name, g, fluid, origin in [
+                ("uniform_adapter", ug, fluid_u, (0.0, 0.0)),
+                ("amr_api", ag, fluid_a, origin_a),
+            ]:
+                data[name] = [cv(g, fluid, origin, r, nu) for r in c["contours"]]
+                aperture[name] = measured_aperture_flux(g, probes[name])
+                rho, _, _ = moments(g)
+                if not np.isfinite(rho[fluid]).all() or rho[fluid].min() <= 0:
+                    raise RuntimeError("Nonfinite/nonpositive fluid density")
+            ox, oy = np.rint(2 * origin_a).astype(int)
+            ay, ax = fluid_a.shape
+            _, _, uv = moments(ug[:, oy : oy + ay, ox : ox + ax])
+            _, _, av = moments(ag)
+            error = float(
+                np.linalg.norm((av - uv)[:, fluid_a])
+                / max(np.linalg.norm(uv[:, fluid_a]), 1e-30)
+            )
+            if recorder is not None:
+                recorder.capture(step, ug, levels, data, aperture, error)
+                recorder.records[-1]["raw_cumulative"] = np.array(
+                    [
+                        [cumulative_raw[name][k].copy() for k in masks]
+                        for name in ("uniform_adapter", "amr_api")
+                    ]
+                )
+            return data, error, aperture
+
+        previous, _, previous_aperture = capture(0)
+        previous_sample_step = 0
+        initial = previous
+        integrated = {name: np.zeros((2, 3)) for name in previous}
+        integrated_aperture = {name: np.zeros(2) for name in previous}
+        accumulated = {name: {k: np.zeros(4) for k in masks} for name in previous}
+        rows = []
+        start = time.monotonic()
+        print(
+            json.dumps(
+                dict(
+                    status="READY",
+                    case=c,
+                    links=count,
+                    actual_amr_jet_api=True,
+                    uniform_reference="private optional-input validation adapter",
+                    q_range=[float(b.tables["q"].min()), float(b.tables["q"].max())],
+                    fallback_count=int(np.count_nonzero(b.tables["fallback"])),
+                )
+            ),
+            flush=True,
+        )
+        for step in range(1, c["steps"] + 1):
+            if time.monotonic() - start > c.get("post_ready_budget_seconds", 1800):
+                raise RuntimeError(
+                    "Declared post-READY budget exceeded; partial report retained"
+                )
+            samples.clear()
+            amr.step()
+            uniform.step(2)
+            if len(samples) != 2:
+                raise RuntimeError("Uniform interval must contain two fine hops")
+            ur = sum(a.astype(float) * [0.25, 0.25, 0.125, 0.25] for a in samples)
+            ar = np.empty((count, 4), np.float64)
+            with ctx.activate():
+                cuda.memcpy_dtoh(ar, b.interval_integral)
+            if not np.isfinite(ar).all() or not np.isfinite(ur).all():
+                raise RuntimeError("Nonfinite raw exchange in completed base interval")
+            obs = amr.read_raw_exchange(0)
+            if obs["time_start"] != step - 1 or obs["time_end"] != step:
+                raise RuntimeError("Raw interval time mismatch")
+            np.testing.assert_allclose(
+                ar.sum(axis=0)[:2], obs["total"]["impulse"], atol=1e-13, rtol=1e-13
+            )
+            for name, rr in [("uniform_adapter", ur), ("amr_api", ar)]:
+                for k, mask in masks.items():
+                    subtotal = rr[mask].sum(axis=0)
+                    accumulated[name][k] += subtotal
+                    cumulative_raw[name][k] += subtotal
+            if step not in scheduled:
+                continue
+            current, error, current_aperture = capture(step)
+            sample_dt = step - previous_sample_step
+            for name in current:
+                integrated_aperture[name] += (
+                    0.5 * sample_dt * (previous_aperture[name] + current_aperture[name])
+                )
+                integrated[name] += (
+                    0.5
+                    * sample_dt
+                    * (
+                        np.array([x[1] for x in previous[name]])
+                        + np.array([x[1] for x in current[name]])
+                    )
+                )
+            previous = current
+            previous_aperture = current_aperture
+            previous_sample_step = step
+            if step % c["window"] == 0 or step == c["steps"]:
+                end0 = rows[-1]["step_end"] if rows else 0
+                duration = step - end0
+                row = dict(
+                    step_start=end0,
+                    step_end=step,
+                    t_star=step * u / c["diameter"],
+                    fine_velocity_relative_l2=error,
+                    elapsed_seconds=time.monotonic() - start,
+                    measurements={},
+                )
+                for name in current:
+                    delta = np.array([x[0] for x in current[name]]) - np.array(
+                        [x[0] for x in initial[name]]
+                    )
+                    mass_cv = delta[:, 0] + integrated[name][:, 0]
+                    reaction_cv = -delta[:, 1:] - integrated[name][:, 1:]
+                    ledger = accumulated[name]
+                    total = ledger["total"]
+                    normal = 0.5 * u * u * c["diameter"]
+                    row["measurements"][name] = dict(
+                        raw_interval={k: v.tolist() for k, v in ledger.items()},
+                        cv_mass_into_fluid=mass_cv.tolist(),
+                        cv_reaction_impulse=reaction_cv.tolist(),
+                        mass_balance_residual=(mass_cv - total[3]).tolist(),
+                        reaction_balance_residual=(reaction_cv - total[:2]).tolist(),
+                        reaction_residual_coefficient=(
+                            (reaction_cv - total[:2]) / duration / normal
+                        ).tolist(),
+                        raw_reaction_coefficient=(
+                            total[:2] / duration / normal
+                        ).tolist(),
+                        cv_reaction_coefficient=(
+                            reaction_cv / duration / normal
+                        ).tolist(),
+                        contour_reaction_spread_coefficient=(
+                            (reaction_cv[1] - reaction_cv[0]) / duration / normal
+                        ).tolist(),
+                        measured_arc_mass_rates=(
+                            integrated_aperture[name] / duration
+                        ).tolist(),
+                        arc_max_sample_distance=[v[-1] for v in probes[name]],
+                        nominal_volume_fluxes=[
+                            2 * np.pi / 12 * radius * c["jet_speed"],
+                            -2 * np.pi / 12 * radius * c["jet_speed"],
+                        ],
+                        aperture_mass_rates={
+                            k: float(ledger[k][3] / duration)
+                            for k in ("blowing", "suction", "wall", "total")
+                        },
+                    )
+                    integrated[name].fill(0)
+                    integrated_aperture[name].fill(0)
+                    for v in accumulated[name].values():
+                        v.fill(0)
+                rows.append(row)
+                initial = current
+                if recorder is not None:
+                    recorder.flush()
+                payload = dict(
+                    case=c,
+                    source_revision=PIN,
+                    **provenance,
+                    status="PILOT_DIAGNOSTIC_ONLY",
+                    windows=rows,
+                    physics_validation_pass=False,
+                    precision="FP32 evolution, FP64 host diagnostics",
+                )
+                (output / f"re{c['re_d']}_s{c['scale']}.json").write_text(
+                    json.dumps(payload, indent=2) + "\n"
+                )
+                print(
+                    json.dumps(
+                        dict(
+                            status="PROGRESS",
+                            step=step,
+                            steps=c["steps"],
+                            elapsed_seconds=row["elapsed_seconds"],
+                            fine_velocity_relative_l2=error,
+                        )
+                    ),
+                    flush=True,
+                )
+        return payload
+
+
+def main():
+    p = argparse.ArgumentParser(description=__doc__)
+    p.add_argument("--repo", type=Path, required=True)
+    p.add_argument("--output", type=Path, required=True)
+    p.add_argument("--scale", type=int, choices=[1, 2, 4], default=1)
+    p.add_argument("--re", type=int, choices=[100, 400], default=100)
+    p.add_argument(
+        "--steps", type=int, default=800, help="Base s=1 steps; scales as s squared"
+    )
+    p.add_argument("--allow-re400", action="store_true")
+    p.add_argument(
+        "--continuation",
+        action="store_true",
+        help="Fixed measurement campaign: Re100, scale1/2, 3200 reference steps",
+    )
+    p.add_argument("--plan", action="store_true")
+    a = p.parse_args()
+    if a.steps <= 0 or a.steps > 3200:
+        p.error("steps must be 1..3200")
+    if a.re == 400 and not a.allow_re400:
+        p.error("Re400 requires explicit conditional-campaign opt-in")
+    if a.continuation:
+        if a.scale not in (1, 2) or a.re != 100 or a.steps != 3200:
+            p.error("--continuation requires --re 100 --scale 1/2 --steps 3200")
+    c = case(a.scale, a.re, a.steps)
+    if a.continuation:
+        from jet_cv_measurements import sample_steps, factors
+
+        c.update(
+            measurement_schema=1,
+            sample_count=len(sample_steps(c)),
+            startup_sample_dt_reference=1,
+            later_sample_dt_reference=4,
+            dense_startup_end_reference=100,
+            reference_conversion=factors(a.scale),
+            post_ready_budget_seconds=1200 if a.scale == 1 else 5400,
+            runtime_estimate_seconds=[210, 500] if a.scale == 1 else [900, 3200],
+        )
+    plan = dict(
+        case=c,
+        required_revision=PIN,
+        required_free_device_mib=1024,
+        estimated_peak_device_mib=128 * a.scale**2,
+        estimated_output_mib=64 if a.continuation else 2,
+        max_steps=c["steps"],
+        gpu_launch_owner="root; this runner does not select or reserve a GPU",
+    )
+    if a.plan:
+        print(json.dumps(plan, indent=2))
+        return
+    if (
+        not os.environ.get("CUDA_VISIBLE_DEVICES")
+        or "," in os.environ["CUDA_VISIBLE_DEVICES"]
+    ):
+        p.error("Root must select one verified idle GPU via CUDA_VISIBLE_DEVICES")
+    repo = a.repo.resolve()
+    source = verify(repo)
+    provenance = dict(
+        source_experiment=source["experiment"], **checkout_provenance(repo)
+    )
+    sys.path.insert(0, str(repo / "src"))
+    import CelerisLab
+
+    if Path(CelerisLab.__file__).resolve() != repo / "src/CelerisLab/__init__.py":
+        raise RuntimeError("Wrong source import")
+    from CelerisLab.cuda.context import CudaContext
+    import pycuda.driver as cuda
+
+    a.output.mkdir(parents=True, exist_ok=False)
+    (a.output / "plan.json").write_text(json.dumps(plan, indent=2) + "\n")
+    with CudaContext(0) as ctx:
+        with ctx.activate():
+            free, total = cuda.mem_get_info()
+        if free < max(1024, 2 * plan["estimated_peak_device_mib"]) * 1024**2:
+            raise RuntimeError("Insufficient free device memory")
+        run_case(c, ctx, a.output, provenance, continuation=a.continuation)
+    print(
+        json.dumps(
+            dict(status="COMPLETE_DIAGNOSTIC_ONLY", physics_validation_pass=False)
+        ),
+        flush=True,
+    )
+
+
+if __name__ == "__main__":
+    main()

tests/validation/run_refined_dfg.py

@@ -0,0 +1,631 @@
+"""Frozen matched SRT/MRT DFG comparison. GPU execution requires explicit root admission."""
+
+import argparse
+import csv
+import hashlib
+import json
+import math
+import os
+from pathlib import Path
+import sys
+import time
+from tempfile import TemporaryDirectory
+
+import numpy as np
+
+HERE = Path(__file__).resolve().parent
+DEFAULT_SOURCE = HERE.parents[1]
+BOUNDS = (40, 24, 320, 140)
+BASE_BLOCK = (2, 881)  # Exact 166 x 881 domain; no geometric padding.
+FINE_BLOCK = (8, 8)  # 240 x 568 including four fine halo cells.
+SPEED = 0.01
+
+
+def helpers(source):
+    sys.path.insert(0, str(source / "src"))
+    sys.path.insert(0, str(source / "tests/validation"))
+    import CelerisLab
+
+    if (
+        Path(CelerisLab.__file__).resolve()
+        != source.resolve() / "src/CelerisLab/__init__.py"
+    ):
+        raise RuntimeError("Wrong CelerisLab import")
+    import re100_helpers as h
+    import periodic
+
+    if h.ROOT.resolve() != source.resolve():
+        raise RuntimeError("Helper module loaded from wrong source")
+    return h, periodic
+
+
+def config(h, diameter, collision="SRT"):
+    cfg, body, p = h.make_config(diameter, SPEED, collision, "neq_extrap", "zou_he_local")
+    cfg["method"]["outlet"]["srt_neq_damp"] = 0.5
+    return cfg, body, p
+
+
+def plan(h, collision="SRT"):
+    _, _, c = config(h, 40, collision)
+    _, _, f = config(h, 80, collision)
+    x0, y0, x1, y1 = BOUNDS
+    fine_active = 4 * (x1 - x0) * (y1 - y0)
+    coarse_uncovered = c["nx"] * c["ny"] - fine_active // 4
+    return dict(
+        status="prepared_not_run",
+        collision=collision,
+        root_gpu_gates_required=True,
+        root=c,
+        uniform=f,
+        bounds_root_edge=list(BOUNDS),
+        physical_patch_m=[
+            (x0 - 0.5) * c["dx"],
+            (y0 - 1) * c["dx"],
+            (x1 - 0.5) * c["dx"],
+            (y1 - 1) * c["dx"],
+        ],
+        body_root_edge=[80.5, 81.0],
+        radius_root=20.0,
+        root_block=list(BASE_BLOCK),
+        fine_block=list(FINE_BLOCK),
+        root_wall_edges=[1.0, 165.0],
+        root_inlet_outlet_edges=[0.5, 880.5],
+        physical_domain_m=[2.2, 0.41],
+        physical_center_m=[0.2, 0.2],
+        physical_diameter_m=0.1,
+        physical_mean_speed_m_s=1.0,
+        physical_viscosity_m2_s=0.001,
+        finest_dx_m=f["dx"],
+        finest_dt_s=f["dt"],
+        approximate_leaf_cell_ratio=(coarse_uncovered + fine_active)
+        / (f["nx"] * f["ny"]),
+        approximate_collision_work_ratio=(coarse_uncovered + 2 * fine_active)
+        / (2 * f["nx"] * f["ny"]),
+        cost_caveat="Geometry-only estimates include ghosts/solids; exclude block halos, transfers, observation and launch overhead. Not measured speedup.",
+        sampling_offset="For origin (38,22), fine physical x=(j+75.5)*dx80 versus uniform i*dx80: half-fine-cell x offset. Fine physical y=(k+42.5)*dx80 equals uniform (l-.5)*dx80 at l=k+43: y samples align. Body, wall and source planes coincide.",
+        outlet=(
+            "SRT pressure NEQ, all nine fneq damped by .5, no backflow clamp; noncanonical DFG outlet"
+            if collision == "SRT" else
+            "MRT pressure NEQ, unknowns q2/6/8 extrapolated; no backflow clamp; noncanonical DFG outlet"
+        ),
+        modes={"pilot_tstar": 5.0, "full_tstar": 300.0},
+        observation_dense_tstar=0.02,
+    )
+
+
+def macros(g):
+    rho = np.sum(g, axis=0, dtype=np.float64)
+    ux = (g[1] - g[2] + g[5] - g[6] + g[7] - g[8]) / rho
+    uy = (g[3] - g[4] + g[5] - g[6] - g[7] + g[8]) / rho
+    return rho, ux, uy
+
+
+def leaf_mass(arrays, dx, masks):
+    """Finite public leaves only; covered coarse, solid and halos are excluded."""
+    if any(not np.isfinite(g[:, mask]).all() for g, mask in zip(arrays, masks)):
+        raise ValueError("Nonfinite physical leaf population")
+    return sum(
+        float(np.sum(np.sum(g, axis=0, dtype=np.float64)[mask], dtype=np.float64))
+        * (dx / 2**level) ** 2
+        for level, (g, mask) in enumerate(zip(arrays, masks))
+    )
+
+
+def initial_state(shape, origin, spacing, height=164.0, speed=SPEED, root=False):
+    """Same physical parabolic initial condition as uniform InitTubeFlow_v2."""
+    y = origin[1] + spacing * (np.arange(shape[0]) + 0.5)
+    u = 6 * speed * ((y - 1) / height) * (1 - (y - 1) / height)
+    velocity = np.broadcast_to(u[:, None], shape).copy()
+    if root:
+        velocity[[0, -1], :] = 0
+        velocity[:, [0, -1]] = 0
+    cx = np.array([0, 1, -1, 0, 0, 1, -1, 1, -1])[:, None, None]
+    w = np.array([4 / 9] + [1 / 9] * 4 + [1 / 36] * 4)[:, None, None]
+    cu = cx * velocity
+    return w * (1 + 3 * cu + 4.5 * cu * cu - 1.5 * velocity * velocity)
+
+
+def host_geometry(source, h, collision="SRT"):
+    """Production donor/transfer and complete cut-link checks, without CUDA."""
+    from CelerisLab.config import load_lbm_config
+    from CelerisLab.lbm.refinement import FixedPatch2D, prepare_fine_boundary
+    from CelerisLab.lbm.blocks import BlockLayout2D
+    from CelerisLab.lbm.channel import (
+        validate_channel_config,
+        validate_channel_transfers,
+    )
+    from CelerisLab.lbm.refinement_tables import blocked_base_tables
+    from CelerisLab.body.geometry.circle import CircleGeometry
+    from CelerisLab.body.objects import SimObject
+
+    cfgdict, _, _ = config(h, 40, collision)
+    with TemporaryDirectory() as tmp:
+        p = Path(tmp) / "cfg.json"
+        p.write_text(json.dumps(cfgdict))
+        cfg = load_lbm_config(str(p))
+    pair = FixedPatch2D.from_config(cfg, BOUNDS)
+    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]
+    validate_channel_transfers(pair, layout, targets)
+    BlockLayout2D.build(*pair.fine_shape, *FINE_BLOCK, periodic=(False, False))
+    body = SimObject(
+        0, geometry=CircleGeometry(80.5, 81.0, 20.0), center=(80.5, 81.0), radius=20.0
+    )
+    host = prepare_fine_boundary(cfg, pair, (body,), _shifted_blocked=True)
+    return cfg, pair, body, host
+
+
+def verify_source(source, collision="SRT"):
+    if collision not in ("SRT", "MRT"):
+        raise ValueError("Matched collision must be SRT or MRT")
+    name = "source_manifest.json" if collision == "SRT" else "source_manifest_mrt.json"
+    expected = json.loads((HERE / "data/refined_dfg" / name).read_text())
+    for name, digest in expected["sha256"].items():
+        if hashlib.sha256((source / name).read_bytes()).hexdigest() != digest:
+            raise RuntimeError("Source mismatch: " + name)
+    return expected
+
+
+def run_case(a, h, periodic, case):
+    from CelerisLab.simulation import Simulation
+    from CelerisLab.cuda.context import CudaContext
+    from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+    from CelerisLab.lbm.descriptors import FLUID
+    from contextlib import ExitStack
+
+    refined = case == "refined"
+    cfgdict, bodydict, p = config(h, 40 if refined else 80, a.collision)
+    cap = 5.0 if a.mode == "pilot" else 300.0
+    max_steps = math.ceil(cap * p["diameter"] / SPEED)
+    rows = []
+    healthy = True
+    step_seconds = 0.0
+    started = time.monotonic()
+    check_deadline(a.deadline)
+    dest = a.output / case
+    dest.mkdir()
+    with ExitStack() as stack:
+        if refined:
+            cfg, pair, body, host = host_geometry(a.source, h, a.collision)
+            context = stack.enter_context(CudaContext(0))
+            backend = stack.enter_context(
+                FixedPatchStepper(
+                    cfg,
+                    BOUNDS,
+                    context=context,
+                    coarse_boundary="channel",
+                    base_block=BASE_BLOCK,
+                    fine_block=FINE_BLOCK,
+                    fine_objects=(body,),
+                    observe_raw_exchange=True,
+                )
+            )
+            from CelerisLab.lbm.channel import channel_masks
+            from CelerisLab.lbm.refinement import ACTIVE
+            from CelerisLab.lbm.descriptors import SOLID
+
+            masks = [
+                channel_masks(pair.coarse_shape)[0] & ~pair.covered,
+                (pair.fine_kind == ACTIVE)
+                & ((host.flags.reshape(pair.fine_shape) & SOLID) == 0),
+            ]
+            backend.initialize(
+                initial_state(pair.coarse_shape, (0.0, 0.0), 1.0, root=True),
+                initial_state(pair.fine_shape, pair.fine_origin, 0.5),
+            )
+            probe_p = dict(
+                p,
+                center=[
+                    (80.5 - pair.fine_origin[0]) / 0.5 - 0.5,
+                    (81.0 - pair.fine_origin[1]) / 0.5 - 0.5,
+                ],
+                radius=40.0,
+            )
+        else:
+            tmp = Path(stack.enter_context(TemporaryDirectory()))
+            (tmp / "lbm.json").write_text(json.dumps(cfgdict))
+            (tmp / "body.json").write_text(json.dumps(bodydict))
+            sim = stack.enter_context(
+                Simulation(str(tmp / "lbm.json"), str(tmp / "body.json"), device_id=0)
+            )
+            sim.initialize()
+            probe_p = p
+            fluid = (sim.field.flag.reshape(p["ny"], p["nx"]) & FLUID) != 0
+        startup = time.monotonic() - started
+        check_deadline(a.deadline)
+        begin = 0
+        with (dest / "history.csv").open("w") as out:
+            writer = None
+            while begin < max_steps:
+                check_deadline(a.deadline)
+                width = (
+                    0.02
+                    if a.mode == "pilot" or begin * SPEED / p["diameter"] >= 240
+                    else 0.5
+                )
+                n = min(max(1, round(width * p["diameter"] / SPEED)), max_steps - begin)
+                impulse = np.zeros(2)
+                tick = time.monotonic()
+                if refined:
+                    for _ in range(n):
+                        check_deadline(a.deadline)
+                        backend.step(1)
+                        exchange = backend.read_raw_exchange(
+                            0, level=1, window="last_base_step"
+                        )
+                        if (
+                            exchange["units"] != "base_lattice_per_unit_depth"
+                            or exchange["sample_count"] != 2
+                            or exchange["time_end"] - exchange["time_start"] != 1
+                        ):
+                            raise RuntimeError("Unexpected force observation contract")
+                        impulse += exchange["total"]["impulse"]
+                        check_deadline(a.deadline)
+                else:
+                    done = 0
+                    while done < n:
+                        check_deadline(a.deadline)
+                        batch = min(100, n - done)
+                        sim.run(batch)
+                        impulse += sim.read_force(0).astype(np.float64) * batch
+                        done += batch
+                        check_deadline(a.deadline)
+                step_seconds += time.monotonic() - tick
+                if refined:
+                    arrays = backend.download()
+                    fields = [macros(g) for g in arrays]
+                    root_rho, root_ux, _ = fields[0]
+                    rho, ux, uy = fields[1]
+                    fluid = masks[1]
+                    if any(
+                        not np.isfinite(g[:, mask]).all()
+                        for g, mask in zip(arrays, masks)
+                    ):
+                        healthy = False
+                        break
+                    mass = leaf_mass(arrays, p["dx"], masks)
+                    vals = [
+                        tuple(v[mask] for v in field)
+                        for mask, field in zip(masks, fields)
+                    ]
+                else:
+                    m = sim.get_macroscopic()
+                    rho, ux, uy = [
+                        np.asarray(m[k], np.float64).reshape(p["ny"], p["nx"])
+                        for k in ("rho", "ux", "uy")
+                    ]
+                    root_rho, root_ux = rho, ux
+                    mass = float(rho[fluid].sum()) * p["dx"] ** 2
+                    vals = [(rho[fluid], ux[fluid], uy[fluid])]
+                healthy = (
+                    all(np.isfinite(v).all() for vs in vals for v in vs)
+                    and np.isfinite(impulse).all()
+                )
+                if not healthy:
+                    break
+                den = SPEED**2 * p["diameter"] / 2
+                row = dict(
+                    step=begin + n,
+                    tstar=(begin + n) * SPEED / p["diameter"],
+                    tstar_force=(begin + n / 2) * SPEED / p["diameter"],
+                    cd=float(impulse[0] / n / den),
+                    cl=float(impulse[1] / n / den),
+                    mass=mass,
+                    rho_min=float(min(vs[0].min() for vs in vals)),
+                    rho_max=float(max(vs[0].max() for vs in vals)),
+                    max_mach=float(
+                        max(np.sqrt(3 * (vs[1] ** 2 + vs[2] ** 2)).max() for vs in vals)
+                    ),
+                    flux_in=float(np.sum(root_rho[1:-1, 2] * root_ux[1:-1, 2]))
+                    * p["dx"]
+                    / SPEED,
+                    flux_out=float(np.sum(root_rho[1:-1, -3] * root_ux[1:-1, -3]))
+                    * p["dx"]
+                    / SPEED,
+                    **h.pressure_probes(rho, fluid, probe_p),
+                )
+                check_deadline(a.deadline)
+                ensure_finite(row)
+                rows.append(row)
+                if writer is None:
+                    writer = csv.DictWriter(out, fieldnames=row)
+                    writer.writeheader()
+                writer.writerow(row)
+                out.flush()
+                healthy = (
+                    row["max_mach"] <= 0.2
+                    and row["rho_min"] >= 0.5
+                    and row["rho_max"] <= 1.5
+                )
+                begin += n
+                if not healthy:
+                    break
+                if len(rows) == 1:
+                    ready = dict(
+                        case=case,
+                        collision=a.collision,
+                        **a.provenance,
+                        mode=a.mode,
+                        step=begin,
+                        tstar=row["tstar"],
+                        runner_sha256=runner_hash(),
+                        all_runtime_and_observation_paths_exercised=True,
+                    )
+                    (dest / "ready.json").write_text(
+                        json.dumps(ready, indent=2, allow_nan=False)
+                    )
+                    print("READY " + json.dumps(ready, allow_nan=False), flush=True)
+    result = dict(
+        case=case,
+        collision=a.collision,
+        **a.provenance,
+        mode=a.mode,
+        history=rows,
+        healthy=bool(healthy),
+        steps=begin,
+        startup_seconds=startup,
+        step_seconds=step_seconds,
+        wall_seconds=time.monotonic() - started,
+        force="raw stationary-body link momentum reaction, per-unit-depth; average over entire observation interval; normalized in each base lattice",
+        mass_units="rho_reference * m^2 per unit depth",
+        flux_units="rho_reference * m^2/s per unit depth",
+        pressure="postcollision end-sample; physical pressure difference / (rho_reference * 1 m/s squared)",
+        suite_validated=False,
+        **periodic.assess(rows),
+    )
+    if not healthy:
+        result["periodic"] = False
+    reference = json.loads(
+        (a.source / "tests/validation/data/dfg_re100/reference_cycle.json").read_text()
+    )
+    if result["cycles"]:
+        result["reference_comparison"] = periodic.compare(
+            result["cycles"][-1], reference
+        )
+    result["status"] = (
+        "invalid_state"
+        if not healthy
+        else (
+            "pilot_complete"
+            if a.mode == "pilot"
+            else ("periodic" if result["periodic"] else "incomplete_periodicity")
+        )
+    )
+    (dest / "summary.json").write_text(json.dumps(result, indent=2, allow_nan=False))
+    return result
+
+
+def runner_hash():
+    return hashlib.sha256(Path(__file__).read_bytes()).hexdigest()
+
+
+def ensure_finite(value):
+    if isinstance(value, dict):
+        for item in value.values():
+            ensure_finite(item)
+    elif isinstance(value, list):
+        for item in value:
+            ensure_finite(item)
+    elif isinstance(value, (float, int)) and not isinstance(value, bool):
+        if not math.isfinite(value):
+            raise ValueError("Nonfinite JSON value")
+
+
+def read_finite_json(path):
+    def reject(value):
+        raise ValueError("Nonfinite JSON constant: " + value)
+
+    value = json.loads(path.read_text(), parse_constant=reject)
+    ensure_finite(value)
+    return value
+
+
+def check_deadline(deadline):
+    if time.monotonic() >= deadline:
+        raise TimeoutError(
+            "Campaign wall-time ceiling reached; no continuation or acceptance"
+        )
+
+
+def verify_pilot(path, manifest, collision="SRT"):
+    if path is None:
+        raise ValueError(
+            "Full comparison requires --pilot-results from a completed matched pilot"
+        )
+    previous = read_finite_json(path / "plan.json")
+    if previous.get("collision") != collision:
+        raise ValueError("Pilot collision differs")
+    if previous["source_manifest"] != manifest or previous["bounds_root_edge"] != list(
+        BOUNDS
+    ):
+        raise ValueError("Pilot source or frozen geometry differs")
+    if previous.get("runner_sha256") != runner_hash():
+        raise ValueError("Pilot runner differs")
+    calibrated_seconds = 0.0
+    for case, diameter in (("uniform", 80), ("refined", 40)):
+        summary = read_finite_json(path / case / "summary.json")
+        if (
+            summary.get("collision") != collision
+            or summary.get("case") != case
+            or summary.get("mode") != "pilot"
+            or summary.get("status") != "pilot_complete"
+            or summary.get("healthy") is not True
+        ):
+            raise ValueError("Mislabeled or unhealthy pilot receipt")
+        rows = summary.get("history")
+        interval = round(0.02 * diameter / SPEED)
+        total = round(5 * diameter / SPEED)
+        if (
+            not isinstance(rows, list)
+            or len(rows) != 250
+            or summary.get("steps") != total
+        ):
+            raise ValueError("Incomplete pilot history or step count")
+        for i, row in enumerate(rows):
+            step = (i + 1) * interval
+            if type(row.get("step")) is not int or row["step"] != step:
+                raise ValueError("Nonmonotone or incomplete pilot step history")
+            for key, expected in (
+                ("tstar", step * SPEED / diameter),
+                ("tstar_force", (step - interval / 2) * SPEED / diameter),
+            ):
+                value = row.get(key)
+                if type(value) not in (float, int) or not math.isclose(
+                    value, expected, rel_tol=0, abs_tol=1e-12
+                ):
+                    raise ValueError("Incorrect pilot observation time")
+            for key in (
+                "cd",
+                "cl",
+                "dp",
+                "mass",
+                "rho_min",
+                "rho_max",
+                "max_mach",
+                "flux_in",
+                "flux_out",
+                "rho_front",
+                "rho_back",
+            ):
+                if type(row.get(key)) not in (int, float):
+                    raise ValueError("Missing numeric pilot observation")
+            if not (
+                0 < row["mass"]
+                and 0.5 <= row["rho_min"] <= row["rho_max"] <= 1.5
+                and 0 <= row["max_mach"] <= 0.2
+            ):
+                raise ValueError("Unhealthy pilot observations")
+        ready = read_finite_json(path / case / "ready.json")
+        if (
+            ready.get("collision") != collision
+            or ready.get("case") != case
+            or ready.get("mode") != "pilot"
+            or ready.get("step") != interval
+            or ready.get("runner_sha256") != runner_hash()
+            or ready.get("all_runtime_and_observation_paths_exercised") is not True
+        ):
+            raise ValueError("Invalid pilot READY receipt")
+        startup = summary.get("startup_seconds")
+        wall = summary.get("wall_seconds")
+        stepping = summary.get("step_seconds")
+        if (
+            any(type(v) not in (float, int) for v in (startup, wall, stepping))
+            or not 0 <= startup < wall
+            or not 0 < stepping <= wall
+        ):
+            raise ValueError("Invalid pilot calibration timings")
+        # Full physical duration is 60 times pilot; dense pilot sampling makes
+        # this conservative for the sparser early full observations. Not a guarantee.
+        calibrated_seconds += startup + 60 * (wall - startup)
+    return 3 * calibrated_seconds
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--collision", choices=["SRT", "MRT"], default="SRT")
+    parser.add_argument("--source", type=Path, default=DEFAULT_SOURCE)
+    parser.add_argument("--output", type=Path, required=True)
+    parser.add_argument("--mode", choices=["plan", "pilot", "full"], default="plan")
+    parser.add_argument("--gpu-gates-approved", action="store_true")
+    parser.add_argument("--pilot-results", type=Path)
+    parser.add_argument(
+        "--max-wall-seconds",
+        type=float,
+        help="Required explicit campaign wall-time ceiling for GPU execution",
+    )
+    a = parser.parse_args()
+    a.source = a.source.resolve()
+    manifest = verify_source(a.source, a.collision)
+    h, periodic = helpers(a.source)
+    from source_provenance import checkout_provenance
+
+    a.provenance = checkout_provenance(a.source)
+    calibration = (
+        verify_pilot(a.pilot_results, manifest, a.collision)
+        if a.mode == "full" else None
+    )
+    if a.mode != "plan" and (
+        a.max_wall_seconds is None
+        or not math.isfinite(a.max_wall_seconds)
+        or a.max_wall_seconds <= 0
+    ):
+        raise ValueError("GPU execution requires finite positive --max-wall-seconds")
+    budget = (
+        min(a.max_wall_seconds, calibration)
+        if calibration is not None
+        else a.max_wall_seconds
+    )
+    if a.output.exists() and any(a.output.iterdir()):
+        raise ValueError("Output must be empty/new")
+    a.output.mkdir(parents=True, exist_ok=True)
+    specification = plan(h, a.collision)
+    specification.update(
+        source_manifest=manifest,
+        **a.provenance,
+        gates=periodic.GATES,
+        reference_gates=periodic.REFERENCE_GATES,
+        runner_sha256=runner_hash(),
+        wall_time_budget_seconds=budget,
+        full_pilot_calibration_seconds=calibration,
+    )
+    host_geometry(a.source, h, a.collision)
+    (a.output / "plan.json").write_text(json.dumps(specification, indent=2))
+    if a.mode == "plan":
+        return 0
+    if not a.gpu_gates_approved or not os.environ.get("CUDA_VISIBLE_DEVICES"):
+        raise RuntimeError(
+            "GPU execution blocked: root GPU gates approval and explicit device selection required"
+        )
+    results = []
+    a.deadline = time.monotonic() + budget
+    for case in ("uniform", "refined"):
+        try:
+            result = run_case(a, h, periodic, case)
+        except TimeoutError as exc:
+            failure = dict(
+                status="wall_time_limit",
+                collision=a.collision,
+                **a.provenance,
+                case=case,
+                matched_case_pass=False,
+                error=str(exc),
+            )
+            (a.output / "comparison.json").write_text(
+                json.dumps(failure, indent=2, allow_nan=False)
+            )
+            print(json.dumps(failure, allow_nan=False), flush=True)
+            return 3
+        results.append(result)
+        if not result["healthy"]:
+            break
+    comparison = dict(
+        collision=a.collision,
+        **a.provenance,
+        status="pilot_only" if a.mode == "pilot" else "incomplete",
+        matched_case_pass=False,
+    )
+    if len(results) == 2 and all(r["healthy"] and r["periodic"] for r in results):
+        comparison.update(
+            status="periodic_comparison",
+            refined_vs_uniform=periodic.compare(
+                results[1]["cycles"][-1], results[0]["cycles"][-1]
+            ),
+        )
+        comparison["matched_case_pass"] = comparison["refined_vs_uniform"][
+            "reference_agreement"
+        ]
+    (a.output / "comparison.json").write_text(
+        json.dumps(comparison, indent=2, allow_nan=False)
+    )
+    return (
+        0
+        if len(results) == 2
+        and all(r["healthy"] for r in results)
+        and (a.mode == "pilot" or comparison["matched_case_pass"])
+        else 2
+    )
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())

tests/validation/run_uniform_les_flow.py

@@ -0,0 +1,321 @@
+"""Bounded SRT/DB uniform-grid wall-bounded shear validation, Horsea only.
+
+Uses production Simulation/collision/streaming. Test-only raw population/flag
+adapter follows tests/integration/test_numerical_gpu.py; no periodic-y API.
+"""
+
+import argparse
+import hashlib
+import json
+import os
+from pathlib import Path
+import subprocess
+import sys
+from tempfile import TemporaryDirectory
+import time
+import numpy as np
+from les_flow_oracle import H, U, NU, DELTA, TIMES, solve, amplitude, relative
+
+HERE = Path(__file__).resolve().parent
+HEADER = "src/CelerisLab/lbm/kernels/operators/turbulence_smag.cuh"
+LIMITS = dict(
+    linear_relative_l2=0.003,
+    les_relative_l2=0.005,
+    sgs_increment_relative_error=0.15,
+    oracle_grid_relative=0.00008,
+    mass_drift=0.00005,
+    max_transverse=0.00001,
+    x_invariance=0.000001,
+    initial_relative_l2=0.00001,
+)
+
+
+def decode_macros(data, w, cx, cy):
+    """Reject invalid physical populations/density/macros before reporting norms."""
+    with np.errstate(over="ignore", invalid="ignore", divide="ignore"):
+        f = data.astype(np.float64) + w[:, None, None].astype(np.float64)
+        if not np.isfinite(f).all():
+            raise RuntimeError("Nonfinite populations")
+        f = f[:, 1:-1, :]
+        rho = f.sum(axis=0)
+        if not np.isfinite(rho).all() or np.any(rho <= 0):
+            raise RuntimeError("Nonpositive or nonfinite fluid density")
+        vx = np.einsum("q,qyx->yx", cx, f) / rho
+        vy = np.einsum("q,qyx->yx", cy, f) / rho
+        profile = vx.mean(axis=1)
+        if not (
+            np.isfinite(vx).all()
+            and np.isfinite(vy).all()
+            and np.isfinite(profile).all()
+        ):
+            raise RuntimeError("Nonfinite decoded velocity/profile")
+        metrics = (
+            float(rho.mean()),
+            float(np.max(np.abs(vy))),
+            float(np.max(np.ptp(vx, axis=1))),
+        )
+        if not np.isfinite(metrics).all():
+            raise RuntimeError("Nonfinite decoded diagnostics")
+    return profile, *metrics
+
+
+def source_guard(root, variant):
+    manifest = json.loads((HERE / "data/les_flow/source_pins.json").read_text())
+    for name, sha in manifest["sources"].items():
+        expected = (
+            manifest["candidate_header_sha256"]
+            if name == HEADER and variant == "corrected"
+            else sha
+        )
+        if hashlib.sha256((root / name).read_bytes()).hexdigest() != expected:
+            raise RuntimeError("Source pin mismatch: " + name)
+    return dict(
+        base_commit=manifest["base_commit"],
+        production_origin_commit=manifest["production_origin_commit"],
+        source_experiment=manifest["source_experiment"],
+        checkout_commit=subprocess.check_output(
+            ["git", "-C", str(root), "rev-parse", "HEAD"], text=True
+        ).strip(),
+        source_file_count=len(manifest["sources"]),
+        header_sha256=hashlib.sha256((root / HEADER).read_bytes()).hexdigest(),
+        variant=variant,
+    )
+
+
+def main():
+    ap = argparse.ArgumentParser(description=__doc__)
+    ap.add_argument("--root", type=Path, required=True)
+    ap.add_argument("--variant", choices=["baseline", "corrected"], required=True)
+    ap.add_argument("--output", type=Path, required=True)
+    ap.add_argument(
+        "--device", type=int, default=0, help="Logical device in CUDA_VISIBLE_DEVICES"
+    )
+    ap.add_argument("--run-gpu", action="store_true")
+    args = ap.parse_args()
+    if not args.run_gpu:
+        ap.error("Explicit --run-gpu is required; authorized Horsea executor only")
+    if not os.environ.get("CUDA_VISIBLE_DEVICES"):
+        ap.error("Explicit CUDA_VISIBLE_DEVICES required")
+    if args.output.exists():
+        ap.error("Output exists; choose a new evidence path")
+    args.output.parent.mkdir(parents=True, exist_ok=True)
+    root = args.root.resolve()
+    pins = source_guard(root, args.variant)
+    # Refuse installed-package shadowing; import only after explicit GPU opt-in.
+    sys.path.insert(0, str(root / "src"))
+    import CelerisLab
+
+    if Path(CelerisLab.__file__).resolve() != root / "src/CelerisLab/__init__.py":
+        raise RuntimeError("Wrong CelerisLab import")
+    import pycuda.driver as cuda
+    from CelerisLab.simulation import Simulation
+    from CelerisLab.lbm.descriptors import D2Q9_EX, D2Q9_EY, FLUID, SOLID, BC_WALL
+
+    cx = np.asarray(D2Q9_EX)
+    cy = np.asarray(D2Q9_EY)
+    w = np.asarray([4 / 9] + [1 / 9] * 4 + [1 / 36] * 4, dtype=np.float32)
+    h = int(H)
+    nx = 8
+    ny = h + 2
+    y = np.arange(h) + 0.5
+    report = dict(
+        source=pins,
+        package=CelerisLab.__file__,
+        cuda_visible_devices=os.environ["CUDA_VISIBLE_DEVICES"],
+        logical_device=args.device,
+        geometry=dict(nx=nx, ny=ny, halfway_wall_height=H),
+        parameters=dict(
+            nu=NU,
+            U=U,
+            delta_lattice=DELTA,
+            tau0=0.5 + 3 * NU,
+            re_height=U * H / NU,
+            ma=U * np.sqrt(3),
+        ),
+        limits=LIMITS,
+        criterion="standard |S|=sqrt(2 S:S); no baseline-failure masking",
+        cases=[],
+    )
+
+    def save():
+        args.output.write_text(json.dumps(report, indent=2, allow_nan=False) + "\n")
+
+    started = time.monotonic()
+    for cs in (0.0, 0.16, 1.0):
+        cfg = json.loads((root / "src/CelerisLab/configs/config_lbm.json").read_text())
+        cfg["grid"].update(nx=nx, ny=ny, nz=1, lattice_model="D2Q9")
+        cfg["physics"].update(viscosity=NU, velocity=0.0, rho=1.0)
+        cfg["method"].update(
+            collision="SRT",
+            streaming="double_buffer",
+            store_precision="FP32",
+            ddf_shifting=True,
+            y_wall_bc="bounce_back",
+        )
+        cfg["method"]["les"].update(enabled=True, cs=cs, closed_form=True)
+        with TemporaryDirectory(prefix="uniform-les-flow-") as temp:
+            lbm = Path(temp) / "lbm.json"
+            bodies = Path(temp) / "bodies.json"
+            lbm.write_text(json.dumps(cfg))
+            bodies.write_text('{"objects":[]}')
+            sim = Simulation(str(lbm), str(bodies), device_id=args.device)
+        try:
+            sim.initialize()
+            # Exactly the existing integration-test spatial postcollision contract.
+            # y-wall ghosts stay zero velocity; x edge fluid gives modulo x streaming.
+            ux = np.zeros((ny, nx))
+            ux[1:-1] = U * np.sin(np.pi * y / H)[:, None]
+            cu = cx[:, None, None] * ux[None, :, :]
+            physical = w[:, None, None] * (
+                1 + 3 * cu + 4.5 * cu * cu - 1.5 * ux[None, :, :] ** 2
+            )
+            shifted = np.ascontiguousarray(
+                physical - w[:, None, None], dtype=np.float32
+            )
+
+            def install_initial():
+                with sim.ctx.activate():
+                    cuda.memcpy_htod(sim.field.ddf_gpu, shifted)
+                    cuda.memcpy_htod(sim.field.temp_gpu, shifted)
+                sim.stepper._step_count = (
+                    1  # count metadata; physical elapsed time is 0
+                )
+                sim.field.invalidate_host_ddf_cache()
+                flags = sim.field.flag.reshape(ny, nx)
+                flags[:] = FLUID
+                flags[[0, -1], :] = SOLID | BC_WALL
+                sim.field.upload_flags()
+                sim.field.update_params(fx=0.0, fy=0.0, fz=0.0)
+
+            def observe():
+                data = np.empty((9, ny, nx), np.float32)
+                with sim.ctx.activate():
+                    cuda.memcpy_dtoh(data, sim.field.ddf_gpu)
+                return decode_macros(data, w, cx, cy)
+
+            install_initial()
+            observe()  # reject invalid state before any advancing launch
+            sim.run(1)  # exercise the actual advancing kernel before claiming READY
+            observe()  # also require valid density/macros after the warm-up launch
+            install_initial()  # restore both (possibly swapped) DB buffers, flags and count
+            initial, rho0, _, _ = observe()
+            initial_error = relative(initial, U * np.sin(np.pi * y / H))
+            if (
+                not np.isfinite(initial_error)
+                or initial_error > LIMITS["initial_relative_l2"]
+            ):
+                raise RuntimeError("Initial state adapter failed")
+            with sim.ctx.activate():
+                device_name = cuda.Context.get_device().name()
+                cuda.Context.synchronize()
+            print(
+                json.dumps(
+                    dict(
+                        status="READY",
+                        cs=cs,
+                        device=device_name,
+                        variant=args.variant,
+                        source=pins,
+                        warmup_steps=1,
+                        initial_state_restored=True,
+                        initial_relative_l2=initial_error,
+                        max_steps=TIMES[-1],
+                        mode="SRT double_buffer shifted FP32",
+                    )
+                ),
+                flush=True,
+            )
+            yf, ref = solve(cs, n=256)
+            yc, coarse = solve(cs, n=128)
+            _, wrong = solve(cs * 2 ** (-0.25), n=256) if cs else (yf, ref)
+            case = dict(
+                cs=cs,
+                device=device_name,
+                warmup_steps=1,
+                initial_state_restored=True,
+                configuration=cfg,
+                initial_relative_l2=initial_error,
+                observations=[],
+            )
+            report["cases"].append(case)
+            elapsed = 0
+            for target in TIMES:
+                sim.run(target - elapsed)
+                elapsed = target
+                profile, mass, transverse, xspread = observe()
+                standard = np.interp(y, yf, ref[target])
+                analytic = (
+                    U * np.sin(np.pi * y / H) * np.exp(-NU * (np.pi / H) ** 2 * target)
+                )
+                oracle_grid = relative(np.interp(yc, yf, ref[target]), coarse[target])
+                reference = standard if cs else analytic
+                error = relative(profile, reference)
+                obs = dict(
+                    elapsed_steps=elapsed,
+                    completed_count=sim.stepper.step_count,
+                    relative_l2=error,
+                    oracle_grid_relative=oracle_grid,
+                    mass_drift=abs(mass - rho0),
+                    max_transverse=transverse,
+                    x_invariance=xspread,
+                    amplitude=amplitude(y, profile),
+                    reference_amplitude=amplitude(y, reference),
+                    wrong_convention_relative_l2=relative(
+                        profile, np.interp(y, yf, wrong[target])
+                    ),
+                    profile=profile.tolist(),
+                    reference_profile=reference.tolist(),
+                )
+                obs["passed"] = bool(
+                    error <= LIMITS["les_relative_l2" if cs else "linear_relative_l2"]
+                    and oracle_grid <= LIMITS["oracle_grid_relative"]
+                    and abs(mass - rho0) <= LIMITS["mass_drift"]
+                    and transverse <= LIMITS["max_transverse"]
+                    and xspread <= LIMITS["x_invariance"]
+                )
+                case["observations"].append(obs)
+                save()
+                print(
+                    json.dumps(
+                        {k: v for k, v in obs.items() if not k.endswith("profile")}
+                    ),
+                    flush=True,
+                )
+        finally:
+            sim.close()
+    # Cancel the common molecular lattice error using the actual Cs=0 run.
+    increments = []
+    for zero, les in zip(
+        report["cases"][0]["observations"], report["cases"][-1]["observations"]
+    ):
+        measured = zero["amplitude"] - les["amplitude"]
+        expected = zero["reference_amplitude"] - les["reference_amplitude"]
+        increments.append(
+            dict(
+                elapsed_steps=zero["elapsed_steps"],
+                measured=measured,
+                expected=expected,
+                relative_error=abs(measured - expected) / abs(expected),
+            )
+        )
+    report.update(sgs_increments=increments, seconds=time.monotonic() - started)
+    # Last two times are the declared signal gate; earliest is transient diagnostic.
+    report["passed"] = bool(
+        all(o["passed"] for c in report["cases"] for o in c["observations"])
+        and all(
+            v["relative_error"] <= LIMITS["sgs_increment_relative_error"]
+            for v in increments[1:]
+        )
+    )
+    save()
+    print(
+        json.dumps(
+            dict(status="DONE", passed=report["passed"], output=str(args.output))
+        ),
+        flush=True,
+    )
+    return 0 if report["passed"] else 1
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())

tests/validation/source_provenance.py

@@ -0,0 +1,15 @@
+"""Record runtime checkout identity separately from the reviewed source origin."""
+
+import subprocess
+
+
+def checkout_provenance(repo):
+    """Require a Git checkout; never label changed-source runs as historical evidence."""
+    revision = subprocess.check_output(
+        ["git", "-C", str(repo), "rev-parse", "HEAD"], text=True
+    ).strip()
+    dirty = subprocess.check_output(
+        ["git", "-C", str(repo), "status", "--porcelain", "--untracked-files=normal"],
+        text=True,
+    )
+    return dict(checkout_revision=revision, checkout_dirty=bool(dirty))

tests/validation/taylor_green/assess.py

@@ -0,0 +1,124 @@
+"""Frozen assessment of the complete 16-case matrix, independent of GPU launch."""
+
+import argparse
+import json
+import math
+from pathlib import Path
+
+from oracle import TIMES, LIMITS, U
+
+# Numerical floors reflect finite-Mach/FP32 and equilibrium startup, not observed runs.
+FLOORS = dict(velocity_l2=0.002, pressure_l2=0.02, energy_relative=0.004)
+BASELINE_FACTOR = 2.0
+REFINEMENT_FACTOR = 0.9
+
+
+def assess(records):
+    indexed = {}
+    for r in records:
+        key = (r["re_L"], r["n"], r["depth"])
+        if key in indexed:
+            raise ValueError(f"duplicate case {key}")
+        indexed[key] = r
+    expected = {(re, n, 1) for re in (100, 1000) for n in (32, 64, 128, 256)}
+    expected |= {(re, n, d) for re in (100, 1000) for n in (32, 64) for d in (2, 3)}
+    missing = sorted(expected - indexed.keys())
+    checks = {
+        "complete_matrix": not missing,
+        "no_extra_cases": set(indexed) <= expected,
+    }
+    identities = {json.dumps(r["source"], sort_keys=True) for r in records}
+    checks["same_source_and_harness"] = len(identities) == 1
+    checks["clean_source"] = bool(records) and all(
+        not r["source"]["dirty"] for r in records
+    )
+    for key, r in indexed.items():
+        label = str(key)
+        rows = r["rows"]
+        checks[label + " complete"] = r["completed"] and len(rows) == len(TIMES)
+        checks[label + " scales"] = (
+            r["U"] == U and abs(r["viscosity"] - U * r["n"] / r["re_L"]) < 1e-15
+        )
+        checks[label + " times"] = len(rows) == len(TIMES) and all(
+            abs(v["tstar"] - t) < 1e-12 for v, t in zip(rows, TIMES)
+        )
+        checks[label + " absolute"] = bool(rows) and all(
+            0 <= row[k] <= limit for row in rows for k, limit in LIMITS.items()
+        )
+    comparisons = []
+
+    def worst(key, metric):
+        values = [r[metric] for r in indexed.get(key, {}).get("rows", [])[1:]]
+        return max(values) if values and all(math.isfinite(v) for v in values) else None
+
+    # Diagnostics survive failed screens and partial pilot matrices. Missing or
+    # nonfinite operands are explicit nulls; they cannot acquire a passing gate.
+    for re in (100, 1000):
+        for depth in (2, 3):
+            for n in (32, 64):
+                if (re, n, depth) not in indexed:
+                    continue
+                for metric, floor in FLOORS.items():
+                    amr = worst((re, n, depth), metric)
+                    coarse = worst((re, n, 1), metric)
+                    fine = worst((re, n * 2 ** (depth - 1), 1), metric)
+                    ok = (
+                        amr is not None
+                        and coarse is not None
+                        and amr <= max(BASELINE_FACTOR * coarse, floor)
+                    )
+                    comparisons.append(
+                        dict(
+                            re_L=re,
+                            n=n,
+                            depth=depth,
+                            metric=metric,
+                            amr=amr,
+                            uniform_coarse=coarse,
+                            uniform_finest=fine,
+                            amr_to_finest=(
+                                None
+                                if amr is None or fine is None
+                                else amr / max(fine, 1e-30)
+                            ),
+                            passed=ok,
+                        )
+                    )
+                    checks[f"baseline {re}/{n}/{depth}/{metric}"] = ok
+            for metric, floor in FLOORS.items():
+                low = worst((re, 32, depth), metric)
+                high = worst((re, 64, depth), metric)
+                checks[f"refinement {re}/{depth}/{metric}"] = (
+                    low is not None
+                    and high is not None
+                    and high <= max(REFINEMENT_FACTOR * low, floor)
+                )
+        for metric, floor in FLOORS.items():
+            low = worst((re, 32, 1), metric)
+            high = worst((re, 256, 1), metric)
+            checks[f"uniform refinement {re}/{metric}"] = (
+                low is not None
+                and high is not None
+                and high <= max(REFINEMENT_FACTOR * low, floor)
+            )
+    return dict(
+        passed=all(checks.values()),
+        missing=missing,
+        checks=checks,
+        comparisons=comparisons,
+        floors=FLOORS,
+        baseline_factor=BASELINE_FACTOR,
+        refinement_factor=REFINEMENT_FACTOR,
+        scope="2D periodic smooth-decay AMR verification only; no turbulence, LES, wall or Jet certification",
+    )
+
+
+if __name__ == "__main__":
+    p = argparse.ArgumentParser(description=__doc__)
+    p.add_argument("files", type=Path, nargs="+")
+    p.add_argument("--output", type=Path, required=True)
+    a = p.parse_args()
+    result = assess([json.loads(path.read_text()) for path in a.files])
+    a.output.write_text(json.dumps(result, indent=2, allow_nan=False) + "\n")
+    print(json.dumps(result, indent=2))
+    raise SystemExit(0 if result["passed"] else 1)

tests/validation/taylor_green/oracle.py

@@ -0,0 +1,103 @@
+"""CPU-only 2D Taylor–Green reference and composite-leaf diagnostics."""
+
+import numpy as np
+
+CX = np.array([0, 1, -1, 0, 0, 1, -1, 1, -1])
+CY = np.array([0, 0, 0, 1, -1, 1, -1, -1, 1])
+W = np.array([4 / 9] + [1 / 9] * 4 + [1 / 36] * 4)
+U = 0.03
+HORIZON = 0.48  # t* = t U/L; exactly 16*N base steps
+TIMES = [0.0, 0.12, 0.24, 0.36, 0.48]
+# Frozen screening ceilings; baseline-relative qualification is in assess.py.
+LIMITS = dict(
+    velocity_l2=0.05, pressure_l2=0.15, energy_relative=0.05, mass_relative=1e-5
+)
+
+
+def exact(x, y, t, length, viscosity, amplitude=U):
+    k = 2 * np.pi / length
+    a = amplitude * np.exp(-2 * viscosity * k * k * t)
+    # TensorMesh phase: negative cos(x)sin(y), positive sin(x)cos(y).
+    ux = -a * np.cos(k * x) * np.sin(k * y)
+    uy = a * np.sin(k * x) * np.cos(k * y)
+    p = -0.25 * a * a * (np.cos(2 * k * x) + np.cos(2 * k * y))
+    return ux, uy, p  # kinematic pressure, mean zero
+
+
+def centers(shape, level=0):
+    # Campaign deliberately uses level-wide union layouts, all origins zero.
+    y, x = np.indices(shape, dtype=float)
+    h = 2.0 ** (-level)
+    return h * (x + 0.5), h * (y + 0.5)
+
+
+def initial(shape, level, length, viscosity):
+    x, y = centers(shape, level)
+    ux, uy, p = exact(x, y, 0.0, length, viscosity)
+    rho = 1 + 3 * p  # cs^2=1/3, rho0=1
+    cu = CX[:, None, None] * ux + CY[:, None, None] * uy
+    # Exact FP64 weights; public shifted initializer encodes before FP32 rounding.
+    return (
+        W[:, None, None]
+        * rho
+        * (1 + 3 * cu + 4.5 * cu * cu - 1.5 * (ux * ux + uy * uy))
+    )
+
+
+def rectangles(n, depth):
+    if n < 32 or n % 32 or depth not in (2, 3):
+        raise ValueError("AMR requires N >= 32 divisible by 32 and depth 2 or 3")
+    # Broad off-axis rectangle cuts nonzero velocity and both gradients.
+    out = [[(n // 8, n // 8, 7 * n // 8, 3 * n // 4)]]
+    if depth == 3:
+        # L1 index coordinates; global physical bounds (3N/8,N/4,3N/4,5N/8).
+        out.append([(3 * n // 4, n // 2, 3 * n // 2, 5 * n // 4)])
+    return out
+
+
+def metrics(arrays, masks, t, length, viscosity, initial_mass=None):
+    accum = []
+    for level, (g, mask) in enumerate(zip(arrays, masks)):
+        if g.shape != (9, *mask.shape):
+            raise ValueError("population/mask shape mismatch")
+        values = g[:, mask].astype(float)
+        if not np.isfinite(values).all():
+            raise FloatingPointError("nonfinite leaf populations")
+        rho = values.sum(axis=0)
+        if (rho <= 0).any():
+            raise FloatingPointError("nonpositive leaf density")
+        ux, uy = CX @ values / rho, CY @ values / rho
+        x, y = centers(mask.shape, level)
+        ex, ey, ep = exact(x[mask], y[mask], t, length, viscosity)
+        area = np.full(rho.size, 0.25**level)
+        accum.append((rho, ux, uy, (rho - 1) / 3, ex, ey, ep, area))
+    rho, ux, uy, p, ex, ey, ep, area = [np.concatenate(v) for v in zip(*accum)]
+    volume = area.sum()
+    if not np.isclose(volume, length**2, rtol=0, atol=1e-10):
+        raise ValueError("leaf partition does not cover domain once")
+    mass = float(area @ rho)
+    # Remove one global pressure gauge, never one gauge per level.
+    p = p - area @ p / volume
+    ep = ep - area @ ep / volume
+    energy = float(area @ (ux * ux + uy * uy) / 2 / volume)
+    exact_energy = U * U / 4 * np.exp(-4 * viscosity * (2 * np.pi / length) ** 2 * t)
+    return dict(
+        tstar=t * U / length,
+        mass=mass,
+        mass_relative=(
+            0.0 if initial_mass is None else abs(mass - initial_mass) / initial_mass
+        ),
+        velocity_l2=float(
+            np.sqrt(
+                area @ ((ux - ex) ** 2 + (uy - ey) ** 2) / (area @ (ex * ex + ey * ey))
+            )
+        ),
+        pressure_l2=float(np.sqrt(area @ ((p - ep) ** 2) / (area @ (ep * ep)))),
+        energy=energy,
+        exact_energy=exact_energy,
+        energy_relative=abs(energy - exact_energy) / exact_energy,
+        rho_min=float(rho.min()),
+        rho_max=float(rho.max()),
+        max_speed=float(np.max(np.hypot(ux, uy))),
+        leaf_cells=int(len(rho)),
+    )

tests/validation/taylor_green/run_case.py

@@ -0,0 +1,205 @@
+"""One bounded GPU case. No production edits, plots, checkpoints or field dumps."""
+
+import argparse
+import hashlib
+import json
+import os
+from pathlib import Path
+import subprocess
+import sys
+import time
+
+import numpy as np
+from oracle import U, TIMES, LIMITS, initial, metrics, rectangles
+
+
+def provenance(repo):
+    git = lambda *a: subprocess.check_output(
+        ["git", "-C", str(repo), *a], text=True
+    ).strip()
+    source_hash = hashlib.sha256()
+    for path in sorted((repo / "src").rglob("*")):
+        if path.is_file() and "__pycache__" not in path.parts and path.suffix != ".pyc":
+            source_hash.update(
+                str(path.relative_to(repo)).encode() + b"\0" + path.read_bytes() + b"\0"
+            )
+    return dict(
+        revision=git("rev-parse", "HEAD"),
+        dirty=bool(git("status", "--porcelain")),
+        source_tree_sha256=source_hash.hexdigest(),
+        harness_sha256={
+            p.name: hashlib.sha256(p.read_bytes()).hexdigest()
+            for p in Path(__file__).parent.glob("*.py")
+        },
+    )
+
+
+def config(n, re):
+    from CelerisLab.config import LBMConfig
+
+    cfg = LBMConfig(
+        nx=n,
+        ny=n,
+        viscosity=U * n / re,
+        velocity=0.0,
+        rho=1.0,
+        lattice_model="D2Q9",
+        collision="SRT",
+        streaming="double_buffer",
+        store_precision="FP32",
+        ddf_shifting=True,
+        les_enabled=False,
+    )
+    cfg.validate()
+    return cfg
+
+
+def workload(masks, base_steps):
+    cells = [int(np.count_nonzero(m)) for m in masks]
+    hops = [base_steps * 2**level for level in range(len(masks))]
+    updates = [count * hop for count, hop in zip(cells, hops)]
+    return dict(
+        active_cells_by_level=cells,
+        level_hops=hops,
+        active_collision_updates_by_level=updates,
+        active_collision_updates=sum(updates),
+        root_intervals=base_steps,
+        excludes="ghost streaming, halo copies, transfers, initialization and diagnostics",
+    )
+
+
+def main():
+    p = argparse.ArgumentParser(description=__doc__)
+    p.add_argument("--repo", type=Path, required=True)
+    p.add_argument("--re", type=int, choices=[100, 1000], required=True)
+    p.add_argument("--n", type=int, choices=[32, 64, 128, 256], required=True)
+    p.add_argument("--depth", type=int, choices=[1, 2, 3], required=True)
+    p.add_argument("--output", type=Path, required=True)
+    args = p.parse_args()
+    if not os.environ.get("CUDA_VISIBLE_DEVICES"):
+        p.error("Explicit physical CUDA_VISIBLE_DEVICES selection required")
+    repo = args.repo.resolve()
+    sys.path.insert(0, str(repo / "src"))
+    import CelerisLab
+
+    if Path(CelerisLab.__file__).resolve() != repo / "src/CelerisLab/__init__.py":
+        raise RuntimeError("Wrong imported source checkout")
+    from CelerisLab.cuda.context import CudaContext
+    from CelerisLab.lbm.blocks import BlockLayout2D
+    from CelerisLab.lbm.block_stepper import PeriodicBlockStepper
+    from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+    from CelerisLab.lbm.refinement import ACTIVE
+    from CelerisLab.lbm.hierarchy import build_hierarchy, leaf_partition
+
+    cfg = config(args.n, args.re)
+    refinements = [] if args.depth == 1 else rectangles(args.n, args.depth)
+    h = build_hierarchy(cfg, refinements)
+    np.testing.assert_array_equal(leaf_partition(h), np.ones((args.n, args.n)))
+    masks = [k == ACTIVE for k in h.kinds]
+    report = dict(
+        schema=1,
+        source=provenance(repo),
+        n=args.n,
+        re_L=args.re,
+        depth=args.depth,
+        refinements=refinements,
+        U=U,
+        mach=U * np.sqrt(3),
+        viscosity=cfg.viscosity,
+        tau_by_level=[0.5 + 3 * c.viscosity for c in h.configs],
+        initialization="physical FP64 equilibrium G at actual cell centers; rho=1+3p",
+        cuda_visible_devices=os.environ["CUDA_VISIBLE_DEVICES"],
+        device_id=0,
+        limits=LIMITS,
+        workload=workload(masks, round(TIMES[-1] * args.n / U)),
+        rows=[],
+        completed=False,
+    )
+    args.output.parent.mkdir(parents=True, exist_ok=True)
+    save = lambda: args.output.write_text(
+        json.dumps(report, indent=2, allow_nan=False) + "\n"
+    )
+    save()
+    started = time.monotonic()
+    with CudaContext(0) as ctx:
+        if args.depth == 1:
+            layout = BlockLayout2D.build(args.n, args.n, 8, 8, periodic=(True, True))
+            run = PeriodicBlockStepper(cfg, layout, context=ctx)
+        else:
+            run = FixedPatchStepper(
+                cfg,
+                context=ctx,
+                coarse_boundary="periodic",
+                refinements=refinements,
+                level_blocks=[(8, 8)] * args.depth,
+            )
+        with run:
+            run.initialize(
+                *[
+                    initial(m.shape, level, args.n, cfg.viscosity)
+                    for level, m in enumerate(masks)
+                ]
+            )
+            initial_mass = None
+            for tstar in TIMES:
+                step = round(tstar * args.n / U)
+                run.step(step - run.step_count)
+                arrays = (run.download(),) if args.depth == 1 else run.download()
+                row = metrics(arrays, masks, step, args.n, cfg.viscosity, initial_mass)
+                row["step"] = step
+                if initial_mass is None:
+                    initial_mass = row["mass"]
+                report["rows"].append(row)
+                save()
+                print(
+                    json.dumps(
+                        dict(
+                            event="PROGRESS",
+                            n=args.n,
+                            re_L=args.re,
+                            depth=args.depth,
+                            **row,
+                        )
+                    ),
+                    flush=True,
+                )
+                if step == 0:
+                    # Exercise real kernels/halos/transfers on every level before READY.
+                    # This is the first measured trajectory step, not discarded warmup.
+                    run.step(1)
+                    probe = (run.download(),) if args.depth == 1 else run.download()
+                    metrics(probe, masks, 1, args.n, cfg.viscosity, initial_mass)
+                    report["ready_after_step"] = 1
+                    save()
+                    print(
+                        json.dumps(
+                            dict(
+                                event="READY",
+                                step=1,
+                                n=args.n,
+                                re_L=args.re,
+                                depth=args.depth,
+                                workload=report["workload"],
+                            )
+                        ),
+                        flush=True,
+                    )
+    report.update(
+        completed=True,
+        seconds=time.monotonic() - started,
+        screening_pass=all(
+            row[key] <= limit for row in report["rows"] for key, limit in LIMITS.items()
+        ),
+    )
+    save()
+    print(
+        json.dumps(
+            {k: report[k] for k in ("n", "re_L", "depth", "screening_pass", "seconds")}
+        ),
+        flush=True,
+    )
+    return 0 if report["screening_pass"] else 1
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())

tests/validation/taylor_green/test_oracle.py

@@ -0,0 +1,170 @@
+"""Standalone CPU oracle/geometry checks. Set PYTHONPATH=<checkout>/src."""
+
+import unittest
+import numpy as np
+from oracle import CX, CY, U, exact, centers, initial, rectangles, metrics
+from run_case import config
+
+
+class OracleTests(unittest.TestCase):
+    def test_phase_pressure_and_decay(self):
+        u, v, p = exact(0.0, 8.0, 0.0, 32.0, 0.01)
+        self.assertAlmostEqual(u, -U)
+        self.assertAlmostEqual(v, 0.0)
+        self.assertAlmostEqual(p, 0.0)
+        self.assertAlmostEqual(exact(0.0, 0.0, 0.0, 32.0, 0.01)[2], -U * U / 2)
+        a = np.array(exact(3.0, 5.0, 0.0, 32.0, 0.01))
+        b = np.array(exact(3.0, 5.0, 20.0, 32.0, 0.01))
+        f = np.exp(-2 * 0.01 * (2 * np.pi / 32) ** 2 * 20)
+        np.testing.assert_allclose(b, a * np.array([f, f, f * f]), atol=1e-17)
+
+    def test_navier_stokes_residual(self):
+        # Independent finite-difference derivatives catch sign/phase mistakes.
+        x = np.array([2.7, 8.1, 13.3])
+        y = np.array([3.2, 7.4, 21.6])
+        t = 10.0
+        h = 0.002
+        n = 32.0
+        nu = 0.01
+        f = lambda x, y, t: np.array(exact(x, y, t, n, nu))
+        z = f(x, y, t)
+        dx = (f(x + h, y, t) - f(x - h, y, t)) / (2 * h)
+        dy = (f(x, y + h, t) - f(x, y - h, t)) / (2 * h)
+        dt = (f(x, y, t + h) - f(x, y, t - h)) / (2 * h)
+        lap = (
+            f(x + h, y, t) + f(x - h, y, t) + f(x, y + h, t) + f(x, y - h, t) - 4 * z
+        ) / h**2
+        self.assertLess(np.max(np.abs(dx[0] + dy[1])), 1e-11)
+        residual = (
+            dt[:2]
+            + z[0] * dx[:2]
+            + z[1] * dy[:2]
+            + np.array([dx[2], dy[2]])
+            - nu * lap[:2]
+        )
+        self.assertLess(np.max(np.abs(residual)), 1e-10)
+
+    def test_equilibrium_moments_and_metrics(self):
+        n = 32
+        nu = 0.01
+        g = initial((n, n), 0, n, nu)
+        rho = g.sum(axis=0)
+        x, y = centers((n, n))
+        u, v, p = exact(x, y, 0, n, nu)
+        np.testing.assert_allclose(rho, 1 + 3 * p, atol=1e-15)
+        np.testing.assert_allclose(np.einsum("q,qyx->yx", CX, g) / rho, u, atol=1e-15)
+        np.testing.assert_allclose(np.einsum("q,qyx->yx", CY, g) / rho, v, atol=1e-15)
+        m = metrics((g,), [np.ones((n, n), bool)], 0, n, nu)
+        for k in ("velocity_l2", "pressure_l2", "energy_relative"):
+            self.assertLess(m[k], 1e-11)
+        with self.assertRaises(FloatingPointError):
+            g[:, 2, 2] = np.nan
+            metrics((g,), [np.ones((n, n), bool)], 0, n, nu)
+
+    def test_geometry_cell_centers_and_scaling(self):
+        from CelerisLab.lbm.hierarchy import build_hierarchy, leaf_partition
+        from CelerisLab.lbm.refinement import ACTIVE
+
+        for n in (32, 64):
+            for re in (100, 1000):
+                for depth in (2, 3):
+                    cfg = config(n, re)
+                    h = build_hierarchy(cfg, rectangles(n, depth))
+                    np.testing.assert_array_equal(leaf_partition(h), np.ones((n, n)))
+                    for level, cfg_l in enumerate(h.configs):
+                        self.assertAlmostEqual(
+                            cfg_l.viscosity, cfg.viscosity * 2**level
+                        )
+                        self.assertLessEqual(cfg_l.omega, cfg_l.omega_max)
+                        x, y = centers(h.kinds[level].shape, level)
+                        self.assertEqual(x[0, 0], 0.5 / 2**level)
+                        self.assertEqual(x[0, -1], n - 0.5 / 2**level)
+                    masks = [k == ACTIVE for k in h.kinds]
+                    g = [
+                        initial(k.shape, l, n, cfg.viscosity)
+                        for l, k in enumerate(h.kinds)
+                    ]
+                    result = metrics(g, masks, 0, n, cfg.viscosity)
+                    self.assertLess(result["velocity_l2"], 1e-12)
+                    self.assertLess(result["pressure_l2"], 1e-10)
+                    self.assertLess(result["energy_relative"], 0.002)
+
+    def test_complete_assessment_and_corruptions(self):
+        from copy import deepcopy
+        from assess import assess
+        from oracle import TIMES
+
+        records = []
+        for re in (100, 1000):
+            cases = [(n, 1) for n in (32, 64, 128, 256)]
+            cases += [(n, d) for n in (32, 64) for d in (2, 3)]
+            for n, d in cases:
+                rows = [
+                    dict(
+                        tstar=t,
+                        velocity_l2=0.0,
+                        pressure_l2=0.0,
+                        energy_relative=0.0,
+                        mass_relative=0.0,
+                    )
+                    for t in TIMES
+                ]
+                records.append(
+                    dict(
+                        re_L=re,
+                        n=n,
+                        depth=d,
+                        U=U,
+                        viscosity=U * n / re,
+                        source=dict(revision="synthetic", dirty=False),
+                        completed=True,
+                        rows=rows,
+                    )
+                )
+        self.assertTrue(assess(records)["passed"])
+        bad = deepcopy(records)
+        bad[-1]["rows"][-1]["pressure_l2"] = 0.5
+        failed = assess(bad)
+        self.assertFalse(failed["passed"])
+        self.assertEqual(len(failed["comparisons"]), 24)
+        partial = assess(records[0:1] + records[4:5])
+        self.assertFalse(partial["passed"])
+        self.assertEqual(len(partial["comparisons"]), 3)
+        self.assertIsNone(partial["comparisons"][0]["uniform_finest"])
+
+        bad = deepcopy(records)
+        bad[-1]["rows"][-1]["pressure_l2"] = float("nan")
+        self.assertFalse(assess(bad)["passed"])
+        bad = deepcopy(records)
+        bad[-1]["source"]["dirty"] = True
+        self.assertFalse(assess(bad)["passed"])
+        bad = deepcopy(records)
+        bad[-1]["rows"][-1]["tstar"] = 0.49
+        self.assertFalse(assess(bad)["passed"])
+        with self.assertRaises(ValueError):
+            assess(records + [records[0]])
+
+    def test_workload(self):
+        from run_case import workload
+        from CelerisLab.lbm.hierarchy import build_hierarchy
+        from CelerisLab.lbm.refinement import ACTIVE
+
+        cases = [(n, 1) for n in (32, 64, 128, 256)]
+        cases += [(n, d) for n in (32, 64) for d in (2, 3)]
+        costs = []
+        for n, d in cases:
+            h = build_hierarchy(config(n, 100), [] if d == 1 else rectangles(n, d))
+            costs.append(workload([k == ACTIVE for k in h.kinds], 16 * n))
+        self.assertEqual(sum(c["active_collision_updates"] for c in costs), 384270336)
+        self.assertEqual(sum(c["root_intervals"] for c in costs), 10752)
+        self.assertEqual(costs[-1]["active_collision_updates"], 50987008)
+
+    def test_incomplete_assessment_fails(self):
+        from assess import assess
+
+        self.assertFalse(assess([])["passed"])
+        self.assertEqual(len(assess([])["missing"]), 16)
+
+
+if __name__ == "__main__":
+    unittest.main()

tests/validation/time_refined_dfg.py

@@ -0,0 +1,326 @@
+"""Measurement-only wall timing for the matched MRT refined DFG case. No numerical edits.
+
+Variants over EQUAL simulated duration (one refined base step = two uniform D80 steps):
+  A  refined: step(1) + read_raw_exchange per base step  -- runner-like force-read workload
+  B  refined: step(1) only, no host reads        -- timing-only, no force history
+  C  refined: one step(N) call                   -- timing-only; ledger keeps final base interval only
+  E  uniform: sim.run(<=100) + read_force                -- runner-like force-read workload
+  E' uniform: one sim.run(2N) call                -- timing-only; at most a whole-window force
+                                                    average (different summation order), no history
+A/E read forces but discard them: no retained force history, no interval aggregation/validation,
+and no macro/probe/CSV cost, so they are not full benchmark throughput.
+observe_raw_exchange stays True for A/B/C: B/C omit host reads only.
+Status lines: CONSTRUCTED = backend/simulation built (JIT done); WARM_READY = first warm-up and
+health check passed for that case.
+Parity compares A/B/C with each other and E/E' with each other, never uniform with refined.
+A-B and B-C are measured comparative differences, not exact causal decompositions.
+Each repeat starts from the same fresh initial state with the same warm-up; construction,
+compilation, parity and health checks are outside timed regions. The actual stepping stream is
+synchronized explicitly immediately before the timer starts and after the timed loop. Wall time
+(perf_counter) is the only timing source; no CUDA-event span is reported as kernel time. An outer
+--max-wall-seconds deadline is checked between variants and inside per-step loops.
+"""
+
+import argparse
+import json
+import math
+import os
+from pathlib import Path
+import statistics
+import sys
+import time
+from tempfile import TemporaryDirectory
+
+import numpy as np
+
+HERE = Path(__file__).resolve().parent
+sys.path.insert(0, str(HERE))
+import run_refined_dfg as r  # noqa: E402  (published runner: same setup functions)
+
+VARIANTS = {
+    "A": dict(case="refined", force_history_valid=False, force_reads=True,
+              force_note="runner-like force-read workload; values discarded, no aggregation"),
+    "B": dict(case="refined", force_history_valid=False, timing_only=True,
+              force_note="no host reads"),
+    "C": dict(case="refined", force_history_valid=False, timing_only=True,
+              force_note="raw ledger retains only the final base interval"),
+    "E": dict(case="uniform", force_history_valid=False, force_reads=True,
+              force_note="runner-like force-read workload; values discarded, no aggregation"),
+    "E'": dict(case="uniform", force_history_valid=False, timing_only=True,
+               force_note="at most a whole-window accumulated average; different summation order"),
+}
+
+
+def schedule(n_base, warmup_base):
+    """Steps per case for equal simulated duration (refined dt_root = 2 dt_uniform)."""
+    if n_base <= 0 or warmup_base < 0:
+        raise ValueError("n_base must be positive and warmup_base nonnegative")
+    return dict(refined=(warmup_base, n_base), uniform=(2 * warmup_base, 2 * n_base),
+                tstar_window=n_base * r.SPEED / 40.0)
+
+
+# Existing run_refined_dfg safety limits (rho in [.5, 1.5], sampled max Mach <= .2).
+RHO_LIMITS, MAX_MACH = (0.5, 1.5), 0.2
+
+
+def same_state(reference, values):
+    """Bitwise populations and identical NaN masks; returns a short description or None.
+
+    Matching NaN is accepted only as a storage mask (inactive/ghost/solid slots); physical
+    leaf health is checked separately. Any infinity in either state is a mismatch.
+    """
+    reference, values = list(reference), list(values)
+    if len(reference) != len(values):
+        return f"level count {len(reference)} != {len(values)}"
+    for level, (a, b) in enumerate(zip(reference, values)):
+        a, b = np.asarray(a), np.asarray(b)
+        if a.shape != b.shape:
+            return f"level {level}: shape {a.shape} != {b.shape}"
+        if np.isinf(a).any() or np.isinf(b).any():
+            return f"level {level}: infinite population"
+        if not np.array_equal(np.isnan(a), np.isnan(b)):
+            return f"level {level}: NaN mask differs"
+        finite = np.isfinite(a)
+        if a.dtype != b.dtype or not np.array_equal(a[finite].view(np.uint8), b[finite].view(np.uint8)):
+            return f"level {level}: finite bits differ"
+    return None
+
+
+def uniform_populations_finite(ddf, fluid):
+    """Simulation.get_ddf() may be flat (9*ny*nx); view it as (Q, cells) before masking."""
+    ddf, fluid = np.asarray(ddf), np.asarray(fluid, bool).ravel()
+    if fluid.size == 0 or ddf.size % fluid.size:
+        raise RuntimeError(f"DDF size {ddf.size} incompatible with {fluid.size} cells")
+    return bool(np.isfinite(ddf.reshape(ddf.size // fluid.size, fluid.size)[:, fluid]).all())
+
+
+def physical_health(fields):
+    """fields: (rho, ux, uy) arrays already restricted to physical leaves/fluid."""
+    if not fields or any(not np.isfinite(v).all() or v.size == 0 for f in fields for v in f):
+        raise RuntimeError("nonfinite or empty physical macroscopic field")
+    rho_min = float(min(f[0].min() for f in fields))
+    rho_max = float(max(f[0].max() for f in fields))
+    max_mach = float(max(np.sqrt(3 * (f[1] ** 2 + f[2] ** 2)).max() for f in fields))
+    if not (RHO_LIMITS[0] <= rho_min <= rho_max <= RHO_LIMITS[1] and max_mach <= MAX_MACH):
+        raise RuntimeError(f"health limits violated: rho [{rho_min}, {rho_max}], Mach {max_mach}")
+    return dict(rho_min=rho_min, rho_max=rho_max, max_mach=max_mach)
+
+
+def refined_health(backend, masks):
+    arrays = backend.download()  # physical FP64 populations; non-leaves NaN by contract
+    if any(not np.isfinite(g[:, m]).all() for g, m in zip(arrays, masks)):
+        raise RuntimeError("nonfinite physical leaf population")
+    return physical_health([tuple(v[m] for v in r.macros(g)) for g, m in zip(arrays, masks)])
+
+
+def uniform_health(sim, fluid):
+    m = sim.get_macroscopic()
+    return physical_health([tuple(np.asarray(m[k], np.float64).reshape(fluid.shape)[fluid]
+                                  for k in ("rho", "ux", "uy"))])
+
+
+def check_deadline(deadline):
+    if time.monotonic() >= deadline:
+        raise TimeoutError("timing harness exceeded --max-wall-seconds")
+
+
+def sync_refined(backend):
+    with backend.ctx.activate():
+        backend._stream.synchronize()
+
+
+def sync_uniform(sim):
+    with sim.ctx.activate():
+        sim.stream.synchronize()
+
+
+def refined_loop(backend, variant, n, deadline):
+    if variant == "C":
+        backend.step(n)
+        return
+    for i in range(n):
+        backend.step(1)
+        if variant == "A":
+            backend.read_raw_exchange(0, level=1, window="last_base_step")
+        if i % 100 == 99:
+            check_deadline(deadline)
+
+
+def uniform_loop(sim, variant, n, deadline):
+    if variant == "E'":
+        sim.run(n)
+        return
+    done = 0
+    while done < n:
+        batch = min(100, n - done)
+        sim.run(batch)
+        sim.read_force(0)
+        done += batch
+        check_deadline(deadline)
+
+
+def time_refined(a, h, sched):
+    from contextlib import ExitStack
+    from CelerisLab.cuda.context import CudaContext
+    from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+    from CelerisLab.lbm.channel import channel_masks
+    from CelerisLab.lbm.refinement import ACTIVE
+    from CelerisLab.lbm.descriptors import SOLID
+
+    warm, n = sched["refined"]
+    rows, reference = [], {}
+    with ExitStack() as stack:
+        cfg, pair, body, host = r.host_geometry(a.source, h, "MRT")
+        context = stack.enter_context(CudaContext(0))
+        backend = stack.enter_context(FixedPatchStepper(
+            cfg, r.BOUNDS, context=context, coarse_boundary="channel", base_block=r.BASE_BLOCK,
+            fine_block=r.FINE_BLOCK, fine_objects=(body,), observe_raw_exchange=True))
+        masks = [channel_masks(pair.coarse_shape)[0] & ~pair.covered,  # as run_refined_dfg
+                 (pair.fine_kind == ACTIVE) & ((host.flags.reshape(pair.fine_shape) & SOLID) == 0)]
+        fresh = (r.initial_state(pair.coarse_shape, (0.0, 0.0), 1.0, root=True),
+                 r.initial_state(pair.fine_shape, pair.fine_origin, 0.5))
+        print(json.dumps(dict(status="CONSTRUCTED", case="refined")), flush=True)
+        for repeat in range(a.repeats):
+            for variant in ("A", "B", "C"):
+                check_deadline(a.deadline)
+                backend.initialize(*fresh)  # resets step count; same fresh state every time
+                refined_loop(backend, variant, warm, a.deadline)
+                warm_health = refined_health(backend, masks)  # untimed
+                if repeat == 0 and variant == "A":
+                    print(json.dumps(dict(status="WARM_READY", case="refined")), flush=True)
+                sync_refined(backend)
+                start = time.perf_counter()
+                refined_loop(backend, variant, n, a.deadline)
+                sync_refined(backend)
+                seconds = time.perf_counter() - start
+                end_health = refined_health(backend, masks)  # untimed
+                values = [np.array(v, copy=True) for v in backend._download_raw().values]
+                if not all(np.isfinite(v[:, m]).all() for v, m in zip(values, backend._state)):
+                    raise RuntimeError(f"refined {variant} repeat {repeat}: nonfinite leaf populations")
+                mismatch = None
+                if repeat == 0 and variant == "A":
+                    reference["A"] = values
+                else:
+                    mismatch = same_state(reference["A"], values)
+                rows.append(dict(variant=variant, repeat=repeat, seconds=seconds,
+                                 base_steps=n, step_count=backend.step_count,
+                                 health_after_warmup=warm_health, health_after_window=end_health,
+                                 parity_vs_A=mismatch or "identical", **VARIANTS[variant]))
+                if mismatch:
+                    raise RuntimeError(f"refined {variant} repeat {repeat}: {mismatch}")
+    return rows
+
+
+def time_uniform(a, h, sched):
+    from CelerisLab.simulation import Simulation
+    from CelerisLab.lbm.descriptors import FLUID
+
+    warm, n = sched["uniform"]
+    rows, reference = [], {}
+    cfgdict, bodydict, params = r.config(h, 80, "MRT")
+    with TemporaryDirectory() as tmp:
+        tmp = Path(tmp)
+        (tmp / "lbm.json").write_text(json.dumps(cfgdict))
+        (tmp / "body.json").write_text(json.dumps(bodydict))
+        with Simulation(str(tmp / "lbm.json"), str(tmp / "body.json"), device_id=0) as sim:
+            sim.initialize()
+            sim.snapshot()  # fresh state; restore keeps clock and streaming parity
+            fluid = (sim.get_flags() & FLUID) != 0
+            ny, nx = params["ny"], params["nx"]
+            fluid = fluid.reshape(ny, nx)
+            print(json.dumps(dict(status="CONSTRUCTED", case="uniform")), flush=True)
+            for repeat in range(a.repeats):
+                for variant in ("E", "E'"):
+                    check_deadline(a.deadline)
+                    sim.restore()
+                    if sim.stepper.step_count % 2:
+                        raise RuntimeError("uniform restore parity changed between variants")
+                    uniform_loop(sim, variant, warm, a.deadline)
+                    warm_health = uniform_health(sim, fluid)  # untimed
+                    if repeat == 0 and variant == "E":
+                        print(json.dumps(dict(status="WARM_READY", case="uniform")), flush=True)
+                    sync_uniform(sim)
+                    start = time.perf_counter()
+                    uniform_loop(sim, variant, n, a.deadline)
+                    sync_uniform(sim)
+                    seconds = time.perf_counter() - start
+                    end_health = uniform_health(sim, fluid)  # untimed
+                    values = [sim.get_ddf()]
+                    if not uniform_populations_finite(values[0], fluid):
+                        raise RuntimeError(f"uniform {variant} repeat {repeat}: nonfinite fluid populations")
+                    mismatch = None
+                    if repeat == 0 and variant == "E":
+                        reference["E"] = values
+                    else:
+                        mismatch = same_state(reference["E"], values)
+                    rows.append(dict(variant=variant, repeat=repeat, seconds=seconds,
+                                     uniform_steps=n, parity_vs_E=mismatch or "identical",
+                                     health_after_warmup=warm_health, health_after_window=end_health,
+                                     **VARIANTS[variant]))
+                    if mismatch:
+                        raise RuntimeError(f"uniform {variant} repeat {repeat}: {mismatch}")
+    return rows
+
+
+def summarize(rows, tstar_window):
+    summary = {}
+    for variant in VARIANTS:
+        seconds = [row["seconds"] for row in rows if row["variant"] == variant]
+        if seconds:
+            median = statistics.median(seconds)
+            summary[variant] = dict(median_seconds=median, seconds=seconds,
+                                    seconds_per_tstar=median / tstar_window,
+                                    **VARIANTS[variant])
+    if {"A", "B", "C"} <= summary.keys():
+        summary["A_minus_B_seconds"] = summary["A"]["median_seconds"] - summary["B"]["median_seconds"]
+        summary["B_minus_C_seconds"] = summary["B"]["median_seconds"] - summary["C"]["median_seconds"]
+    if {"A", "E"} <= summary.keys():
+        summary["A_over_E"] = summary["A"]["median_seconds"] / summary["E"]["median_seconds"]
+    if {"C", "E'"} <= summary.keys():
+        summary["C_over_Eprime"] = summary["C"]["median_seconds"] / summary["E'"]["median_seconds"]
+    summary["interpretation"] = ("Measured comparative differences under GPU variability; "
+                                 "not exact causal decompositions or performance claims.")
+    return summary
+
+
+def main():
+    p = argparse.ArgumentParser(description=__doc__)
+    p.add_argument("--source", type=Path, default=r.DEFAULT_SOURCE)
+    p.add_argument("--output", type=Path, required=True)
+    p.add_argument("--base-steps", type=int, default=2000)
+    p.add_argument("--warmup-base-steps", type=int, default=200)
+    p.add_argument("--repeats", type=int, default=3)
+    p.add_argument("--max-wall-seconds", type=float, help="Required outer timeout for GPU runs")
+    p.add_argument("--plan", action="store_true")
+    a = p.parse_args()
+    if a.repeats < 1 or a.base_steps > 20000:
+        p.error("repeats must be >= 1 and base-steps <= 20000")
+    a.source = a.source.resolve()
+    manifest = r.verify_source(a.source, "MRT")
+    sched = schedule(a.base_steps, a.warmup_base_steps)
+    plan = dict(collision="MRT", manifest_commit=manifest["commit"], schedule=sched,
+                repeats=a.repeats, variants=VARIANTS, timer="time.perf_counter wall",
+                excluded="construction/JIT, warm-up, state download and parity checks")
+    if a.plan:
+        print(json.dumps(plan, indent=2))
+        return 0
+    if a.max_wall_seconds is None or not math.isfinite(a.max_wall_seconds) or a.max_wall_seconds <= 0:
+        p.error("GPU timing requires a finite positive --max-wall-seconds")
+    a.deadline = time.monotonic() + a.max_wall_seconds
+    if not os.environ.get("CUDA_VISIBLE_DEVICES") or "," in os.environ["CUDA_VISIBLE_DEVICES"]:
+        p.error("Root must select exactly one verified idle GPU via CUDA_VISIBLE_DEVICES")
+    if a.output.exists() and any(a.output.iterdir()):
+        p.error("Output must be empty/new")
+    a.output.mkdir(parents=True, exist_ok=True)
+    h, _ = r.helpers(a.source)
+    rows = time_uniform(a, h, sched) + time_refined(a, h, sched)
+    result = dict(plan=plan, rows=rows, summary=summarize(rows, sched["tstar_window"]))
+    if not all(math.isfinite(row["seconds"]) for row in rows):
+        raise RuntimeError("Nonfinite timing")
+    (a.output / "timing.json").write_text(json.dumps(result, indent=2))
+    print(json.dumps(result["summary"], indent=2), flush=True)
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())

tests/validation/time_window_dfg.py

@@ -0,0 +1,183 @@
+"""Measurement-only A-versus-window timing for the matched shifted-MRT refined DFG setup.
+
+Candidate source (window ledger) is recorded by checkout revision, dirty flag and file hashes;
+the frozen refined-DFG manifests/results are NOT used or modified (they pin pre-window source).
+Equal simulated duration, same fresh state and warm-up per variant, one observation interval =
+80 root base steps (the runner's dense-phase interval):
+  A   per base step: step(1) + read_raw_exchange, summed per interval (current runner workload)
+  W   per interval: start_exchange_window(); step(80); read_exchange_window()
+  W1  per interval: start_exchange_window(); 80 x step(1); read_exchange_window()
+All variants keep observe_raw_exchange=True; the backend is built once with
+observe_window_exchange=True (a closed window adds no launch). Populations (bitwise) and
+per-interval total impulses are compared OUTSIDE timing; the impulse tolerance is a
+conservative engineering tolerance (see INTERVAL_RTOL); the tighter substep-scaled engineering
+tolerance is exercised by the capability test tests/validation/multigrid/test_exchange_window.py.
+Wall time (perf_counter) with explicit stream synchronization is the only timing source.
+"""
+
+import argparse
+import hashlib
+import json
+import math
+import os
+from pathlib import Path
+import statistics
+import sys
+import time
+
+import numpy as np
+
+HERE = Path(__file__).resolve().parent
+sys.path.insert(0, str(HERE))
+import time_refined_dfg as t  # noqa: E402  (reuses run_refined_dfg setup via t.r)
+
+r = t.r
+INTERVAL = 80
+VARIANTS = ("A", "W", "W1")
+# Interval impulses differ only by FP64 summation order (per link over time vs per step over
+# links). Engineering tolerance: 1e-9 of the interval's summed absolute per-step totals.
+INTERVAL_RTOL = 1e-9
+CANDIDATE_FILES = ("src/CelerisLab/lbm/refinement_stepper.py", "src/CelerisLab/lbm/raw_exchange.py",
+                   "src/CelerisLab/lbm/refinement_tables.py", "src/CelerisLab/lbm/kernels/blocks.cu")
+
+
+def components(obs):
+    o = obs["total"]
+    return np.array([o["impulse"][0], o["impulse"][1], o["angular_impulse"], o["mass_into_fluid"]])
+
+
+def candidate_identity(source):
+    from source_provenance import checkout_provenance
+    return dict(checkout_provenance(source),
+                sha256={f: hashlib.sha256((source / f).read_bytes()).hexdigest()
+                        for f in CANDIDATE_FILES})
+
+
+def run_variant(backend, variant, n_intervals, deadline):
+    """Returns per-interval totals and, for A, per-interval summed |per-step totals|."""
+    totals, scales = [], []
+    for _ in range(n_intervals):
+        if variant == "A":
+            impulse, scale = np.zeros(4), np.zeros(4)
+            for _ in range(INTERVAL):
+                backend.step(1)
+                c = components(backend.read_raw_exchange(0, level=1, window="last_base_step"))
+                impulse += c
+                scale += np.abs(c)
+            totals.append(impulse)
+            scales.append(scale)
+        else:
+            backend.start_exchange_window()
+            if variant == "W":
+                backend.step(INTERVAL)
+            else:
+                for _ in range(INTERVAL):
+                    backend.step(1)
+            totals.append(components(backend.read_exchange_window(0)))
+        t.check_deadline(deadline)
+    return np.array(totals), (np.array(scales) if scales else None)
+
+
+def compare_intervals(reference, scales, totals):
+    """Max violation ratio; <= 1 passes. Scale includes a floor so zero cancels cannot divide."""
+    bound = INTERVAL_RTOL * scales + 1e-300
+    return float(np.max(np.abs(totals - reference) / bound))
+
+
+def main():
+    p = argparse.ArgumentParser(description=__doc__)
+    p.add_argument("--source", type=Path, default=r.DEFAULT_SOURCE)
+    p.add_argument("--output", type=Path, required=True)
+    p.add_argument("--base-steps", type=int, default=2000)
+    p.add_argument("--warmup-base-steps", type=int, default=200)
+    p.add_argument("--repeats", type=int, default=3)
+    p.add_argument("--max-wall-seconds", type=float)
+    p.add_argument("--plan", action="store_true")
+    a = p.parse_args()
+    if a.base_steps <= 0 or a.base_steps % INTERVAL or a.base_steps > 20000 or a.repeats < 1:
+        p.error(f"base-steps must be a positive multiple of {INTERVAL} (<= 20000); repeats >= 1")
+    a.source = a.source.resolve()
+    identity = candidate_identity(a.source)
+    plan = dict(collision="MRT", shifted_storage="as run_refined_dfg MRT config",
+                candidate_source=identity, frozen_manifests_used=False, interval_base_steps=INTERVAL,
+                base_steps=a.base_steps, warmup_base_steps=a.warmup_base_steps, repeats=a.repeats,
+                variants=VARIANTS, tstar_window=a.base_steps * r.SPEED / 40.0,
+                timer="time.perf_counter wall with explicit stream sync",
+                excluded="construction/JIT, warm-up, health, downloads, parity/impulse comparison")
+    if a.plan:
+        print(json.dumps(plan, indent=2))
+        return 0
+    if a.max_wall_seconds is None or not math.isfinite(a.max_wall_seconds) or a.max_wall_seconds <= 0:
+        p.error("GPU timing requires a finite positive --max-wall-seconds")
+    if not os.environ.get("CUDA_VISIBLE_DEVICES") or "," in os.environ["CUDA_VISIBLE_DEVICES"]:
+        p.error("Root must select exactly one verified idle GPU via CUDA_VISIBLE_DEVICES")
+    if a.output.exists() and any(a.output.iterdir()):
+        p.error("Output must be empty/new")
+    a.output.mkdir(parents=True, exist_ok=True)
+    deadline = time.monotonic() + a.max_wall_seconds
+    h, _ = r.helpers(a.source)
+
+    from contextlib import ExitStack
+    from CelerisLab.cuda.context import CudaContext
+    from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+    from CelerisLab.lbm.channel import channel_masks
+    from CelerisLab.lbm.refinement import ACTIVE
+    from CelerisLab.lbm.descriptors import SOLID
+
+    rows, ref = [], {}
+    n_intervals = a.base_steps // INTERVAL
+    with ExitStack() as stack:
+        cfg, pair, body, host = r.host_geometry(a.source, h, "MRT")
+        context = stack.enter_context(CudaContext(0))
+        backend = stack.enter_context(FixedPatchStepper(
+            cfg, r.BOUNDS, context=context, coarse_boundary="channel", base_block=r.BASE_BLOCK,
+            fine_block=r.FINE_BLOCK, fine_objects=(body,), observe_raw_exchange=True,
+            observe_window_exchange=True))
+        masks = [channel_masks(pair.coarse_shape)[0] & ~pair.covered,
+                 (pair.fine_kind == ACTIVE) & ((host.flags.reshape(pair.fine_shape) & SOLID) == 0)]
+        fresh = (r.initial_state(pair.coarse_shape, (0.0, 0.0), 1.0, root=True),
+                 r.initial_state(pair.fine_shape, pair.fine_origin, 0.5))
+        print(json.dumps(dict(status="CONSTRUCTED", case="refined")), flush=True)
+        for repeat in range(a.repeats):
+            for variant in VARIANTS:
+                t.check_deadline(deadline)
+                backend.initialize(*fresh)
+                for _ in range(a.warmup_base_steps):
+                    backend.step(1)
+                warm = t.refined_health(backend, masks)
+                if repeat == 0 and variant == "A":
+                    print(json.dumps(dict(status="WARM_READY", case="refined")), flush=True)
+                t.sync_refined(backend)
+                start = time.perf_counter()
+                totals, scales = run_variant(backend, variant, n_intervals, deadline)
+                t.sync_refined(backend)
+                seconds = time.perf_counter() - start
+                end = t.refined_health(backend, masks)
+                values = [np.array(v, copy=True) for v in backend._download_raw().values]
+                if variant == "A" and repeat == 0:
+                    ref.update(values=values, totals=totals, scales=scales)
+                    parity, ratio = "identical", 0.0
+                else:
+                    parity = t.same_state(ref["values"], values) or "identical"
+                    ratio = compare_intervals(ref["totals"], ref["scales"], totals)
+                rows.append(dict(variant=variant, repeat=repeat, seconds=seconds,
+                                 intervals=n_intervals, population_parity_vs_A=parity,
+                                 interval_impulse_violation_ratio=ratio,
+                                 health_after_warmup=warm, health_after_window=end))
+                if parity != "identical" or not ratio <= 1:
+                    raise RuntimeError(f"{variant} repeat {repeat}: parity {parity}, impulse ratio {ratio}")
+    summary = {v: dict(median_seconds=statistics.median(x["seconds"] for x in rows if x["variant"] == v),
+                       seconds=[x["seconds"] for x in rows if x["variant"] == v]) for v in VARIANTS}
+    summary.update(W_over_A=summary["W"]["median_seconds"] / summary["A"]["median_seconds"],
+                   W1_over_A=summary["W1"]["median_seconds"] / summary["A"]["median_seconds"],
+                   interpretation="Measured comparative timings under GPU variability; W also "
+                                  "batches step calls, W1 isolates the read pattern. Not a "
+                                  "benchmark-throughput or causal claim.")
+    (a.output / "window_timing.json").write_text(json.dumps(dict(plan=plan, rows=rows, summary=summary),
+                                                           indent=2))
+    print(json.dumps(summary, indent=2), flush=True)
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())