step(n) docstring: raw-exchange ledger holds the final interval only

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

README.md

@@ -123,7 +123,9 @@ with CudaContext(0) as ctx, FixedPatchStepper(cfg, (8, 6, 24, 18), context=ctx,
 ```
 
 Scope and limits:
-- D2Q9, SRT, FP32, double buffer, unshifted, no LES or forcing; fixed 2:1 patch(es), two levels (or one nested third).
+- The examples below use D2Q9, SRT, FP32, double buffer, unshifted, no LES or
+  forcing. Fixed 2:1 hierarchies can use level lists of arbitrary depth; the
+  legacy two/three-level forms are shown below.
 - `coarse_boundary="periodic"` (default) or `"slot_reservoir"` (bottom slot, top
   reservoir; patch needs `y0 >= 4`, `y1 <= ny-4`). Periodic mode ignores the wall/slot config fields.
 - Object centers/radii use the coarse **cell-edge** frame (cell `I` spans `[I, I+1)`),
@@ -135,16 +137,20 @@ Scope and limits:
   divide the WHOLE fine allocation, i.e. `2*(y1-y0)+8` by `2*(x1-x0)+8` including the
   four-cell ghost band. It gives results identical to dense storage and is
   infrastructure for future block AMR, not a speedup.
-- Several separated patches: pass a sequence of `(x0, y0, x1, y1)`
-  rectangles as `bounds`, at least 4 coarse cells apart, with both `fine_block`
-  and `base_block`. The fine level is then one level-wide grid `(2*ny, 2*nx)`:
+- Several regions: pass a sequence of `(x0, y0, x1, y1)` rectangles as
+  `bounds`, with both `fine_block` and `base_block`. Transfers are built once
+  from their union; adjacent, overlapping and concave unions are accepted when
+  their actual dependencies close. There is no fixed rectangle-gap rule.
+  The fine level is then one level-wide grid `(2*ny, 2*nx)`:
   `run.patch.fine_shape` gives it, `initialize`/`download` use it, and only the
-  patch windows (active region plus 4-cell ghost band, origin `(2*y0-4, 2*x0-4)`)
-  are stored; each window must align with `fine_block`. Order of rectangles does
-  not matter. `fine_objects`/`jets` work as for one patch: each circle (same
-  coarse cell-edge frame) must lie wholly inside one patch's active region with
+  grid-aligned tiles intersecting the union plus its four-cell ghost band are stored.
+  Block sizes must divide the level-wide grid; rectangle edges need not align
+  with tiles. Order of rectangles does not matter. `fine_objects`/`jets` work
+  as for one patch: each circle (same coarse cell-edge frame) must lie wholly
+  inside the union's active region with
   its boundary stencil clear of transfer cells; a jet may not select links of two
-  different bodies, even across patches. No wall-touching patches or nesting yet.
+  different bodies, even across regions. Wall-touching patches remain unsupported.
+  For nesting, use the level lists or legacy child forms below.
 
   ```python
   cfg = LBMConfig(nx=32, ny=16, viscosity=.1)
@@ -158,7 +164,7 @@ Scope and limits:
       coarse_g, fine_g = run.download()  # coarse: NaN on every covered cell;
                                          # fine: NaN on every non-ACTIVE cell
   ```
-- Three levels (one nested chain): add `child=(x0, y0, x1, y1)` in
+- Legacy three-level scalar chain: add `child=(x0, y0, x1, y1)` in
   **L1 window cells** (the patch's fine grid including its 4-cell ghost band) and
   `child_block`, with `fine_block` and `base_block`. L2 must sit inside L1's
   ACTIVE cells, away from L1's ghost band and transfer footprint (checked before
@@ -192,14 +198,30 @@ Scope and limits:
       run.step(4)
       coarse_g, l1_g, l2_g = run.download()  # solids/covered/nonleaves are NaN
   ```
-- Not available: per-link surface exchange/force readout in blocked mode,
-  checkpoints, moving bodies, L1 bodies/Jets in a three-level chain,
-  several patches per level with nesting, dynamic refinement, `Simulation`
-  integration. Validation so far is storage/operator equivalence and short runs,
-  not accuracy, convergence or long-time stability. The new L2-body GPU gates
-  (layouts, masks, amplitude reversal, legacy equivalence) and independent
-  every-substep numerical replay remain required acceptance checks; CPU tests
-  alone do not establish GPU numerical correctness.
+- Level lists: `refinements[l]` describes level `l+1` in allocated level-`l`
+  integer cells; use a rectangle list for a union, or a scalar rectangle for a
+  single window (a scalar child requires a scalar parent). `level_blocks`,
+  `level_objects` and `level_jets` each have one entry per level, including L0.
+  These lists are mutually exclusive with the legacy patch keywords. Three or
+  more levels require all-blocked storage and bodies/Jets only on the finest
+  level. Every adjacent pair must preserve inherited ACTIVE coverage, valid
+  donors/receivers, surviving pulls and the upstream transfer footprint.
+  Level `l` takes `2**l` steps per base interval. This describes implemented
+  topology and scheduling, not accuracy or convergence qualification.
+- Blocked body levels can enable `observe_raw_exchange=True` and call
+  `read_raw_exchange(body_id, level=...)` after a completed base interval.
+  It reduces unique cut-link observations into wall/aperture/total force,
+  torque and mass exchange in base-lattice units per unit depth, for the last
+  base interval (default) or `window="last_substep"`.
+  `observe_window_exchange=True`, alongside `observe_raw_exchange=True`, enables
+  `start_exchange_window()` / `read_exchange_window(body_id, level=...)` across
+  completed base intervals. These are raw boundary-reaction diagnostics, not a general
+  stress or nozzle-load accuracy claim; dense body layouts are unsupported.
+- Not available: checkpoints, moving bodies, L1 bodies/Jets in a three-level
+  chain, automatic dynamic refinement, `Simulation` integration. Qualification
+  depends on the tested configuration; see [the validation scope](docs/validation/2d_scope.md).
+  CPU topology/schedule tests alone do not establish GPU numerical correctness,
+  general convergence or long-time stability.
 
 ## Installation
 

docs/development/projected_pair_explosion.md

@@ -0,0 +1,57 @@
+# Experimental projected G explosion
+
+`FixedPatchStepper(..., explosion="projected_linear")` changes only the existing
+G-to-G explosion reconstruction. The omitted/default `"constant"` selector retains
+the same kernel, launch arguments and device allocations. Projected mode adds two
+compact tables per adjacent pair; recursion, collision, BFL, restriction, buffer
+ownership and synchronized publication are unchanged. `remap_child` preserves it.
+
+Runtime eligibility is periodic-root, fully blocked D2Q9/SRT/FP32, body/Jet-free,
+without LES; the fixed-patch backend has no forcing interface. Shifted and unshifted
+storage use the same reconstruction. One generic builder handles scalar windows,
+unique union rows and all adjacent hierarchy pairs without depth/rectangle caps.
+Donors are original explosion donors intersected with owned, uncovered ACTIVE
+cells. An axis uses centered differences only if both neighbors are legal;
+otherwise both addresses repeat the center. Nested layout periodicity is explicit.
+No covered cells, holes, inactive cells or stale halo copies enter the stencil.
+
+The 33-line kernel is unchanged from our `519473a` experiment and retains its
+mathematical attribution. Host tests cover union decomposition, concavity, narrow
+gaps, grazing packet paths, offsets, sparse/nested owner addresses, three/four-level
+source timing, unchanged constant dispatch, remapping and shifted/unshifted actual
+kernel arithmetic compiled as host C++. These are not CUDA or flow validation.
+
+## Reviewed bounded evidence
+
+The `1aa8110` CUDA gate, with production exactly `a135252`, passed exact-dev
+omitted/default and explicit-constant equivalence, per-population reconstruction,
+scalar/union/four-level scheduling, alternate layouts, and changed remapping in
+both storage modes. Operator maximum error was 2.98023224e-8; the eight-step
+concave replay maximum was 1.58037614e-7. Whole-step mass/momentum remain
+finite diagnostics, not an FP32 conservation certificate. Reviewed evidence:
+`7aa66ff0542175092753ea099f40b903ece0527e`.
+
+The fresh four-case `2cf9853` N128/256 Taylor–Green repeat retained all `11901a7`
+settings and gates. All 20 main scalar rows and 260 early-pressure samples matched
+the older receipts exactly; this is not field-array equality. Projected worst
+velocity was 0.001995299260844 versus 0.002 (only 0.235% margin), pressure
+0.012157174904932 and energy 0.001184362645162. Constant velocity
+0.004418076615223 failed its gate, so the aggregate verdict remains false.
+The projected early sampled pressure peak was 0.021621453573334; it remains a
+descriptive sampled maximum. Reviewed evidence: `90849a55cad4a0da9b72747da7bfe6eb75962a44`.
+
+The old N64 projected failures remain relevant. These results support only this
+bounded plumbing gate and short periodic Re100 smooth-decay experiment. They do
+not establish union/multilevel physics accuracy, convergence, long-time stability,
+body/Jet, boundary, forcing, MRT or LES suitability. Constant remains the default.
+
+## Separate LES source profile
+
+The two older LES source manifests and the historical default are unchanged and
+correctly reject this different source tree. The existing LES runner now offers
+`--experiment a135252`, pinning all 101 actual committed corrected source files.
+Its explicit baseline-header override describes only a synthetic isolated reverse
+control, never the committed a135252 header. The unchanged corrected-only uniform
+LES regression (Cs=0/.16/1) is the remaining integration gate; no LES result is
+inherited from the projected evidence. The separate historical jet-CV `channel.py`
+pin mismatch predates this port and is not addressed here.

docs/fixed_three_level_multipatch.md

@@ -1,8 +1,11 @@
-# Fixed three-level separated patches
+# Fixed three-level rectangle unions
 
-`FixedPatchStepper` supports exactly three static, blocked levels with separated
-rectangle sets on either adjacent pair. This is not general AMR: there is no
-per-region timestep, moving geometry, dynamic refinement or L1 body support.
+This page describes the legacy three-level entry point of `FixedPatchStepper`,
+with rectangle unions on either adjacent pair. The `refinements` level-list
+entry point also supports arbitrary depth; three or more levels require
+all-blocked storage and bodies/Jets only on the finest level. This is not
+general AMR: there is no per-region timestep, moving geometry, automatic
+dynamic refinement or L1 body support in the three-level form.
 All regions share one current/spare buffer pair per level. The existing schedule
 advances L0 once, L1 twice and L2 four times per base step, with aggregate transfers.
 No CUDA kernel changes are involved.
@@ -30,12 +33,15 @@ are physically equivalent to parent set `[(4,4,12,12)]` with child set
 `[(8,8,16,16)]`. Reusing the same numbers after changing the parent frame can
 move the physical region.
 
-Both adjacent pairs require a Chebyshev gap of at least four in their own
-parent-cell units. The whole child union must preserve inherited ACTIVE
-coverage, uncovered ACTIVE donors/receivers, surviving pulls, and the upstream
+Both adjacent pairs are built from their unions, without a fixed rectangle-gap
+rule. Adjacent, overlapping and concave unions are accepted only when their
+actual transfer dependencies close. The whole child union must preserve
+inherited ACTIVE coverage, uncovered ACTIVE donors/receivers, surviving pulls, and the upstream
 restriction footprint. Sparse holes remain unowned; inherited INNER remains
-stream-only and OUTER remains invalid. Block sizes must align with every
-support edge, including its four-cell ghost band.
+stream-only and OUTER remains invalid. Block sizes must divide their allocated
+level grids. Sparse storage owns every grid-aligned tile intersecting the union
+plus its four-cell ghost band; rectangle and support edges need not align with
+tile edges.
 
 ## Example
 
@@ -65,6 +71,12 @@ CPU geometry, table, pre-CUDA rejection, schedule and cleanup tests:
 PYTHONPATH=src python -m unittest discover -s tests/unit -p 'test_refinement*.py' -v
 ```
 
+`tests/unit/test_union_refinement.py` covers union dependencies and tile
+ownership; `tests/unit/test_hierarchy.py` covers composed geometry through four
+levels. `tests/unit/test_raw_exchange.py` checks the four-level host schedule
+with CUDA mocks. These establish topology/control-flow coverage, not numerical
+qualification at every depth or interface shape.
+
 Explicit GPU gates are in
 `tests/validation/multigrid/test_refinement_three_level_multipatch_gpu.py`.
 They cover physical-frame equivalence, twins, asymmetric refinement, multiple

docs/synchronized_leaf_remap.md

@@ -1,11 +1,15 @@
 # Synchronized leaf G remap and manual continuation
 
 This is a CPU state-transfer component, **not a runtime adaptive solver**.
-It changes only a separated L2 rectangle set in a three-level, body-free,
-periodic hierarchy. L0/L1 topology, relaxation parameters, direction order,
+It changes only an L2 rectangle union in a three-level periodic hierarchy.
+The body-free case is shown first; stationary L2 bodies/Jets follow the bounded
+support contract below. L0/L1 topology, relaxation parameters, direction order,
 physical frame and completed time remain unchanged. Existing geometry checks
-validate child support/nesting and separation. The CPU helper itself adds no CUDA, scheduler, allocator, boundary handling,
-collision or ghost refill. A bounded runtime adapter is described below.
+validate child support/nesting and actual transfer dependencies, without a fixed
+rectangle-gap rule. The three-level remap limit is unchanged by the stepper's
+arbitrary-depth level-list support. The CPU helper itself adds no CUDA,
+scheduler, allocator, boundary handling, collision or ghost refill. A bounded
+runtime adapter is described below.
 
 ```python
 from CelerisLab.lbm.refinement import FixedPatchSet2D
@@ -29,11 +33,13 @@ origin is retained. Physical lower edges are `origin + spacing*(x,y)` on L0,
 `origin + spacing*(parent.fine_origin + (x,y)/4)` on L2.
 
 Factory options `coarse_boundary` and `objects` must declare `"periodic"` and
-an empty body collection (all levels). LBMConfig's channel settings do not
-select this operator's boundary mode. Geometry metadata is factory-created;
+an empty collection; L2 bodies use `child_objects` as described below.
+LBMConfig's channel settings do not select this operator's boundary mode.
+Geometry metadata is factory-created;
 manual mutation/forging of metadata is unsupported. Caller certifies the
-actual state has no bodies/forcing, that all device work and old readers are
-quiescent, and that inputs are the **current owned-leaf postcollision G**.
+actual state matches the declared boundary and has no forcing, that all device
+work and old readers are quiescent, and that inputs are the **current owned-leaf
+postcollision G**.
 Count/phase stamps check consistency, not freshness or active-buffer identity.
 
 `current_g` is three NumPy Q-major arrays `(9,ny,nx)`. For blocked inputs,
@@ -73,9 +79,10 @@ PYTHONPATH=src python -m unittest discover -s tests/unit -p 'test_synchronized_r
 
 `replacement = current.remap_child(rectangles)` returns a new ready
 `FixedPatchStepper` at the **same absolute step_count**. It supports only
-body-free periodic three-level runs whose existing L2 uses
-`child_rectangles_in_parent`. Scalar or union L1 parents are supported; scalar
-`child=` is deliberately rejected. L0/L1, relaxation and block sizes stay fixed.
+periodic three-level runs with no L0/L1 bodies and an existing L2
+`FixedPatchSet2D`; stationary L2 bodies/Jets follow the support contract below.
+Scalar or union L1 parents are supported; scalar `child=` is deliberately
+rejected. L0/L1, relaxation and block sizes stay fixed.
 No tagger, scheduler, in-place replacement, or general rollback is provided.
 
 The method holds the stepper lock, validates geometry, synchronizes its stream,
@@ -113,13 +120,14 @@ Per-population conservation is not an advancement target under collision.
 ### Bounded pinned-body extension
 
 `replacement = current.remap_child(rectangles, child_objects=original_objects)`
-also accepts stationary Circles in separated L2 rectangles. Every body’s
-**complete containing rectangle must remain identical**. Several bodies may
-share one pinned rectangle; there need not be any fluid-only rectangles.
-Separated fluid-only rectangles may be added, removed or moved. Existing
-nesting, four-cell gap, full-circle containment, reverse cut-link coverage and two-hop BFL
-transfer-footprint checks still apply. This is manual remapping, not adaptive
-body tracking or runtime body editing.
+also accepts stationary Circles inside the L2 union. Each body's **physical
+boundary, complete BFL support and current per-link velocity must remain
+identical**. Rectangles are not pinned: they may be split, merged or resized,
+and other fluid coverage may change, including within the same connected union,
+provided that physical support is unchanged. Existing nesting, union transfer
+dependencies, full-circle containment, reverse cut-link
+coverage and two-hop BFL transfer-footprint checks still apply. This is manual
+remapping, not adaptive body tracking or runtime body editing.
 
 Supply all original root edge-frame definitions in their construction order.
 Do not pass `current.child_boundary.objects`: those are projected fine-local copies.
@@ -128,8 +136,8 @@ state, remapping reprojects the supplied definition onto the old topology and
 compares exact ordered circle centers/radii, flags, physical cut links, body
 IDs, q, lever arms, hits, BFL support,
 legacy SurfaceActuation specification and current per-link velocity against
-the existing boundary. Omission, changed geometry/actuation and changed body
-rectangle fail without device reads or allocations; reordered bodies fail
+the existing boundary. Omission, changed geometry/actuation or changed physical
+boundary/BFL support fail without device reads or allocations; reordered bodies fail
 too. Root object IDs are still assigned by list position, as in the constructor.
 
 Jets are taken from the live L2 boundary, including the most recent successful
@@ -141,12 +149,16 @@ boundary and velocities must match exactly.
 typed `FineBoundary2D`; it derives L2 fluid leaves as ACTIVE minus SOLID. No
 caller mask is accepted. Solids, ghosts and holes are never donors and remain
 NaN, physical-fluid overlap is exact, and exposed L1 still restricts all four
-old fluid children. The helper checks the old/new body pins and boundary too.
-Its immutable `pinned_rectangles` is a canonical, deduplicated tuple, empty for body-free layouts; rectangle input
-order does not affect it. Scalar-parent window offsets remain part of the
-coordinate conversion.
+old fluid children. The helper checks old/new physical boundary and velocity
+equality too. Its immutable `pinned_rectangles` is legacy containment metadata,
+not a remap eligibility constraint; it may be empty for a seam-spanning body.
+Scalar-parent window offsets remain part of the coordinate conversion.
 
-Focused CPU mocks: `tests/unit/test_pinned_body_remap.py`. GPU declarations:
+Focused CPU mocks: `tests/unit/test_pinned_body_remap.py`.
+`tests/unit/test_union_body_support.py` also checks seam-spanning circles,
+rectangle repartition/resize, unchanged physical support and live-Jet
+preservation. These host checks do not establish GPU evolution accuracy.
+GPU declarations:
 `tests/validation/multigrid/test_pinned_body_remap_gpu.py`, covering no-op
 next-step bitwise identity, moved-fluid A→B→A, current-amplitude preservation,
 physical boundary/velocity equality, source independence and absolute counts.

src/CelerisLab/lbm/kernels/blocks.cu

@@ -63,6 +63,39 @@ extern "C" __global__ void BlockExplode(
         to[index_f(t % TOTAL_CELLS, q)] = from[q * coarse_cells + d % coarse_cells];
 }
 
+// Experimental Schornbaum--Ruede (arXiv:1508.07982, Algorithm 2)
+// and Chen et al. (doi:10.1016/j.physa.2005.09.036) population interpolation.
+// Independent implementation of the mathematical stencil; no upstream source copied.
+// Stencil cells are owner cores; a disabled axis repeats its center twice.
+extern "C" __global__ void BlockExplodeProjected(
+    const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
+    unsigned int count, unsigned long coarse_cells,
+    const uint64_t* neighbors, const float* offsets)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const uint64_t t = target[row], d = donor[row];
+    fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
+    for (unsigned int q = 0; q < NQ; ++q) {
+        float v[4];
+        for (unsigned int a = 0; a < 4; ++a) {
+            const uint64_t n = neighbors[4ul * row + a];
+            v[a] = coarse[(n / coarse_cells) * NQ * coarse_cells
+                          + q * coarse_cells + n % coarse_cells];
+        }
+        float gx = .5f * (v[1] - v[0]), gy = .5f * (v[3] - v[2]);
+        const float ex = d_cx[q], ey = d_cy[q];
+        if (q != 0) {
+            const float along = (ex * gx + ey * gy) / (ex * ex + ey * ey);
+            gx -= ex * along;
+            gy -= ey * along;
+        }
+        to[index_f(t % TOTAL_CELLS, q)] =
+            coarse[(d / coarse_cells) * NQ * coarse_cells + q * coarse_cells + d % coarse_cells]
+            + offsets[2ul * row] * gx + offsets[2ul * row + 1] * gy;
+    }
+}
+
 // Targets are receiver-local SOURCE slots for the immediately following pull.
 // Keep child summation order identical to PatchCoarseStep; no collision here.
 extern "C" __global__ void BlockCoalesce(

src/CelerisLab/lbm/refinement_stepper.py

@@ -22,7 +22,7 @@ from .block_stepper import MAX_GRID_DIM
 from .hierarchy import build_hierarchy, is_rectangle
 from .refinement import ACTIVE, HALO, FixedPatch2D, FixedPatchSet2D, prepare_fine_boundary
 from .refinement_tables import (blocked_base_tables, blocked_fine_tables,  # noqa: F401 (re-export)
-                                blocked_kind, _prepare_blocked_boundary,
+                                blocked_kind, projected_explosion_tables, _prepare_blocked_boundary,
                                 BLOCKED_BOUNDARY_DEVICE_COLUMNS)
 from .descriptors import SOLID
 from .raw_exchange import summarize_exchange
@@ -41,6 +41,8 @@ class _Pair:
     n_ghosts: int
     n_rows: int
     parent_cells: int    # parent block window cells (explode donor stride)
+    neighbors: object = None  # projected-only owner-core cells; constant allocates neither
+    offsets: object = None
 
 
 @dataclass
@@ -206,6 +208,12 @@ class FixedPatchStepper:
     halo -> BFL -> step four times per base step. Runtime child Jet changes
     require ``set_jet_amplitude(index, amplitude, level=2)``.
 
+    ``explosion="projected_linear"`` selects experimental transverse linear
+    reconstruction on each adjacent pair. Default "constant" is unchanged.
+    Projected mode requires periodic root, all-blocked SRT levels without
+    bodies/Jets or LES; this backend supplies no forcing. Geometry is generic,
+    but union/multilevel GPU and matched-physics qualification remain pending.
+
     Level lists (any depth): ``refinements[l]`` gives level l+1 in level-l
     integer cells (scalar window or rectangle list, see lbm.hierarchy);
     ``level_blocks``/``level_objects``/``level_jets`` hold one entry per level,
@@ -219,7 +227,7 @@ 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_window_exchange=False):
+                 observe_raw_exchange=False, observe_window_exchange=False, explosion="constant"):
         if HALO != 4:
             raise RuntimeError("refinement.cu requires a four-cell fine halo")
         if coarse_boundary not in ("periodic", "slot_reservoir", "channel"):
