From 1fa498c7d7ba8857c72c9740a0d6fed086e3f528 Mon Sep 17 00:00:00 2001 From: Claude Date: Wed, 21 Oct 2026 01:24:27 +0000 Subject: [PATCH] kernels: shared block-window to global coordinate helper BlockStep (slot-reservoir wall rows) and BlockChannelStep computed the global cell of a block-window cell with the same four lines. block_global_xy (core/layout.cuh, 2D, __forceinline__) now holds them; both kernels call it with their own domain width. Same integer operations: no numerical, ABI or scheduling change. The slot-reservoir kernel layout is unchanged (no split). Co-Authored-By: Claude Opus 5.5 --- src/CelerisLab/lbm/kernels/blocks.cu | 6 ++---- src/CelerisLab/lbm/kernels/core/layout.cuh | 12 ++++++++++++ src/CelerisLab/lbm/kernels/step/block_channel.cu | 6 ++---- 3 files changed, 16 insertions(+), 8 deletions(-) diff --git a/src/CelerisLab/lbm/kernels/blocks.cu b/src/CelerisLab/lbm/kernels/blocks.cu index 141d23a..b8c7ef4 100644 --- a/src/CelerisLab/lbm/kernels/blocks.cu +++ b/src/CelerisLab/lbm/kernels/blocks.cu @@ -168,10 +168,8 @@ extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int hal stream_pull_load(k, f, in + base, j); if (cell_kind == KIND_ACTIVE) { if (d_block_wall.nx != 0) { - const unsigned int bx = NX - 2 * halo, by = NY - 2 * halo; - const unsigned int nbx = d_block_wall.nx / bx; - const unsigned int gx = (blockIdx.z % nbx) * bx + x - halo; - const unsigned int gy = (blockIdx.z / nbx) * by + y - halo; + unsigned int gx, gy; + block_global_xy(x, y, halo, d_block_wall.nx, gx, gy); if (gy == 1) reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity( gx, d_block_wall.nx, d_block_wall.slot_width, d_block_wall.amplitude)); diff --git a/src/CelerisLab/lbm/kernels/core/layout.cuh b/src/CelerisLab/lbm/kernels/core/layout.cuh index 56a6d61..3b0f76a 100644 --- a/src/CelerisLab/lbm/kernels/core/layout.cuh +++ b/src/CelerisLab/lbm/kernels/core/layout.cuh @@ -53,6 +53,18 @@ __device__ __forceinline__ void index_from_thread(unsigned int& x, k = (unsigned long)y * NX + x; } +// Global cell of window cell (x, y) in block blockIdx.z of a row-major, +// full-domain block layout domain_nx cells wide. NX/NY are the local strides. +__device__ __forceinline__ void block_global_xy(unsigned int x, unsigned int y, + unsigned int halo, unsigned int domain_nx, + unsigned int& gx, unsigned int& gy) +{ + const unsigned int bx = NX - 2 * halo, by = NY - 2 * halo; + const unsigned int nbx = domain_nx / bx; + gx = (blockIdx.z % nbx) * bx + x - halo; + gy = (blockIdx.z / nbx) * by + y - halo; +} + // --------------------------------------------------------------------------- // Neighbor list (periodic BCs via modular arithmetic) // --------------------------------------------------------------------------- diff --git a/src/CelerisLab/lbm/kernels/step/block_channel.cu b/src/CelerisLab/lbm/kernels/step/block_channel.cu index 67a613d..d31bd0d 100644 --- a/src/CelerisLab/lbm/kernels/step/block_channel.cu +++ b/src/CelerisLab/lbm/kernels/step/block_channel.cu @@ -15,10 +15,8 @@ extern "C" __global__ void BlockChannelStep( unsigned int x, y; unsigned long k; index_from_thread(x, y, k); if (x < halo || y < halo || x >= NX - halo || y >= NY - halo) return; - const unsigned int bx = NX - 2 * halo, by = NY - 2 * halo; - const unsigned int nbx = d_block_channel.nx / bx; - const unsigned int gx = (blockIdx.z % nbx) * bx + x - halo; - const unsigned int gy = (blockIdx.z / nbx) * by + y - halo; + unsigned int gx, gy; + block_global_xy(x, y, halo, d_block_channel.nx, gx, gy); const unsigned long base = (unsigned long)blockIdx.z * NQ * TOTAL_CELLS; const unsigned long cell = (unsigned long)blockIdx.z * TOTAL_CELLS + k; float f[NQ];