Blocked steps: one block_collide extension point; drop unused HaloPlan, collide_trt_no_force, MIN_PATCH_GAP

SRT keeps its two explicit arithmetic variants (level: unclamped no-force,
shifted direct-momentum feq; channel: uniform OneStep arithmetic), selected
at compile time by BLOCK_ROOT_BC; text moved verbatim, PTX parity pending.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

src/CelerisLab/lbm/kernels/blocked.cu

@@ -34,9 +34,7 @@
 #if BLOCK_ROOT_BC == 1
 #include "boundary/block_wall.cuh"
 #endif
-#if BLOCK_ROOT_BC != 2
 #include "operators/block_collision.cuh"
-#endif
 
 // Kernels, in definition order: transfers, boundary staging, then the one full
 // step selected for this module (BlockStep, or BlockChannelStep for a channel root).

src/CelerisLab/lbm/kernels/core/params.cuh

@@ -44,16 +44,4 @@ struct RigidBodyControl2D {
     int   control_mode;           // 0=force, 1=velocity, 2=displacement
 };
 
-// ============================================================================
-// Halo plan (multi-GPU / AMR 预留)
-// ============================================================================
-struct HaloPlan {
-    int       face_id;
-    int       peer_gpu;
-    int       level;
-    size_t    count;
-    int*      send_idx;
-    int*      recv_idx;
-};
-
 #endif // CELERIS_CORE_PARAMS_CUH

src/CelerisLab/lbm/kernels/operators/block_collision.cuh

@@ -1,4 +1,10 @@
-// Collision used by periodic/fine and slot/reservoir BlockStep.
+// The one collision extension point of the blocked steps (BlockStep, BlockChannelStep).
+// Compile-time selection: MRT shares one path; SRT keeps two explicit arithmetic
+// variants, unchanged from their original kernels:
+//   level  (BLOCK_ROOT_BC 0/1): unclamped omega, no forcing term; shifted storage
+//          builds feq from the retained drho and direct momentum (numerical variant);
+//   channel (BLOCK_ROOT_BC 2): matches uniform OneStep SRT arithmetic (omega clamp,
+//          zero Guo forcing term, compute_feq).
 #ifndef CELERIS_OPERATORS_BLOCK_COLLISION_CUH
 #define CELERIS_OPERATORS_BLOCK_COLLISION_CUH
 
@@ -11,7 +17,13 @@
 
 __device__ __forceinline__ void block_collide(float* f) {
 #if COLLISION_MODEL == 2
-    collide_mrt_unforced(f);  // SRT path below unchanged.
+    collide_mrt_unforced(f);
+#elif BLOCK_ROOT_BC == 2
+    float rho, ux, uy, feq[NQ], forcing[NQ] = {};
+    compute_rho_u(f, rho, ux, uy);
+    compute_feq(rho, ux, uy, feq);
+    const float omega = fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega));
+    collide_srt(f, feq, forcing, omega);
 #else
     float rho, ux, uy, feq[NQ];
     compute_rho_u(f, rho, ux, uy);

src/CelerisLab/lbm/kernels/operators/collision_trt.cuh

@@ -57,25 +57,4 @@ __device__ __forceinline__ void collide_trt(float* __restrict__ f,
     }
 }
 
-__device__ __forceinline__ void collide_trt_no_force(float* __restrict__ f,
-                                                     const float* __restrict__ feq,
-                                                     float omega)
-{
-    const float wp = omega;
-    const float wm = compute_omega_minus(wp);
-
-    f[0] = f[0] - wp * (f[0] - feq[0]);
-
-    #pragma unroll
-    for (int i = 1; i < NQ; i += 2) {
-        const int ib = i + 1;
-        const float fi = f[i], fb = f[ib];
-        const float feqi = feq[i], feqb = feq[ib];
-        const float delta_s = 0.5f * ((fi - feqi) + (fb - feqb));
-        const float delta_a = 0.5f * ((fi - feqi) - (fb - feqb));
-        f[i]  = fi  - wp * delta_s - wm * delta_a;
-        f[ib] = fb  - wp * delta_s + wm * delta_a;
-    }
-}
-
 #endif // CELERIS_OPERATORS_COLLISION_TRT_CUH

src/CelerisLab/lbm/kernels/step/block_channel.cu

@@ -48,18 +48,7 @@ extern "C" __global__ void BlockChannelStep(
         } else {
             apply_wall_bb_d2q9_rows(gy, d_block_channel.ny - 2, f, in + base, k);
         }
-        if (!wall && !outlet) {
-#if COLLISION_MODEL == 2
-            collide_mrt_unforced(f);
-#else
-            // Match uniform SRT arithmetic, including its shifted equilibrium path.
-            float rho, ux, uy, feq[NQ], forcing[NQ] = {};
-            compute_rho_u(f, rho, ux, uy);
-            compute_feq(rho, ux, uy, feq);
-            const float omega = fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega));
-            collide_srt(f, feq, forcing, omega);
-#endif
-        }
+        if (!wall && !outlet) block_collide(f);
     }
     stream_pull_store(k, f, out + base);
 }

src/CelerisLab/lbm/refinement.py

@@ -151,7 +151,6 @@ class FixedPatch2D:
             raise ValueError("levels are not synchronized: fine_steps must equal 2*coarse_steps")
 
 
-MIN_PATCH_GAP = 4  # Legacy exported constant; union transfers impose no gap policy.
 
 
 @dataclass(frozen=True, eq=False)