@@ -248,6 +256,15 @@ 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
+        if explosion not in ("constant", "projected_linear"):
+            raise ValueError("explosion must be constant or projected_linear")
+        if explosion == "projected_linear" and (
+                coarse_boundary != "periodic" or any(b is None for b in level_blocks)
+                or any(level_objects) or any(level_jets)
+                or cfg.collision != "SRT" or cfg.les_enabled):
+            raise ValueError("projected explosion requires periodic all-blocked SRT levels "
+                             "without bodies/Jets or LES")
+        self._explosion = explosion
         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")
@@ -479,6 +496,11 @@ class FixedPatchStepper:
                        self._upload(children), len(ghosts), len(targets),
                        int(np.prod(pair.coarse_shape if parent.layout is None
                                    else parent.layout.local_shape)))
+        if self._explosion == "projected_linear":
+            neighbors, offsets = projected_explosion_tables(
+                pair, parent.layout, self.hierarchy.kinds[index])
+            result.explode = child.module.get_function("BlockExplodeProjected")
+            result.neighbors, result.offsets = self._upload(neighbors), self._upload(offsets)
         if child.layout is not None:
             child.kind = self._upload(blocked_kind(child.layout, self.hierarchy.kinds[index + 1]))
         if index == 0 and base is not None:
@@ -653,7 +675,7 @@ class FixedPatchStepper:
                 child_rectangles_in_parent=new.child.bounds,
                 child_objects=child_objects, child_jets=jets,
                 observe_raw_exchange=self._observe_raw_exchange,
-                observe_window_exchange=self._observe_window)
+                observe_window_exchange=self._observe_window, explosion=self._explosion)
             try:
                 if self._ddf_shifting:
                     candidate._initialize_raw(_StoredPopulations(transferred.g, transferred.shifted))
