kernels (K3): cell kinds in core/, capability guard first, self-contained block collision

- core/cell_kind.cuh holds KIND_ACTIVE/INNER/OUTER (was a constexpr in
  blocks.cu that step/block_channel.cu and the channel_root.cpp shim
  relied on or redefined); blocks.cu, block_channel.cu and the shim now
  include it.
- blocks.cu checks the lattice/model #error guard before the
  BLOCK_ROOT_BC guard and mode includes, so an unsupported build reports
  the real reason first.
- operators/block_collision.cuh includes what it uses (params, macro,
  SRT/MRT collision) instead of relying on blocks.cu's include order.
Preprocessor-only: same kernels, same statements, same arithmetic.

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

docs/supported_settings.md

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

src/CelerisLab/cuda/compiler_v2.py

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

src/CelerisLab/lbm/kernels/blocks.cu

@@ -1,5 +1,6 @@
 // CelerisLab/lbm/kernels/blocks.cu
-// One periodic level of fixed-size blocks (lbm/blocks.py layout). Separate
+// One fixed-size blocked level (lbm/blocks.py layout), specialized at compile
+// time for periodic/fine, slot/reservoir or channel stepping. Separate
 // compilation entry: NX, NY are the block window including halo, so index_f
 // and linear_index address one block; blockIdx.z selects the block, whose
 // populations start at b*NQ*TOTAL_CELLS (block-contiguous, Q-major inside).
@@ -12,6 +13,7 @@
 #include "core/params.cuh"
 #include "core/flags.cuh"
 #include "core/obs.cuh"
+#include "core/cell_kind.cuh"
 #include "operators/macro.cuh"
 #include "operators/collision_srt.cuh"
 #if COLLISION_MODEL == 2
@@ -19,19 +21,22 @@
 #endif
 #include "streaming/pull_double_buffer.cuh"
 #include "boundary/curved_boundary.cuh"
-#include "boundary/planar_velocity.cuh"
 
 #if DIM != 2 || NQ != 9 || (COLLISION_MODEL != 0 && COLLISION_MODEL != 2) || STORE_PRECISION != 0 || USE_LES || STREAMING_MODEL != 0
 #error "Blocks require FP32 D2Q9 SRT or MRT double buffers without LES"
 #endif
+#ifndef BLOCK_ROOT_BC
+#error "Blocks require a generated BLOCK_ROOT_BC mode"
+#elif BLOCK_ROOT_BC < 0 || BLOCK_ROOT_BC > 2
+#error "BLOCK_ROOT_BC must be 0, 1 or 2"
+#endif
 
-// Fixed-patch cell kinds, matching refinement.py ACTIVE/INNER_GHOST/OUTER_GHOST.
-constexpr unsigned char KIND_ACTIVE = 0u, KIND_INNER = 1u, KIND_OUTER = 2u;
-
-// Zero-initialized for ordinary blocks/fine levels. Set only for parent walls;
-// NX/NY remain local storage strides, never physical wall coordinates.
-struct BlockWall { unsigned int nx, ny, slot_width; float amplitude; };
-extern "C" { __constant__ BlockWall d_block_wall; }
+#if BLOCK_ROOT_BC == 1
+#include "boundary/block_wall.cuh"
+#endif
+#if BLOCK_ROOT_BC != 2
+#include "operators/block_collision.cuh"
+#endif
 
 // Halo refresh: block cell ids b*TOTAL_CELLS+c, sources are always core cells.
 extern "C" __global__ void BlockHaloCopy(
@@ -147,6 +152,7 @@ extern "C" __global__ void BlockAccumulateExchange(
     integral[k + 3] += h * h * (double)raw[k + 3];
 }
 
+#if BLOCK_ROOT_BC != 2
 // Pull/collide/store core cells only; halo cells of out are left untouched.
 // flags (block cell ids, nullable): solid cores store NaN and never collide.
 // kind (block cell ids, nullable = all ACTIVE): OUTER skips the pull and
@@ -167,34 +173,18 @@ extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int hal
         compute_neighbors(x, y, j);  // Core neighbors stay inside the window: no wrap.
         stream_pull_load(k, f, in + base, j);
         if (cell_kind == KIND_ACTIVE) {
-            if (d_block_wall.nx != 0) {
-                unsigned int gx, gy;
-                block_global_xy(x, y, halo, d_block_wall.nx, gx, gy);
-                if (gy == 1)
-                    reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity(
-                        gx, d_block_wall.nx, d_block_wall.slot_width, d_block_wall.amplitude));
-                else if (gy == d_block_wall.ny - 2) reconstruct_top_reservoir(f, RHO);
-            }
-#if COLLISION_MODEL == 2
-            collide_mrt_unforced(f);  // SRT path below unchanged.
-#else
-            float rho, ux, uy, feq[NQ];
-            compute_rho_u(f, rho, ux, uy);
-#if USE_DDF_SHIFTING
-            // Preserve the density perturbation before adding the unit reference.
-            const float drho = f[0]+f[1]+f[2]+f[3]+f[4]+f[5]+f[6]+f[7]+f[8];
-            const float jx = f[1]-f[2]+f[5]-f[6]+f[7]-f[8];
-            const float jy = f[3]-f[4]+f[5]-f[6]-f[7]+f[8];
-            compute_feq_from_momentum(drho, jx, jy, ux, uy, feq);
-#else
-            compute_feq(rho, ux, uy, feq);
-#endif
-            collide_srt_no_force(f, feq, d_params.omega);
+#if BLOCK_ROOT_BC == 1
+            block_wall(f, x, y, halo);
 #endif
+            block_collide(f);
         }
     }
     stream_pull_store(k, f, out + base);
 }
 
