refinement_stepper: compatibility views are kept on purpose (comment only)

The views and the PeriodicPatchStepper alias stay for the existing tests
and docs (about 130 references); the comment no longer says REMOVE AFTER.
No code change.

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

docs/supported_settings.md

@@ -1,6 +1,7 @@
-# Supported settings: 2D Jet + fixed multilevel refinement (candidate c95a7ba)
+# Supported settings: 2D Jet + fixed multilevel refinement (candidate 888c8038)
 
-Every entry names the guard that enforces it (file:line at c95a7ba) and, where one
+Every entry names the guard that enforces it (file and function or kernel guard,
+at source revision 888c80381c006e976be2af8a452288729fc17c1b) and, where one
 exists, an existing test that exercises it. A setting not listed here is not
 claimed. Observations at the end are tied to the exact case in which they were
 measured; they are not general bounds.
@@ -9,12 +10,12 @@ measured; they are not general bounds.
 
 | Setting | Supported | Guard |
 |---|---|---|
-| Dimension / lattice | 2D D2Q9 only (3D refinement is future work) | `lbm/refinement.py:65`, `lbm/kernels/blocks.cu:25`, `lbm/kernels/refinement.cu:26` |
+| Dimension / lattice | 2D D2Q9 only (3D refinement is future work) | `lbm/refinement.py` `FixedPatch2D.from_config`; `#error` guards in `lbm/kernels/blocks.cu` and `lbm/kernels/refinement.cu` |
 | Precision / streaming | FP32 storage, double buffer | same guards |
-| Collision on refinement levels | SRT or MRT; TRT and LES rejected | `lbm/refinement.py:65` (re-checked for fine boundaries at `lbm/refinement.py:457`) |
-| Shifted (ddf_shifting) storage | only when every level is blocked; dense or mixed rejected | `lbm/refinement_stepper.py:198`, `:469`; `lbm/kernels/refinement.cu:26` (dense patch kernel unshifted only) |
-| Fine relaxation | 1/tau_fine must lie inside the configured omega guards | `lbm/refinement.py:80` |
-| Uniform blocked backend `PeriodicBlockStepper` | doubly periodic, SRT only, stationary passive circles | `lbm/block_stepper.py:84`, `:91`, `:45`-`:67` |
+| Collision on refinement levels | SRT or MRT; TRT and LES rejected | `lbm/refinement.py` `FixedPatch2D.from_config` (re-checked for fine boundaries in `prepare_fine_boundary`) |
+| Shifted (ddf_shifting) storage | only when every level is blocked; dense or mixed rejected | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__`, `FixedPatchStepper._check_storage`; `#error` guard in `lbm/kernels/refinement.cu` (dense patch kernel unshifted only) |
+| Fine relaxation | 1/tau_fine must lie inside the configured omega guards | `lbm/refinement.py` `FixedPatch2D.from_config` |
+| Uniform blocked backend `PeriodicBlockStepper` | doubly periodic, SRT only, stationary passive circles | `lbm/block_stepper.py` `PeriodicBlockStepper.__init__`, `_static_circles` |
 
 Tests: `tests/unit/test_refinement_mrt.py`, `tests/validation/multigrid/test_refinement_mrt_gpu.py`,
 `tests/unit/test_shifted_storage.py`, `tests/unit/test_shifted_block_operators.py`.