@@ -746,6 +768,7 @@ class FixedPatchStepper:
         pair, level = self._pairs[index - 1], self._levels[index]
         pair.explode(parent_buffer, level.current, pair.ghosts, pair.donors, np.uint32(pair.n_ghosts),
                      np.uint64(pair.parent_cells),
+                     *((pair.neighbors, pair.offsets) if pair.neighbors is not None else ()),
                      grid=((pair.n_ghosts + self._block[0] - 1) // self._block[0], 1), **kw)
         self._hop(index, level.current, level.spare, **kw)
         self._hop(index, level.spare, level.current, **kw)
@@ -799,27 +822,41 @@ class FixedPatchStepper:
                             grid=((pairs + self._block[0] - 1) // self._block[0], 1), **kw)
 
     def step(self, n=1):
-        """Advance G→G by n base intervals: level l takes 2**l hops per interval."""
+        """Advance G→G by n base intervals: level l takes 2**l hops per interval.
+
+        All n intervals are enqueued in order on the one stream, and the host waits
+        once, before returning; step(0) launches and waits for nothing. On return
+        the device work is complete, step_count and an open window cover all n
+        intervals, and the raw-exchange ledger holds the final interval only.
+
+        Asynchronous device errors surface at that single wait, so they are not
+        attributed to the interval that caused them; call step(1) repeatedly to
+        localize one (same launches, same bits). On any failure the instance is
+        failed, the exchange ledger and window are invalid, and step_count keeps
+        its value from entry to the failing call.
+        """
         if isinstance(n, bool) or not isinstance(n, Integral) or n < 0:
             raise ValueError("n must be a nonnegative integer")
         with self._lock, self.ctx.activate():
             self._require("ready")
+            if n == 0:
+                return
             root = self._levels[0]
             kw = dict(block=self._block, stream=self._stream)
             try:
                 self._status = "advancing"
+                self._exchange_valid = False
                 for _ in range(n):
-                    if self._observe_raw_exchange:
-                        self._exchange_valid = False
+                    if self._observe_raw_exchange:  # stream-ordered: holds the last interval only
                         for b in self._level_boundaries:
                             if b is not None:
                                 cuda.memset_d8_async(b.interval_integral, 0,
                                                      len(b.tables["cell"]) * 32, self._stream)
                     self._hop(0, root.current, root.spare, **kw)
-                    self._stream.synchronize()
-                    root.current, root.spare = root.spare, root.current
-                    self._step_count += 1
-                    self._exchange_valid = self._observe_raw_exchange
+                    root.current, root.spare = root.spare, root.current  # host handles only
+                self._stream.synchronize()
+                self._step_count += n
+                self._exchange_valid = self._observe_raw_exchange
                 self._status = "ready"
             except BaseException:
                 self._window_valid = False

src/CelerisLab/lbm/refinement_tables.py

@@ -66,6 +66,50 @@ def blocked_base_tables(patch, layout, order, wall_aware=False, *, composed_kind
             blocked_kind(layout, dense_kind))
 
 
+def projected_explosion_tables(pair, parent_layout, composed_parent_kind):
+    """Experimental transverse linear explosion, in the pair's existing row order.
+
+    Schornbaum & Rüde (2016), arXiv:1508.07982, Algorithm 2, following
+    Chen et al., doi:10.1016/j.physa.2005.09.036. Independently implemented
+    from the mathematics: centered gradients, perpendicular projection and
+    child offsets +/-1/4. The kernel applies the projection per population.
+
+    Only original explosion donors that are owned, uncovered ACTIVE cells
+    may supply gradients. Both axis neighbors must be legal; otherwise repeat
+    the center twice for zero slope. Owner-core addresses avoid stale halos.
+    The same rule applies to scalar windows, unions and every adjacent pair;
+    nested layouts retain their own per-axis (usually nonperiodic) topology.
+    """
+    ny, nx = pair.coarse_shape
+    kind = np.asarray(composed_parent_kind)
+    if parent_layout.shape != pair.coarse_shape or kind.shape != pair.coarse_shape:
+        raise ValueError("projected explosion parent layout/kind has the wrong shape")
+    donor = pair.ghost_coarse.astype(np.intp)
+    legal = np.zeros(ny * nx, bool)
+    legal[donor] = True
+    legal &= (parent_layout.owned.ravel() & ~pair.covered.ravel()
+              & (kind.ravel() == ACTIVE))
+    if not legal[donor].all():
+        raise ValueError("projected explosion donors need owned uncovered parent ACTIVE cells")
+    y, x = np.divmod(donor, nx)
+    neighbors = np.repeat(donor[:, None], 4, axis=1)
+    for axis, (dx, dy) in enumerate(((1, 0), (0, 1))):
+        minus = ((y - dy) % ny) * nx + (x - dx) % nx
+        plus = ((y + dy) % ny) * nx + (x + dx) % nx
+        valid = legal[minus] & legal[plus]
+        if not parent_layout.periodic[axis]:
+            coordinate, extent = (x, nx) if axis == 0 else (y, ny)
+            valid &= (coordinate > 0) & (coordinate < extent - 1)
+        neighbors[valid, 2 * axis] = minus[valid]
+        neighbors[valid, 2 * axis + 1] = plus[valid]
+    fy, fx = np.divmod(pair.ghost_fine, pair.fine_shape[1])
+    ox, oy = pair.fine_origin
+    offsets = np.column_stack((ox + (fx + .5) / 2 - (x + .5),
+                               oy + (fy + .5) / 2 - (y + .5)))
+    return (np.ascontiguousarray(parent_layout.owner.ravel()[neighbors], dtype=np.uint64),
+            np.ascontiguousarray(offsets, dtype=np.float32))
+
+
 # Required device columns plus optional actuation; observation arms are uploaded
 # separately only when enabled. body_id/aperture remain host reduction metadata.
 BLOCKED_BOUNDARY_DEVICE_COLUMNS = ("block", "cell", "dir", "q", "fallback", "surface_velocity")

tests/unit/test_les_flow_runner.py

@@ -43,12 +43,12 @@ class GuardTests(unittest.TestCase):
 
     def test_current_corrected_pins(self):
         self.assertEqual(
-            source_guard(ROOT, "corrected", "dev8403956")["base_commit"],
-            "8403956a9afc0b0bb89265cec374392013d89016",
+            source_guard(ROOT, "corrected", "a135252")["base_commit"],
+            "a135252ae78b47b57d7e77d19001700b3bedac8b",
         )
 
     def test_corrected_and_tamper_guard(self):
-        pins = json.loads((HERE / "data/les_flow/source_pins_dev8403956.json").read_text())
+        pins = json.loads((HERE / "data/les_flow/source_pins_a135252.json").read_text())
         with TemporaryDirectory() as directory:
             root = Path(directory)
             for name in pins["sources"]:
@@ -59,24 +59,45 @@ class GuardTests(unittest.TestCase):
                 "run_uniform_les_flow.subprocess.check_output",
                 return_value="isolated-fixture\n",
             ):
-                source_guard(root, "corrected", "dev8403956")
+                source_guard(root, "corrected", "a135252")
                 with self.assertRaisesRegex(RuntimeError, "Source pin mismatch"):
-                    source_guard(root, "baseline", "dev8403956")
-                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", "dev8403956")
+                    source_guard(root, "baseline", "a135252")
+                for name in ("kernels/step/one_step_double.cu", "kernels/blocks.cu",
+                             "refinement_stepper.py", "refinement_tables.py"):
+                    target = root / "src/CelerisLab/lbm" / name
+                    original = target.read_bytes()
+                    target.write_bytes(original + b"\n// tamper\n")
+                    with self.assertRaisesRegex(RuntimeError, "Source pin mismatch"):
+                        source_guard(root, "corrected", "a135252")
+                    target.write_bytes(original)
 
     def test_historical_default_and_manifest_unchanged(self):
-        manifest = HERE / "data/les_flow/source_pins.json"
-        self.assertEqual(hashlib.sha256(manifest.read_bytes()).hexdigest(),
-                         "87ca033619da1fc02bee822bfc5681b518a92a155c3c3d1dbbe5265c96066b39")
-        for variant in ("baseline", "corrected"):
-            with self.assertRaisesRegex(RuntimeError, "Source pin mismatch"):
-                source_guard(ROOT, variant)
+        historical = {
+            "historical": ("source_pins.json", "87ca033619da1fc02bee822bfc5681b518a92a155c3c3d1dbbe5265c96066b39"),
+            "dev8403956": ("source_pins_dev8403956.json", "58187908485f2af7c78677d53de98015878372444265ce9e667683b8ef5d410f"),
+        }
+        for experiment, (name, expected) in historical.items():
+            manifest = HERE / "data/les_flow" / name
+            self.assertEqual(hashlib.sha256(manifest.read_bytes()).hexdigest(), expected)
+            for variant in ("baseline", "corrected"):
+                with self.assertRaisesRegex(RuntimeError, "Source pin mismatch"):
+                    source_guard(ROOT, variant, experiment)
+        with self.assertRaisesRegex(RuntimeError, "Source pin mismatch"):
+            source_guard(ROOT, "corrected")  # Omission still uses the historical profile.
+
+    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")
 
     def test_named_baseline_accepts_only_reversed_header(self):
-        pins = json.loads((HERE / "data/les_flow/source_pins_dev8403956.json").read_text())
+        pins = json.loads((HERE / "data/les_flow/source_pins_a135252.json").read_text())
         with TemporaryDirectory() as directory:
             root = Path(directory)
             for name in pins["sources"]:
@@ -85,9 +106,9 @@ class GuardTests(unittest.TestCase):
                 shutil.copyfile(ROOT / name, dest)
             subprocess.run(["git", "apply", "--reverse", str(HERE / "les/data/smagorinsky_normalization.patch")], cwd=root, check=True)
             with patch("run_uniform_les_flow.subprocess.check_output", return_value="baseline-fixture\n"):
-                source_guard(root, "baseline", "dev8403956")
+                source_guard(root, "baseline", "a135252")
                 with self.assertRaisesRegex(RuntimeError, "Source pin mismatch"):
-                    source_guard(root, "corrected", "dev8403956")
+                    source_guard(root, "corrected", "a135252")
 
     def test_gpu_requires_opt_in_before_import(self):
         with TemporaryDirectory() as directory:

tests/unit/test_projected_explosion.py

@@ -0,0 +1,394 @@
+"""CPU stencil and actual closed-interface packet budgets; no CUDA claims."""
+import numpy as np
+import unittest
+from pathlib import Path
+from dataclasses import replace
+import shutil
+import subprocess
+import tempfile
+
+from CelerisLab.config import LBMConfig
+from CelerisLab.lbm.blocks import BlockLayout2D
+from CelerisLab.lbm.descriptors import D2Q9_EX, D2Q9_EY
+from CelerisLab.lbm.refinement import FixedPatch2D, FixedPatchSet2D, ACTIVE, OUTER_GHOST
+from CelerisLab.lbm.refinement_tables import projected_explosion_tables
+from CelerisLab.lbm.hierarchy import build_hierarchy
+from tests.unit.test_hierarchy import CFG, FOUR
+
+E = np.array([D2Q9_EX, D2Q9_EY]).T
+W = np.array([4/9] + [1/9]*4 + [1/36]*4)
+
+
+def setup(union=False, bounds=(4, 4, 12, 12)):
+    cfg = LBMConfig(nx=16, ny=16, viscosity=.1)
+    p = (FixedPatchSet2D.from_config(cfg, [bounds]) if union else
+         FixedPatch2D.from_config(cfg, bounds))
+    layout = BlockLayout2D.build(16, 16, 4, 4, periodic=(True, True))
+    neighbors, offsets = projected_explosion_tables(p, layout, np.where(p.covered, OUTER_GHOST, ACTIVE))
+    inverse = np.full(layout.nblocks * np.prod(layout.local_shape), -1, dtype=int)
+    inverse[layout.owner.ravel()] = np.arange(256)
+    return p, inverse[neighbors], offsets
+
+
+def reconstruct(c, p, neighbors, offsets, project=True):
+    # Independent NumPy evaluation of the documented stencil; neighbors have
+    # been decoded from actual blocked owner IDs, not recomputed geometry.
+    values = c.reshape(9, -1)[:, neighbors]
+    g = .5 * (values[:, :, 1::2] - values[:, :, ::2])
+    if project:
+        g[1:] -= (np.einsum('qra,qa->qr', g[1:], E[1:]) /
+                  np.sum(E[1:]**2, axis=1)[:, None])[:, :, None] * E[1:, None, :]
+    out = c.reshape(9, -1)[:, p.ghost_coarse] + np.einsum('qra,ra->qr', g, offsets)
+    return out, g
+
+
+def check_legal_footprint_mean_and_characteristic(union, bounds):
+    p, neighbors, offsets = setup(union, bounds)
+    assert set(neighbors.ravel()) <= set(p.ghost_coarse)
+    assert not p.covered.ravel()[neighbors].any()
+    assert (neighbors >= 0).all()
+    c = np.random.default_rng(72).normal(size=(9,16,16))
+    c[:, p.covered] = np.nan  # poison every forbidden donor
+    out, g = reconstruct(c, p, neighbors, offsets)
+    assert np.isfinite(out).all()
+    np.testing.assert_allclose(np.einsum('qra,qa->qr', g, E), 0, atol=5e-16)
+    for d in np.unique(p.ghost_coarse):
+        rows = p.ghost_coarse == d
+        assert rows.sum() == 4
+        np.testing.assert_allclose(out[:, rows].mean(axis=1), c.reshape(9,-1)[:,d], atol=5e-16)
+        for q in range(1,9):
+            delta = offsets[rows, None, :] - offsets[None, rows, :]
+            # Parallel displacements have zero 2D cross product.
+            parallel = np.isclose(delta[...,0]*E[q,1]-delta[...,1]*E[q,0], 0)
+            diff = out[q, rows][:,None]-out[q, rows][None,:]
+            np.testing.assert_allclose(diff[parallel], 0, atol=5e-16)
+    const = np.broadcast_to(W[:,None,None], (9,16,16)).copy()
+    constant, _ = reconstruct(const,p,neighbors,offsets)
+    np.testing.assert_array_equal(constant, np.broadcast_to(W[:,None], constant.shape))
+
+
+def collide(f, tau):
+    rho = f.sum(axis=0)
+    u = np.einsum('qa,q...->a...', E, f) / rho
+    cu = np.einsum('qa,a...->q...', E, u)
+    eq = W[:,None] * rho * (1 + 3*cu + 4.5*cu**2 - 1.5*np.sum(u*u,axis=0))
+    return f + (eq-f)/tau
+
+
+def totals(c, f, p):
+    q = c[:,~p.covered].sum(axis=1)+.25*f[:,p.fine_kind==ACTIVE].sum(axis=1)
+    return np.r_[q.sum(), E.T@q]
+
+
+def check_full_two_hop_closed_corner_budget(union, collision, bounds):
+    p, neighbors, offsets = setup(union,bounds)
+    rng=np.random.default_rng(17)
+    c=W[:,None,None]*(1+rng.normal(0,.002,(9,16,16)))
+    f=W[:,None,None]*(1+rng.normal(0,.002,(9,*p.fine_shape)))
+    start=totals(c,f,p)
+    # Count exactly the eight coarse diagonal links whose source is also
+    # uncovered, but whose midpoint lies on the closed rectangular interface.
+    q, k=np.divmod(p.coalesce_coarse.astype(np.int64),256)
+    y,x=np.divmod(k,16)
+    grazing=~p.covered[y-E[q,1],x-E[q,0]]
+    assert grazing.sum()==8
+    c[:,p.covered]=np.nan
+    for _ in range(8):
+        f.reshape(9,-1)[:,p.ghost_fine]=reconstruct(c,p,neighbors,offsets)[0]
+        for _ in range(2):
+            f=np.array([np.roll(f[q],(ey,ex),(0,1)) for q,(ex,ey) in enumerate(E)])
+            f[:,p.fine_kind==OUTER_GHOST]=np.nan
+            if collision:
+                f[:,p.fine_kind==ACTIVE]=collide(f[:,p.fine_kind==ACTIVE],p.tau_fine)
+        out=np.array([np.roll(c[q],(ey,ex),(0,1)) for q,(ex,ey) in enumerate(E)])
+        packets=f.reshape(-1)[p.coalesce_fine].mean(axis=1)
+        assert np.isfinite(packets[grazing]).all()
+        out.reshape(-1)[p.coalesce_coarse]=packets
+        if collision:out[:,~p.covered]=collide(out[:,~p.covered],p.tau_coarse)
+        out[:,p.covered]=np.nan
+        c=out
+        np.testing.assert_allclose(totals(c,f,p),start,rtol=0,atol=2e-12)
+
+
+def check_projection_is_distinguished_by_characteristic_not_global_mass():
+    p,n,o=setup()
+    c=np.random.default_rng(8).normal(size=(9,16,16))
+    _,g=reconstruct(c,p,n,o,project=False)
+    assert np.max(abs(np.einsum('qra,qa->qr',g,E))) > .1
+
+
+
+
+class ProjectedExplosionTests(unittest.TestCase):
+    def test_stencil(self):
+        for union in (False, True):
+            for bounds in ((4,4,12,12),(2,3,9,10),(2,2,3,3)):
+                with self.subTest(union=union,bounds=bounds):
+                    check_legal_footprint_mean_and_characteristic(union,bounds)
+
+    def test_closed_interface(self):
+        for union in (False,True):
+            for collision in (False,True):
+                for bounds in ((4,4,12,12),(2,3,14,12),(2,2,3,3)):
+                    with self.subTest(union=union,collision=collision,bounds=bounds):
+                        check_full_two_hop_closed_corner_budget(union,collision,bounds)
+
+    def test_projection(self):
+        check_projection_is_distinguished_by_characteristic_not_global_mass()
+
+    def test_affine_and_shift_invariance(self):
+        p,n,o=setup()
+        y,x=np.indices((16,16))
+        c=np.array([.7+q*.01+x*.03+y*.05 for q in range(9)])
+        out,g=reconstruct(c,p,n,o)
+        available=n[:,::2]!=n[:,1::2]
+        gradient=np.array([.03,.05])*available
+        expected=np.broadcast_to(gradient,(9,*gradient.shape)).copy()
+        expected[1:]-=(np.einsum('qra,qa->qr',expected[1:],E[1:])/
+                       np.sum(E[1:]**2,axis=1)[:,None])[:,:,None]*E[1:,None,:]
+        np.testing.assert_allclose(g,expected,atol=2e-16)
+        shifted,_=reconstruct(c-W[:,None,None],p,n,o)
+        np.testing.assert_allclose(out-W[:,None],shifted,atol=5e-16)
+
+
+class ProjectedKernelTests(unittest.TestCase):
+    @unittest.skipUnless(shutil.which("g++"), "g++ is required for actual-kernel CPU arithmetic")
+    def test_actual_kernel_fp32_shifted_and_unshifted(self):
+        p, neighbors, offsets = setup(True)
+        parent = BlockLayout2D.build(16,16,4,4,periodic=(True,True))
+        child = p.fine_layout((8,4))  # deliberately unequal parent/child strides
+        ids, offsets = projected_explosion_tables(p,parent,np.where(p.covered,OUTER_GHOST,ACTIVE))
+        cells = int(np.prod(parent.local_shape))
+        targets = child.owner.ravel()[p.ghost_fine]
+        donors = parent.owner.ravel()[p.ghost_coarse]
+        source = (Path(__file__).resolve().parents[2] /
+                  "src/CelerisLab/lbm/kernels/blocks.cu").read_text()
+        kernel = source.split('extern "C" __global__ void BlockExplodeProjected(',1)[1]
+        kernel = 'void BlockExplodeProjected(' + kernel.split('// Targets are receiver-local SOURCE slots',1)[0]
+        def array(name, values, kind):
+            v = np.asarray(values).ravel()
+            text = ','.join(str(int(x)) if kind=='uint64_t' else float(x).hex()+'f' for x in v)
+            return f"{kind} {name}[] = {{{text}}};\n"
+        for shifted in (False,True):
+            c=np.random.default_rng(11).normal(0,.001,(9,16,16))
+            if not shifted:c+=W[:,None,None]
+            c=c.astype(np.float32)
+            packed=np.full(parent.nblocks*9*cells,np.nan,np.float32)
+            for q in range(9):packed[parent.slot(parent.owner.ravel(),q)]=c[q].ravel()
+            # All halo/covered inputs are poisoned; kernel must use legal cores.
+            for q in range(9):packed[parent.slot(parent.owner.ravel()[p.covered.ravel()],q)]=np.nan
+            # Hex initializer needs a portable expression for poisoned NaNs.
+            data=array('coarse',packed,'float').replace('nanf','NAN')
+            code = '#include <cstdint>\n#include <cmath>\n#include <cstdio>\nusing fpxx=float;\n'
+            code += f'constexpr unsigned NQ=9, TOTAL_CELLS={int(np.prod(child.local_shape))};\n'
+            code += 'struct Id {unsigned x;}; Id blockIdx{0},blockDim{1},threadIdx{0};\n'
+            code += 'unsigned index_f(unsigned k,unsigned q){return q*TOTAL_CELLS+k;}\n'
+            code += 'int d_cx[]={0,1,-1,0,0,1,-1,1,-1},d_cy[]={0,0,0,1,-1,1,-1,-1,1};\n'+kernel
+            code += data+array('target',targets,'uint64_t')+array('donor',donors,'uint64_t')
+            code += array('neighbor',ids,'uint64_t')+array('offset',offsets,'float')
+            code += f'float fine[{child.nblocks*9*int(np.prod(child.local_shape))}]={{}};\n'
+            code += f'int main(){{for(unsigned r=0;r<{len(targets)};++r){{threadIdx.x=r;BlockExplodeProjected(coarse,fine,target,donor,{len(targets)},{cells},neighbor,offset);}}'
+            code += f'for(unsigned q=0;q<9;++q)for(unsigned r=0;r<{len(targets)};++r){{auto t=target[r];printf("%.9g\\n",fine[(t/TOTAL_CELLS)*9*TOTAL_CELLS+q*TOTAL_CELLS+t%TOTAL_CELLS]);}}}}'
+            with tempfile.TemporaryDirectory() as folder:
+                src=Path(folder)/'kernel.cpp'; binary=Path(folder)/'kernel'
+                src.write_text(code)
+                subprocess.run(['g++','-std=c++17','-O2',str(src),'-o',str(binary)],check=True,capture_output=True)
+                result=np.fromstring(subprocess.check_output([str(binary)],text=True),sep=' ').reshape(9,-1)
+            expected,_=reconstruct(c.astype(float),p,neighbors,offsets)
+            np.testing.assert_allclose(result,expected,atol=6e-8,rtol=2e-7)
+
+
+UNIONS = [
+    [(8, 8, 24, 20)],
+    [(8, 8, 16, 20), (16, 8, 24, 20)],       # abutting
+    [(8, 8, 18, 20), (14, 8, 24, 20)],       # overlap
+    [(8, 8, 16, 20), (17, 8, 24, 20)],       # gap 1
+    [(8, 8, 16, 20), (18, 8, 24, 20)],       # gap 2
+    [(8, 8, 16, 20), (19, 8, 24, 20)],       # gap 3
+    [(8, 8, 20, 14), (8, 14, 14, 22)],       # L concavity
+    [(4, 4, 10, 10), (18, 16, 26, 26)],      # separated
+    [(4, 4, 5, 7), (4, 5, 7, 7)],           # known concave exit/re-entry path
+]
+
+
+def decode(layout, neighbors):
+    inverse = np.full(layout.hole, -1, dtype=int)
+    inverse[layout.owner[layout.owned]] = np.flatnonzero(layout.owned)
+    return inverse[neighbors]
+
+
+def check_pair_stencil(pair, layout, kind):
+    ids, offsets = projected_explosion_tables(pair, layout, kind)
+    neighbors = decode(layout, ids)
+    legal = set(map(int, pair.ghost_coarse))
+    assert len(np.unique(pair.ghost_fine)) == len(pair.ghost_fine)
+    assert set(neighbors.ravel()) <= legal
+    assert layout.owned.ravel()[neighbors].all()
+    assert (kind.ravel()[neighbors] == ACTIVE).all()
+    assert not pair.covered.ravel()[neighbors].any()
+    ny, nx = pair.coarse_shape
+    for row, donor in enumerate(pair.ghost_coarse):
+        y, x = divmod(int(donor), nx)
+        for axis, (dx, dy) in enumerate(((1, 0), (0, 1))):
+            ends = [(x - dx, y - dy), (x + dx, y + dy)]
+            expected = []
+            for xx, yy in ends:
+                outside = not (0 <= xx < nx and 0 <= yy < ny)
+                if outside and not layout.periodic[axis]:
+                    break
+                cell = (yy % ny) * nx + xx % nx
+                if cell not in legal:
+                    break
+                expected.append(cell)
+            if len(expected) != 2:
+                expected = [int(donor)] * 2
+            np.testing.assert_array_equal(neighbors[row, 2*axis:2*axis+2], expected)
+        fy, fx = divmod(int(pair.ghost_fine[row]), pair.fine_shape[1])
+        expected = [pair.fine_origin[0]+(fx+.5)/2-(x+.5),
+                    pair.fine_origin[1]+(fy+.5)/2-(y+.5)]
+        np.testing.assert_array_equal(offsets[row], expected)
+    assert set(offsets.ravel()) == {-.25, .25}
+    assert ids.dtype == np.uint64 and ids.flags.c_contiguous
+    assert offsets.dtype == np.float32 and offsets.flags.c_contiguous
+    return neighbors, offsets
+
+
+class GenericProjectedPairs(unittest.TestCase):
+    def test_single_window_and_singleton_union_exact_stencils(self):
+        for bounds in ((4,4,12,12), (2,3,9,10), (2,2,3,3)):
+            a, an, ao = setup(False, bounds)
+            b, bn, bo = setup(True, bounds)
+            np.testing.assert_array_equal(a.ghost_coarse, b.ghost_coarse)
+            np.testing.assert_array_equal(an, bn)
+            np.testing.assert_array_equal(ao, bo)
+            # Coordinates differ between scalar and union arrays; physical rows don't.
+            ay, ax = np.divmod(a.ghost_fine, a.fine_shape[1])
+            by, bx = np.divmod(b.ghost_fine, b.fine_shape[1])
+            np.testing.assert_array_equal(ax + 2*a.fine_origin[0], bx)
+            np.testing.assert_array_equal(ay + 2*a.fine_origin[1], by)
+
+    def test_union_decomposition_permutation_and_topologies(self):
+        cfg = LBMConfig(nx=30, ny=30, viscosity=.1)
+        layout = BlockLayout2D.build(30,30,6,5,periodic=(True,True))
+        stencils = []
+        for rectangles in UNIONS:
+            pair = FixedPatchSet2D.from_config(cfg, rectangles)
+            kind = np.where(pair.covered, OUTER_GHOST, ACTIVE)
+            n, o = check_pair_stencil(pair, layout, kind)
+            other = FixedPatchSet2D.from_config(cfg, rectangles[::-1])
+            on, oo = projected_explosion_tables(other, layout, kind)
+            np.testing.assert_array_equal(n, decode(layout, on))
+            np.testing.assert_array_equal(o, oo)
+            stencils.append((pair.ghost_fine, pair.ghost_coarse, n, o))
+            c = np.random.default_rng(92).normal(size=(9,30,30))
+            c[:, pair.covered] = np.nan
+            values, gradients = reconstruct(c, pair, n, o)
+            self.assertTrue(np.isfinite(values).all())
+            np.testing.assert_allclose(np.einsum('qra,qa->qr', gradients, E),0,atol=1e-15)
+            # q0 keeps its full gradient; no division by zero or forced zero.
+            expected_q0 = .5*(c[0].ravel()[n[:,1::2]]-c[0].ravel()[n[:,::2]])
+            np.testing.assert_array_equal(gradients[0], expected_q0)
+            self.assertGreater(np.max(abs(gradients[0])), .1)
+            for donor in np.unique(pair.ghost_coarse):
+                rows = pair.ghost_coarse == donor
+                self.assertEqual(rows.sum(),4)
+                np.testing.assert_allclose(values[:,rows].mean(axis=1),c.reshape(9,-1)[:,donor],atol=1e-15)
+                for q in range(1,9):
+                    delta = o[rows,None,:]-o[None,rows,:]
+                    parallel = delta[...,0]*E[q,1] == delta[...,1]*E[q,0]
+                    diff = values[q,rows][:,None]-values[q,rows][None,:]
+                    np.testing.assert_allclose(diff[parallel],0,atol=1e-15)
+        for split in stencils[1:3]:
+            for expected, actual in zip(stencils[0],split):
+                np.testing.assert_array_equal(actual,expected)
+
+    def test_nested_scalar_origins_sparse_holes_and_unequal_strides(self):
+        cases = [(CFG, FOUR), (LBMConfig(nx=16,ny=16,viscosity=.02),
+                               [(4,4,12,12),(8,8,16,16),(8,8,16,16)])]
+        for cfg, refs in cases:
+            hierarchy = build_hierarchy(cfg,refs)
+            for index,pair in enumerate(hierarchy.pairs):
+                if index == 0:
+                    layout = BlockLayout2D.build(cfg.ny,cfg.nx,4,4,periodic=(True,True))
+                else:
+                    prior = hierarchy.pairs[index-1]
+                    layout = (prior.fine_layout((4,8)) if isinstance(prior,FixedPatchSet2D)
+                              else BlockLayout2D.build(*prior.fine_shape,4,8,periodic=(False,False)))
+                kind = hierarchy.kinds[index]
+                n,o = check_pair_stencil(pair,layout,kind)
+                c = np.random.default_rng(index).normal(size=(9,*pair.coarse_shape))
+                c[:,~layout.owned | (kind != ACTIVE)] = np.nan
+                self.assertTrue(np.isfinite(reconstruct(c,pair,n,o)[0]).all())
+                if index and isinstance(pair,FixedPatchSet2D):
+                    self.assertTrue((~layout.owned).any())
+                self.assertEqual(layout.periodic, (index == 0, index == 0))
+
+    def test_periodicity_is_per_axis_and_missing_pairs_repeat_center(self):
+        cfg=LBMConfig(nx=16,ny=16,viscosity=.1)
+        p=FixedPatch2D.from_config(cfg,(2,2,14,14))
+        kind=np.where(p.covered,OUTER_GHOST,ACTIVE)
+        results=[]
+        for periodic in ((False,False),(True,False),(False,True),(True,True)):
+            layout=BlockLayout2D.build(16,16,4,4,periodic=periodic)
+            n,_=check_pair_stencil(p,layout,kind)
+            results.append(n)
+            y,x=np.divmod(p.ghost_coarse,16)
+            for axis,coordinate in enumerate((x,y)):
+                edge=(coordinate==0)|(coordinate==15)
+                if not periodic[axis]:
+                    np.testing.assert_array_equal(n[edge,2*axis:2*axis+2],
+                        np.repeat(p.ghost_coarse[edge,None],2,axis=1))
+        self.assertTrue((results[0]!=results[1]).any())
+        self.assertTrue((results[0]!=results[2]).any())
+        self.assertTrue((results[1]!=results[3]).any())
+
+    def test_invalid_parent_donors_reject_before_owner_lookup(self):
+        p,_,_=setup()
+        layout=BlockLayout2D.build(16,16,4,4)
+        kind=np.where(p.covered,OUTER_GHOST,ACTIVE)
+        for nonactive in (1,2,255):
+            bad=kind.copy(); bad.ravel()[p.ghost_coarse[0]]=nonactive
+            with self.assertRaisesRegex(ValueError,'ACTIVE'):
+                projected_explosion_tables(p,layout,bad)
+        with self.assertRaisesRegex(ValueError,'shape'):
+            projected_explosion_tables(p,layout,kind[:1])
+        hole_layout=BlockLayout2D.build(16,16,4,4,origins=[(0,0)])
+        with self.assertRaisesRegex(ValueError,'owned'):
+            projected_explosion_tables(p,hole_layout,kind)
+
+    def test_generic_two_hop_mass_and_momentum_budget(self):
+        cfg=LBMConfig(nx=30,ny=30,viscosity=.1)
+        layout=BlockLayout2D.build(30,30,6,5)
+        grazing_counts=set()
+        for rectangles in UNIONS:
+            pair=FixedPatchSet2D.from_config(cfg,rectangles)
+            ids,o=projected_explosion_tables(pair,layout,np.where(pair.covered,OUTER_GHOST,ACTIVE))
+            n=decode(layout,ids)
+            q,k=np.divmod(pair.coalesce_coarse.astype(np.int64),900)
+            y,x=np.divmod(k,30)
+            grazing=~pair.covered[(y-E[q,1])%30,(x-E[q,0])%30]
+            grazing_counts.add(int(grazing.sum()))
+            self.assertTrue(grazing.any())
+            for collision in (False,True):
+                rng=np.random.default_rng(17)
+                c=W[:,None,None]*(1+rng.normal(0,.002,(9,30,30)))
+                f=W[:,None,None]*(1+rng.normal(0,.002,(9,*pair.fine_shape)))
+                start=totals(c,f,pair)
+                c[:,pair.covered]=np.nan
+                for _ in range(4):
+                    f.reshape(9,-1)[:,pair.ghost_fine]=reconstruct(c,pair,n,o)[0]
+                    for _ in range(2):
+                        f=np.array([np.roll(f[q],(ey,ex),(0,1)) for q,(ex,ey) in enumerate(E)])
+                        f[:,pair.fine_kind==OUTER_GHOST]=np.nan
+                        if collision:
+                            f[:,pair.fine_kind==ACTIVE]=collide(f[:,pair.fine_kind==ACTIVE],pair.tau_fine)
+                    out=np.array([np.roll(c[q],(ey,ex),(0,1)) for q,(ex,ey) in enumerate(E)])
+                    packets=f.reshape(-1)[pair.coalesce_fine].mean(axis=1)
+                    self.assertTrue(np.isfinite(packets[grazing]).all())
+                    out.reshape(-1)[pair.coalesce_coarse]=packets
+                    if collision:
+                        out[:,~pair.covered]=collide(out[:,~pair.covered],pair.tau_coarse)
+                    out[:,pair.covered]=np.nan
+                    c=out
+                    np.testing.assert_allclose(totals(c,f,pair),start,rtol=0,atol=4e-12)
+        self.assertGreater(len(grazing_counts),1)  # not a hard-coded eight-corner oracle

tests/unit/test_projected_pair_gate.py

@@ -0,0 +1,162 @@
+"""Host-only checks for the opt-in CUDA gate; never imports or compiles CUDA."""
+from copy import deepcopy
+import importlib.util
+from pathlib import Path
+import unittest
+from unittest.mock import patch
+
+import numpy as np
+
+from CelerisLab.config import LBMConfig
+from CelerisLab.lbm.blocks import BlockLayout2D
+from CelerisLab.lbm.hierarchy import build_hierarchy
+from CelerisLab.lbm.refinement import ACTIVE, OUTER_GHOST, FixedPatch2D, FixedPatchSet2D
+from CelerisLab.lbm.refinement_tables import projected_explosion_tables
+from tests.unit.test_hierarchy import CFG, FOUR
+from tests.unit.test_projected_explosion import decode, reconstruct, totals
+
+PATH = Path(__file__).resolve().parents[1] / 'validation/multigrid/check_projected_pairs.py'
+spec = importlib.util.spec_from_file_location('projected_pair_gate', PATH)
+gate = importlib.util.module_from_spec(spec)
+spec.loader.exec_module(gate)
+
+
+class ProjectedGateOracle(unittest.TestCase):
+    def check_pair(self, pair, layout, kind):
+        ids, offsets = projected_explosion_tables(pair, layout, kind)
+        neighbors = decode(layout, ids)
+        coarse = np.random.default_rng(731).normal(size=(9, *pair.coarse_shape))
+        legal = np.zeros(pair.coarse_shape, bool)
+        legal.ravel()[pair.ghost_coarse] = True
+        coarse[:, ~legal] = np.nan
+        expected, _ = reconstruct(coarse, pair, neighbors, offsets)
+        # The coordinate oracle must not accidentally share production tables.
+        with patch('CelerisLab.lbm.refinement_tables.projected_explosion_tables',
+                   side_effect=AssertionError('oracle used production stencil')):
+            actual = gate.independent_explosion(coarse, pair, layout, kind)
+        np.testing.assert_allclose(actual, expected, rtol=0, atol=2e-15)
+        center = coarse.reshape(9, -1)[:, pair.ghost_coarse]
+        self.assertTrue(np.isfinite(actual).all())
+        self.assertGreater(np.max(abs(actual[0] - center[0])), .01)
+        self.assertGreater(np.max(abs(actual[1:] - center[1:])), .01)
+        return coarse, neighbors, offsets, actual
+
+    def test_scalar_union_nested_sparse_and_unequal_strides(self):
+        def fine_layout(pair, block):
+            return (pair.fine_layout(block) if isinstance(pair, FixedPatchSet2D) else
+                    BlockLayout2D.build(*pair.fine_shape, *block, periodic=(False,False)))
+        small = LBMConfig(nx=16, ny=16, viscosity=.02)
+        fixtures = [(small, [(4,4,12,12), (8,8,16,16), (8,8,16,16)]),
+                    (CFG, FOUR), (small, [[(4,4,5,7), (4,5,7,7)]])]
+        for cfg, refs in fixtures:
+            hierarchy = build_hierarchy(cfg, refs)
+            for index, pair in enumerate(hierarchy.pairs):
+                block, child_block = ((4,4), (4,8)) if index % 2 == 0 else ((4,8), (4,4))
+                layout = (BlockLayout2D.build(cfg.ny, cfg.nx, *block) if index == 0
+                          else fine_layout(hierarchy.pairs[index-1], block))
+                with self.subTest(refs=refs, pair=index):
+                    self.assertNotEqual(np.prod(layout.local_shape),
+                                        np.prod(fine_layout(pair, child_block).local_shape))
+                    if index and isinstance(pair, FixedPatchSet2D):
+                        self.assertTrue((~layout.owned).any())
+                    self.check_pair(pair, layout, hierarchy.kinds[index])
+
+    def test_axis_periodicity_and_mean_preserving_mutants(self):
+        pair = FixedPatch2D.from_config(LBMConfig(nx=16,ny=16,viscosity=.1), (2,2,14,14))
+        kind = np.where(pair.covered, OUTER_GHOST, ACTIVE)
+        outputs = []
+        for periodic in ((False,False), (True,False), (False,True), (True,True)):
+            layout = BlockLayout2D.build(16,16,4,4,periodic=periodic)
+            coarse, neighbors, offsets, actual = self.check_pair(pair, layout, kind)
+            outputs.append(actual)
+            donors, first, inverse = np.unique(pair.ghost_coarse, return_index=True, return_inverse=True)
+            wrong = neighbors[first[np.roll(np.arange(len(donors)), 1)]][inverse]
+            for bad_neighbors, bad_offsets in ((neighbors, -offsets), (wrong, offsets)):
+                mutant, _ = reconstruct(coarse, pair, bad_neighbors, bad_offsets)
+                self.assertGreater(np.max(abs(mutant - actual)), .01)
+                for donor in donors:
+                    rows = pair.ghost_coarse == donor
+                    self.assertEqual(rows.sum(), 4)
+                    np.testing.assert_allclose(mutant[:,rows].mean(axis=1),
+                                               coarse.reshape(9,-1)[:,donor], rtol=0, atol=1e-15)
+            bad_kind = kind.copy(); bad_kind.ravel()[pair.ghost_coarse[0]] = OUTER_GHOST
+            with self.assertRaisesRegex(AssertionError, 'illegal center donor'):
+                gate.independent_explosion(coarse, pair, layout, bad_kind)
+        for i, j in ((0,1), (0,2), (1,3), (2,3)):
+            self.assertGreater(np.max(abs(outputs[i] - outputs[j])), .01)
+
+    def test_concave_replay_preserves_existing_fp64_budget(self):
+        cfg, refs, blocks = gate.cases()['concave']
+        hierarchy = build_hierarchy(cfg, refs); pair = hierarchy.pairs[0]
+        layout = BlockLayout2D.build(cfg.ny, cfg.nx, *blocks[0])
+        rng = np.random.default_rng(17)
+        values = tuple(gate.W[:,None,None] * (1 + rng.normal(0,.002,(9,*shape)))
+                       for shape in (pair.coarse_shape, pair.fine_shape))
+        start = totals(*values, pair)
+        values[0][:,pair.covered] = np.nan
+        for _ in range(8):
+            values = gate.replay_two_hops(values, pair, layout, hierarchy.kinds[0])
+            np.testing.assert_allclose(totals(*values, pair), start, rtol=0, atol=4e-12)
+        self.assertEqual(gate.REPLAY_ATOL, 3e-6)  # CUDA replay bound is separate.
+
+
+class ProjectedGateReceipts(unittest.TestCase):
+    def test_digest_exact_finite_bits_and_nan_masks(self):
+        values = np.array([0., -0., np.nan, 1.], dtype=np.float32)
+        self.assertEqual(gate.digest(values), gate.digest(values.copy()))
+        payload = values.copy(); payload.view(np.uint32)[2] = 0x7fc00001
+        self.assertEqual(gate.digest(values), gate.digest(payload))
+        variants = [np.array([-0., -0., np.nan, 1.], np.float32),
+                    np.array([0., np.nan, -0., 1.], np.float32),
+                    np.array([0., -0., np.nan, np.nextafter(np.float32(1), np.float32(2))], np.float32),
+                    values.reshape(2,2), values.astype(np.float64)]
+        for changed in variants:
+            self.assertNotEqual(gate.digest(values), gate.digest(changed))
+        with self.assertRaisesRegex(AssertionError, 'infinite'):
+            gate.digest(np.array([np.inf]))
+
+    def test_snapshot_negative_controls_and_exact_dev_options(self):
+        snapshot = [dict(step=0, levels=[dict(raw='raw',public='public',active=4)],
+                         mass_momentum=[1.,0.,0.],observed_level_hops=[0,0])]
+        gate.compare_snapshots(snapshot, deepcopy(snapshot))
+        diagnostic = deepcopy(snapshot); diagnostic[0]['mass_momentum'][0] = 2.
+        gate.compare_snapshots(snapshot, diagnostic)  # Diagnostics are not bitwise gate fields.
+        mutants = [[]]
+        for key, value in (('raw','changed'), ('public','changed'), ('active',3)):
+            changed = deepcopy(snapshot); changed[0]['levels'][0][key] = value
+            mutants.append(changed)
+        changed = deepcopy(snapshot); changed[0]['step'] = 1; mutants.append(changed)
+        changed = deepcopy(snapshot); changed[0]['observed_level_hops'] = [0,1]; mutants.append(changed)
+        for changed in mutants:
+            with self.assertRaises(AssertionError):
+                gate.compare_snapshots(snapshot, changed)
+        self.assertEqual(gate.default_options('omitted'), {})
+        self.assertEqual(gate.default_options('constant'), {'explosion':'constant'})
+        self.assertEqual(gate.default_options('projected_linear'), {'explosion':'projected_linear'})
+
+    def test_baseline_receipt_rejects_provenance_and_contract_mutants(self):
+        keys = [f'{name}:shifted={shifted}' for name in gate.cases() for shifted in (False,True)]
+        good = dict(completed=True,dry_run=False,mode='baseline',production_revision=gate.BASE,
+                    source=dict(revision=gate.BASE,dirty=False),driver_sha256='driver',
+                    stops=list(gate.STOPS),operator_tolerance=[6e-8,2e-7],
+                    rest_tolerance=[3e-7,1e-6],replay_tolerance=[3e-6,0],default_snapshots=dict.fromkeys(keys, []),
+                    actual_build_sidecars=[{'ptx_sha256':'captured'}])
+        self.assertIs(gate.baseline_receipt(good, 'driver', keys), good)
+        mutants = dict(completed=False,dry_run=True,mode='candidate',production_revision='wrong',
+                       driver_sha256='wrong',stops=[0,1,2],operator_tolerance=[1e-6,2e-7],
+                       rest_tolerance=[3e-7,1e-5],replay_tolerance=[3e-6,1e-7],actual_build_sidecars=[],default_snapshots={})
+        for field, value in mutants.items():
+            bad = deepcopy(good); bad[field] = value
+            with self.subTest(field=field), self.assertRaises(ValueError):
+                gate.baseline_receipt(bad, 'driver', keys)
+        for changes in (dict(revision='wrong'), dict(dirty=True), dict(dirty=None)):
+            bad = deepcopy(good); bad['source'].update(changes)
+            with self.assertRaises(ValueError):
+                gate.baseline_receipt(bad, 'driver', keys)
+        bad = deepcopy(good); bad['default_snapshots']['unexpected'] = []
+        with self.assertRaises(ValueError):
+            gate.baseline_receipt(bad, 'driver', keys)
+
+
+if __name__ == '__main__':
+    unittest.main()

tests/unit/test_projected_pair_stepper.py

@@ -0,0 +1,198 @@
+"""Production host dispatch with CUDA doubles; no CUDA compilation or flow runs."""
+from contextlib import nullcontext
+from dataclasses import replace
+from unittest.mock import Mock
+import unittest
+
+import numpy as np
+
+from CelerisLab.config import LBMConfig
+from tests.unit.test_hierarchy import CFG, FOUR
+from tests.unit.test_refinement_three_level_multipatch import host_module
+
+
+def backend(**options):
+    host = host_module()
+    calls, allocations, storage = [], [], {}
+    context = Mock(activate=nullcontext, sm_arch='sm_70')
+    host.cuda.mem_alloc.side_effect = lambda size: allocate(size)
+
+    def allocate(size):
+        pointer = Mock(name=f'buffer{len(allocations)}', size=size)
+        allocations.append(pointer)
+        return pointer
+
+    def load(_):
+        module = Mock()
+        module.get_global.return_value = (Mock(), 20)
+        module.get_function.side_effect = lambda name: Mock(
+            name=name, side_effect=lambda *a, **kw: calls.append((name, a, kw)))
+        return module
+
+    host.compiler.load_module.side_effect = load
+    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)])
+    cfg = options.pop('cfg', CFG)
+    refinements = options.pop('refinements', FOUR)
+    blocks = options.pop('level_blocks', [(4, 4)] * (len(refinements) + 1))
+    run = host.FixedPatchStepper(cfg, context=context, refinements=refinements,
+                                level_blocks=blocks, **options)
+    run.initialize(*(np.full(level.shape, 1/9, np.float32) for level in run._levels))
+    return run, host, calls, allocations, storage
+
+
+def trace(run, calls, allocations, storage):
+    """Stable pointer identities expose allocations/arguments, not just kernel names."""
+    pointer = {id(value): index for index, value in enumerate(allocations)}
+    launches = [(name, tuple(('ptr', pointer[id(a)]) if id(a) in pointer else int(a)
+                            for a in args), kw['grid'], kw['block'])
+                for name, args, kw in calls]
+    uploads = [(index, value.dtype.str, value.shape, value.tobytes())
+               for index, ptr in enumerate(allocations) if (value := storage.get(id(ptr))) is not None]
+    return [p.size for p in allocations], uploads, launches
+
+
+class ProjectedPairDispatch(unittest.TestCase):
+    def test_omitted_and_constant_have_identical_allocations_uploads_and_launches(self):
+        results = []
+        for options in ({}, {'explosion': 'constant'}):
+            run, host, calls, allocations, storage = backend(**options)
+            self.assertEqual(run._explosion, 'constant')
+            self.assertTrue(all(p.neighbors is None and p.offsets is None for p in run._pairs))
+            run.step()
+            explosion = [args for name, args, _ in calls if name == 'BlockExplode']
+            self.assertEqual(len(explosion), 7)
+            self.assertTrue(all(len(args) == 6 for args in explosion))
+            self.assertFalse(any(name == 'BlockExplodeProjected' for name, _, _ in calls))
+            results.append(trace(run, calls, allocations, storage))
+            run.close()
+        self.assertEqual(*results)
+
+    def test_projected_adds_exactly_two_tables_per_adjacent_pair(self):
+        for depth in (2, 3, 4):
+            constant, _, _, ca, _ = backend(refinements=FOUR[:depth-1])
+            projected, _, calls, pa, storage = backend(
+                refinements=FOUR[:depth-1], explosion='projected_linear')
+            self.assertEqual(len(pa)-len(ca), 2*(depth-1))
+            for pair, geometry in zip(projected._pairs, projected.hierarchy.pairs):
+                self.assertEqual(storage[id(pair.neighbors)].shape, (len(geometry.ghost_fine), 4))
+                self.assertEqual(storage[id(pair.offsets)].shape, (len(geometry.ghost_fine), 2))
+                self.assertEqual(pair.neighbors.size, len(geometry.ghost_fine)*4*8)
+                self.assertEqual(pair.offsets.size, len(geometry.ghost_fine)*2*4)
+            projected.step()
+            explosions = [args for name, args, _ in calls if name == 'BlockExplodeProjected']
+            self.assertEqual(len(explosions), 2**(depth-1)-1)
+            self.assertTrue(all(len(args) == 8 for args in explosions))
+            self.assertFalse(any(name == 'BlockExplode' for name, _, _ in calls))
+            for args in explosions:
+                pair = next(p for p in projected._pairs if args[2] is p.ghosts)
+                self.assertIs(args[6], pair.neighbors)
+                self.assertIs(args[7], pair.offsets)
+                self.assertEqual(args[5], pair.parent_cells)
+            constant.close(); projected.close()
+
+    def test_three_and_four_level_schedule_has_fresh_midpoint_sources(self):
+        for depth in (3, 4):
+            for explosion in ('constant', 'projected_linear'):
+                run, host, calls, allocations, storage = backend(
+                    refinements=FOUR[:depth-1], explosion=explosion)
+                before = len(allocations)
+                events = []
+                for index, pair in enumerate(run._pairs):
+                    pair.explode.side_effect = lambda source, *a, index=index, **kw: events.append(('explode', index, source))
+                    pair.coalesce.side_effect = lambda child, source, *a, index=index, **kw: events.append(('restrict', index, source))
+                for index, level in enumerate(run._levels):
+                    level.kernel.side_effect = lambda source, *a, index=index, **kw: events.append(('step', index, source))
+                    level.halo_copy.side_effect = lambda source, *a, index=index, **kw: events.append(('halo', index, source))
+                expected = []
+
+                def hop(index, source):
+                    if index+1 < depth:
+                        expected.append(('explode', index, source))
+                        child = run._levels[index+1]
+                        hop(index+1, child.current)
+                        hop(index+1, child.spare)
+                    expected.append(('halo', index, source))
+                    if index+1 < depth:
+                        expected.append(('restrict', index, source))
+                    expected.append(('step', index, source))
+
+                hop(0, run._levels[0].current)
+                run.step()
+                self.assertEqual(events, expected)
+                self.assertEqual([level.kernel.call_count for level in run._levels], [2**i for i in range(depth)])
+                self.assertEqual(len(allocations), before)
+                self.assertEqual(run.step_count, 1)
+                for index in range(1, depth-1):
+                    feeds = [source for name, pair, source in events if name == 'explode' and pair == index]
+                    self.assertEqual(feeds, [run._levels[index].current, run._levels[index].spare] * (2**(index-1)))
+                run.close()
+
+    def test_scalar_windows_follow_the_same_pair_path(self):
+        run, _, _, _, _ = backend(cfg=LBMConfig(nx=16,ny=16,viscosity=.02),
+            refinements=[(4,4,12,12),(8,8,16,16),(8,8,16,16)],
+            level_blocks=[(4,4),(4,8),(8,4),(4,4)], explosion='projected_linear')
+        self.assertTrue(all(pair.neighbors is not None for pair in run._pairs))
+        self.assertEqual([pair.parent_cells for pair in run._pairs], [36,60,60])
+        run.step()
+        self.assertEqual([level.kernel.call_count for level in run._levels], [1,2,4,8])
+        run.close()
+
+    def test_physics_gate_precedes_context_and_allocations(self):
+        host = host_module()
+        context = Mock()
+        options = dict(refinements=FOUR[:1], level_blocks=[(4,4)]*2,
+                       explosion='projected_linear', context=context)
+        for cfg, changes in [(CFG, {'explosion':'unknown'}),
+                             (CFG, {'coarse_boundary':'slot_reservoir'}),
+                             (CFG, {'coarse_boundary':'channel'}),
+                             (CFG, {'level_blocks':[None,None], 'refinements':[(6,6,26,26)]}),
+                             (CFG, {'level_objects':[(),(object(),)]}),
+                             (CFG, {'level_jets':[(),(object(),)]}),
+                             (replace(CFG,collision='MRT'), {}),
+                             (replace(CFG,les_enabled=True), {})]:
+            with self.subTest(changes=changes, cfg=cfg), self.assertRaisesRegex(ValueError,'explosion'):
+                host.FixedPatchStepper(cfg, **dict(options, **changes))
+        context.activate.assert_not_called()
+        host.cuda.mem_alloc.assert_not_called()
+        # The legacy constant path still allows MRT and dense storage.
+        for options in ({'cfg':replace(CFG,collision='MRT')},
+                        {'level_blocks':[None,None], 'refinements':[(6,6,26,26)]}):
+            run, _, _, _, _ = backend(**options)
+            run.close()
+
+    def test_remap_preserves_selector_tables_and_constant_default(self):
+        for explosion in ('constant','projected_linear'):
+            run, host, _, _, storage = backend(refinements=FOUR[:2], explosion=explosion)
+            run._step_count = 7
+            old = run.download()
+            clone = run.remap_child(FOUR[1])
+            self.assertEqual(clone._explosion, explosion)
+            self.assertEqual(clone.step_count, 7)
+            for pair in clone._pairs:
+                self.assertEqual(pair.neighbors is None, explosion == 'constant')
+            for expected, actual in zip(old, clone.download()):
+                np.testing.assert_array_equal(actual,expected)
+            clone.step()
+            self.assertEqual((run.step_count,clone.step_count), (7,8))
+            clone.close(); run.close()
+
+    def test_changed_remap_geometry_keeps_projected_in_both_representations(self):
+        changed = [(22,22,42,42), (72,68,76,76)]
+        for shifted in (False, True):
+            run, _, _, _, _ = backend(cfg=replace(CFG,ddf_shifting=shifted),
+                refinements=FOUR[:2], explosion='projected_linear')
+            run._step_count = 3
+            before = run.download()
+            clone = run.remap_child(changed)
+            self.assertEqual(clone._explosion, 'projected_linear')
+            self.assertEqual(clone.hierarchy.pairs[1].bounds, tuple(changed))
+            self.assertNotEqual(clone._pairs[1].n_ghosts, run._pairs[1].n_ghosts)
+            for old, new in zip(run._pairs, clone._pairs):
+                self.assertIsNot(old.neighbors, new.neighbors)
+                self.assertIsNotNone(new.neighbors)
+            for expected, actual in zip(before, run.download()):
+                np.testing.assert_array_equal(actual, expected)
+            clone.step()
+            self.assertEqual((run.step_count, clone.step_count), (3, 4))
+            clone.close(); run.close()

