From 43e226c59a4986d699863df03219b9dd17112eae Mon Sep 17 00:00:00 2001 From: Claude Date: Thu, 22 Oct 2026 00:35:07 +0000 Subject: [PATCH] test_projected_explosion: give the host harness the real block address helpers K4 made BlockExplodeProjected address storage through block_ddf and block_slot (core/layout.cuh). This host test compiles the kernel text extracted from blocks.cu, so it now also extracts those two helper definitions from core/layout.cuh (device qualifiers -> inline) instead of failing to compile. Assertions and tolerances unchanged. Co-Authored-By: Claude Opus 5.5 --- tests/unit/test_projected_explosion.py | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/tests/unit/test_projected_explosion.py b/tests/unit/test_projected_explosion.py index 3c16634..3768d16 100644 --- a/tests/unit/test_projected_explosion.py +++ b/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 \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'