diff --git a/src/CelerisLab/lbm/kernels/blocks.cu b/src/CelerisLab/lbm/kernels/blocks.cu index 8e3764f..141d23a 100644 --- a/src/CelerisLab/lbm/kernels/blocks.cu +++ b/src/CelerisLab/lbm/kernels/blocks.cu @@ -178,11 +178,7 @@ extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int hal else if (gy == d_block_wall.ny - 2) reconstruct_top_reservoir(f, RHO); } #if COLLISION_MODEL == 2 - // Uniform MRT sequence (collide_dispatch, zero forcing); SRT path below unchanged. - float rho, ux, uy; - compute_rho_u(f, rho, ux, uy); - collide_mrt_no_force(f, rho, ux, uy, - fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega))); + collide_mrt_unforced(f); // SRT path below unchanged. #else float rho, ux, uy, feq[NQ]; compute_rho_u(f, rho, ux, uy); diff --git a/src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh b/src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh index efc59c1..80d4637 100644 --- a/src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh +++ b/src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh @@ -114,6 +114,16 @@ __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/src/CelerisLab/lbm/kernels/refinement.cu b/src/CelerisLab/lbm/kernels/refinement.cu index 93c79b9..f496b10 100644 --- a/src/CelerisLab/lbm/kernels/refinement.cu +++ b/src/CelerisLab/lbm/kernels/refinement.cu @@ -37,10 +37,7 @@ __device__ unsigned int incoming_counts[NQ * TOTAL_CELLS]; __device__ __forceinline__ void patch_collide(float* f, unsigned long k) { #if COLLISION_MODEL == 2 - float rho, ux, uy; - compute_rho_u(f, rho, ux, uy); - collide_mrt_no_force(f, rho, ux, uy, - fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega))); + collide_mrt_unforced(f); #else float rho, ux, uy, feq[NQ]; compute_rho_u(f, rho, ux, uy); diff --git a/src/CelerisLab/lbm/kernels/step/block_channel.cu b/src/CelerisLab/lbm/kernels/step/block_channel.cu index 17e90fb..67a613d 100644 --- a/src/CelerisLab/lbm/kernels/step/block_channel.cu +++ b/src/CelerisLab/lbm/kernels/step/block_channel.cu @@ -55,11 +55,7 @@ extern "C" __global__ void BlockChannelStep( } if (!wall && !outlet) { #if COLLISION_MODEL == 2 - // Match uniform MRT arithmetic (collide_dispatch, zero forcing). - float rho, ux, uy; - compute_rho_u(f, rho, ux, uy); - collide_mrt_no_force(f, rho, ux, uy, - fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega))); + collide_mrt_unforced(f); #else // Match uniform SRT arithmetic, including its shifted equilibrium path. float rho, ux, uy, feq[NQ], forcing[NQ] = {}; diff --git a/tests/unit/test_refinement_mrt.py b/tests/unit/test_refinement_mrt.py index 5f33782..be4d6f5 100644 --- a/tests/unit/test_refinement_mrt.py +++ b/tests/unit/test_refinement_mrt.py @@ -37,8 +37,13 @@ class RefinementMRTHost(unittest.TestCase): self.assertIn("STORE_PRECISION != 0", source) self.assertIn('#if COLLISION_MODEL == 2\n#include "operators/collision_mrt.cuh"\n#endif', source) self.assertIn("USE_DDF_SHIFTING ||", refinement) # dense/mixed shifted stays rejected + # 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, channel, refinement): - self.assertEqual(source.count("collide_mrt_no_force("), 1) + 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) if __name__ == "__main__": unittest.main()