@@ -23,17 +24,17 @@ Tests: `tests/unit/test_refinement_mrt.py`, `tests/validation/multigrid/test_ref
 
 | Setting | Supported | Guard |
 |---|---|---|
-| Depth | at least one refinement level; one entry per level in `level_blocks`/`level_objects`/`level_jets` | `lbm/refinement_stepper.py:72`, `:75` |
-| Single L1 rectangle | two-coarse-cell margin to the domain edge | `lbm/refinement.py:73` |
-| Several L1 rectangles (union) | require `level_blocks[0]` and `level_blocks[1]`; transfers must close (dependency checks, no fixed gap rule) | `lbm/refinement_stepper.py:84`; `lbm/refinement.py:283`-`:295` |
-| Nested levels | depth >= 3 must be all blocked; a blocked level requires blocked levels below; a scalar child needs a scalar parent | `lbm/refinement_stepper.py:82`, `:87`; `lbm/hierarchy.py:54` |
-| Nested transfer support | child only over inherited ACTIVE cells, donors/receivers uncovered ACTIVE, no child over the parent restriction footprint | `lbm/refinement.py:391`-`:407` |
-| Tile shape | positive `(by, bx)`; the level allocation must be divisible by the tile (no partial tiles yet) | `lbm/refinement_stepper.py:78`, `lbm/refinement.py:246` (shape); `lbm/blocks.py:134` (divisibility) |
-| Coarse boundary | `periodic`, `slot_reservoir` or `channel` | `lbm/refinement_stepper.py:177` |
-| `slot_reservoir` | on-site velocity bottom slot, reservoir top, velocity 0, 2 <= slot width < nx, patch rows 4..ny-4 | `lbm/refinement_stepper.py:205`, `:212`; `lbm/refinement.py:546` |
-| `channel` | all levels blocked; Zou-He inlet, pressure-NEQ outlet, stationary bounce-back walls; transfers away from the closure | `lbm/channel.py:22`, `:26`, `:43`, `:54` |
-| Explosion | `constant` (default); `projected_linear` only periodic, all-blocked SRT, no bodies/Jets, no LES (experimental) | `lbm/refinement_stepper.py:188`, `:193` |
-| L2 runtime remap (`remap_child`) | periodic three-level L2 sets only, no L0/L1 bodies, identical pinned boundary | `lbm/refinement_stepper.py:557`, `:567`-`:581` |
+| Depth | at least one refinement level; one entry per level in `level_blocks`/`level_objects`/`level_jets` | `lbm/refinement_stepper.py` `_check_levels` |
+| Single L1 rectangle | two-coarse-cell margin to the domain edge | `lbm/refinement.py` `FixedPatch2D.from_config` |
+| Several L1 rectangles (union) | require `level_blocks[0]` and `level_blocks[1]`; transfers must close (dependency checks, no fixed gap rule) | `lbm/refinement_stepper.py` `_check_levels`; `lbm/refinement.py` `_check_parent_pull_closure` |
+| Nested levels | depth >= 3 must be all blocked; a blocked level requires blocked levels below; a scalar child needs a scalar parent | `lbm/refinement_stepper.py` `_check_levels`; `lbm/hierarchy.py` `build_hierarchy` |
+| Nested transfer support | child only over inherited ACTIVE cells, donors/receivers uncovered ACTIVE, no child over the parent restriction footprint | `lbm/refinement.py` `_nested_kind` |
+| Tile shape | positive `(by, bx)`; the level allocation must be divisible by the tile (no partial tiles yet) | `lbm/refinement_stepper.py` `_check_levels`, `lbm/refinement.py` `FixedPatchSet2D.fine_layout` (shape); `lbm/blocks.py` `BlockLayout2D.build` (divisibility) |
+| Coarse boundary | `periodic`, `slot_reservoir` or `channel` | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__` |
+| `slot_reservoir` | on-site velocity bottom slot, reservoir top, velocity 0, 2 <= slot width < nx, patch rows 4..ny-4 | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__`; `lbm/refinement.py` `prepare_fine_boundary` |
+| `channel` | all levels blocked; Zou-He inlet, pressure-NEQ outlet, stationary bounce-back walls; transfers away from the closure | `lbm/channel.py` `validate_channel_config`, `validate_channel_transfers` |
+| Explosion | `constant` (default); `projected_linear` only periodic, all-blocked SRT, no bodies/Jets, no LES (experimental) | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__` |
+| L2 runtime remap (`remap_child`) | periodic three-level L2 sets only, no L0/L1 bodies, identical pinned boundary | `lbm/refinement_stepper.py` `FixedPatchStepper.remap_child` |
 
 Tests: `tests/unit/test_refinement_three_level_multipatch.py`, `tests/validation/multigrid/test_refinement_multipatch.py`,
 `test_refinement_three_level.py`, `test_refinement_three_level_multipatch_gpu.py`, `test_refinement_blocked_base.py`,
@@ -44,14 +45,14 @@ Tests: `tests/unit/test_refinement_three_level_multipatch.py`, `tests/validation
 
 | Setting | Supported | Guard |
 |---|---|---|
-| Bodies | stationary circles without sensor/force-region roles, entirely inside the patch union, fully resolved, disjoint, stencil away from the two-hop transfer footprint | `lbm/refinement.py:474`-`:542` |
-| Level | not on L0; only on the finest level of a nested hierarchy | `lbm/refinement_stepper.py:80`, `:89` |
-| Jet condition | velocity only; `pressure` is reserved and raises NotImplementedError | `body/jets.py:148`, `:150` |
-| Jet geometry | each Jet selects resolved links of exactly one stationary body; no Jet/Jet or Jet/SurfaceActuation overlap; abs(amplitude) < 1 lattice | `body/jets.py:155`, `:175`, `:177`, `:255`, `:261`, `:263` |
-| Jets with shifted storage | only in the all-blocked stepper with SRT and no LES | `body/jets.py:249` |
-| SurfaceActuation | unshifted storage only | `body/objects.py:155` |
-| Raw boundary exchange (force) | blocked body levels only; the window ledger requires `observe_raw_exchange` | `lbm/refinement_stepper.py:264`, `:269` |
-| Checkpoint with Jets | refused | `simulation.py:687` |
+| Bodies | stationary circles without sensor/force-region roles, entirely inside the patch union, fully resolved, disjoint, stencil away from the two-hop transfer footprint | `lbm/refinement.py` `prepare_fine_boundary` |
+| Level | not on L0; only on the finest level of a nested hierarchy | `lbm/refinement_stepper.py` `_check_levels` |
+| Jet condition | velocity only; `pressure` is reserved and raises NotImplementedError | `body/jets.py` `Jet.__post_init__` |
+| Jet geometry | each Jet selects resolved links of exactly one stationary body; no Jet/Jet or Jet/SurfaceActuation overlap; abs(amplitude) < 1 lattice | `body/jets.py` `Jet.__post_init__`, `jet_link_shapes`, `_surface_links` |
+| Jets with shifted storage | only in the all-blocked stepper with SRT and no LES | `body/jets.py` `_surface_links` |
+| SurfaceActuation | unshifted storage only | `body/objects.py` `SimObject.validate_surface_actuation` |
+| Raw boundary exchange (force) | blocked body levels only; the window ledger requires `observe_raw_exchange` | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__` |
+| Checkpoint with Jets | refused | `simulation.py` `Simulation.save_checkpoint` |
 
 Raw exchange is the cut-link momentum reaction (wall + aperture), not a validated
 nozzle thrust or stress load. Jet amplitude is a lattice velocity, not a flow rate.

src/CelerisLab/lbm/kernels/blocks.cu

