Frank14f/CelerisLab
1
0 Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

test_projected_explosion: give the host harness the real block address helpers

Browse Source 
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 <noreply@anthropic.com>
This commit is contained in:
ClaudeandClaude Opus 5.5 committed 2026-10-22 00:35:07 +00:00
1 parent b47af12c9d
commit 43e226c59a
1 file changed
+7 -1
Download Patch File Download Diff File Expand all files Collapse all files
tests/unit/test_projected_explosion.py
+7 -1
View File
@@ -163,6 +163,12 @@ class ProjectedKernelTests(unittest.TestCase):
donors = parent.owner.ravel()[p.ghost_coarse] donors = parent.owner.ravel()[p.ghost_coarse]
source = (Path(__file__).resolve().parents[2] / source = (Path(__file__).resolve().parents[2] /
"src/CelerisLab/lbm/kernels/blocks.cu").read_text() "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 = source.split('extern "C" __global__ void BlockExplodeProjected(',1)[1]
kernel = 'void BlockExplodeProjected(' + kernel.split('// Targets are receiver-local SOURCE slots',1)[0] kernel = 'void BlockExplodeProjected(' + kernel.split('// Targets are receiver-local SOURCE slots',1)[0]
def array(name, values, kind): 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 += 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 += '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 += '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 += data+array('target',targets,'uint64_t')+array('donor',donors,'uint64_t')
code += array('neighbor',ids,'uint64_t')+array('offset',offsets,'float') code += array('neighbor',ids,'uint64_t')+array('offset',offsets,'float')
code += f'float fine[{child.nblocks*9*int(np.prod(child.local_shape))}]={{}};\n' code += f'float fine[{child.nblocks*9*int(np.prod(child.local_shape))}]={{}};\n'