tests/unit/test_step_sync.py

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

tests/validation/data/les_flow/source_pins_a135252.json

@@ -0,0 +1,110 @@
+{
+  "base_commit": "a135252ae78b47b57d7e77d19001700b3bedac8b",
+  "production_origin_commit": "a135252ae78b47b57d7e77d19001700b3bedac8b",
+  "source_experiment": "Explicit a135252 uniform SRT LES regression. Sources pin the committed corrected tree; baseline is a synthetic isolated reverse-header control, not the committed a135252 source. Historical profiles and receipts remain unchanged.",
+  "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": "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": "668e5cf14e88bdcfa281c2fa6c75c0f8f8cdbbfb8ba2ca03064ae02ef7ad0057",
+    "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": "a3db46c2e7d27530b2d1436e65cdbed0b72cc7e047bfbeef2a640877271b51aa",
+    "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": "9031a196006a0f74dee4028e6f3bb673bc19ef7959f9f04dde9e01f93293aeb3",
+    "src/CelerisLab/lbm/refinement.py": "100d5e46ac894f6008b61364cdad5d37c2ca7ed949d8d71405572d6366526417",
+    "src/CelerisLab/lbm/refinement_stepper.py": "536cd1d3482591e4ff57c4ae8e888c19e31701687d287aa509129df5b2a43be3",
+    "src/CelerisLab/lbm/refinement_tables.py": "0cbc0f29aabd0ab4ba1f912bd4d0ac2b430246584cb6ddf13dd28e10582accf3",
+    "src/CelerisLab/lbm/stepper.py": "0dcb7cd7ae87e33e3dc67a38b6aa43042fbe673e37d9ad165342ee1d695933ce",
+    "src/CelerisLab/lbm/synchronized_remap.py": "7400fc4beaffd4649d6a9c524d38ce5fd6fa7eeb32eb5332e4030a4b6648f197",
+    "src/CelerisLab/simulation.py": "0a03904dc6de88667d7af699005a289adf8b09f0976cb0d099b041f4830a8f4c"
+  },
+  "candidate_header_sha256": "a3db46c2e7d27530b2d1436e65cdbed0b72cc7e047bfbeef2a640877271b51aa",
+  "baseline_header_sha256": "ff7f000297ee0817e255ad9fcb6a13b647e3ddd5b206518d35f38d2895a9f501"
+}