@@ -13,6 +13,7 @@
 #include "core/params.cuh"
 #include "core/flags.cuh"
 #include "core/obs.cuh"
+#include "core/cell_kind.cuh"
 #include "operators/macro.cuh"
 #include "operators/collision_srt.cuh"
 #if COLLISION_MODEL == 2
@@ -20,6 +21,10 @@
 #endif
 #include "streaming/pull_double_buffer.cuh"
 #include "boundary/curved_boundary.cuh"
+
+#if DIM != 2 || NQ != 9 || (COLLISION_MODEL != 0 && COLLISION_MODEL != 2) || STORE_PRECISION != 0 || USE_LES || STREAMING_MODEL != 0
+#error "Blocks require FP32 D2Q9 SRT or MRT double buffers without LES"
+#endif
 #ifndef BLOCK_ROOT_BC
 #error "Blocks require a generated BLOCK_ROOT_BC mode"
 #elif BLOCK_ROOT_BC < 0 || BLOCK_ROOT_BC > 2
@@ -33,160 +38,12 @@
 #include "operators/block_collision.cuh"
 #endif
 
-#if DIM != 2 || NQ != 9 || (COLLISION_MODEL != 0 && COLLISION_MODEL != 2) || STORE_PRECISION != 0 || USE_LES || STREAMING_MODEL != 0
-#error "Blocks require FP32 D2Q9 SRT or MRT double buffers without LES"
-#endif
-
-// Fixed-patch cell kinds, matching refinement.py ACTIVE/INNER_GHOST/OUTER_GHOST.
-constexpr unsigned char KIND_ACTIVE = 0u, KIND_INNER = 1u, KIND_OUTER = 2u;
-
-// Halo refresh: block cell ids b*TOTAL_CELLS+c, sources are always core cells.
-extern "C" __global__ void BlockHaloCopy(
-    fpxx* f, const uint64_t* dst, const uint64_t* src, unsigned int count)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const uint64_t d = dst[row], s = src[row];
-    fpxx* to = f + (d / TOTAL_CELLS) * NQ * TOTAL_CELLS;
-    const fpxx* from = f + (s / TOTAL_CELLS) * NQ * TOTAL_CELLS;
-    for (unsigned int q = 0; q < NQ; ++q)
-        to[index_f(d % TOTAL_CELLS, q)] = from[index_f(s % TOTAL_CELLS, q)];
-}
-
-// Fixed-patch explosion into blocked fine storage: target = owner block cell
-// id of a patch-ghost cell; donors are parent block cell ids. Dense parent
-// storage is the one-window case with coarse_cells = full domain size.
-extern "C" __global__ void BlockExplode(
-    const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
-    unsigned int count, unsigned long coarse_cells)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const uint64_t t = target[row];
-    fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
-    const uint64_t d = donor[row];
-    const fpxx* from = coarse + (d / coarse_cells) * NQ * coarse_cells;
-    for (unsigned int q = 0; q < NQ; ++q)
-        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(
-    const fpxx* fine, fpxx* coarse, const uint64_t* target, const uint64_t* children,
-    unsigned int count)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const uint64_t* s = children + 4ul * row;
-    coarse[target[row]] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
-}
-
-// Static cut links, receiver-local staging: cell is the receiver's local core
-// cell, so the existing BFL writes the solid-source window image that this
-// block's next pull reads (no halo copy between this kernel and BlockStep).
-// surface_velocity (nullable, 2 floats per link): prescribed fluid velocity on
-// fixed geometry, same convention as CurvedBoundaryKernel. Optional raw rows
-// (jx,jy,torque,dm) are unique physical links, never body-observation slots.
-extern "C" __global__ void BlockCurvedBoundary(
-    fpxx* f, const uint32_t* block, const uint32_t* cell, const unsigned char* dir,
-    const float* q, const unsigned char* fallback, unsigned int count,
-    const float* surface_velocity, const float* rx, const float* ry, float* raw)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    float Uw = 0.0f, Vw = 0.0f;
-    if (surface_velocity != nullptr) {
-        const float u = surface_velocity[2 * row], v = surface_velocity[2 * row + 1];
-        if (u != 0.0f || v != 0.0f) { Uw = u; Vw = v; }
-    }
-    apply_bouzidi_link(dir[row], q[row], cell[row], f + (unsigned long)block[row] * NQ * TOTAL_CELLS,
-                       Uw, Vw, rx ? rx[row] : 0.0f, ry ? ry[row] : 0.0f,
-                       fallback[row], nullptr, -1, raw ? raw + 4ul * row : nullptr,
-                       raw ? raw + 4ul * row + 3 : nullptr);
-}
-
-// Same-stream BFL -> this row-owned accumulation -> pull step. No atomics,
-// no cross-row reduction and no history allocation proportional to substeps.
-extern "C" __global__ void BlockAccumulateExchange(
-    const float* raw, double* integral, unsigned int count, double h, double dt)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const unsigned long k = 4ul * row;
-    const double impulse = h * dt;
-    integral[k] += impulse * (double)raw[k];
-    integral[k + 1] += impulse * (double)raw[k + 1];
-    integral[k + 2] += h * impulse * (double)raw[k + 2];
-    integral[k + 3] += h * h * (double)raw[k + 3];
-}
-
+// Kernels, in definition order: transfers, boundary staging, then the one full
+// step selected for this module (BlockStep, or BlockChannelStep for a channel root).
+#include "step/block_transfer.cu"
+#include "step/block_boundary.cu"
 #if BLOCK_ROOT_BC != 2
