From 78c0f1811bf8859a92fe19466b0b6e385b53c184 Mon Sep 17 00:00:00 2001 From: Claude Date: Thu, 22 Oct 2026 01:32:40 +0000 Subject: [PATCH] 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 --- src/CelerisLab/lbm/kernels/blocks.cu | 15 ++++++++------- src/CelerisLab/lbm/kernels/core/layout.cuh | 15 --------------- tests/unit/test_projected_explosion.py | 8 +------- 3 files changed, 9 insertions(+), 29 deletions(-) diff --git a/src/CelerisLab/lbm/kernels/blocks.cu b/src/CelerisLab/lbm/kernels/blocks.cu index 21f61b3..44fc8c2 100644 --- a/src/CelerisLab/lbm/kernels/blocks.cu +++ b/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; } } diff --git a/src/CelerisLab/lbm/kernels/core/layout.cuh b/src/CelerisLab/lbm/kernels/core/layout.cuh index 1fce767..3b0f76a 100644 --- a/src/CelerisLab/lbm/kernels/core/layout.cuh +++ b/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 -__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) // --------------------------------------------------------------------------- diff --git a/tests/unit/test_projected_explosion.py b/tests/unit/test_projected_explosion.py index 3768d16..3c16634 100644 --- a/tests/unit/test_projected_explosion.py +++ b/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 \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'