tests/validation/multigrid/check_projected_pairs.py

@@ -0,0 +1,425 @@
+"""Bounded exact-source CUDA plumbing gate; never a flow-accuracy certificate.
+
+Run this driver externally against exact dev with --mode baseline first, then
+against the experimental source with --mode candidate --baseline <receipt>.
+Only scalar receipts/hashes and actual compiler sidecars are retained. Whole-step
+mass/momentum are diagnostics: no verified FP32 conservation bound is claimed.
+"""
+import argparse
+from contextlib import contextmanager
+from dataclasses import replace
+import hashlib
+import json
+import os
+import shutil
+from pathlib import Path
+import subprocess
+import sys
+
+import numpy as np
+
+BASE = 'ed9d8c937d446b854d6963bd92d718bb4fb98b96'
+PRODUCTION = 'a135252ae78b47b57d7e77d19001700b3bedac8b'
+STOPS = (0, 1, 2, 8)
+OPERATOR_ATOL, OPERATOR_RTOL = 6e-8, 2e-7
+REST_ATOL, REST_RTOL = 3e-7, 1e-6
+REPLAY_ATOL = 3e-6
+E = np.array([(0,0),(1,0),(-1,0),(0,1),(0,-1),(1,1),(-1,-1),(1,-1),(-1,1)])
+W = np.array([4/9]+[1/9]*4+[1/36]*4)
+
+
+def digest(array):
+    """All finite bits (including signed zero) and exact NaN mask; no NaN payloads."""
+    array = np.asarray(array)
+    if np.isinf(array).any():
+        raise AssertionError('infinite population')
+    finite = np.isfinite(array)
+    h = hashlib.sha256()
+    h.update(str((array.shape, array.dtype.str)).encode())
+    h.update(np.ascontiguousarray(finite).tobytes())
+    h.update(np.ascontiguousarray(array[finite]).tobytes())
+    return h.hexdigest()
+
+
+def independent_explosion(coarse, pair, layout, kind):
+    """Coordinate oracle; never reads the production neighbor/offset tables."""
+    ny, nx = pair.coarse_shape
+    legal = np.zeros((ny,nx), bool)
+    legal.ravel()[pair.ghost_coarse] = True
+    legal &= layout.owned & ~pair.covered & (kind == 0)
+    if not legal.ravel()[pair.ghost_coarse].all():
+        raise AssertionError('illegal center donor')
+    out = np.empty((9,len(pair.ghost_fine)))
+    flat = np.asarray(coarse, np.float64).reshape(9,-1)
+    for row, (target, donor) in enumerate(zip(pair.ghost_fine,pair.ghost_coarse)):
+        y,x = divmod(int(donor),nx)
+        gradient = np.zeros((9,2))
+        for axis, (dx,dy) in enumerate(((1,0),(0,1))):
+            ends = [(x-dx,y-dy),(x+dx,y+dy)]
+            if not layout.periodic[axis] and any(not (0<=xx<nx and 0<=yy<ny) for xx,yy in ends):
+                continue
+            ends = [(xx%nx,yy%ny) for xx,yy in ends]
+            if all(legal[yy,xx] for xx,yy in ends):
+                (xm,ym),(xp,yp) = ends
+                gradient[:,axis] = (flat[:,yp*nx+xp]-flat[:,ym*nx+xm])/2
+        for q in range(1,9):
+            gradient[q] -= E[q]*np.dot(E[q],gradient[q])/np.dot(E[q],E[q])
+        fy,fx = divmod(int(target),pair.fine_shape[1])
+        offset = np.array([pair.fine_origin[0]+(fx+.5)/2-(x+.5),
+                           pair.fine_origin[1]+(fy+.5)/2-(y+.5)])
+        out[:,row] = flat[:,int(donor)] + gradient@offset
+    return out
+
+
+def default_options(selector):
+    # Exact dev has no explosion keyword. Baseline and omitted paths pass none.
+    return {} if selector == 'omitted' else {'explosion':selector}
+
+
+def snapshot(run, raw):
+    from CelerisLab.lbm.refinement import ACTIVE, OUTER_GHOST
+    stored, public = raw(run), run.download()
+    hashes, totals = [], np.zeros(3)
+    for level,(values,visible,kind) in enumerate(zip(stored,public,run.hierarchy.kinds)):
+        active = kind == ACTIVE
+        if not np.isfinite(visible[:,active]).all() or not np.isnan(visible[:,~active]).all():
+            raise AssertionError('public ACTIVE/invalid mask mismatch')
+        if not np.isnan(values[:,kind==OUTER_GHOST]).all():
+            raise AssertionError('OUTER owner core is not poisoned')
+        population = visible[:,active].sum(axis=1,dtype=np.float64)*(.25**level)
+        totals += np.r_[population.sum(), E.T@population]
+        hashes.append(dict(raw=digest(values), public=digest(visible), active=int(active.sum())))
+    if not np.isfinite(totals).all():
+        raise AssertionError('nonfinite mass/momentum diagnostic')
+    return dict(step=run.step_count, levels=hashes, mass_momentum=totals.tolist())
+
+
+def initialize(run, rest=False):
+    arrays, origin, spacing = [], np.zeros(2), 1.
+    for index,level in enumerate(run._levels):
+        if index:
+            origin += spacing*np.asarray(run.hierarchy.pairs[index-1].fine_origin)
+            spacing /= 2
+        y,x = np.indices(level.shape[1:],dtype=np.float64)
+        x,y = origin[0]+spacing*(x+.5),origin[1]+spacing*(y+.5)
+        q = np.arange(9)[:,None,None]
+        values = np.broadcast_to(W[:,None,None],level.shape).copy()
+        if not rest:
+            # q-distinct smooth non-equilibrium G; same physical frame on all levels.
+            values += 1e-4*np.sin(.17*x+.23*y+q)
+        arrays.append(values)
+    run.initialize(*arrays)
+
+
+def operator_case(run, index):
+    import pycuda.driver as cuda
+    parent,child = run._levels[index:index+2]
+    geometry,pair = run.hierarchy.pairs[index],run._pairs[index]
+    c = np.random.default_rng(19+index).normal(0,.001,(9,*geometry.coarse_shape))
+    if not run._ddf_shifting:
+        c += W[:,None,None]
+    c = c.astype(np.float32)
+    legal = np.zeros(geometry.coarse_shape,bool)
+    legal.ravel()[geometry.ghost_coarse] = True
+    c[:,~legal] = np.nan
+    packed = parent.layout.to_blocked(c)
+    output = np.full(child.layout.blocked_shape,np.nan,np.float32)
+    cuda.memcpy_htod(parent.current,packed)
+    cuda.memcpy_htod(child.current,output)
+    calls, hop = [],run._hop
+    try:
+        run._hop = lambda *a,**kw: calls.append(a)
+        run._advance(index+1,parent.current,block=run._block,stream=run._stream)
+        run._stream.synchronize()
+    finally:
+        run._hop = hop
+    if calls != [(index+1,child.current,child.spare),(index+1,child.spare,child.current)]:
+        raise AssertionError('explosion changed two-hop launch contract')
+    cuda.memcpy_dtoh(output,child.current)
+    target = child.layout.owner.ravel()[geometry.ghost_fine]
+    slots = np.array([child.layout.slot(target,q) for q in range(9)])
+    actual = output.ravel()[slots]
+    expected = independent_explosion(c,geometry,parent.layout,run.hierarchy.kinds[index])
+    constant = c.reshape(9,-1)[:,geometry.ghost_coarse]
+    if not np.isfinite(actual).all() or np.max(abs(expected-constant)) <= 1e-5:
+        raise AssertionError('operator must exercise finite nonzero projected slopes')
+    np.testing.assert_allclose(actual,expected,atol=OPERATOR_ATOL,rtol=OPERATOR_RTOL)
+    untouched = np.ones(output.size,bool); untouched[slots.ravel()] = False
+    if not np.isnan(output.ravel()[untouched]).all():
+        raise AssertionError('explosion wrote outside its targets')
+    return dict(pair=index,rows=len(target),parent_stride=pair.parent_cells,
+                child_stride=int(np.prod(child.layout.local_shape)),
+                max_error_by_q=np.max(abs(actual-expected),axis=1).tolist(),
+                nonzero_projection=True,poisoned_nondonors=True,untouched_preserved=True)
+
+
+def replay_two_hops(values, pair, layout, kind):
+    """Existing FP64 two-hop algebra, using the independent coordinate oracle."""
+    from tests.unit.test_projected_explosion import collide
+    coarse,fine = (a.copy() for a in values)
+    fine.reshape(9,-1)[:,pair.ghost_fine] = independent_explosion(coarse,pair,layout,kind)
+    for _ in range(2):
+        fine = np.array([np.roll(fine[q],(ey,ex),(0,1)) for q,(ex,ey) in enumerate(E)])
+        fine[:,pair.fine_kind==2] = np.nan
+        active = pair.fine_kind==0
+        fine[:,active] = collide(fine[:,active],pair.tau_fine)
+    out = np.array([np.roll(coarse[q],(ey,ex),(0,1)) for q,(ex,ey) in enumerate(E)])
+    out.reshape(-1)[pair.coalesce_coarse] = fine.reshape(-1)[pair.coalesce_fine].mean(axis=1)
+    out[:,~pair.covered] = collide(out[:,~pair.covered],pair.tau_coarse)
+    out[:,pair.covered] = np.nan
+    return out,fine
+
+
+def run_case(ctx, case, shifted, selector, raw, projected=False, replay=False):
+    from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+    cfg,refs,blocks = case
+    cfg = replace(cfg,velocity=0,ddf_shifting=shifted)
+    with FixedPatchStepper(cfg,context=ctx,refinements=refs,level_blocks=blocks,
+                           **default_options(selector)) as run:
+        operator = [operator_case(run,i) for i in range(len(run._pairs))] if projected else []
+        counts = [0]*len(run._levels)
+        for index,level in enumerate(run._levels):
+            def counted(*a, index=index, kernel=level.kernel, **kw):
+                counts[index] += 1
+                return kernel(*a,**kw)
+            level.kernel = counted
+        def check_counts(stop):
+            if counts != [stop*2**index for index in range(len(counts))]:
+                raise AssertionError('actual level hop counts differ from acoustic schedule')
+        rest = []
+        if projected:
+            initialize(run,rest=True)
+            for stop in STOPS:
+                run.step(stop-run.step_count)
+                check_counts(stop)
+                snapshot(run,raw)  # actual public and raw invalid masks at every stop
+                maximum = 0.
+                for values,kind in zip(run.download(),run.hierarchy.kinds):
+                    active = values[:,kind==0]
+                    np.testing.assert_allclose(active,np.broadcast_to(W[:,None],active.shape),
+                                               atol=REST_ATOL,rtol=REST_RTOL)
+                    maximum = max(maximum,float(np.max(abs(active-W[:,None]))))
+                rest.append(dict(step=stop,max_absolute_error=maximum))
+        counts[:] = [0]*len(counts)
+        initialize(run)
+        reference = tuple(a.astype(np.float64) for a in run.download()) if replay else None
+        snapshots,replay_errors,previous = [],[],0
+        for stop in STOPS:
+            run.step(stop-run.step_count)
+            if replay:
+                for _ in range(stop-previous):
+                    reference = replay_two_hops(reference,run.hierarchy.pairs[0],
+                                               run._levels[0].layout,run.hierarchy.kinds[0])
+                errors = []
+                for actual,expected,kind in zip(run.download(),reference,run.hierarchy.kinds):
+                    np.testing.assert_allclose(actual[:,kind==0],expected[:,kind==0],
+                                               atol=REPLAY_ATOL,rtol=0)
+                    errors.append(np.max(abs(actual[:,kind==0]-expected[:,kind==0]),axis=1).tolist())
+                replay_errors.append(dict(step=stop,max_error_by_level_q=errors))
+            check_counts(stop)
+            snapshots.append(dict(snapshot(run,raw),observed_level_hops=counts.copy())); previous = stop
+        change = (np.array(snapshots[-1]['mass_momentum'])-
+                  np.array(snapshots[0]['mass_momentum'])).tolist()
+        return dict(operator=operator,snapshots=snapshots,rest_fixed_point=rest,
+                    replay_max_errors=replay_errors,mass_momentum_change_diagnostic=change,
+                    conservation_status='diagnostic_only_no_FP32_bound')
+
+
+def compare_snapshots(actual,expected):
+    if len(actual) != len(expected):
+        raise AssertionError('snapshot count differs')
+    for got,want in zip(actual,expected):
+        if (got['step'] != want['step'] or got['levels'] != want['levels']
+                or got.get('observed_level_hops') != want.get('observed_level_hops')):
+            raise AssertionError('owner bits, public values or masks differ')
+
+
+def cases():
+    from CelerisLab.config import LBMConfig
+    from tests.unit.test_hierarchy import CFG, FOUR
+    scalar = (LBMConfig(nx=16,ny=16,viscosity=.02),
+              [(4,4,12,12),(8,8,16,16),(8,8,16,16)],
+              [(4,4),(4,8),(6,4),(4,6)])
+    four = (CFG,FOUR,[(6,8),(8,6),(12,8),(16,12)])
+    concave = (LBMConfig(nx=30,ny=30,viscosity=.1),
+               [[(4,4,5,7),(4,5,7,7)]],[(6,5),(4,6)])
+    return dict(scalar=scalar,four=four,concave=concave)
+
+
+def baseline_receipt(receipt, driver_hash, keys):
+    if (receipt.get('completed') is not True or receipt.get('dry_run') is not False
+            or receipt.get('mode') != 'baseline' or receipt.get('production_revision') != BASE
+            or receipt.get('source',{}).get('revision') != BASE
+            or receipt.get('source',{}).get('dirty') is not False
+            or receipt.get('driver_sha256') != driver_hash
+            or receipt.get('stops') != list(STOPS)
+            or receipt.get('operator_tolerance') != [OPERATOR_ATOL,OPERATOR_RTOL]
+            or receipt.get('rest_tolerance') != [REST_ATOL,REST_RTOL]
+            or receipt.get('replay_tolerance') != [REPLAY_ATOL,0]
+            or set(receipt.get('default_snapshots',{})) != set(keys)
+            or not receipt.get('actual_build_sidecars')):
+        raise ValueError('baseline receipt does not match the frozen exact-dev gate')
+    return receipt
+
+
+def remap_case(ctx, shifted, raw):
+    from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
+    cfg,refs,blocks = cases()['four']
+    with FixedPatchStepper(replace(cfg,velocity=0,ddf_shifting=shifted),context=ctx,
+            refinements=refs[:2],level_blocks=blocks[:3],explosion='projected_linear') as run:
+        initialize(run); run.step(1)
+        before = snapshot(run,raw)
+        with run.remap_child([(22,22,42,42),(72,68,76,76)]) as clone:
+            if clone._explosion != 'projected_linear' or any(p.neighbors is None for p in clone._pairs):
+                raise AssertionError('remap lost projected selection')
+            if clone.step_count != 1: raise AssertionError('remap lost synchronized count')
+            initial = snapshot(clone,raw)
+            clone.step(1)
+            after = snapshot(clone,raw)
+            if clone.step_count != 2: raise AssertionError('remap did not continue count')
+        compare_snapshots([snapshot(run,raw)],[before])
+        return dict(shifted=shifted,source_unchanged=True,initial=initial,continued=after)
+
+
+@contextmanager
+def capture_builds(report, save, baseline=None):
+    from CelerisLab.cuda import compiler_v2 as compiler
+    original = compiler.load_module
+    def load(path):
+        path = Path(path)
+        metadata = json.loads(Path(str(path)+'.json').read_text())
+        if metadata['ptx_sha256'] != hashlib.sha256(path.read_bytes()).hexdigest():
+            raise AssertionError('actual PTX and build sidecar disagree')
+        if baseline and (metadata['arch'],metadata['flags']) not in [
+                (m['arch'],m['flags']) for m in baseline['actual_build_sidecars']]:
+            raise AssertionError('baseline build architecture or flags differ')
+        report['actual_build_sidecars'].append(metadata)
+        save()  # capture AMR metadata before production cleanup removes scratch
+        return original(str(path))
+    compiler.load_module = load
+    try:
+        yield
+    finally:
+        compiler.load_module = original
+
+
+def preflight(repo, revision, mode):
+    git = lambda *a: subprocess.check_output(['git','-C',str(repo),*a],text=True).strip()
+    if git('rev-parse','HEAD') != revision or git('status','--porcelain'):
+        raise ValueError('requires the exact expected clean checkout')
+    source = BASE if mode == 'baseline' else PRODUCTION
+    if mode == 'baseline' and revision != BASE:
+        raise ValueError('baseline must be exact ed9d8c9')
+    if subprocess.call(['git','-C',str(repo),'diff','--quiet',source,'--','src']):
+        raise ValueError('production source differs from the frozen source revision')
+    return source
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument('--repo',type=Path,required=True)
+    parser.add_argument('--expected-revision',required=True)
+    parser.add_argument('--mode',choices=('baseline','candidate'),required=True)
+    parser.add_argument('--baseline',type=Path)
+    parser.add_argument('--output',type=Path,required=True)
+    parser.add_argument('--run-gpu',action='store_true')
+    parser.add_argument('--dry-run',action='store_true')
+    args = parser.parse_args()
+    if args.output.exists(): parser.error('output already exists')
+    selected = os.environ.get('CUDA_VISIBLE_DEVICES','')
+    if not args.dry_run and (not args.run_gpu or not selected or ',' in selected or selected=='-1'):
+        parser.error('--run-gpu and exactly one explicit CUDA_VISIBLE_DEVICES are required')
+    repo = args.repo.resolve()
+    source = preflight(repo,args.expected_revision,args.mode)
+    baseline = None
+    if args.mode == 'candidate':
+        if args.baseline is None: parser.error('candidate requires --baseline receipt')
+        baseline = json.loads(args.baseline.read_text())
+    sys.path[:0] = [str(repo/'src'),str(repo),str(repo/'tests/validation/multigrid'),
+                   str(repo/'tests/validation/taylor_green')]
+    import CelerisLab
+    if Path(CelerisLab.__file__).resolve() != repo/'src/CelerisLab/__init__.py':
+        raise RuntimeError('wrong imported checkout')
+    from run_case import provenance
+    driver_hash = hashlib.sha256(Path(__file__).read_bytes()).hexdigest()
+    keys = [f'{name}:shifted={shifted}' for name in cases() for shifted in (False,True)]
+    if baseline is not None: baseline_receipt(baseline,driver_hash,keys)
+    report = dict(mode=args.mode,production_revision=source,source=provenance(repo),
+                  driver_sha256=driver_hash,
+                  device_selection=selected,completed=False,dry_run=args.dry_run,
+                  operator_tolerance=[OPERATOR_ATOL,OPERATOR_RTOL],
+                  rest_tolerance=[REST_ATOL,REST_RTOL],replay_tolerance=[REPLAY_ATOL,0],stops=STOPS,
+                  conservation_status='diagnostic_only_no_FP32_bound',
+                  actual_build_sidecars=[],default_snapshots={},default_equivalence=[],
+                  projected_operator=[],union_multilevel_smoke=[],changed_remap=[])
+    args.output.parent.mkdir(parents=True,exist_ok=True)
+    def save():
+        text = json.dumps(report,indent=2,allow_nan=False)+'\n'
+        if len(text.encode()) > 8*1024*1024: raise RuntimeError('receipt exceeds 8 MiB cap')
+        temporary = args.output.with_suffix(args.output.suffix+'.tmp')
+        temporary.write_text(text); os.replace(temporary,args.output)
+    save()
+    if args.dry_run:
+        print('Source preflight only; no CUDA imports or execution.'); return
+    try:
+        from CelerisLab.cuda.context import CudaContext
+        from test_refinement_blocked_base import raw
+        with CudaContext(0) as ctx, capture_builds(report,save,baseline):
+            import pycuda.driver as cuda
+            report['device'] = dict(name=ctx.device.name(),pci_bus_id=ctx.device.pci_bus_id(),
+                                    capability=list(ctx.device.compute_capability()),
+                                    driver=cuda.get_driver_version(),runtime=list(cuda.get_version()))
+            nvcc = shutil.which('nvcc') or '/usr/local/cuda/bin/nvcc'
+            report['nvcc_version'] = subprocess.check_output([nvcc,'--version'],text=True).strip()
+            if baseline and (report['device'] != baseline.get('device')
+                             or report['nvcc_version'] != baseline.get('nvcc_version')
+                             or selected != baseline.get('device_selection')):
+                raise ValueError('baseline GPU/device selection or toolchain differs')
+            for name,case in cases().items():
+                for shifted in (False,True):
+                    key = f'{name}:shifted={shifted}'
+                    report['active_case'] = f'default:{key}'
+                    result = run_case(ctx,case,shifted,'omitted',raw)
+                    report['default_snapshots'][key] = result['snapshots']
+                    if baseline:
+                        compare_snapshots(result['snapshots'],baseline['default_snapshots'][key])
+                        explicit = run_case(ctx,case,shifted,'constant',raw)
+                        compare_snapshots(explicit['snapshots'],result['snapshots'])
+                        report['default_equivalence'].append(dict(case=key,passed=True))
+                    save()
+            if baseline:
+                for name,case in cases().items():
+                    for shifted in (False,True):
+                        report['active_case'] = f'projected:{name}:shifted={shifted}'
+                        result = run_case(ctx,case,shifted,'projected_linear',raw,True,name=='concave')
+                        cfg,refs,blocks = case
+                        if name == 'four':
+                            variant_refs = [list(reversed(rects)) for rects in refs]
+                        elif name == 'concave':
+                            variant_refs = [[(4,5,7,7),(4,4,5,5)]]  # same union, disjoint decomposition
+                        else:
+                            variant_refs = refs
+                        variant = (cfg,variant_refs,[(2,2)]*len(blocks))
+                        alternate = run_case(ctx,variant,shifted,'projected_linear',raw,True)
+                        compare_snapshots(alternate['snapshots'],result['snapshots'])
+                        report['projected_operator'].append(dict(case=name,shifted=shifted,rows=result['operator']))
+                        report['union_multilevel_smoke'].append(dict(case=name,shifted=shifted,result=result,
+                                                                   alternate_bits_identical=True))
+                        save()
+                for shifted in (False,True):
+                    report['active_case'] = f'changed-remap:shifted={shifted}'
+                    report['changed_remap'].append(remap_case(ctx,shifted,raw)); save()
+                expected_builds = {(s['arch'],tuple(s['flags'])) for s in baseline['actual_build_sidecars']}
+                actual_builds = {(s['arch'],tuple(s['flags'])) for s in report['actual_build_sidecars']}
+                if actual_builds != expected_builds:
+                    raise AssertionError('baseline build architecture or compiler flags differ')
+            report['completed'] = True
+            save()
+    except Exception as error:
+        report['error'] = f'{type(error).__name__}: {error}'
+        save(); raise
+    print(json.dumps(dict(completed=True,mode=args.mode,receipt=str(args.output))))
+
+
+if __name__ == '__main__':
+    main()