-// Pull/collide/store core cells only; halo cells of out are left untouched.
-// flags (block cell ids, nullable): solid cores store NaN and never collide.
-// kind (block cell ids, nullable = all ACTIVE): OUTER skips the pull and
-// stores NaN, INNER pulls without collision, ACTIVE pulls and collides.
-extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int halo,
-                                     const uint16_t* flags, const unsigned char* kind) {
-    unsigned int x, y; unsigned long k;
-    index_from_thread(x, y, k);
-    if (x < halo || y < halo || x >= NX - halo || y >= NY - halo) return;
-    const unsigned long base = (unsigned long)blockIdx.z * NQ * TOTAL_CELLS;
-    const unsigned long cell = (unsigned long)blockIdx.z * TOTAL_CELLS + k;
-    const unsigned char cell_kind = kind != nullptr ? kind[cell] : KIND_ACTIVE;
-    float f[NQ];
-    if (cell_kind == KIND_OUTER || (flags != nullptr && is_solid(flags[cell]))) {
-        for (int q = 0; q < NQ; ++q) f[q] = NAN;
-    } else {
-        unsigned long j[NQ];
-        compute_neighbors(x, y, j);  // Core neighbors stay inside the window: no wrap.
-        stream_pull_load(k, f, in + base, j);
-        if (cell_kind == KIND_ACTIVE) {
-#if BLOCK_ROOT_BC == 1
-            block_wall(f, x, y, halo);
-#endif
-            block_collide(f);
-        }
-    }
-    stream_pull_store(k, f, out + base);
-}
-
-#endif
-
-// Root channel is a replacement full step, selected once by the host.
-#if BLOCK_ROOT_BC == 2
+#include "step/block_step.cu"
+#else
 #include "step/block_channel.cu"
 #endif

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

@@ -15,10 +15,13 @@
 // D2Q9 y-wall half-way bounce-back for pull double-buffer streaming.
 // ---------------------------------------------------------------------------
 #if LATTICE_MODEL == LATTICE_D2Q9
-__device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y,
-                                                  float* __restrict__ f,
-                                                  const fpxx* __restrict__ fi_in,
-                                                  unsigned long k)
+// Rows 1 and top_row (= physical ny - 2) next to walls at rows 0 and ny - 1.
+// y is the physical row; k addresses the node in fi_in (a block window may
+// have NX/NY strides different from the physical extent).
+__device__ inline void apply_wall_bb_d2q9_rows(unsigned int y, unsigned int top_row,
+                                               float* __restrict__ f,
+                                               const fpxx* __restrict__ fi_in,
+                                               unsigned long k)
 {
     if (y == 1) {
         // Directions sourced from y=0 wall in pull step: +y, +x+y, -x+y.
@@ -26,14 +29,22 @@ __device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y,
         f[5] = load_ddf(fi_in, index_f(k, 6u));
         f[8] = load_ddf(fi_in, index_f(k, 7u));
     }
-    else if (y == (unsigned int)(NY - 2)) {
-        // Directions sourced from y=NY-1 wall in pull step: -y, -x-y, +x-y.
+    else if (y == top_row) {
+        // Directions sourced from the y = ny - 1 wall in pull step: -y, -x-y, +x-y.
         f[4] = load_ddf(fi_in, index_f(k, 3u));
         f[6] = load_ddf(fi_in, index_f(k, 5u));
         f[7] = load_ddf(fi_in, index_f(k, 8u));
     }
 }
 
+__device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y,
+                                                  float* __restrict__ f,
+                                                  const fpxx* __restrict__ fi_in,
+                                                  unsigned long k)
+{
+    apply_wall_bb_d2q9_rows(y, (unsigned int)(NY - 2), f, fi_in, k);
+}
+
 // ---------------------------------------------------------------------------
 // D2Q9 y-wall free-slip (specular reflection) for pull double-buffer streaming.
 //

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

@@ -0,0 +1,16 @@
+// CelerisLab – core/cell_kind.cuh
+// Fixed-patch cell kinds stored per block cell, matching lbm/refinement.py
+// ACTIVE/INNER_GHOST/OUTER_GHOST: OUTER is invalid (stores NaN, no pull),
+// INNER streams without collision, ACTIVE pulls and collides.
+#ifndef CELERIS_CORE_CELL_KIND_CUH
+#define CELERIS_CORE_CELL_KIND_CUH
+
+constexpr unsigned char KIND_ACTIVE = 0u, KIND_INNER = 1u, KIND_OUTER = 2u;
+
+// NaN populations mark a cell without valid state (OUTER ghost or solid core).
+__device__ __forceinline__ void fill_invalid(float* f)
+{
+    for (int q = 0; q < NQ; ++q) f[q] = NAN;
+}
+
+#endif

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

@@ -2,6 +2,13 @@
 #ifndef CELERIS_OPERATORS_BLOCK_COLLISION_CUH
 #define CELERIS_OPERATORS_BLOCK_COLLISION_CUH
 
