From a4fc46d849c854b3fba794e04208650bad4fe548 Mon Sep 17 00:00:00 2001 From: Claude Date: Wed, 28 Oct 2026 10:55:02 +0000 Subject: [PATCH] block_collide: one unforced uniform-equivalent branch; drop collide_mrt_unforced The blocked collision had three arms (MRT wrapper, channel SRT, level SRT) and a one-caller MRT wrapper in collision_mrt.cuh. Now block_collide computes rho/u once and has two compile-time families: - uniform-equivalent (MRT any level, SRT channel root): collide_dispatch arithmetic with zero force, no LES: one clamped omega, collide_mrt_no_force or collide_srt with its zero Guo term (unchanged); - level SRT (BLOCK_ROOT_BC 0/1): unclamped, collide_srt_no_force, shifted momentum feq exception kept visible. No runtime d_params.fx/fy is read. Uniform kernels unchanged (only the unused inline wrapper is removed from collision_mrt.cuh). Source-text test updated to the new single call sites. Co-Authored-By: Claude Opus 5.5 --- .../lbm/kernels/operators/block_collision.cuh | 30 +++++++++++-------- .../lbm/kernels/operators/collision_mrt.cuh | 10 ------- tests/unit/test_refinement_mrt.py | 20 ++++++------- 3 files changed, 27 insertions(+), 33 deletions(-) diff --git a/src/CelerisLab/lbm/kernels/operators/block_collision.cuh b/src/CelerisLab/lbm/kernels/operators/block_collision.cuh index 9c436a9..a0e6234 100644 --- a/src/CelerisLab/lbm/kernels/operators/block_collision.cuh +++ b/src/CelerisLab/lbm/kernels/operators/block_collision.cuh @@ -1,10 +1,12 @@ +// CelerisLab – operators/block_collision.cuh // 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). +// Unforced by construction: FixedPatchStepper supplies no forcing, so no runtime +// d_params.fx/fy branch is read. Two compile-time arithmetic families: +// uniform-equivalent (MRT on every level; SRT on the channel root, BLOCK_ROOT_BC 2): +// collide_dispatch arithmetic with zero force and no LES: clamped omega; +// SRT keeps the zero Guo forcing term of collide_srt; +// level SRT (BLOCK_ROOT_BC 0/1): unclamped omega, no forcing term; shifted storage +// builds feq from the retained drho and direct momentum (numerical variant). #ifndef CELERIS_OPERATORS_BLOCK_COLLISION_CUH #define CELERIS_OPERATORS_BLOCK_COLLISION_CUH @@ -16,17 +18,19 @@ #endif __device__ __forceinline__ void block_collide(float* f) { -#if COLLISION_MODEL == 2 - collide_mrt_unforced(f); -#elif BLOCK_ROOT_BC == 2 - float rho, ux, uy, feq[NQ], forcing[NQ] = {}; + float rho, ux, uy; compute_rho_u(f, rho, ux, uy); - compute_feq(rho, ux, uy, feq); +#if COLLISION_MODEL == 2 || BLOCK_ROOT_BC == 2 const float omega = fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega)); + #if COLLISION_MODEL == 2 + collide_mrt_no_force(f, rho, ux, uy, omega); + #else + float feq[NQ], forcing[NQ] = {}; + compute_feq(rho, ux, uy, feq); collide_srt(f, feq, forcing, omega); + #endif #else - float rho, ux, uy, feq[NQ]; - compute_rho_u(f, rho, ux, uy); + float feq[NQ]; #if USE_DDF_SHIFTING // Preserve the density perturbation before adding the unit reference. const float drho = f[0]+f[1]+f[2]+f[3]+f[4]+f[5]+f[6]+f[7]+f[8]; diff --git a/src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh b/src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh index 80d4637..efc59c1 100644 --- a/src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh +++ b/src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh @@ -114,16 +114,6 @@ __device__ __forceinline__ void collide_mrt_no_force(float* __restrict__ g, collide_mrt(g, rho, ux, uy, Fin, omega); } -// Unforced MRT with the clamped device omega, as in collide_dispatch with zero -// forcing. Shared by the blocked, channel-root and dense-patch collisions. -__device__ __forceinline__ void collide_mrt_unforced(float* __restrict__ g) -{ - float rho, ux, uy; - compute_rho_u(g, rho, ux, uy); - collide_mrt_no_force(g, rho, ux, uy, - fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega))); -} - #elif LATTICE_MODEL == LATTICE_D3Q19 __device__ __forceinline__ void collide_mrt(float* __restrict__ g, diff --git a/tests/unit/test_refinement_mrt.py b/tests/unit/test_refinement_mrt.py index 02e2614..5a61a0d 100644 --- a/tests/unit/test_refinement_mrt.py +++ b/tests/unit/test_refinement_mrt.py @@ -32,22 +32,22 @@ class RefinementMRTHost(unittest.TestCase): "collide_srt_no_force(f, feq, d_params.omega);"): self.assertIn(text, block_collision) self.assertIn("collide_srt(f, feq, forcing, omega);", block_collision) - self.assertIn("#elif BLOCK_ROOT_BC == 2", block_collision) + self.assertIn("#if COLLISION_MODEL == 2 || BLOCK_ROOT_BC == 2", block_collision) self.assertEqual(channel.count("block_collide(f);"), 1) - self.assertNotIn("collide_mrt_unforced(f);", channel) + for text in ("collide_mrt_unforced", "collide_dispatch(", "d_params.fx", "d_params.fy"): + self.assertNotIn(text, channel + "\n" + block_collision) for source in (blocks,): self.assertIn("(COLLISION_MODEL != 0 && COLLISION_MODEL != 2)", source) self.assertIn("USE_LES", source) self.assertIn("STORE_PRECISION != 0", source) self.assertIn('#if COLLISION_MODEL == 2\n#include "operators/collision_mrt.cuh"\n#endif', source) - # One shared unforced MRT body; each blocked/patch path calls it exactly once. - mrt = (KERNELS / "operators/collision_mrt.cuh").read_text() - self.assertEqual(mrt.count("void collide_mrt_unforced("), 1) - for source in (blocks + "\n" + block_collision, channel + "\n" + block_collision): - self.assertEqual(source.count("collide_mrt_unforced(f);"), 1) - self.assertNotIn("collide_mrt_no_force(", source) - self.assertNotIn("compute_rho_u(f, rho, ux, uy);\n collide_mrt", source) - self.assertNotRegex(source, r"compute_rho_u\(f, rho, ux, uy\);\s*collide_mrt") + # One unforced MRT call, in block_collide only; the steps never collide directly. + self.assertEqual(block_collision.count("collide_mrt_no_force(f, rho, ux, uy, omega);"), 1) + self.assertEqual(block_collision.count("compute_rho_u(f, rho, ux, uy);"), 1) + self.assertEqual(block_collision.count("fmaxf(OMEGA_COLLISION_MIN, d_params.omega)"), 1) + for source in (blocks, channel): + self.assertNotIn("collide_mrt", source) + self.assertNotIn("collide_srt", source) if __name__ == "__main__": unittest.main()