From 93e692a40fa7105d9660d1cd678209f871e284ab Mon Sep 17 00:00:00 2001
From: Claude <claudeworker@noreply.invalid>
Date: Thu, 22 Oct 2026 00:32:07 +0000
Subject: [PATCH] kernels (K5): reuse the existing bounce_back_swap instead of
 a new copy

operators/helpers.cuh already had bounce_back_swap (the uniform steps use
it). It moves verbatim, #pragma unroll included, into its own
operators/bounce_back_swap.cuh, which helpers.cuh includes; the channel
root includes that header and calls it. bounce_back_swap_d2q9 from the
previous commit is removed. bounce_back.cuh keeps only the half-way rows
helper.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
---
 .../lbm/kernels/boundary/bounce_back.cuh           |  8 --------
 .../lbm/kernels/operators/bounce_back_swap.cuh     | 14 ++++++++++++++
 src/CelerisLab/lbm/kernels/operators/helpers.cuh   |  7 +------
 src/CelerisLab/lbm/kernels/step/block_channel.cu   |  3 ++-
 4 files changed, 17 insertions(+), 15 deletions(-)
 create mode 100644 src/CelerisLab/lbm/kernels/operators/bounce_back_swap.cuh

diff --git a/src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh b/src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh
index 1eaac1b..e0a977e 100644
--- a/src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh
+++ b/src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh
@@ -45,14 +45,6 @@ __device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y,
     apply_wall_bb_d2q9_rows(y, (unsigned int)(NY - 2), f, fi_in, k);
 }
 
-// Full-way swap of every opposite pair (paired order q, q+1) at a wall node.
-__device__ __forceinline__ void bounce_back_swap_d2q9(float* f)
-{
-    for (int q = 1; q < NQ; q += 2) {
-        const float t = f[q]; f[q] = f[q + 1]; f[q + 1] = t;
-    }
-}
-
 // ---------------------------------------------------------------------------
 // D2Q9 y-wall free-slip (specular reflection) for pull double-buffer streaming.
 //
diff --git a/src/CelerisLab/lbm/kernels/operators/bounce_back_swap.cuh b/src/CelerisLab/lbm/kernels/operators/bounce_back_swap.cuh
new file mode 100644
index 0000000..439ea73
--- /dev/null
+++ b/src/CelerisLab/lbm/kernels/operators/bounce_back_swap.cuh
@@ -0,0 +1,14 @@
+// CelerisLab – operators/bounce_back_swap.cuh
+// Full-way bounce-back at a wall node: swap every opposite pair (paired
+// order q, q+1). Needs only NQ; shared by the uniform and blocked steps.
+#ifndef CELERIS_OPERATORS_BOUNCE_BACK_SWAP_CUH
+#define CELERIS_OPERATORS_BOUNCE_BACK_SWAP_CUH
+
+__device__ __forceinline__ void bounce_back_swap(float* f) {
+    #pragma unroll
+    for (int i = 1; i < NQ; i += 2) {
+        float t = f[i]; f[i] = f[i+1]; f[i+1] = t;
+    }
+}
+
+#endif
diff --git a/src/CelerisLab/lbm/kernels/operators/helpers.cuh b/src/CelerisLab/lbm/kernels/operators/helpers.cuh
index 95bee56..a21106e 100644
--- a/src/CelerisLab/lbm/kernels/operators/helpers.cuh
+++ b/src/CelerisLab/lbm/kernels/operators/helpers.cuh
@@ -5,12 +5,7 @@
 #ifndef CELERIS_OPERATORS_HELPERS_CUH
 #define CELERIS_OPERATORS_HELPERS_CUH
 
-__device__ __forceinline__ void bounce_back_swap(float* f) {
-    #pragma unroll
-    for (int i = 1; i < NQ; i += 2) {
-        float t = f[i]; f[i] = f[i+1]; f[i+1] = t;
-    }
-}
+#include "bounce_back_swap.cuh"
 
 __device__ __forceinline__ void collide_dispatch(
     float* f, float rho_n,
diff --git a/src/CelerisLab/lbm/kernels/step/block_channel.cu b/src/CelerisLab/lbm/kernels/step/block_channel.cu
index 4efa5be..7390633 100644
--- a/src/CelerisLab/lbm/kernels/step/block_channel.cu
+++ b/src/CelerisLab/lbm/kernels/step/block_channel.cu
@@ -5,6 +5,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 "../operators/bounce_back_swap.cuh"
 #include "../boundary/bounce_back.cuh"
 #include "../boundary/inlet/common.cuh"
 #include "../boundary/inlet/zou_he_local.cuh"
@@ -35,7 +36,7 @@ extern "C" __global__ void BlockChannelStep(
         const bool wall = gy == 0 || gy == d_block_channel.ny - 1;
         const bool outlet = gx == d_block_channel.nx - 1;
         if (wall) {
-            bounce_back_swap_d2q9(f);  // Wall owns all four corners; ghosts never collide.
+            bounce_back_swap(f);  // Wall owns all four corners; ghosts never collide.
         } else if (gx == 0) {
             apply_zou_he_left_velocity_inlet_d2q9(
                 f, inlet_target_u((float)gy, d_block_channel.ny, d_block_channel.velocity), 0.0f);