Sah04验证部分通过，临时修复部分NEQ inlet解决高阻塞流动。kan99b验证初期。

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

@@ -6,17 +6,16 @@
 // CurvedBoundaryKernel: per-link geometry is passed as cl_rx/y/z and runtime
 // angular state is read from action[body_id * OBS_BODY_SLOT_FLOATS + ...].
 // For DIM==2, wall velocity uses rigid-body rotation: Uw=-omega*ry, Vw=omega*rx.
+//
+// Important timing:
+//   CurvedBoundaryKernel runs BEFORE the main pull step. It writes corrected
+//   source populations into obstacle nodes so the pull loader receives valid
+//   incoming distributions on the same step.
 // ============================================================================
 
 #ifndef CELERIS_STEP_AUX_KERNELS_CU
 #define CELERIS_STEP_AUX_KERNELS_CU
 
-// ============================================================================
-// CurvedBoundaryKernel — per-link Bouzidi curved BC (n_curved threads)
-//
-// One thread per cut link.  Reads from fi_out (current step, post-collision)
-// and fi_in (previous step) to compute Bouzidi linear interpolation.
-// ============================================================================
 __device__ __forceinline__ float action_omega(const float* action, int body_id)
 {
     if (action == nullptr || body_id < 0) return 0.0f;
@@ -25,8 +24,7 @@ __device__ __forceinline__ float action_omega(const float* action, int body_id)
 }
 
 __global__ void CurvedBoundaryKernel(
-    fpxx* fi_out,
-    const fpxx* fi_in,
+    fpxx* fi,
     const unsigned int* cl_fluid_idx,
     const unsigned char* cl_dir,
     const float* cl_q,
@@ -42,34 +40,30 @@ __global__ void CurvedBoundaryKernel(
     unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
     if (tid >= n_curved) return;
 
-    unsigned long k_f = (unsigned long)cl_fluid_idx[tid];
-    unsigned int dir = (unsigned int)cl_dir[tid];
-    float q = cl_q[tid];
-    int bid = cl_body_id[tid];
-    float rx = cl_rx[tid];
-    float ry = cl_ry[tid];
-    unsigned char fallback_class = cl_fallback_class[tid];
-    float omega = action_omega(action, bid);
+    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
-    float Uw = -omega * ry;
-    float Vw =  omega * rx;
+    const float Uw = -omega * ry;
+    const float Vw =  omega * rx;
     (void)cl_rz;
-    apply_bouzidi_link(dir, q, k_f, fi_out, fi_in, Uw, Vw, rx, ry, fallback_class, obs, bid);
+    apply_bouzidi_link(dir, q, k_f, fi, Uw, Vw, rx, ry, fallback_class, obs, bid);
 #elif DIM == 3
-    float rz = cl_rz[tid];
-    // Placeholder 3D angular contract: read scalar omega_z from action tail.
-    float Uw = -omega * ry;
-    float Vw =  omega * rx;
-    float Ww = 0.0f;
-    apply_bouzidi_link(dir, q, k_f, fi_out, fi_in, Uw, Vw, Ww,
+    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
 }
 
-// ============================================================================
-// SensorKernel — compact-list sensor readout (n_sensor threads)
-// ============================================================================
 __global__ void SensorKernel(
     const fpxx* fi,
     const unsigned short* flag,