Merge K4 helper withdrawal into K5

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

src/CelerisLab/lbm/kernels/blocks.cu

@@ -45,8 +45,8 @@ extern "C" __global__ void BlockHaloCopy(
     unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
     if (row >= count) return;
     const uint64_t d = dst[row], s = src[row];
-    fpxx* to = block_ddf(f, d);
-    const fpxx* from = block_ddf(f, s);
+    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)];
 }
@@ -61,9 +61,9 @@ extern "C" __global__ void BlockExplode(
     unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
     if (row >= count) return;
     const uint64_t t = target[row];
-    fpxx* to = block_ddf(fine, t);
+    fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
     const uint64_t d = donor[row];
-    const fpxx* from = block_ddf(coarse, d, coarse_cells);
+    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];
 }
@@ -80,12 +80,13 @@ extern "C" __global__ void BlockExplodeProjected(
     unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
     if (row >= count) return;
     const uint64_t t = target[row], d = donor[row];
-    fpxx* to = block_ddf(fine, t);
+    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[block_slot(n, coarse_cells, q)];
+            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];
@@ -95,7 +96,7 @@ extern "C" __global__ void BlockExplodeProjected(
             gy -= ey * along;
         }
         to[index_f(t % TOTAL_CELLS, q)] =
-            coarse[block_slot(d, coarse_cells, q)]
+            coarse[(d / coarse_cells) * NQ * coarse_cells + q * coarse_cells + d % coarse_cells]
             + offsets[2ul * row] * gx + offsets[2ul * row + 1] * gy;
     }
 }

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

@@ -65,21 +65,6 @@ __device__ __forceinline__ void block_global_xy(unsigned int x, unsigned int y,
     gy = (blockIdx.z / nbx) * by + y - halo;
 }
 
-// Block-contiguous storage: block b holds NQ*cells values (Q-major inside) and
-// block cell id b*cells + c names cell c of block b. cells is TOTAL_CELLS for a
-// block window, or the whole domain for a dense (one-window) parent.
-template <typename T>
-__device__ __forceinline__ T* block_ddf(T* f, uint64_t block_cell, unsigned long cells = TOTAL_CELLS)
-{
-    return f + (block_cell / cells) * NQ * cells;
-}
-
-// Offset of population q of block cell block_cell from the start of storage.
-__device__ __forceinline__ uint64_t block_slot(uint64_t block_cell, unsigned long cells, unsigned int q)
-{
-    return (block_cell / cells) * NQ * cells + q * cells + block_cell % cells;
-}
-
 // ---------------------------------------------------------------------------
 // Neighbor list (periodic BCs via modular arithmetic)
 // ---------------------------------------------------------------------------

tests/unit/test_projected_explosion.py

@@ -163,12 +163,6 @@ class ProjectedKernelTests(unittest.TestCase):
         donors = parent.owner.ravel()[p.ghost_coarse]
         source = (Path(__file__).resolve().parents[2] /
                   "src/CelerisLab/lbm/kernels/blocks.cu").read_text()
-        # The kernel addresses storage through the real core/layout.cuh helpers.
-        layout = (Path(__file__).resolve().parents[2] /
-                  "src/CelerisLab/lbm/kernels/core/layout.cuh").read_text()
-        start = layout.index("template <typename T>\n__device__ __forceinline__ T* block_ddf(")
-        end = layout.index("}\n", layout.index("__device__ __forceinline__ uint64_t block_slot(")) + 2
-        helpers = layout[start:end].replace("__device__ __forceinline__ ", "inline ")
         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):
@@ -189,7 +183,7 @@ class ProjectedKernelTests(unittest.TestCase):
             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'+helpers+kernel
+            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'