Merge K4 test harness fix into K5

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

tests/unit/test_projected_explosion.py

@@ -163,6 +163,12 @@ 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):
@@ -183,7 +189,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'+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'+helpers+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'