tests/validation/run_uniform_les_flow.py

@@ -62,6 +62,7 @@ def decode_macros(data, w, cx, cy):
 EXPERIMENTS = {
     "historical": "source_pins.json",
     "dev8403956": "source_pins_dev8403956.json",
+    "a135252": "source_pins_a135252.json",
 }
 
 
@@ -69,11 +70,10 @@ def source_guard(root, variant, experiment="historical"):
     manifest_path = HERE / "data/les_flow" / EXPERIMENTS[experiment]
     manifest = json.loads(manifest_path.read_text())
     for name, sha in manifest["sources"].items():
-        expected = (
-            manifest["candidate_header_sha256"]
-            if name == HEADER and variant == "corrected"
-            else sha
-        )
+        expected = sha
+        if name == HEADER:
+            expected = (manifest["candidate_header_sha256"] if variant == "corrected"
+                        else manifest.get("baseline_header_sha256", sha))
         if hashlib.sha256((root / name).read_bytes()).hexdigest() != expected:
             raise RuntimeError("Source pin mismatch: " + name)
     identity = dict(

tests/validation/taylor_green/assess_projected_pairs.py

@@ -0,0 +1,208 @@
+"""New-source four-case repeat of 11901a7; frozen oracle and gates are unchanged."""
+
+import argparse
+import json
+import math
+from pathlib import Path
+
+from assess import BASELINE_FACTOR, FLOORS
+from oracle import LIMITS, TIMES, U, rectangles
+from run_case import (INITIALIZATION, sample_steps, NUMERICAL_REVISION, SOURCE_SHA256,
+                      BUILD_CAPTURE_SHA256, REFERENCE_HARNESS)
+
+# Host-generated snapshots of a135252 at the frozen two local configurations.
+# These describe expected inputs; actual sidecars are captured before cleanup.
+BUILD_SOURCES = {
+    128: "803177416b7599ab89cc5afd35cf4a28a673b2f4bb8e8a6a1abc0b6d77610873",
+    256: "17ec5e60f87b840d3d13b8778bc8a3aecf485469c5a061af46c5cc19b73396cf",
+}
+CASES = ((100, 128, 1, "constant"), (100, 256, 1, "constant"),
+         (100, 128, 2, "constant"), (100, 128, 2, "projected_linear"))
+NUMERICS = dict(lattice_model="D2Q9", collision="SRT", store_precision="FP32",
+                ddf_shifting=True, streaming="double_buffer", les_enabled=False,
+                boundary="periodic", block_shape=[8, 8],
+                sample_phase="postcollision G at root synchronization")
+
+
+def finite(value):
+    return isinstance(value, (int, float)) and not isinstance(value, bool) and math.isfinite(value)
+
+
+def finite_tree(value):
+    if isinstance(value, dict):
+        return all(finite_tree(v) for v in value.values())
+    if isinstance(value, list):
+        return all(finite_tree(v) for v in value)
+    return not isinstance(value, float) or math.isfinite(value)
+
+
+def at_steps(rows, steps, n):
+    return len(rows) == len(steps) and all(
+        type(r.get("step")) is int and r["step"] == step and finite(r.get("tstar"))
+        and abs(r["tstar"] - step * U / n) <= 3e-17
+        for r, step in zip(rows, steps)
+    )
+
+
+def maximum(rows, metric):
+    values = [r.get(metric) for r in rows]
+    return max(values) if values and all(finite(v) and v >= 0 for v in values) else None
+
+
+def valid_builds(record, n, depth):
+    builds = record.get("actual_build_sidecars", [])
+    if not isinstance(builds, list) or len(builds) != depth:
+        return False
+    for level, build in enumerate(builds):
+        if not isinstance(build, dict):
+            return False
+        size = n * 2**level
+        cfg = build.get("config", {})
+        expected = dict(nx=10, ny=10, nz=1, lattice_model="D2Q9", collision="SRT",
+                        store_precision="FP32", streaming="double_buffer", ddf_shifting=True,
+                        les_enabled=False, velocity=0.0, viscosity=U*size/100, rho=1.0,
+                        omega=1/(.5+3*U*size/100), threads_per_block=256)
+        ptx = build.get("ptx_sha256", "")
+        if (build.get("source_sha256") != BUILD_SOURCES[size]
+                or build.get("arch") != "sm_70" or build.get("flags") != ["--use_fast_math"]
+                or build.get("n_objects") != 0 or not isinstance(cfg, dict)
+                or any(cfg.get(k) != value for k, value in expected.items())
+                or not isinstance(ptx, str) or len(ptx) != 64
+                or any(c not in "0123456789abcdef" for c in ptx)):
+            return False
+    return True
+
+
+def assess(records, expected_revision=None):
+    indexed, duplicates = {}, []
+    for record in records:
+        key = tuple(record.get(k) for k in ("re_L", "n", "depth", "explosion"))
+        if key in indexed:
+            duplicates.append(str(key))
+        else:
+            indexed[key] = record
+    missing = [str(k) for k in CASES if k not in indexed]
+    checks = dict(complete_matrix=not missing, no_extra_cases=set(indexed) <= set(CASES),
+                  no_duplicate_cases=not duplicates, finite_receipts=finite_tree(records))
+    sources = [r.get("source", {}) for r in records]
+    checks["same_source_and_harness"] = bool(sources) and len({
+        json.dumps(s, sort_keys=True) for s in sources}) == 1
+    checks["clean_a135252_numerical_source"] = bool(sources) and all(
+        s.get("dirty") is False and s.get("source_tree_sha256") == SOURCE_SHA256
+        and bool(s.get("revision")) and bool(s.get("harness_sha256"))
+        and s.get("build_capture_sha256") == BUILD_CAPTURE_SHA256
+        and (expected_revision is None or s.get("revision") == expected_revision) for s in sources)
+    runtimes = [r.get("runtime_identity") for r in records]
+    checks["same_recorded_runtime"] = bool(runtimes) and all(
+        isinstance(v, dict) and all(v.get(k) for k in
+            ("device", "nvcc_path", "nvcc_version", "python", "numpy", "pycuda"))
+        and isinstance(v["device"], dict) and v["device"].get("capability") == [7, 0]
+        and all(v["device"].get(k) for k in ("name", "pci_bus_id", "driver", "runtime"))
+        for v in runtimes
+    ) and len({json.dumps(v, sort_keys=True) for v in runtimes}) == 1
+    devices = [r.get("cuda_visible_devices") for r in records]
+    checks["same_explicit_device"] = bool(devices) and all(
+        isinstance(d, str) and d and "," not in d and d != "-1" for d in devices
+    ) and len(set(devices)) == 1
+    enabled = [r.get("early_pressure", {}).get("enabled") for r in records]
+    checks["same_early_sampling"] = bool(enabled) and all(type(v) is bool for v in enabled) and len(set(enabled)) == 1
+    shared_valid = all(checks.values())
+    complete, case_valid, early = {}, {}, []
+    for key in CASES:
+        if key not in indexed:
+            continue
+        r = indexed[key]
+        _, n, depth, scheme = key
+        label = f"{n}/{depth}/{scheme}"
+        rows = r.get("rows", [])
+        main_steps, early_steps = sample_steps(n, True)
+        complete[key] = r.get("completed") is True and at_steps(rows, main_steps, n)
+        checks[label + " complete_times"] = complete[key]
+        checks[label + " scales"] = (
+            r.get("U") == U and finite(r.get("viscosity"))
+            and abs(r["viscosity"] - U * n / 100) < 1e-15
+            and r.get("tau_by_level") == [0.5 + 3 * U * n / 100 * 2**l for l in range(depth)]
+            and r.get("mach") == U * math.sqrt(3)
+        )
+        checks[label + " actual_builds"] = valid_builds(r, n, depth)
+        checks[label + " reference"] = (r.get("numerical_revision") == NUMERICAL_REVISION
+                                         and r.get("reference_harness") == REFERENCE_HARNESS)
+        checks[label + " numerics"] = r.get("numerics") == NUMERICS and r.get("initialization") == INITIALIZATION
+        runtime = r.get("runtime_source", {})
+        package = Path(runtime.get("checkout", "")) / "src/CelerisLab"
+        checks[label + " runtime_source"] = bool(runtime.get("checkout")) and all(
+            runtime.get(k) == str(package / tail) for k, tail in
+            (("module", "__init__.py"), ("compiler", "cuda/compiler_v2.py"), ("kernel_root", "lbm/kernels")))
+        expected_patch = [] if depth == 1 else rectangles(n, depth)
+        checks[label + " patch"] = json.dumps(r.get("refinements")) == json.dumps(expected_patch)
+        cells = [n*n] if depth == 1 else [8704, 30720]
+        hops = [16*n * 2**l for l in range(depth)]
+        work = r.get("workload", {})
+        checks[label + " workload"] = (
+            work.get("active_cells_by_level") == cells and work.get("level_hops") == hops
+            and work.get("root_intervals") == 16*n
+            and work.get("active_collision_updates_by_level") == [c*h for c, h in zip(cells, hops)]
+            and work.get("active_collision_updates") == sum(c*h for c, h in zip(cells, hops))
+        )
+        checks[label + " absolute"] = bool(rows) and all(
+            finite(row.get(metric)) and 0 <= row[metric] <= limit
+            for row in rows for metric, limit in LIMITS.items())
+        trace = r.get("early_pressure", {})
+        trace_rows = trace.get("rows", [])
+        peak = maximum(trace_rows, "pressure_l2")
+        trace_complete = at_steps(trace_rows, early_steps, n) and peak is not None
+        checks[label + " early_trace"] = (
+            trace_complete and trace.get("sampled_max_pressure_l2") == peak
+            and len(rows) >= 2 and all(trace_rows[i]["pressure_l2"] == rows[j].get("pressure_l2")
+                                      for i, j in ((0, 0), (-1, 1)))
+            if trace.get("enabled") is True else
+            trace.get("enabled") is False and trace_rows == [] and trace.get("sampled_max_pressure_l2") is None
+        )
+        early.append(dict(n=n, depth=depth, explosion=scheme, enabled=trace.get("enabled"),
+                          samples=len(trace_rows), complete=trace_complete,
+                          sampled_max_pressure_l2=peak))
+        case_valid[key] = all(value for name, value in checks.items() if name.startswith(label + " "))
+
+    def worst(key, metric):
+        return maximum(indexed[key]["rows"][1:], metric) if complete.get(key) else None
+
+    comparisons, scheme_results = [], {}
+    for scheme in ("constant", "projected_linear"):
+        key = (100, 128, 2, scheme)
+        for metric, floor in FLOORS.items():
+            amr, coarse, fine = (worst(k, metric) for k in (key, CASES[0], CASES[1]))
+            budget = None if coarse is None else max(BASELINE_FACTOR * coarse, floor)
+            budget = budget if finite(budget) else None
+            ratio = None if amr is None or fine is None else amr / max(fine, 1e-30)
+            passed = amr is not None and budget is not None and amr <= budget
+            checks[f"baseline {scheme}/{metric}"] = passed
+            comparisons.append(dict(explosion=scheme, metric=metric, amr=amr,
+                                    uniform128=coarse, baseline_budget=budget, passed=passed,
+                                    uniform256=fine, amr_to_uniform256=ratio if finite(ratio) else None))
+        baseline_passed = all(checks[f"baseline {scheme}/{metric}"] for metric in FLOORS)
+        valid = shared_valid and all(case_valid.get(k, False) for k in (CASES[0], CASES[1], key))
+        scheme_results[scheme] = dict(passed=valid and baseline_passed,
+                                     valid_case_and_controls=valid, baseline_passed=baseline_passed)
+    return dict(passed=all(checks.values()), missing=missing, duplicates=duplicates, checks=checks,
+                comparisons=comparisons, scheme_results=scheme_results, early_pressure=early,
+                numerical_revision=NUMERICAL_REVISION, source_tree_sha256=SOURCE_SHA256,
+                reference_harness=REFERENCE_HARNESS, expected_revision=expected_revision,
+                floors=FLOORS, baseline_factor=BASELINE_FACTOR,
+                main_statistic="maximum of four post-initial samples at t*=.12,.24,.36,.48",
+                early_statistic="sampled pressure maximum only, t*=k*.001875, k=0..64; not an all-step or continuous maximum",
+                finest_control="uniform256 and AMR/ uniform256 ratios are descriptive, not acceptance gates",
+                scope="2D periodic Re100 smooth-decay diagnostic; no new algorithm, convergence-order or general AMR qualification")
+
+
+if __name__ == "__main__":
+    p = argparse.ArgumentParser(description=__doc__)
+    p.add_argument("files", type=Path, nargs="+")
+    p.add_argument("--output", type=Path, required=True)
+    p.add_argument("--expected-revision", required=True)
+    args = p.parse_args()
+    if args.output.exists():
+        p.error("Output already exists; choose a new path to retain prior evidence")
+    result = assess([json.loads(path.read_text()) for path in args.files], args.expected_revision)
+    args.output.write_text(json.dumps(result, indent=2, allow_nan=False) + "\n")
+    print(json.dumps(result, indent=2, allow_nan=False))
+    raise SystemExit(0 if result["passed"] else 1)

tests/validation/taylor_green/run_case.py

@@ -1,9 +1,10 @@
-"""One bounded GPU case. No production edits, plots, checkpoints or field dumps."""
+"""One bounded GPU case, with explicit opt-in experimental explosion."""
 
 import argparse
 import hashlib
 import json
 import os
+import shutil
 from pathlib import Path
 import subprocess
 import sys
@@ -12,6 +13,35 @@ import time
 import numpy as np
 from oracle import U, TIMES, LIMITS, initial, metrics, rectangles
 
+NUMERICAL_REVISION = "a135252ae78b47b57d7e77d19001700b3bedac8b"
+SOURCE_SHA256 = "a5dbfd804befce868067d79a23e5fd47dce7f505b3f9b1bb2145fb57cb5f2178"
+BUILD_CAPTURE_SHA256 = "3dbaa6547cb0a4cb1541054f3b64ad8b6e18cc70527ac6166e56b944e9bd464a"
+REFERENCE_HARNESS = "11901a7f77557a901e017a171eeae1758d584fa5"
+INITIALIZATION = "physical FP64 equilibrium G at actual cell centers; rho=1+3p"
+
+
+def sample_steps(n, early_pressure=False):
+    """Root synchronization steps; early ladder matches the 707afe6 diagnostic."""
+    main = [round(tstar * n / U) for tstar in TIMES]
+    early = list(range(0, 4 * n + 1, n // 16)) if early_pressure else []
+    return main, early
+
+
+def runtime_source(repo):
+    """Fail before CUDA if Python or the compiler snapshot comes from another tree."""
+    import CelerisLab
+    from CelerisLab.cuda import compiler_v2
+
+    paths = dict(module=Path(CelerisLab.__file__).resolve(),
+                 compiler=Path(compiler_v2.__file__).resolve(),
+                 kernel_root=Path(compiler_v2.kernel_path("blocks.cu")).resolve().parent)
+    package = repo / "src/CelerisLab"
+    expected = dict(module=package / "__init__.py", compiler=package / "cuda/compiler_v2.py",
+                    kernel_root=package / "lbm/kernels")
+    if paths != expected:
+        raise RuntimeError("Wrong imported source or CUDA source/include root")
+    return dict(checkout=str(repo), **{key: str(path) for key, path in paths.items()})
+
 
 def provenance(repo):
     git = lambda *a: subprocess.check_output(
@@ -34,6 +64,46 @@ def provenance(repo):
     )
 
 
+def verify_source(repo, revision):
+    """Exact clean harness commit, fixed production tree, and reviewed capture helper."""
+    import oracle
+    from tests.validation.multigrid import check_projected_pairs as gate
+
+    harness = repo / "tests/validation/taylor_green"
+    if (Path(__file__).resolve() != harness / "run_case.py"
+            or Path(oracle.__file__).resolve() != harness / "oracle.py"):
+        raise RuntimeError("Wrong runner or oracle checkout")
+    helper = repo / "tests/validation/multigrid/check_projected_pairs.py"
+    if (Path(gate.__file__).resolve() != helper
+            or hashlib.sha256(helper.read_bytes()).hexdigest() != BUILD_CAPTURE_SHA256):
+        raise RuntimeError("Wrong build-capture helper")
+    if gate.preflight(repo, revision, "candidate") != NUMERICAL_REVISION:
+        raise RuntimeError("Wrong production revision")
+    result = provenance(repo)
+    if result["source_tree_sha256"] != SOURCE_SHA256:
+        raise RuntimeError("Wrong numerical source fingerprint")
+    result["build_capture_sha256"] = BUILD_CAPTURE_SHA256
+    return result
+
+
+def build_capture(report, save):
+    from tests.validation.multigrid.check_projected_pairs import capture_builds
+    return capture_builds(report, save)
+
+
+def runtime_identity(ctx):
+    import pycuda
+    import pycuda.driver as cuda
+
+    nvcc = shutil.which("nvcc") or "/usr/local/cuda/bin/nvcc"
+    return dict(device=dict(name=ctx.device.name(), pci_bus_id=ctx.device.pci_bus_id(),
+                            capability=list(ctx.device.compute_capability()),
+                            driver=cuda.get_driver_version(), runtime=list(cuda.get_version())),
+                nvcc_path=str(Path(nvcc).resolve()),
+                nvcc_version=subprocess.check_output([nvcc, "--version"], text=True).strip(),
+                python=sys.version, numpy=np.__version__, pycuda=pycuda.VERSION_TEXT)
+
+
 def config(n, re):
     from CelerisLab.config import LBMConfig
 
@@ -71,19 +141,31 @@ def workload(masks, base_steps):
 def main():
     p = argparse.ArgumentParser(description=__doc__)
     p.add_argument("--repo", type=Path, required=True)
+    p.add_argument("--expected-revision", required=True)
+    p.add_argument("--run-gpu", action="store_true")
+    p.add_argument("--dry-run", action="store_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)
+    p.add_argument("--explosion", choices=["constant", "projected_linear"], default="constant")
+    p.add_argument("--early-pressure", action="store_true",
+                   help="also save 65 scalar pressure samples at t*=k*.001875, k=0..64")
     args = p.parse_args()
-    if not os.environ.get("CUDA_VISIBLE_DEVICES"):
-        p.error("Explicit physical CUDA_VISIBLE_DEVICES selection required")
+    if args.output.exists():
+        p.error("Output already exists; choose a new path to retain prior evidence")
+    if args.explosion != "constant" and args.depth != 2:
+        p.error("projected_linear requires depth 2")
+    selected = os.environ.get("CUDA_VISIBLE_DEVICES", "")
+    if not args.dry_run and (not args.run_gpu or not selected or "," in selected or selected == "-1"):
+        p.error("--run-gpu and exactly one explicit CUDA_VISIBLE_DEVICES are 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")
+    sys.path[:0] = [str(repo / "src"), str(repo)]
+    source = verify_source(repo, args.expected_revision)
+    runtime = runtime_source(repo)
+    if args.dry_run:
+        print(json.dumps(dict(source=source, runtime_source=runtime, dry_run=True)))
+        return 0
     from CelerisLab.cuda.context import CudaContext
     from CelerisLab.lbm.blocks import BlockLayout2D
     from CelerisLab.lbm.block_stepper import PeriodicBlockStepper
@@ -96,9 +178,17 @@ def main():
     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]
+    main_steps, early_steps = sample_steps(args.n, args.early_pressure)
     report = dict(
         schema=1,
-        source=provenance(repo),
+        source=source,
+        numerical_revision=NUMERICAL_REVISION,
+        reference_harness=REFERENCE_HARNESS,
+        actual_build_sidecars=[],
+        runtime_identity=None,
+        runtime_source=runtime,
+        historical_baseline_revision="9805383f9c50d867e8555471baf5dbda8333a652",
+        explosion=args.explosion,
         n=args.n,
         re_L=args.re,
         depth=args.depth,
@@ -107,21 +197,35 @@ def main():
         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"],
+        initialization=INITIALIZATION,
+        numerics=dict(lattice_model=cfg.lattice_model, collision=cfg.collision,
+                      store_precision=cfg.store_precision, ddf_shifting=cfg.ddf_shifting,
+                      streaming=cfg.streaming, les_enabled=cfg.les_enabled,
+                      boundary="periodic", block_shape=[8, 8],
+                      sample_phase="postcollision G at root synchronization"),
+        cuda_visible_devices=selected,
         device_id=0,
         limits=LIMITS,
         workload=workload(masks, round(TIMES[-1] * args.n / U)),
         rows=[],
+        early_pressure=dict(enabled=args.early_pressure, rows=[],
+                            sampled_max_pressure_l2=None),
         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"
-    )
+    def save():
+        temporary = args.output.with_suffix(args.output.suffix + ".tmp")
+        text = json.dumps(report, indent=2, allow_nan=False) + "\n"
+        if len(text.encode()) > 1024 * 1024:
+            raise RuntimeError("Scalar receipt exceeds 1 MiB cap")
+        temporary.write_text(text)
+        temporary.replace(args.output)
+
     save()
     started = time.monotonic()
-    with CudaContext(0) as ctx:
+    with CudaContext(0) as ctx, build_capture(report, save):
+        report["runtime_identity"] = runtime_identity(ctx)
+        save()
         if args.depth == 1:
             layout = BlockLayout2D.build(args.n, args.n, 8, 8, periodic=(True, True))
             run = PeriodicBlockStepper(cfg, layout, context=ctx)
@@ -132,6 +236,7 @@ def main():
                 coarse_boundary="periodic",
                 refinements=refinements,
                 level_blocks=[(8, 8)] * args.depth,
+                explosion=args.explosion,
             )
         with run:
             run.initialize(
@@ -141,24 +246,30 @@ def main():
                 ]
             )
             initial_mass = None
-            for tstar in TIMES:
-                step = round(tstar * args.n / U)
+            for step in sorted(set(main_steps + early_steps)):
                 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)
