diff --git a/src/CelerisLab/lbm/kernels/blocked.cu b/src/CelerisLab/lbm/kernels/blocked.cu
index 18f7e42..98c5b22 100644
--- a/src/CelerisLab/lbm/kernels/blocked.cu
+++ b/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).
diff --git a/src/CelerisLab/lbm/kernels/core/params.cuh b/src/CelerisLab/lbm/kernels/core/params.cuh
index a4b2ded..eec3833 100644
--- a/src/CelerisLab/lbm/kernels/core/params.cuh
+++ b/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
diff --git a/src/CelerisLab/lbm/kernels/operators/block_collision.cuh b/src/CelerisLab/lbm/kernels/operators/block_collision.cuh
index 612adf7..9c436a9 100644
--- a/src/CelerisLab/lbm/kernels/operators/block_collision.cuh
+++ b/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);
diff --git a/src/CelerisLab/lbm/kernels/operators/collision_trt.cuh b/src/CelerisLab/lbm/kernels/operators/collision_trt.cuh
index ed28496..e5aa211 100644
--- a/src/CelerisLab/lbm/kernels/operators/collision_trt.cuh
+++ b/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
diff --git a/src/CelerisLab/lbm/kernels/step/block_channel.cu b/src/CelerisLab/lbm/kernels/step/block_channel.cu
index 7390633..664413a 100644
--- a/src/CelerisLab/lbm/kernels/step/block_channel.cu
+++ b/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);
 }
diff --git a/src/CelerisLab/lbm/refinement.py b/src/CelerisLab/lbm/refinement.py
index 37d04b9..7c7e157 100644
--- a/src/CelerisLab/lbm/refinement.py
+++ b/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)