+#include "../core/params.cuh"
+#include "macro.cuh"
+#include "collision_srt.cuh"
+#if COLLISION_MODEL == 2
+#include "collision_mrt.cuh"
+#endif
+
 __device__ __forceinline__ void block_collide(float* f) {
 #if COLLISION_MODEL == 2
     collide_mrt_unforced(f);  // SRT path below unchanged.

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

@@ -0,0 +1,14 @@
+// CelerisLab – operators/bounce_back_swap.cuh
+// Full-way bounce-back at a wall node: swap every opposite pair (paired
+// order q, q+1). Needs only NQ; shared by the uniform and blocked steps.
+#ifndef CELERIS_OPERATORS_BOUNCE_BACK_SWAP_CUH
+#define CELERIS_OPERATORS_BOUNCE_BACK_SWAP_CUH
+
+__device__ __forceinline__ void bounce_back_swap(float* f) {
+    #pragma unroll
+    for (int i = 1; i < NQ; i += 2) {
+        float t = f[i]; f[i] = f[i+1]; f[i+1] = t;
+    }
+}
+
+#endif

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

@@ -5,12 +5,7 @@
 #ifndef CELERIS_OPERATORS_HELPERS_CUH
 #define CELERIS_OPERATORS_HELPERS_CUH
 
-__device__ __forceinline__ void bounce_back_swap(float* f) {
-    #pragma unroll
-    for (int i = 1; i < NQ; i += 2) {
-        float t = f[i]; f[i] = f[i+1]; f[i+1] = t;
-    }
-}
+#include "bounce_back_swap.cuh"
 
 __device__ __forceinline__ void collide_dispatch(
     float* f, float rho_n,

src/CelerisLab/lbm/kernels/refinement.cu

@@ -9,6 +9,7 @@
 #include "core/params.cuh"
 #include "core/flags.cuh"
 #include "core/obs.cuh"
+#include "core/cell_kind.cuh"
 #include "operators/macro.cuh"
 #include "operators/collision_srt.cuh"
 #if COLLISION_MODEL == 2
@@ -35,93 +36,5 @@ __device__ unsigned int incoming_counts[NQ * TOTAL_CELLS];
 }
 #endif
 
-__device__ __forceinline__ void patch_collide(float* f, unsigned long k) {
-#if COLLISION_MODEL == 2
-    collide_mrt_unforced(f);
-#else
-    float rho, ux, uy, feq[NQ];
-    compute_rho_u(f, rho, ux, uy);
-    compute_feq(rho, ux, uy, feq);
-    collide_srt_no_force(f, feq, d_params.omega);
-#endif
-#ifdef CELERIS_REFINEMENT_AUDIT
-    ++collision_counts[k];
-#endif
-}
-
-extern "C" __global__ void PatchExplode(
-    const float* coarse, float* fine, const uint32_t* targets,
-    const uint32_t* donors, unsigned int count, unsigned long coarse_cells)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    for (unsigned int q = 0; q < NQ; ++q)
-        fine[index_f(targets[row], q)] = coarse[q * coarse_cells + donors[row]];
-#ifdef CELERIS_REFINEMENT_AUDIT
-    ++explosion_counts[targets[row]];
-#endif
-}
-
-extern "C" __global__ void PatchFineStep(const float* in, float* out, const uint16_t* flags) {
-    unsigned int x, y; unsigned long k;
-    index_from_thread(x, y, k);
-    if (x >= NX || y >= NY) return;
-    float f[NQ];
-    // Outer halo is invalid after EACH hop, including when reusing the buffer
-    // that held explosion data. Inner halo only streams; it never collides.
-    if (x < 2 || y < 2 || x >= NX - 2 || y >= NY - 2
-            || (flags != nullptr && is_solid(flags[k]))) {
-        for (int q = 0; q < NQ; ++q) f[q] = NAN;
-    } else {
-        unsigned long j[NQ];
-        compute_neighbors(x, y, j);  // These interior targets never wrap.
-        stream_pull_load(k, f, in, j);
-        if (x >= 4 && y >= 4 && x < NX - 4 && y < NY - 4)
-            patch_collide(f, k);
-    }
-    stream_pull_store(k, f, out);
-}
-
-extern "C" __global__ void PatchCoarseStep(
-    const float* in, float* out, const float* fine,
-    const uint64_t* targets, const uint64_t* children, unsigned int count,
-    unsigned int x0, unsigned int y0, unsigned int x1, unsigned int y1,
-    unsigned int wall_aware)
-{
-    unsigned int x, y; unsigned long k;
-    index_from_thread(x, y, k);
-    if (x >= NX || y >= NY) return;
-    // Both buffers were poisoned at initialization; covered cells stay invalid.
-    if (x >= x0 && x < x1 && y >= y0 && y < y1) return;
-    if (wall_aware && (y == 0 || y == NY - 1)) return;
-    unsigned long j[NQ];
-    compute_neighbors(x, y, j);
-    float f[NQ];
-    stream_pull_load(k, f, in, j);
-    if (x >= x0 - 1 && x <= x1 && y >= y0 - 1 && y <= y1) {
-        for (unsigned int q = 1; q < NQ; ++q) {
-            uint64_t slot = index_f(k, q);
-            unsigned int lo = 0, hi = count;
-            while (lo < hi) {
-                unsigned int mid = lo + (hi - lo) / 2;
-                if (targets[mid] < slot) lo = mid + 1;
-                else hi = mid;
-            }
-            if (lo < count && targets[lo] == slot) {
-                const uint64_t* s = children + 4ul * lo;
-                f[q] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
-#ifdef CELERIS_REFINEMENT_AUDIT
-                ++incoming_counts[slot];
-#endif
-            }
-        }
-    }
-    // Includes the eight grazing-corner links with UNCOVERED coarse donors.
-    // All replacement arrivals must be in place before computing rho/u.
-    if (wall_aware) {
-        if (y == 1) reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity(x));
-        else if (y == NY - 2) reconstruct_top_reservoir(f, RHO);
-    }
-    patch_collide(f, k);
-    stream_pull_store(k, f, out);
-}
+// Kernels: patch collision helper, PatchExplode, PatchFineStep, PatchCoarseStep.
+#include "step/patch_step.cu"

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