+                if not all(np.isfinite(value) for value in row.values()):
+                    raise FloatingPointError("nonfinite sample diagnostics")
                 row["step"] = step
                 if initial_mass is None:
                     initial_mass = row["mass"]
-                report["rows"].append(row)
+                if step in main_steps:
+                    report["rows"].append(row)
+                if step in early_steps:
+                    trace = report["early_pressure"]
+                    trace["rows"].append({key: row[key] for key in ("step", "tstar", "pressure_l2")})
+                    trace["sampled_max_pressure_l2"] = max(r["pressure_l2"] for r in trace["rows"])
                 save()
                 print(
                     json.dumps(
                         dict(
-                            event="PROGRESS",
+                            event="PROGRESS" if step in main_steps else "EARLY_PRESSURE",
                             n=args.n,
                             re_L=args.re,
                             depth=args.depth,
-                            **row,
+                            **(row if step in main_steps else report["early_pressure"]["rows"][-1]),
                         )
                     ),
                     flush=True,
@@ -168,7 +279,9 @@ def main():
                     # 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)
+                    probe_metrics = metrics(probe, masks, 1, args.n, cfg.viscosity, initial_mass)
+                    if not all(np.isfinite(value) for value in probe_metrics.values()):
+                        raise FloatingPointError("nonfinite READY diagnostics")
                     report["ready_after_step"] = 1
                     save()
                     print(

tests/validation/taylor_green/test_projected_pairs.py

@@ -0,0 +1,434 @@
+"""Host geometry, receipt-corruption tests and mocked scheduling; no CUDA/flow."""
+
+from contextlib import nullcontext, redirect_stdout
+from copy import deepcopy
+from dataclasses import asdict, replace
+import hashlib
+import io
+import json
+from pathlib import Path
+import subprocess
+import sys
+import tempfile
+from types import SimpleNamespace
+import unittest
+from unittest.mock import patch
+
+import numpy as np
+
+import run_case
+from assess_projected_pairs import (BUILD_SOURCES, CASES, NUMERICS, NUMERICAL_REVISION,
+                                    SOURCE_SHA256, assess)
+from oracle import LIMITS, TIMES, U, rectangles
+
+REPO = Path(__file__).resolve().parents[3]
+OLD_REVISION = "519473a61c7b319183aa54e55b9f097499dde0bf"
+MOCK_REVISION = "f" * 40
+RUNTIME = dict(device=dict(name="host mock", pci_bus_id="mock", capability=[7, 0],
+                           driver=12040, runtime=[12, 4, 0]), nvcc_path="/mock/nvcc",
+               nvcc_version="mock nvcc", python="mock Python", numpy=np.__version__, pycuda="mock")
+
+
+def source():
+    return dict(revision=MOCK_REVISION, dirty=False, source_tree_sha256=SOURCE_SHA256,
+                harness_sha256={"run_case.py": "a" * 64},
+                build_capture_sha256=run_case.BUILD_CAPTURE_SHA256)
+
+
+def build_metadata(size):
+    cfg = replace(run_case.config(size, 100), nx=10, ny=10)
+    cfg.validate()
+    return dict(source_sha256=BUILD_SOURCES[size], config=asdict(cfg), n_objects=0,
+                arch="sm_70", flags=["--use_fast_math"], ptx_sha256="b" * 64)
+
+
+def records(early=False):
+    from CelerisLab.lbm.hierarchy import build_hierarchy
+    from CelerisLab.lbm.refinement import ACTIVE
+
+    out = []
+    for re, n, depth, scheme in CASES:
+        cfg = run_case.config(n, re)
+        refinements = [] if depth == 1 else rectangles(n, depth)
+        h = build_hierarchy(cfg, refinements)
+        main_steps, early_steps = run_case.sample_steps(n, early)
+        out.append(dict(
+            re_L=re, n=n, depth=depth, explosion=scheme, U=U, mach=float(U*np.sqrt(3)),
+            viscosity=cfg.viscosity, tau_by_level=[0.5+3*c.viscosity for c in h.configs],
+            refinements=refinements, numerics=NUMERICS.copy(), initialization=run_case.INITIALIZATION,
+            source=source(), numerical_revision=NUMERICAL_REVISION,
+            reference_harness=run_case.REFERENCE_HARNESS,
+            actual_build_sidecars=[build_metadata(n * 2**level) for level in range(depth)],
+            runtime_identity=deepcopy(RUNTIME), cuda_visible_devices="0",
+            runtime_source=run_case.runtime_source(REPO), completed=True,
+            workload=run_case.workload([k == ACTIVE for k in h.kinds], 16*n),
+            rows=[dict(step=s, tstar=s*U/n, **{k: 0.0 for k in LIMITS}) for s in main_steps],
+            early_pressure=dict(enabled=early, rows=[
+                dict(step=s, tstar=s*U/n, pressure_l2=.01 if s not in main_steps else 0.0) for s in early_steps],
+                sampled_max_pressure_l2=.01 if early else None)))
+    return out
+
+
+class ProjectedPairsTests(unittest.TestCase):
+    def test_exact_geometry_work_and_samples(self):
+        from CelerisLab.lbm.hierarchy import build_hierarchy, leaf_partition
+
+        data = records(True)
+        self.assertEqual(sum(r["workload"]["active_collision_updates"] for r in data), 589299712)
+        for r in data:
+            n, depth = r["n"], r["depth"]
+            h = build_hierarchy(run_case.config(n, 100), r["refinements"])
+            np.testing.assert_array_equal(leaf_partition(h), np.ones((n, n)))
+            main, early = run_case.sample_steps(n, True)
+            self.assertEqual(main, [0, 4*n, 8*n, 12*n, 16*n])
+            self.assertEqual(early, [k*(n//16) for k in range(65)])
+            np.testing.assert_allclose(np.array(early)*U/n, np.arange(65)*.001875, rtol=0, atol=3e-17)
+            if depth == 2:
+                self.assertEqual(r["refinements"], [[(16, 16, 112, 96)]])
+        self.assertTrue(assess(data)["passed"])
+        self.assertTrue(assess(records())["passed"])
+
+    def test_exact_old519_geometry_and_table_parity(self):
+        from CelerisLab.lbm.blocks import BlockLayout2D
+        from CelerisLab.lbm.hierarchy import build_hierarchy
+        from CelerisLab.lbm.refinement_tables import projected_explosion_tables
+
+        # Frozen from executing exact 519473a61c7b319183aa54e55b9f097499dde0bf
+        # geometry and table definitions, after checking unchanged config,
+        # descriptors and blocks. No historical Git object is needed at runtime.
+        # Historical source SHA256, refinement.py then refinement_tables.py:
+        # 08c2a0557276a98cbcaa459d9bb7a9735d8bbbf64a6114c57af2fda61481a687
+        # dabff948bb30496384d1a52c666b0edbc2085f9c61ed1bbc39f748502629113d
+        frozen = {
+            "covered": ("|b1", (128, 128), "14a1b5be54e3309bea825a00a653105951f0c805622b877a86dd453d65556df3"),
+            "fine_kind": ("|u1", (256, 256), "891ff72c3d60686c62b1f35ffe59d054ac23b8f9d59aeb5782bc941202794cc0"),
+            "ghost_fine": ("<u8", (2880,), "2d0311ac0feb1c0e1e2b6bcf806fe616b8217513de218407267dcd7652287579"),
+            "ghost_coarse": ("<u4", (2880,), "5a3b305a0f841c8d05a13f105fda02b43ef875ab1321f9bca579f2b8ae313233"),
+            "coalesce_coarse": ("<u8", (1060,), "a1d88c60f282153dcf2bb497c59d777968b6db773d061d8fea08163c98155967"),
+            "coalesce_fine": ("<u8", (1060, 4), "08914c5d17c1ec8e2bfd234f581e4b3283de26296ae7d80403385feed7b89532"),
+            "neighbors": ("<u8", (2880, 4), "cfef3a07103241e3dfff24fd65ab967afcf20520f699574b191044d47cf5d1f4"),
+            "offsets": ("<f4", (2880, 2), "c98ba658d5f2cc18605e029fbe1da448ded3fd2436edcaf8ea097fa46aeb1f2d"),
+        }
+        hierarchy = build_hierarchy(run_case.config(128, 100), rectangles(128, 2))
+        pair = hierarchy.pairs[0]
+        self.assertEqual((pair.coarse_shape, pair.fine_shape, pair.fine_origin),
+                         ((128, 128), (256, 256), (0.0, 0.0)))
+        self.assertEqual(len(np.unique(pair.ghost_coarse)), 720)
+        self.assertTrue((hierarchy.kinds[0].ravel()[pair.ghost_coarse] == 0).all())
+        layout = BlockLayout2D.build(128, 128, 8, 8, periodic=(True, True))
+        neighbors, offsets = projected_explosion_tables(pair, layout, hierarchy.kinds[0])
+        arrays = dict(neighbors=neighbors, offsets=offsets)
+        arrays.update({name: getattr(pair, name) for name in frozen if name not in arrays})
+        for name, array in arrays.items():
+            with self.subTest(array=name):
+                self.assertEqual((array.dtype.str, array.shape, hashlib.sha256(array.tobytes()).hexdigest()),
+                                 frozen[name])
+
+    def test_actual_build_capture_rejects_tampered_ptx_and_sidecar(self):
+        from CelerisLab.cuda import compiler_v2
+        from tests.validation.multigrid.check_projected_pairs import capture_builds
+
+        for tamper in (None, "ptx", "sidecar"):
+            with self.subTest(tamper=tamper), tempfile.TemporaryDirectory() as tmp:
+                path = Path(tmp) / "blocks.ptx"
+                path.write_bytes(b"host-only mock PTX")
+                metadata = dict(ptx_sha256=hashlib.sha256(path.read_bytes()).hexdigest())
+                if tamper == "ptx":
+                    path.write_bytes(b"tampered")
+                if tamper == "sidecar":
+                    metadata["ptx_sha256"] = "0" * 64
+                Path(str(path) + ".json").write_text(json.dumps(metadata))
+                report = {"actual_build_sidecars": []}
+                with patch.object(compiler_v2, "load_module", return_value="mock-module") as load:
+                    with capture_builds(report, lambda: None):
+                        if tamper:
+                            with self.assertRaisesRegex(AssertionError, "PTX and build sidecar"):
+                                compiler_v2.load_module(str(path))
+                            load.assert_not_called()
+                            self.assertEqual(report["actual_build_sidecars"], [])
+                        else:
+                            self.assertEqual(compiler_v2.load_module(str(path)), "mock-module")
+                            self.assertEqual(report["actual_build_sidecars"], [metadata])
+                    self.assertIs(compiler_v2.load_module, load)
+
+    def test_frozen_formula_and_finest_control(self):
+        data = records(True)
+        data[0]["rows"][2]["velocity_l2"] = .002
+        data[0]["rows"][0]["velocity_l2"] = .04  # retained absolute gate, excluded from worst
+        for r in data[2:]:
+            for row in r["rows"][1:]:
+                row.update(velocity_l2=.0039, pressure_l2=.02, energy_relative=.004)
+            r["early_pressure"]["rows"][-1]["pressure_l2"] = .02
+            r["early_pressure"]["sampled_max_pressure_l2"] = .02
+        data[1]["rows"][2]["velocity_l2"] = .00001
+        data[3]["early_pressure"]["rows"][7]["pressure_l2"] = .9
+        data[3]["early_pressure"]["sampled_max_pressure_l2"] = .9
+        result = assess(data)
+        self.assertTrue(result["passed"])
+        velocity = result["comparisons"][0]
+        self.assertEqual(velocity["baseline_budget"], .004)
+        self.assertAlmostEqual(velocity["amr_to_uniform256"], 390)
+        self.assertEqual(result["early_pressure"][-1]["sampled_max_pressure_l2"], .9)
+        data[3]["rows"][-1]["pressure_l2"] = .02001
+        result = assess(data)
+        self.assertFalse(result["passed"])
+        self.assertEqual(result["comparisons"][4]["amr"], .02001)
+
+    def test_scheme_verdicts_are_independent(self):
+        data = records()
+        data[2]["rows"][-1]["pressure_l2"] = .03
+        result = assess(data)
+        self.assertFalse(result["passed"])
+        self.assertFalse(result["scheme_results"]["constant"]["passed"])
+        self.assertTrue(result["scheme_results"]["projected_linear"]["passed"])
+
+    def test_corruptions_cannot_pass(self):
+        good = records(True)
+        corruptions = [
+            lambda r: r.update(completed=False),
+            lambda r: r["rows"].pop(),
+            lambda r: r["rows"][-1].update(pressure_l2=float("nan")),
+            lambda r: r["rows"][-1].update(energy_relative=float("inf")),
+            lambda r: r["rows"][-1].update(pressure_l2=1e308),
+            lambda r: r["rows"][-1].update(velocity_l2=-1),
+            lambda r: r["rows"][-1].update(step=2047),
+            lambda r: r["rows"][-1].update(tstar=.4801),
+            lambda r: r["source"].update(dirty=True),
+            lambda r: r["source"].update(source_tree_sha256="wrong"),
+            lambda r: r["source"].update(revision=OLD_REVISION),
+            lambda r: r["source"].pop("build_capture_sha256"),
+            lambda r: r.update(numerical_revision=OLD_REVISION),
+            lambda r: r.update(reference_harness="wrong"),
+            lambda r: r.pop("runtime_identity"),
+            lambda r: r["runtime_identity"].pop("nvcc_version"),
+            lambda r: r["runtime_identity"].update(device="wrong"),
+            lambda r: r["runtime_identity"]["device"].update(capability=[8, 0]),
+            lambda r: r.update(cuda_visible_devices="0,1"),
+            lambda r: r.pop("actual_build_sidecars"),
+            lambda r: r["actual_build_sidecars"].pop(),
+            lambda r: r["actual_build_sidecars"].reverse(),
+            lambda r: r["actual_build_sidecars"][0].update(source_sha256="c" * 64),
+            lambda r: r["actual_build_sidecars"][0].update(arch="sm_80"),
+            lambda r: r["actual_build_sidecars"][0].update(flags=[]),
+            lambda r: r["actual_build_sidecars"][0].update(ptx_sha256="not-a-hash"),
+            lambda r: r["actual_build_sidecars"][0]["config"].update(ddf_shifting=False),
+            lambda r: r["source"]["harness_sha256"].update(extra="mixed"),
+            lambda r: r.update(viscosity=.01),
+            lambda r: r["tau_by_level"].reverse(),
+            lambda r: r["numerics"].update(ddf_shifting=False),
+            lambda r: r.update(initialization="different equilibrium"),
+            lambda r: r.update(refinements=[[(16, 16, 112, 104)]]),
+            lambda r: r["workload"].update(active_collision_updates=0),
+            lambda r: r["runtime_source"].update(kernel_root="/wrong/installed/tree"),
+            lambda r: r["early_pressure"]["rows"].pop(),
+            lambda r: r["early_pressure"]["rows"][8].update(step=65),
+            lambda r: r["early_pressure"]["rows"][-1].update(pressure_l2=.005),
+            lambda r: r["early_pressure"]["rows"][8].update(pressure_l2=float("nan")),
+            lambda r: r["early_pressure"].update(sampled_max_pressure_l2=.5),
+            lambda r: r["early_pressure"].update(enabled=False),
+        ]
+        for mutation in corruptions:
+            bad = deepcopy(good)
+            mutation(bad[-1])
+            with self.subTest(mutation=mutation):
+                result = assess(bad)
+                self.assertFalse(result["passed"])
+                json.dumps(result, allow_nan=False)  # failure summary itself remains serializable
+        self.assertTrue(assess(good, expected_revision=MOCK_REVISION)["passed"])
+        self.assertFalse(assess(good, expected_revision=OLD_REVISION)["passed"])
+        old = deepcopy(good)
+        for record in old:
+            record["source"].update(revision=OLD_REVISION,
+                source_tree_sha256="766e173df81d2881d3871007f15d5e321f5c7205ae578b0f80b43d1e7fe3e220")
+        self.assertFalse(assess(old)["passed"])
+        self.assertFalse(assess(good[:-1])["passed"])
+        self.assertFalse(assess(good + [good[-1]])["passed"])
+        self.assertFalse(assess([])["passed"])
+        incomplete = deepcopy(good)
+        incomplete[-1]["completed"] = False
+        self.assertFalse(assess(incomplete)["comparisons"][3]["passed"])
+
+    def test_pinned_build_sources_match_host_generated_inputs(self):
+        from CelerisLab.cuda import compiler_v2
+
+        for size in (128, 256):
+            cfg = replace(run_case.config(size, 100), nx=10, ny=10)
+            cfg.validate()
+            root = compiler_v2.generate_config(cfg)
+            try:
+                actual = json.loads((Path(root) / ".celeris-build").read_text())
+                expected = build_metadata(size)
+                self.assertEqual(actual, {key: expected[key] for key in
+                                          ("source_sha256", "config", "n_objects")})
+            finally:
+                compiler_v2.cleanup_build_dir(root)
+
+    def test_source_preflight_requires_clean_exact_revision_and_tree(self):
+        import oracle
+        from tests.validation.multigrid import check_projected_pairs as gate
+
+        for module in (run_case, oracle):
+            with self.subTest(imported_module=module.__name__), \
+                 patch.object(module, "__file__", "/wrong/checkout/" + Path(module.__file__).name), \
+                 patch.object(gate, "preflight") as preflight:
+                with self.assertRaisesRegex(RuntimeError, "Wrong runner or oracle checkout"):
+                    run_case.verify_source(REPO, MOCK_REVISION)
+                preflight.assert_not_called()
+        for expected, status, diff, fingerprint in (
+                (MOCK_REVISION, "", 0, SOURCE_SHA256),
+                (OLD_REVISION, "", 0, SOURCE_SHA256),
+                (MOCK_REVISION, " M tracked.py", 0, SOURCE_SHA256),
+                (MOCK_REVISION, "", 1, SOURCE_SHA256),
+                (MOCK_REVISION, "", 0, "wrong")):
+            provenance = dict(source(), source_tree_sha256=fingerprint)
+            with self.subTest(expected=expected, status=status, diff=diff, fingerprint=fingerprint), \
+                 patch.object(gate.subprocess, "check_output", side_effect=[MOCK_REVISION, status]), \
+                 patch.object(gate.subprocess, "call", return_value=diff), \
+                 patch.object(run_case, "provenance", return_value=provenance):
+                if expected != MOCK_REVISION or status or diff or fingerprint != SOURCE_SHA256:
+                    with self.assertRaises((ValueError, RuntimeError)):
+                        run_case.verify_source(REPO, expected)
+                else:
+                    self.assertEqual(run_case.verify_source(REPO, expected), source())
+
+    def test_explicit_opt_in_selector_and_dry_run_without_cuda(self):
+        with tempfile.TemporaryDirectory() as tmp:
+            output = Path(tmp) / "case.json"
+            base = ["run_case.py", "--repo", str(REPO), "--re", "100", "--n", "128",
+                    "--depth", "1", "--output", str(output)]
+            pinned = base + ["--expected-revision", MOCK_REVISION]
+            for argv, device in ((base + ["--run-gpu"], "0"), (pinned, "0"),
+                    (pinned + ["--run-gpu"], ""), (pinned + ["--run-gpu"], "0,1"),
+                    (pinned + ["--run-gpu"], "-1"),
+                    (pinned + ["--run-gpu", "--explosion", "projected"], "0"),
+                    (pinned + ["--run-gpu", "--explosion", "projected_linear"], "0")):
+                with self.subTest(argv=argv, device=device), patch.object(sys, "argv", argv), \
+                     patch.dict("os.environ", CUDA_VISIBLE_DEVICES=device), \
+                     patch.object(run_case, "verify_source") as verify, patch("sys.stderr", io.StringIO()):
+                    with self.assertRaises(SystemExit) as caught:
+                        run_case.main()
+                    self.assertEqual(caught.exception.code, 2)
+                    verify.assert_not_called()
+                    self.assertFalse(output.exists())
+            stdout = io.StringIO()
+            cuda_modules = {name: None for name in ("pycuda", "CelerisLab.cuda.context",
+                           "CelerisLab.lbm.block_stepper", "CelerisLab.lbm.refinement_stepper")}
+            with patch.object(sys, "argv", pinned + ["--dry-run"]), \
+                 patch.dict(sys.modules, cuda_modules), patch.dict("os.environ", CUDA_VISIBLE_DEVICES=""), \
+                 patch.object(run_case, "verify_source", return_value=source()) as verify, \
+                 patch.object(run_case, "runtime_identity", side_effect=AssertionError("CUDA reached")), \
+                 redirect_stdout(stdout):
+                self.assertEqual(run_case.main(), 0)
+            verify.assert_called_once_with(REPO, MOCK_REVISION)
+            self.assertTrue(json.loads(stdout.getvalue())["dry_run"])
+            self.assertFalse(output.exists())
+
+    def test_runtime_rejects_installed_kernel_tree(self):
+        from CelerisLab.cuda import compiler_v2
+        with patch.object(compiler_v2, "_KERNEL_DIR", "/other/CelerisLab/lbm/kernels"):
+            with self.assertRaisesRegex(RuntimeError, "CUDA source/include root"):
+                run_case.runtime_source(REPO)
+
+    def test_existing_output_is_never_overwritten(self):
+        with tempfile.TemporaryDirectory() as tmp:
+            output = Path(tmp) / "failed.json"
+            output.write_text('{"completed": false}\n')
+            argv = ["run_case.py", "--repo", str(REPO), "--re", "100", "--n", "128",
+                    "--depth", "1", "--output", str(output),
+                    "--expected-revision", MOCK_REVISION, "--run-gpu"]
+            with patch.object(sys, "argv", argv), patch("sys.stderr", io.StringIO()):
+                with self.assertRaises(SystemExit) as caught:
+                    run_case.main()
+            self.assertEqual(caught.exception.code, 2)
+            self.assertEqual(output.read_text(), '{"completed": false}\n')
+            result = subprocess.run([sys.executable, str(Path(__file__).with_name("assess_projected_pairs.py")),
+                                     "unused.json", "--expected-revision", MOCK_REVISION,
+                                     "--output", str(output)], capture_output=True)
+            self.assertEqual(result.returncode, 2)
+            self.assertEqual(output.read_text(), '{"completed": false}\n')
+
+    def test_mocked_runner_default_early_and_failure(self):
+        # Replace the GPU context and steppers before import. No CUDA module or flow is used.
+        for early, depth, fail in ((False, 1, None), (False, 2, None), (True, 2, None),
+                                   (True, 2, "resource"), (True, 2, "nonfinite"),
+                                   (True, 2, "probe_nonfinite")):
+            calls, selectors = [], []
+
+            class FakeRun:
+                def __init__(self, *args, **kwargs):
+                    self.step_count = 0
+                    if "coarse_boundary" in kwargs:
+                        selectors.append(kwargs["explosion"])
+
+                def __enter__(self):
+                    return self
+
+                def __exit__(self, *args):
+                    pass
+
+                def initialize(self, *arrays):
+                    pass
+
+                def step(self, count):
+                    self.assertion = count >= 0
+                    if not self.assertion:
+                        raise AssertionError("backward step")
+                    if fail == "resource" and self.step_count >= 16:
+                        raise MemoryError("mock resource failure")
+                    self.step_count += count
+                    calls.append(self.step_count)
+
+                def download(self):
+                    return None
+
+            modules = {
+                "CelerisLab.cuda.context": SimpleNamespace(CudaContext=FakeRun),
+                "CelerisLab.lbm.block_stepper": SimpleNamespace(PeriodicBlockStepper=FakeRun),
+                "CelerisLab.lbm.refinement_stepper": SimpleNamespace(FixedPatchStepper=FakeRun),
+            }
+
+            def metric(arrays, masks, step, n, nu, mass):
+                return dict(step=step, tstar=step*U/n, mass=16384.,
+                            **{k: float("nan") if ((fail == "nonfinite" and step == 16)
+                               or (fail == "probe_nonfinite" and step == 1))
+                               and k == "energy_relative" else 0.0 for k in LIMITS})
+
+            with tempfile.TemporaryDirectory() as tmp:
+                output = Path(tmp) / "case.json"
+                argv = ["run_case.py", "--repo", str(REPO), "--re", "100", "--n", "128",
+                        "--depth", str(depth), "--output", str(output),
+                        "--expected-revision", MOCK_REVISION, "--run-gpu"]
+                if early:
+                    argv += ["--early-pressure", "--explosion", "projected_linear"]
+                with patch.dict(sys.modules, modules), patch.object(sys, "argv", argv), \
+                     patch.dict("os.environ", CUDA_VISIBLE_DEVICES="mock"), \
+                     patch.object(run_case, "verify_source", return_value=source()) as verify, \
+                     patch.object(run_case, "runtime_identity", return_value=deepcopy(RUNTIME)), \
+                     patch.object(run_case, "build_capture", return_value=nullcontext()), \
+                     patch.object(run_case, "metrics", side_effect=metric), redirect_stdout(io.StringIO()):
+                    if fail:
+                        with self.assertRaises(MemoryError if fail == "resource" else FloatingPointError):
+                            run_case.main()
+                    else:
+                        self.assertEqual(run_case.main(), 0)
+                receipt = json.loads(output.read_text())
+                verify.assert_called_once_with(REPO, MOCK_REVISION)
+                self.assertEqual(selectors, ["projected_linear" if early else "constant"] if depth == 2 else [])
+                self.assertEqual(receipt["source"], source())
+                self.assertEqual(receipt["runtime_identity"], RUNTIME)
+                self.assertEqual(receipt["completed"], not fail)
+                self.assertEqual(calls[:2], [0, 1])  # READY probe is part of this trajectory
+                if not fail:
+                    self.assertEqual([r["tstar"] for r in receipt["rows"]], TIMES)
+                    self.assertEqual(len(receipt["early_pressure"]["rows"]), 65 if early else 0)
+                    self.assertEqual(calls[-1], 2048)
+                else:
+                    self.assertEqual(len(receipt["rows"]), 1)
+                    self.assertEqual(len(receipt["early_pressure"]["rows"]),
+                                     {"resource": 3, "nonfinite": 2, "probe_nonfinite": 1}[fail])
+                    if fail == "probe_nonfinite":
+                        self.assertNotIn("ready_after_step", receipt)
+
+
+if __name__ == "__main__":
+    unittest.main()