From a3df5855a2114e17887c9f6211cb1ef5c3f693b6 Mon Sep 17 00:00:00 2001 From: Claude Date: Thu, 22 Oct 2026 00:26:41 +0000 Subject: [PATCH] kernels (K3): cell kinds in core/, capability guard first, self-contained block collision - core/cell_kind.cuh holds KIND_ACTIVE/INNER/OUTER (was a constexpr in blocks.cu that step/block_channel.cu and the channel_root.cpp shim relied on or redefined); blocks.cu, block_channel.cu and the shim now include it. - blocks.cu checks the lattice/model #error guard before the BLOCK_ROOT_BC guard and mode includes, so an unsupported build reports the real reason first. - operators/block_collision.cuh includes what it uses (params, macro, SRT/MRT collision) instead of relying on blocks.cu's include order. Preprocessor-only: same kernels, same statements, same arithmetic. Co-Authored-By: Claude Opus 5.5 --- src/CelerisLab/lbm/kernels/blocks.cu | 12 +++++------- src/CelerisLab/lbm/kernels/core/cell_kind.cuh | 10 ++++++++++ .../lbm/kernels/operators/block_collision.cuh | 7 +++++++ src/CelerisLab/lbm/kernels/step/block_channel.cu | 1 + tests/unit/cpp/channel_root.cpp | 2 +- 5 files changed, 24 insertions(+), 8 deletions(-) create mode 100644 src/CelerisLab/lbm/kernels/core/cell_kind.cuh diff --git a/src/CelerisLab/lbm/kernels/blocks.cu b/src/CelerisLab/lbm/kernels/blocks.cu index 2a9ef94..2ff9f89 100644 --- a/src/CelerisLab/lbm/kernels/blocks.cu +++ b/src/CelerisLab/lbm/kernels/blocks.cu @@ -13,6 +13,7 @@ #include "core/params.cuh" #include "core/flags.cuh" #include "core/obs.cuh" +#include "core/cell_kind.cuh" #include "operators/macro.cuh" #include "operators/collision_srt.cuh" #if COLLISION_MODEL == 2 @@ -20,6 +21,10 @@ #endif #include "streaming/pull_double_buffer.cuh" #include "boundary/curved_boundary.cuh" + +#if DIM != 2 || NQ != 9 || (COLLISION_MODEL != 0 && COLLISION_MODEL != 2) || STORE_PRECISION != 0 || USE_LES || STREAMING_MODEL != 0 +#error "Blocks require FP32 D2Q9 SRT or MRT double buffers without LES" +#endif #ifndef BLOCK_ROOT_BC #error "Blocks require a generated BLOCK_ROOT_BC mode" #elif BLOCK_ROOT_BC < 0 || BLOCK_ROOT_BC > 2 @@ -33,13 +38,6 @@ #include "operators/block_collision.cuh" #endif -#if DIM != 2 || NQ != 9 || (COLLISION_MODEL != 0 && COLLISION_MODEL != 2) || STORE_PRECISION != 0 || USE_LES || STREAMING_MODEL != 0 -#error "Blocks require FP32 D2Q9 SRT or MRT double buffers without LES" -#endif - -// Fixed-patch cell kinds, matching refinement.py ACTIVE/INNER_GHOST/OUTER_GHOST. -constexpr unsigned char KIND_ACTIVE = 0u, KIND_INNER = 1u, KIND_OUTER = 2u; - // Halo refresh: block cell ids b*TOTAL_CELLS+c, sources are always core cells. extern "C" __global__ void BlockHaloCopy( fpxx* f, const uint64_t* dst, const uint64_t* src, unsigned int count) diff --git a/src/CelerisLab/lbm/kernels/core/cell_kind.cuh b/src/CelerisLab/lbm/kernels/core/cell_kind.cuh new file mode 100644 index 0000000..16a0c61 --- /dev/null +++ b/src/CelerisLab/lbm/kernels/core/cell_kind.cuh @@ -0,0 +1,10 @@ +// CelerisLab – core/cell_kind.cuh +// Fixed-patch cell kinds stored per block cell, matching lbm/refinement.py +// ACTIVE/INNER_GHOST/OUTER_GHOST: OUTER is invalid (stores NaN, no pull), +// INNER streams without collision, ACTIVE pulls and collides. +#ifndef CELERIS_CORE_CELL_KIND_CUH +#define CELERIS_CORE_CELL_KIND_CUH + +constexpr unsigned char KIND_ACTIVE = 0u, KIND_INNER = 1u, KIND_OUTER = 2u; + +#endif diff --git a/src/CelerisLab/lbm/kernels/operators/block_collision.cuh b/src/CelerisLab/lbm/kernels/operators/block_collision.cuh index 56e0d26..612adf7 100644 --- a/src/CelerisLab/lbm/kernels/operators/block_collision.cuh +++ b/src/CelerisLab/lbm/kernels/operators/block_collision.cuh @@ -2,6 +2,13 @@ #ifndef CELERIS_OPERATORS_BLOCK_COLLISION_CUH #define CELERIS_OPERATORS_BLOCK_COLLISION_CUH +#include "../core/params.cuh" +#include "macro.cuh" +#include "collision_srt.cuh" +#if COLLISION_MODEL == 2 +#include "collision_mrt.cuh" +#endif + __device__ __forceinline__ void block_collide(float* f) { #if COLLISION_MODEL == 2 collide_mrt_unforced(f); // SRT path below unchanged. diff --git a/src/CelerisLab/lbm/kernels/step/block_channel.cu b/src/CelerisLab/lbm/kernels/step/block_channel.cu index 1ff1cd2..e570a8c 100644 --- a/src/CelerisLab/lbm/kernels/step/block_channel.cu +++ b/src/CelerisLab/lbm/kernels/step/block_channel.cu @@ -4,6 +4,7 @@ // Root-only blocked counterpart of OneStep's D2Q9/SRT channel closure. // NX/NY are LOCAL strides; all physical boundary decisions use global coordinates. +#include "../core/cell_kind.cuh" #include "../boundary/inlet/common.cuh" #include "../boundary/inlet/zou_he_local.cuh" #include "../boundary/outlet/pressure_neq.cuh" diff --git a/tests/unit/cpp/channel_root.cpp b/tests/unit/cpp/channel_root.cpp index 6385d28..2374044 100644 --- a/tests/unit/cpp/channel_root.cpp +++ b/tests/unit/cpp/channel_root.cpp @@ -41,7 +41,7 @@ inline float __fdiv_rn(float a,float b){return a/b;} #include "boundary/inlet_outlet.cuh" #include "step/one_step_double.cu" #else -constexpr unsigned char KIND_OUTER=2; +#include "core/cell_kind.cuh" #include "step/block_channel.cu" #endif int main(){