feat(body): add force_region, fix sensor API, reorganize docs

- Add force_region object type: local Guo forcing via sparse compact list
- ForceRegionSoA container, ForceRegionKernel, stepper dispatch
- add_body("force_region", ...) + set_force(id, fx, fy) API
- Fix read_sensor(normalize=...) not being passed from Simulation layer
- Fix force_region incorrectly entering curved cut-link path (P0 blocker)
- Clean up module boundaries: body/__init__ no longer imports from lbm
- Circluar import fix: common/streakline <-> pathline
- Package data globs fixed for recursive kernel files
- Version unified to 0.3.0
- Performance analysis: pycuda launch overhead vs GPU compute at various grid sizes
- Nsight Systems + Nsight Compute profiling data and report
- Documentation reorganized under docs/ (audit, validation_specs)
- README overhaul: multi-body examples, validated benchmarks, force_region docs

Co-authored-by: Cursor <cursoragent@cursor.com>

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

@@ -98,4 +98,93 @@ __global__ void SensorKernel(
     (void)flag;
 }
 
+// ---------------------------------------------------------------------------
+// ForceRegionKernel — operator-split local Guo forcing correction
+//
+// 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.
+// ---------------------------------------------------------------------------
+__device__ __forceinline__ float action_force_x(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 + 0u];
+}
+
+__device__ __forceinline__ float action_force_y(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 + 1u];
+}
+
+__global__ void ForceRegionKernel(
+    fpxx* fi,
+    const unsigned short* flag,
+    const unsigned int* fr_cells,
+    const int* fr_obj_id,
+    const float* action,
+    unsigned int n_cells)
+{
+    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
+    if (tid >= n_cells) return;
+
+    unsigned long k = (unsigned long)fr_cells[tid];
+    int bid = fr_obj_id[tid];
+    float fx = action_force_x(action, bid);
+    float fy = action_force_y(action, bid);
+    if (fx == 0.0f && fy == 0.0f) return;
+
+    // Only apply force to fluid cells — skip obstacles, walls, boundaries
+    if (!is_fluid(flag[k])) return;
+
+    float f[NQ];
+    #pragma unroll
+    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);
+
+    // Guo velocity correction: u* = u + F / (2 * rho)
+    float rho2 = 0.5f / rho_n;
+    ux = fmaf(fx, rho2, ux);
+    uy = fmaf(fy, rho2, uy);
+
+    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];
+#elif DIM == 3
+    float rho_n, ux, uy, uz;
+    compute_rho_u(f, rho_n, ux, uy, uz);
+
+    float rho2 = 0.5f / rho_n;
+    ux = fmaf(fx, rho2, ux);
+    uy = fmaf(fy, rho2, uy);
+    float Fin[NQ];
+    // fz placeholder: 0.0f (z-force not plumbed through action slots yet)
+    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];
+#endif
+
+    #pragma unroll
+    for (int i = 0; i < NQ; i++)
+        store_ddf(fi, index_f(k, (unsigned int)i), f[i]);
+}
+
 #endif // CELERIS_STEP_AUX_KERNELS_CU