@@ -0,0 +1,42 @@
+// CelerisLab/lbm/kernels/step/block_boundary.cu
+// Receiver-local curved-boundary staging and its raw-exchange accumulation.
+// Part of the blocks.cu compilation entry; relies on its includes and guards.
+
+// Static cut links, receiver-local staging: cell is the receiver's local core
+// cell, so the existing BFL writes the solid-source window image that this
+// block's next pull reads (no halo copy between this kernel and BlockStep).
+// surface_velocity (nullable, 2 floats per link): prescribed fluid velocity on
+// fixed geometry, same convention as CurvedBoundaryKernel. Optional raw rows
+// (jx,jy,torque,dm) are unique physical links, never body-observation slots.
+extern "C" __global__ void BlockCurvedBoundary(
+    fpxx* f, const uint32_t* block, const uint32_t* cell, const unsigned char* dir,
+    const float* q, const unsigned char* fallback, unsigned int count,
+    const float* surface_velocity, const float* rx, const float* ry, float* raw)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    float Uw = 0.0f, Vw = 0.0f;
+    if (surface_velocity != nullptr) {
+        const float u = surface_velocity[2 * row], v = surface_velocity[2 * row + 1];
+        if (u != 0.0f || v != 0.0f) { Uw = u; Vw = v; }
+    }
+    apply_bouzidi_link(dir[row], q[row], cell[row], f + (unsigned long)block[row] * NQ * TOTAL_CELLS,
+                       Uw, Vw, rx ? rx[row] : 0.0f, ry ? ry[row] : 0.0f,
+                       fallback[row], nullptr, -1, raw ? raw + 4ul * row : nullptr,
+                       raw ? raw + 4ul * row + 3 : nullptr);
+}
+
+// Same-stream BFL -> this row-owned accumulation -> pull step. No atomics,
+// no cross-row reduction and no history allocation proportional to substeps.
+extern "C" __global__ void BlockAccumulateExchange(
+    const float* raw, double* integral, unsigned int count, double h, double dt)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const unsigned long k = 4ul * row;
+    const double impulse = h * dt;
+    integral[k] += impulse * (double)raw[k];
+    integral[k + 1] += impulse * (double)raw[k + 1];
+    integral[k + 2] += h * impulse * (double)raw[k + 2];
+    integral[k + 3] += h * h * (double)raw[k + 3];
+}

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

@@ -4,6 +4,9 @@
 
 // Root-only blocked counterpart of OneStep's D2Q9/SRT channel closure.
 // NX/NY are LOCAL strides; all physical boundary decisions use global coordinates.
+#include "../core/cell_kind.cuh"
+#include "../operators/bounce_back_swap.cuh"
+#include "../boundary/bounce_back.cuh"
 #include "../boundary/inlet/common.cuh"
 #include "../boundary/inlet/zou_he_local.cuh"
 #include "../boundary/outlet/pressure_neq.cuh"
