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(){