Preserve D2 repairs and add bounded jet and fixed-patch kernels

src/CelerisLab/lbm/kernels/step/aux_kernels.cu

@@ -15,15 +15,7 @@
 //     (pre-streaming semantics = matching double-buffer fi_in[k, dir]).
 //   SensorKernel / ForceRegionKernel run after EsoPullStep, using
 //     t = step_count+1 for post-streaming physical-value semantics.
-//   Wall boundary functions (bounce_back.cuh) use load_physical_dir_esopull
-//     inside EsoPullStep (post-streaming, matching load_f_esopull output).
-//
-// NOTE: The wall boundary helpers use post-streaming reads (unlike Bouzidi)
-//       because they operate inside EsoPullStep on f[] already reconstructed
-//       by load_f_esopull.  Replacing post-streaming f[i] with post-streaming
-//       f[opp(i)] is the correct half-way bounce-back for EsoPull.
-//
-// Verified: Kan99b K2 (60k steps, MRT) — bit-identical to double-buffer.
+//   Y-wall donors are staged by PrepareEsoPullBoundaries before any backing writes.
 // ============================================================================
 
 #ifndef CELERIS_STEP_AUX_KERNELS_CU
@@ -32,54 +24,7 @@
 #include "../streaming/esopull_semantic_helpers.cuh"
 #include "../boundary/esopull_curved_helpers.cuh"
 
-__device__ __forceinline__ float action_omega(const float* action, int body_id)
-{
-    if (action == nullptr || body_id < 0) return 0.0f;
-    unsigned int base = (unsigned int)body_id * (unsigned int)OBS_BODY_SLOT_FLOATS;
-    return action[base + (unsigned int)(OBS_BODY_SLOT_FLOATS - 1)];
-}
-
-__global__ void CurvedBoundaryKernel(
-    fpxx* fi,
-    const unsigned int* cl_fluid_idx,
-    const unsigned char* cl_dir,
-    const float* cl_q,
-    const float* cl_rx,
-    const float* cl_ry,
-    const float* cl_rz,
-    const unsigned char* cl_fallback_class,
-    const int* cl_body_id,
-    const float* action,
-    float* obs,
-    unsigned int n_curved)
-{
-    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
-    if (tid >= n_curved) return;
-
-    const unsigned long k_f = (unsigned long)cl_fluid_idx[tid];
-    const unsigned int dir = (unsigned int)cl_dir[tid];
-    const float q = cl_q[tid];
-    const int bid = cl_body_id[tid];
-    const float rx = cl_rx[tid];
-    const float ry = cl_ry[tid];
-    const unsigned char fallback_class = cl_fallback_class[tid];
-    const float omega = action_omega(action, bid);
-
-#if DIM == 2
-    const float Uw = -omega * ry;
-    const float Vw =  omega * rx;
-    (void)cl_rz;
-    apply_bouzidi_link(dir, q, k_f, fi, Uw, Vw, rx, ry, fallback_class, obs, bid);
-#elif DIM == 3
-    const float rz = cl_rz[tid];
-    const float Uw = -omega * ry;
-    const float Vw =  omega * rx;
-    const float Ww = 0.0f;
-    apply_bouzidi_link(dir, q, k_f, fi, Uw, Vw, Ww,
-                       rx, ry, rz, fallback_class, obs, bid);
-#endif
-}
-// End of CurvedBoundaryKernel
+#include "curved_double.cu"
 
 // ---------------------------------------------------------------------------
 // PrepareEsoPullCurvedKernel -- read-only curved prep for EsoPull
@@ -146,7 +91,12 @@ __global__ void PrepareEsoPullCurvedKernel(
     } else {
         f_reflected = compute_bouzidi_reflection(f_toward, 0.0f, 0.0f, q, fallback);
     }
-    f_reflected += compute_moving_wall_addend(dir, q, fallback, Uw, Vw);
+    float rho_wall = USE_DDF_SHIFTING ? 1.0f : 0.0f;
+    if (Uw != 0.0f || Vw != 0.0f) {
+        for (unsigned int i = 0; i < NQ; ++i)
+            rho_wall += load_physical_dir_pre_streaming(k_f, i, fi, j_f, t);
+    }
+    f_reflected += compute_moving_wall_addend(dir, q, fallback, Uw, Vw, rho_wall);
 
     unsigned long dst_cell;
     unsigned int dst_slot;
@@ -249,14 +199,13 @@ __global__ void SensorKernel(
 }
 
 // ---------------------------------------------------------------------------
-// ForceRegionKernel — operator-split local Guo forcing correction
+// ForceRegionKernel — operator-split exact-difference force-density impulse
 //
-// Runs after the main collide+stream step.  Applies a Guo forcing term to
-// cells within a force_region footprint.  This is an operator-split correction,
-// not inline in OneStep's collide_dispatch.
-//
-// The Guo formula and c_tau = 1 - omega/2 conventions match collide_srt,
-// collide_mrt, and collide_trt exactly.
+// Runs after the main step. The raw source evaluated at u + F/(2 rho)
+// equals feq(rho, u + F/rho) - feq(rho, u) for the quadratic equilibrium.
+// No collision prefactor: collision has already happened. This preserves
+// density and adds exactly F momentum, independently of viscosity. It is
+// a split impulse, not the collision-integrated global Guo discretization.
 // ---------------------------------------------------------------------------
 __device__ __forceinline__ float action_force_x(const float* action, int body_id)
 {
@@ -319,11 +268,9 @@ __global__ void ForceRegionKernel(
     float Fin[NQ];
     compute_guo_forcing(ux, uy, fx, fy, Fin);
 
-    float omega = d_params.omega;
-    float c_tau = 1.0f - 0.5f * omega;
     #pragma unroll
     for (int i = 0; i < NQ; i++)
-        f[i] += c_tau * Fin[i];
+        f[i] += Fin[i];
 #elif DIM == 3
     float rho_n, ux, uy, uz;
     compute_rho_u(f, rho_n, ux, uy, uz);
@@ -334,11 +281,9 @@ __global__ void ForceRegionKernel(
     float Fin[NQ];
     compute_guo_forcing(ux, uy, uz, fx, fy, 0.0f, Fin);
 
-    float omega = d_params.omega;
-    float c_tau = 1.0f - 0.5f * omega;
     #pragma unroll
     for (int i = 0; i < NQ; i++)
-        f[i] += c_tau * Fin[i];
+        f[i] += Fin[i];
 #endif
 
     // Write back with correct semantics