@@ -25,7 +28,7 @@ extern "C" __global__ void BlockChannelStep(
     const unsigned long cell = (unsigned long)blockIdx.z * TOTAL_CELLS + k;
     float f[NQ];
     if (kind != nullptr && kind[cell] == KIND_OUTER) {
-        for (int q = 0; q < NQ; ++q) f[q] = NAN;
+        fill_invalid(f);
     } else {
         unsigned long j[NQ];
         compute_neighbors(x, y, j);
@@ -33,27 +36,17 @@ extern "C" __global__ void BlockChannelStep(
         const bool wall = gy == 0 || gy == d_block_channel.ny - 1;
         const bool outlet = gx == d_block_channel.nx - 1;
         if (wall) {
-            // Wall owns all four corners. Ghosts pull and swap, never collide.
-            for (int q = 1; q < NQ; q += 2) {
-                const float t = f[q]; f[q] = f[q + 1]; f[q + 1] = t;
-            }
+            bounce_back_swap(f);  // Wall owns all four corners; ghosts never collide.
         } else if (gx == 0) {
             apply_zou_he_left_velocity_inlet_d2q9(
                 f, inlet_target_u((float)gy, d_block_channel.ny, d_block_channel.velocity), 0.0f);
         } else if (outlet) {
             // Immutable input donor, receiver-local (also valid across a block seam).
             float donor[NQ];
-            for (int q = 0; q < NQ; ++q)
-                donor[q] = load_ddf(in + base, index_f(k - 1, q));
+            load_node(in + base, k - 1, donor);
             apply_pressure_outlet_d2q9(f, donor);
-        } else if (gy == 1) {
-            f[3] = load_ddf(in + base, index_f(k, 4));
-            f[5] = load_ddf(in + base, index_f(k, 6));
-            f[8] = load_ddf(in + base, index_f(k, 7));
-        } else if (gy == d_block_channel.ny - 2) {
-            f[4] = load_ddf(in + base, index_f(k, 3));
-            f[6] = load_ddf(in + base, index_f(k, 5));
-            f[7] = load_ddf(in + base, index_f(k, 8));
+        } else {
+            apply_wall_bb_d2q9_rows(gy, d_block_channel.ny - 2, f, in + base, k);
         }
         if (!wall && !outlet) {
 #if COLLISION_MODEL == 2

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

@@ -0,0 +1,35 @@
+// CelerisLab/lbm/kernels/step/block_step.cu
+// Generic blocked pull/collide/store for periodic/fine and slot/reservoir levels.
+// Part of the blocks.cu compilation entry; relies on its includes and guards.
+#if BLOCK_ROOT_BC == 2
+#error "BlockStep is not built for the channel root (BLOCK_ROOT_BC == 2)"
+#endif
+
+// Pull/collide/store core cells only; halo cells of out are left untouched.
+// flags (block cell ids, nullable): solid cores store NaN and never collide.
+// kind (block cell ids, nullable = all ACTIVE): OUTER skips the pull and
+// stores NaN, INNER pulls without collision, ACTIVE pulls and collides.
+extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int halo,
+                                     const uint16_t* flags, const unsigned char* kind) {
+    unsigned int x, y; unsigned long k;
+    index_from_thread(x, y, k);
+    if (x < halo || y < halo || x >= NX - halo || y >= NY - halo) return;
+    const unsigned long base = (unsigned long)blockIdx.z * NQ * TOTAL_CELLS;
+    const unsigned long cell = (unsigned long)blockIdx.z * TOTAL_CELLS + k;
+    const unsigned char cell_kind = kind != nullptr ? kind[cell] : KIND_ACTIVE;
+    float f[NQ];
+    if (cell_kind == KIND_OUTER || (flags != nullptr && is_solid(flags[cell]))) {
+        fill_invalid(f);
+    } else {
+        unsigned long j[NQ];
+        compute_neighbors(x, y, j);  // Core neighbors stay inside the window: no wrap.
+        stream_pull_load(k, f, in + base, j);
+        if (cell_kind == KIND_ACTIVE) {
+#if BLOCK_ROOT_BC == 1
+            block_wall(f, x, y, halo);
+#endif
+            block_collide(f);
+        }
+    }
+    stream_pull_store(k, f, out + base);
+}

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

@@ -0,0 +1,78 @@
+// CelerisLab/lbm/kernels/step/block_transfer.cu
+// Same-level halo refresh and 2:1 explosion/coalescence between blocked levels.
+// Part of the blocks.cu compilation entry; relies on its includes and guards.
+
+// Halo refresh: block cell ids b*TOTAL_CELLS+c, sources are always core cells.
+extern "C" __global__ void BlockHaloCopy(
+    fpxx* f, const uint64_t* dst, const uint64_t* src, unsigned int count)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const uint64_t d = dst[row], s = src[row];
+    fpxx* to = f + (d / TOTAL_CELLS) * NQ * TOTAL_CELLS;
+    const fpxx* from = f + (s / TOTAL_CELLS) * NQ * TOTAL_CELLS;
+    for (unsigned int q = 0; q < NQ; ++q)
+        to[index_f(d % TOTAL_CELLS, q)] = from[index_f(s % TOTAL_CELLS, q)];
+}
+
+// Fixed-patch explosion into blocked fine storage: target = owner block cell
+// id of a patch-ghost cell; donors are parent block cell ids. Dense parent
+// storage is the one-window case with coarse_cells = full domain size.
+extern "C" __global__ void BlockExplode(
+    const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
+    unsigned int count, unsigned long coarse_cells)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const uint64_t t = target[row];
+    fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
+    const uint64_t d = donor[row];
+    const fpxx* from = coarse + (d / coarse_cells) * NQ * coarse_cells;
+    for (unsigned int q = 0; q < NQ; ++q)
+        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(
+    const fpxx* fine, fpxx* coarse, const uint64_t* target, const uint64_t* children,
+    unsigned int count)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const uint64_t* s = children + 4ul * row;
+    coarse[target[row]] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
+}

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

@@ -0,0 +1,95 @@
+// CelerisLab/lbm/kernels/step/patch_step.cu
+// Dense fixed-patch collision, explosion and the fine/coarse steps.
+// Part of the refinement.cu compilation entry; relies on its includes, guards
+// and the optional CELERIS_REFINEMENT_AUDIT counters declared there.
+
+__device__ __forceinline__ void patch_collide(float* f, unsigned long k) {
+#if COLLISION_MODEL == 2
+    collide_mrt_unforced(f);
+#else
+    float rho, ux, uy, feq[NQ];
+    compute_rho_u(f, rho, ux, uy);
+    compute_feq(rho, ux, uy, feq);
+    collide_srt_no_force(f, feq, d_params.omega);
+#endif
+#ifdef CELERIS_REFINEMENT_AUDIT
+    ++collision_counts[k];
+#endif
+}
+
+extern "C" __global__ void PatchExplode(
+    const float* coarse, float* fine, const uint32_t* targets,
+    const uint32_t* donors, unsigned int count, unsigned long coarse_cells)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    for (unsigned int q = 0; q < NQ; ++q)
+        fine[index_f(targets[row], q)] = coarse[q * coarse_cells + donors[row]];
+#ifdef CELERIS_REFINEMENT_AUDIT
+    ++explosion_counts[targets[row]];
+#endif
+}
+
+extern "C" __global__ void PatchFineStep(const float* in, float* out, const uint16_t* flags) {
+    unsigned int x, y; unsigned long k;
+    index_from_thread(x, y, k);
+    if (x >= NX || y >= NY) return;
+    float f[NQ];
+    // Outer halo is invalid after EACH hop, including when reusing the buffer
+    // that held explosion data. Inner halo only streams; it never collides.
+    if (x < 2 || y < 2 || x >= NX - 2 || y >= NY - 2
+            || (flags != nullptr && is_solid(flags[k]))) {
+        fill_invalid(f);
+    } else {
+        unsigned long j[NQ];
+        compute_neighbors(x, y, j);  // These interior targets never wrap.
+        stream_pull_load(k, f, in, j);
+        if (x >= 4 && y >= 4 && x < NX - 4 && y < NY - 4)
+            patch_collide(f, k);
+    }
+    stream_pull_store(k, f, out);
+}
+
+extern "C" __global__ void PatchCoarseStep(
+    const float* in, float* out, const float* fine,
+    const uint64_t* targets, const uint64_t* children, unsigned int count,
+    unsigned int x0, unsigned int y0, unsigned int x1, unsigned int y1,
+    unsigned int wall_aware)
+{
+    unsigned int x, y; unsigned long k;
+    index_from_thread(x, y, k);
+    if (x >= NX || y >= NY) return;
+    // Both buffers were poisoned at initialization; covered cells stay invalid.
+    if (x >= x0 && x < x1 && y >= y0 && y < y1) return;
+    if (wall_aware && (y == 0 || y == NY - 1)) return;
+    unsigned long j[NQ];
+    compute_neighbors(x, y, j);
+    float f[NQ];
+    stream_pull_load(k, f, in, j);
+    if (x >= x0 - 1 && x <= x1 && y >= y0 - 1 && y <= y1) {
+        for (unsigned int q = 1; q < NQ; ++q) {
+            uint64_t slot = index_f(k, q);
+            unsigned int lo = 0, hi = count;
+            while (lo < hi) {
+                unsigned int mid = lo + (hi - lo) / 2;
+                if (targets[mid] < slot) lo = mid + 1;
+                else hi = mid;
+            }
+            if (lo < count && targets[lo] == slot) {
+                const uint64_t* s = children + 4ul * lo;
+                f[q] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
+#ifdef CELERIS_REFINEMENT_AUDIT
+                ++incoming_counts[slot];
+#endif
+            }
+        }
+    }
+    // Includes the eight grazing-corner links with UNCOVERED coarse donors.
+    // All replacement arrivals must be in place before computing rho/u.
+    if (wall_aware) {
+        if (y == 1) reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity(x));
+        else if (y == NY - 2) reconstruct_top_reservoir(f, RHO);
+    }
+    patch_collide(f, k);
+    stream_pull_store(k, f, out);
+}