-// Root channel closure; fine and periodic levels keep BlockStep.
+#endif
+
+// Root channel is a replacement full step, selected once by the host.
+#if BLOCK_ROOT_BC == 2
 #include "step/block_channel.cu"
+#endif

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

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

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

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

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

@@ -0,0 +1,10 @@
+// CelerisLab – core/cell_kind.cuh
+// Fixed-patch cell kinds stored per block cell, matching lbm/refinement.py
+// ACTIVE/INNER_GHOST/OUTER_GHOST: OUTER is invalid (stores NaN, no pull),
+// INNER streams without collision, ACTIVE pulls and collides.
+#ifndef CELERIS_CORE_CELL_KIND_CUH
+#define CELERIS_CORE_CELL_KIND_CUH
+
+constexpr unsigned char KIND_ACTIVE = 0u, KIND_INNER = 1u, KIND_OUTER = 2u;
+
+#endif

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

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

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

@@ -1,5 +1,10 @@
+#if BLOCK_ROOT_BC != 2
+#error "BlockChannelStep requires BLOCK_ROOT_BC == 2"
+#endif
+
 // Root-only blocked counterpart of OneStep's D2Q9/SRT channel closure.
 // NX/NY are LOCAL strides; all physical boundary decisions use global coordinates.
+#include "../core/cell_kind.cuh"
 #include "../boundary/inlet/common.cuh"
 #include "../boundary/inlet/zou_he_local.cuh"
 #include "../boundary/outlet/pressure_neq.cuh"

src/CelerisLab/lbm/refinement_stepper.py

@@ -362,7 +362,10 @@ class FixedPatchStepper:
             raise RuntimeError("level omega/tau mismatch")
         host = self._boundary_hosts[index]
         n_objects = len(host.objects) if name == "PatchFineStep" and host is not None else 0
-        build = compiler.generate_config(config, n_objects=n_objects)
+        block_root_bc = 0
+        if index == 0 and layout is not None:
+            block_root_bc = 2 if self._channel else 1 if self._wall_aware else 0
+        build = compiler.generate_config(config, n_objects=n_objects, block_root_bc=block_root_bc)
         try:
             ptx = compiler.compile_kernel(self.ctx.sm_arch, source=source, build_dir=build)
             module = compiler.load_module(ptx)

tests/integration/test_checkpoint_bodies_gpu.py

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

tests/unit/cpp/channel_root.cpp

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

tests/unit/historical_tree.py

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

tests/unit/test_build_contracts.py

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

tests/unit/test_channel_root.py

@@ -97,5 +97,6 @@ def test_evolved_ghost_snapshot_and_physical_download():
         with pytest.raises(ValueError, match="finite"):
             backend.initialize(*arrays)
         assert backend._levels[0].kernel == host.compiler.load_module.return_value.get_function("BlockChannelStep")
+        assert [call.kwargs["block_root_bc"] for call in host.compiler.generate_config.call_args_list] == [2, 0]
     finally:
         backend.close()

tests/unit/test_jet_circle_cv.py

@@ -8,6 +8,7 @@ import numpy as np
 ROOT = Path(__file__).resolve().parents[2]
 REPO = ROOT
 sys.path.insert(0, str(ROOT / "tests/validation"))
+sys.path.insert(0, str(ROOT / "tests/unit"))
 sys.path.insert(0, str(REPO / "src"))
 from run_jet_circle_cv import (
     CX,
@@ -19,6 +20,7 @@ from run_jet_circle_cv import (
     cv_from_macros,
     aperture_probe,
     measured_aperture_flux,
+    PIN,
     verify,
 )
 
@@ -166,7 +168,13 @@ class Host(unittest.TestCase):
         self.assertEqual(a["steps"] * 4, b["steps"])
 
     def test_source_pin(self):
-        verify(REPO)
+        from tempfile import TemporaryDirectory
+        from historical_tree import historical_tree
+
+        with self.assertRaisesRegex(RuntimeError, "Source hash mismatch"):
+            verify(REPO)
+        with TemporaryDirectory() as directory:
+            verify(historical_tree(PIN, directory))
 
 
 if __name__ == "__main__":

tests/unit/test_les_flow_runner.py

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

tests/unit/test_re100.py

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

tests/unit/test_refined_dfg.py

@@ -7,7 +7,10 @@ import numpy as np
 import unittest
 
 sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "validation"))
+sys.path.insert(0, str(Path(__file__).resolve().parent))
 import run_refined_dfg as r
