kernels (K4): withdraw the address helpers; keep only fill_invalid

Horsea's gates showed block_ddf/block_slot change real PTX in
BlockHaloCopy, BlockExplode and BlockExplodeProjected (6/10 modules;
BlockExplodeProjected 38 -> 40 registers) although the results stayed
exact. The stack's bar is byte-identical PTX for naming-only steps, so
those three kernels return to their K3 expressions, the helpers leave
core/layout.cuh, and test_projected_explosion goes back to its K3 form.
fill_invalid stays (its kernels were byte-identical). K4 vs K3 is now
only the NaN-fill helper.

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

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'