src/CelerisLab/lbm/kernels/streaming/pull_double_buffer.cuh

@@ -21,6 +21,14 @@
 //   i.e. the population traveling in direction i at node n
 //   arrived from the node at n - c_i  = neighbor in direction opp(i)
 // ---------------------------------------------------------------------------
+// load_node: all NQ populations of node n as stored (no streaming).
+__device__ __forceinline__ void load_node(const fpxx* __restrict__ fi, unsigned long n,
+                                          float* __restrict__ f)
+{
+    for (int q = 0; q < NQ; ++q)
+        f[q] = load_ddf(fi, index_f(n, q));
+}
+
 __device__ __forceinline__ void stream_pull_load(unsigned long n,
                                                  float* __restrict__ f,
                                                  const fpxx* __restrict__ fi,

src/CelerisLab/lbm/refinement_stepper.py

@@ -949,8 +949,8 @@ class FixedPatchStepper:
             if exc[0] is None:
                 raise
 
-    # Read-only compatibility views over the level lists. Internal code must not use them;
-    # REMOVE AFTER the existing tests/docs migrate to hierarchy/_layouts/_level_* lists.
+    # Read-only compatibility views over the level lists, kept on purpose for the existing
+    # tests and docs. Internal code must not use them; it reads hierarchy/_layouts/_level_*.
     patch = property(lambda self: self.hierarchy.pairs[0])
     child_patch = property(lambda self: self.hierarchy.pairs[1] if self.hierarchy.depth > 2 else None)
     base_layout = property(lambda self: self._layouts[0])

tests/unit/cpp/channel_root.cpp

@@ -41,7 +41,7 @@ inline float __fdiv_rn(float a,float b){return a/b;}
 #include "boundary/inlet_outlet.cuh"
 #include "step/one_step_double.cu"
 #else
-constexpr unsigned char KIND_OUTER=2;
+#include "core/cell_kind.cuh"
 #include "step/block_channel.cu"
 #endif
 int main(){

tests/unit/test_projected_explosion.py

@@ -162,7 +162,7 @@ class ProjectedKernelTests(unittest.TestCase):
         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()
+                  "src/CelerisLab/lbm/kernels/step/block_transfer.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):

tests/unit/test_refinement_mrt.py

@@ -22,12 +22,12 @@ class RefinementMRTHost(unittest.TestCase):
 
     def test_srt_statements_and_guards_preserved(self):
         """MRT is an additive compile-time branch; SRT statements and other guards remain."""
-        blocks = (KERNELS / "blocks.cu").read_text()
+        blocks = (KERNELS / "blocks.cu").read_text() + (KERNELS / "step/block_step.cu").read_text()
         channel = (KERNELS / "step/block_channel.cu").read_text()
         block_collision = (KERNELS / "operators/block_collision.cuh").read_text()
         self.assertIn('#include "operators/block_collision.cuh"', blocks)
         self.assertEqual(blocks.count("block_collide(f);"), 1)
-        refinement = (KERNELS / "refinement.cu").read_text()
+        refinement = (KERNELS / "refinement.cu").read_text() + (KERNELS / "step/patch_step.cu").read_text()
         for text in ("compute_feq_from_momentum(drho, jx, jy, ux, uy, feq);",
                      "compute_feq(rho, ux, uy, feq);",
                      "collide_srt_no_force(f, feq, d_params.omega);"):