+from historical_tree import historical_tree
+from tempfile import TemporaryDirectory
 
 h, periodic = r.helpers(r.DEFAULT_SOURCE)
 
@@ -17,11 +20,24 @@ def test_source_pin():
     assert historical["commit"].startswith("cd3a657")
     mrt_history = json.loads((r.HERE / "data/refined_dfg/source_manifest_mrt.json").read_text())
     assert mrt_history["commit"].startswith("3ea131e")
-    assert r.verify_source(r.DEFAULT_SOURCE, "MRT", "window")["commit"].startswith("a233c90")
+    with unittest.TestCase().assertRaisesRegex(RuntimeError, "Source mismatch"):
+        r.verify_source(r.DEFAULT_SOURCE, "MRT", "window")
     with unittest.TestCase().assertRaisesRegex(RuntimeError, "Source mismatch"):
         r.verify_source(r.DEFAULT_SOURCE, "MRT")
     with unittest.TestCase().assertRaisesRegex(RuntimeError, "Source mismatch"):
         r.verify_source(r.DEFAULT_SOURCE, "SRT")
+    with TemporaryDirectory() as directory:
+        for collision, mode, prefix in (
+            ("SRT", "last_base_step", "cd3a657"),
+            ("MRT", "last_base_step", "3ea131e"),
+            ("MRT", "window", "a233c90"),
+        ):
+            name = {"SRT": "source_manifest.json", "MRT": "source_manifest_mrt.json"}[collision]
+            name = "source_manifest_window.json" if mode == "window" else name
+            commit = json.loads((r.HERE / "data/refined_dfg" / name).read_text())["commit"]
+            assert commit.startswith(prefix)
+            root = historical_tree(commit, Path(directory) / commit)
+            assert r.verify_source(root, collision, mode)["commit"] == commit
 
 
 def test_matched_units_and_planes():

tests/unit/test_refinement_mrt.py

@@ -24,11 +24,14 @@ class RefinementMRTHost(unittest.TestCase):
         """MRT is an additive compile-time branch; SRT statements and other guards remain."""
         blocks = (KERNELS / "blocks.cu").read_text()
         channel = (KERNELS / "step/block_channel.cu").read_text()
+        block_collision = (KERNELS / "operators/block_collision.cuh").read_text()
+        self.assertIn('#include "operators/block_collision.cuh"', blocks)
+        self.assertEqual(blocks.count("block_collide(f);"), 1)
         refinement = (KERNELS / "refinement.cu").read_text()
         for text in ("compute_feq_from_momentum(drho, jx, jy, ux, uy, feq);",
                      "compute_feq(rho, ux, uy, feq);",
                      "collide_srt_no_force(f, feq, d_params.omega);"):
-            self.assertIn(text, blocks)
+            self.assertIn(text, block_collision)
         self.assertIn("collide_srt(f, feq, forcing, omega);", channel)
         self.assertIn("collide_srt_no_force(f, feq, d_params.omega);", refinement)
         for source in (blocks, refinement):
@@ -40,10 +43,11 @@ class RefinementMRTHost(unittest.TestCase):
         # One shared unforced MRT body; each blocked/patch path calls it exactly once.
         mrt = (KERNELS / "operators/collision_mrt.cuh").read_text()
         self.assertEqual(mrt.count("void collide_mrt_unforced("), 1)
-        for source in (blocks, channel, refinement):
+        for source in (blocks + "\n" + block_collision, channel, refinement):
             self.assertEqual(source.count("collide_mrt_unforced(f);"), 1)
             self.assertNotIn("collide_mrt_no_force(", source)
             self.assertNotIn("compute_rho_u(f, rho, ux, uy);\n            collide_mrt", source)
+            self.assertNotRegex(source, r"compute_rho_u\(f, rho, ux, uy\);\s*collide_mrt")
 
 
 if __name__ == "__main__": unittest.main()

tests/unit/test_validation_source_provenance.py

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

tests/validation/multigrid/test_channel_root_gpu.py

@@ -27,7 +27,7 @@ def state(ny, nx):
 
 def compile_module(cfg, context, source):
     cfg.validate()
-    build = compiler.generate_config(cfg)
+    build = compiler.generate_config(cfg, block_root_bc=2 if source == "blocks.cu" else 0)
     try:
         module = compiler.load_module(compiler.compile_kernel(context.sm_arch, source=source, build_dir=build))
     finally:

tests/validation/multigrid/test_refinement_blocked_base.py

@@ -223,6 +223,8 @@ class BlockedBaseCPU(unittest.TestCase):
                     self.assertIs(collide[4], coarse.kind)
                 walls = [a for ptr, a in m["uploads"] if isinstance(ptr, str) and ptr == "d_block_wall"]
                 self.assertEqual(len(walls), 1)
+                self.assertEqual([call.kwargs["block_root_bc"] for call in
+                                  compiler.generate_config.call_args_list], [1, 0])
                 width, height, slot, speed = struct.unpack("=IIIf", walls[0])
                 self.assertEqual((width, height, slot), (NX, NY, 8))
                 self.assertAlmostEqual(speed, .001)

tests/validation/taylor_green/test_projected_pairs.py

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