From 33837697da3b576339b78a813cc610f5da3e52eb Mon Sep 17 00:00:00 2001 From: Claude Date: Thu, 22 Oct 2026 00:31:26 +0000 Subject: [PATCH] kernels (K5): channel root reuses boundary and streaming helpers - boundary/bounce_back.cuh: apply_wall_bb_d2q9_rows(y, top_row, ...) is the existing half-way y-wall correction with an explicit top row; apply_wall_bb_d2q9_y_pull keeps its signature and calls it with NY - 2. bounce_back_swap_d2q9(f) names the full opposite-pair swap. - streaming/pull_double_buffer.cuh: load_node(fi, n, f) loads one node's stored populations. - step/block_channel.cu uses all three instead of its inline copies: the wall swap, the outlet donor gather, and rows 1 / ny-2 (now one call with the physical top row; same y == 1 / y == top_row order and loads). Same statements, branch order, indices and arithmetic. Co-Authored-By: Claude Opus 5.5 --- .../lbm/kernels/boundary/bounce_back.cuh | 31 +++++++++++++++---- .../lbm/kernels/step/block_channel.cu | 19 +++--------- .../kernels/streaming/pull_double_buffer.cuh | 8 +++++ 3 files changed, 38 insertions(+), 20 deletions(-) diff --git a/src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh b/src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh index 69210ba..1eaac1b 100644 --- a/src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh +++ b/src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh @@ -15,10 +15,13 @@ // D2Q9 y-wall half-way bounce-back for pull double-buffer streaming. // --------------------------------------------------------------------------- #if LATTICE_MODEL == LATTICE_D2Q9 -__device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y, - float* __restrict__ f, - const fpxx* __restrict__ fi_in, - unsigned long k) +// Rows 1 and top_row (= physical ny - 2) next to walls at rows 0 and ny - 1. +// y is the physical row; k addresses the node in fi_in (a block window may +// have NX/NY strides different from the physical extent). +__device__ inline void apply_wall_bb_d2q9_rows(unsigned int y, unsigned int top_row, + float* __restrict__ f, + const fpxx* __restrict__ fi_in, + unsigned long k) { if (y == 1) { // Directions sourced from y=0 wall in pull step: +y, +x+y, -x+y. @@ -26,14 +29,30 @@ __device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y, f[5] = load_ddf(fi_in, index_f(k, 6u)); f[8] = load_ddf(fi_in, index_f(k, 7u)); } - else if (y == (unsigned int)(NY - 2)) { - // Directions sourced from y=NY-1 wall in pull step: -y, -x-y, +x-y. + else if (y == top_row) { + // Directions sourced from the y = ny - 1 wall in pull step: -y, -x-y, +x-y. f[4] = load_ddf(fi_in, index_f(k, 3u)); f[6] = load_ddf(fi_in, index_f(k, 5u)); f[7] = load_ddf(fi_in, index_f(k, 8u)); } } +__device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y, + float* __restrict__ f, + const fpxx* __restrict__ fi_in, + unsigned long k) +{ + 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/step/block_channel.cu b/src/CelerisLab/lbm/kernels/step/block_channel.cu index 5791a50..4efa5be 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 "../boundary/bounce_back.cuh" #include "../boundary/inlet/common.cuh" #include "../boundary/inlet/zou_he_local.cuh" #include "../boundary/outlet/pressure_neq.cuh" @@ -34,27 +35,17 @@ 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) { - // Wall owns all four corners. Ghosts pull and swap, never collide. - for (int q = 1; q < NQ; q += 2) { - const float t = f[q]; f[q] = f[q + 1]; f[q + 1] = t; - } + bounce_back_swap_d2q9(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); } else if (outlet) { // Immutable input donor, receiver-local (also valid across a block seam). float donor[NQ]; - for (int q = 0; q < NQ; ++q) - donor[q] = load_ddf(in + base, index_f(k - 1, q)); + load_node(in + base, k - 1, donor); apply_pressure_outlet_d2q9(f, donor); - } else if (gy == 1) { - f[3] = load_ddf(in + base, index_f(k, 4)); - f[5] = load_ddf(in + base, index_f(k, 6)); - f[8] = load_ddf(in + base, index_f(k, 7)); - } else if (gy == d_block_channel.ny - 2) { - f[4] = load_ddf(in + base, index_f(k, 3)); - f[6] = load_ddf(in + base, index_f(k, 5)); - f[7] = load_ddf(in + base, index_f(k, 8)); + } else { + apply_wall_bb_d2q9_rows(gy, d_block_channel.ny - 2, f, in + base, k); } if (!wall && !outlet) { #if COLLISION_MODEL == 2 diff --git a/src/CelerisLab/lbm/kernels/streaming/pull_double_buffer.cuh b/src/CelerisLab/lbm/kernels/streaming/pull_double_buffer.cuh index 413ae3c..6620a62 100644 --- a/src/CelerisLab/lbm/kernels/streaming/pull_double_buffer.cuh +++ b/src/CelerisLab/lbm/kernels/streaming/pull_double_buffer.cuh @@ -21,6 +21,14 @@ // i.e. the population traveling in direction i at node n // arrived from the node at n - c_i = neighbor in direction opp(i) // --------------------------------------------------------------------------- +// load_node: all NQ populations of node n as stored (no streaming). +__device__ __forceinline__ void load_node(const fpxx* __restrict__ fi, unsigned long n, + float* __restrict__ f) +{ + for (int q = 0; q < NQ; ++q) + f[q] = load_ddf(fi, index_f(n, q)); +} + __device__ __forceinline__ void stream_pull_load(unsigned long n, float* __restrict__ f, const fpxx* __restrict__ fi,