Undermind审计前，Cursor重构

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

@@ -0,0 +1,107 @@
+// CelerisLab – step/aux_kernels.cu
+// Auxiliary kernels: CurvedBoundaryKernel, SensorKernel.
+// Launched on compact lists (n_curved / n_sensor threads).
+// Included by kernel_v2.cu.  No standalone compilation.
+//
+// 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.
+// ============================================================================
+
+#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;
+    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_out,
+    const fpxx* fi_in,
+    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;
+
+    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);
+
+#if DIM == 2
+    float Uw = -omega * ry;
+    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);
+#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,
+                       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,
+    const unsigned int* sensor_cells,
+    const int* sensor_obj_id,
+    float* obs,
+    unsigned int n_sensor)
+{
+    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
+    if (tid >= n_sensor) return;
+
+    unsigned long k = (unsigned long)sensor_cells[tid];
+    int id_obj = sensor_obj_id[tid];
+
+    float f[NQ];
+    for (int i = 0; i < NQ; i++)
+        f[i] = load_ddf(fi, index_f(k, (unsigned int)i));
+
+#if DIM == 2
+    float rho_n, ux, uy;
+    compute_rho_u(f, rho_n, ux, uy);
+    atomicAdd(&obs[obs_sensor_index(id_obj, 0)], ux);
+    atomicAdd(&obs[obs_sensor_index(id_obj, 1)], uy);
+#elif DIM == 3
+    float rho_n, ux, uy, uz;
+    compute_rho_u(f, rho_n, ux, uy, uz);
+    atomicAdd(&obs[obs_sensor_index(id_obj, 0)], ux);
+    atomicAdd(&obs[obs_sensor_index(id_obj, 1)], uy);
+    atomicAdd(&obs[obs_sensor_index(id_obj, 2)], uz);
+#endif
+
+    (void)flag;
+}
+
+#endif // CELERIS_STEP_AUX_KERNELS_CU