重构body api，性能分析，项目整理

docs/audit/audit_round2.md

@@ -0,0 +1,209 @@
+## 审计结论（第二轮）
+
+本轮审计的视角与上一轮不同。上一轮聚焦"主链路上是否有显式 bug"（找到了 16 项）。本轮审计的四个维度：
+
+| 维度 | 上一轮 | 本轮 |
+|---|---|---|
+| 主链路正确性 | 找显式 bug | 确认修补正确、无回归 |
+| 代码架构质量 | 少量提及 | 系统评估 Python 层职责分离 |
+| 新代码 | 不存在 | streakline.py 全篇 + test runners |
+| 跨层一致性 | 仅一处 | 系统性检查 config → compiler → kernel → docs 同步 |
+
+核心结论：**上次 16 项修补经逐条审查确认正确，无回归。**
+
+### 第二轮修改完成状态
+
+根据审计发现，已完成以下修改：
+
+| 修改项 | 状态 |
+|---|---|
+| EsoPull 添加 y=1/NY-2 半格 bounce-back 修正（D2Q9 + D3Q19） | ✅ 已实施 |
+| `macro.cuh` diagnostic 函数标注 | ✅ 已标注 `// --- Diagnostic only ---` |
+| `config.py` `omega_max` 默认值 1.99 → 1.96 | ✅ 已修改 |
+| `BC_MOVING`/`BC_PERIODIC` 代码注释说明 | ✅ 已添加 TODO 注释 |
+| Streakline 模块重构：779 行单文件 → `common/streakline/` 子包 (5 文件) | ✅ 已完成 |
+| `run_kan99b_streakline.py` 更新为新 API | ✅ 已完成 |
+| `run_exp_ctrl_matrix_streakline.py` 更新为新 API | ✅ 已完成 |
+| `render_vorticity_field` 移到 `common/render.py` | ✅ 已完成 |
+| `ParticleTrailSet` 移到 `common/pathline.py` | ✅ 已完成 |
+| 向后兼容 shim `common/streakline.py` | ✅ 已创建（带 DeprecationWarning） |
+
+---
+
+## 状态说明
+
+- `[已确认]` 经代码审查确认正确
+- `[无法确认: 需运行验证]` 需数值算例确认，不能单靠读代码定论
+- `[新发现]` 本轮审计首次发现
+- `[已修复]` 已实施修改
+- `[保留说明]` 当前不修，但需在代码或文档中明确限制
+
+---
+
+## 第一轮审计修补确认
+
+逐一确认 16 项"已解决"修补在代码中的真实状态。**全部 [已确认]，无回归。**
+
+| 问题 | 结论 | 关键文件:行号 |
+|---|---|---|
+| `lbm/__init__.py` 导出错误 | 已确认 | 当前导出真实存在的 `add_vortex` |
+| forcing 主链路 + 预因子不一致 | 已确认 — 三模型统一使用 `c_tau = 1-omega/2` | `operators/collision_srt.cuh:15`, `collision_trt.cuh:35`, `collision_mrt.cuh:41/116` |
+| TRT outlet NEQ 重构未补齐 | 已确认 — COMPILE_MODEL==0\|1 时全分布 damped NEQ | `boundary/outlet/pressure_neq.cuh:45-51` (D2Q9) |
+| `add_vortex()` 动量当速度 | 已确认 — `ux = sum(f*cx) / rho_safe` | `lbm/initializers.py:57-58` |
+| Sensor 面积归一化 | 已确认 — ObjectManager 层已提供 | `body/manager.py` |
+| `sync_to_gpu()` 重置非流体 | 已确认 — 已收缩为只覆盖 obstacle interior | `body/manager.py` |
+| curved donor 合法性 | 已确认 — 扩展到实际 domain flags | `body/objects.py` `_donor_is_fluid` |
+| curved Bouzidi 时序 | 已确认 — 步前写入 obstacle source slot | `lbm/stepper.py:70-71` → `step/aux_kernels.cu:14` |
+| `q >= 0.5` 分支读错时间层 | 已确认 — 读 `load_ddf(fi, ...)` 即同一步 post-collision | `boundary/curved_boundary.cuh:73-75` |
+| moving wall 修正未按 q 分支 | 已确认 — 分三路：fallback / q<0.5 / q>=0.5 | `boundary/curved_boundary.cuh:30-43` |
+| 初始化链路 flag 叠加顺序 | 已确认 — obstacle overlay → init kernel preserve → equilibrium | `step/init_flow.cu:48-52`, `lbm/stepper.py:43-54` |
+| `config_body.json` 未进入初始化 | 已确认 — Simulation 已消费 | `simulation.py` |
+| inlet `U0` 语义 | 已确认 — 已补充截面平均速度注释 | `configs/CONFIG.md`, `README` |
+
+---
+
+## 本轮发现与处理
+
+### CUDA Kernel 层
+
+| 严重度 | 问题 | 文件:行号 | 处理 |
+|---|---|---|---|
+| **[新发现]** [高] | **`BC_MOVING` 与 `BC_PERIODIC` 标记未在 step kernel 中分发。** `core/flags.cuh` 定义了两种标记，但 step kernel 中没有 `is_moving()` 或 `is_periodic()` 分支。设了这两种标记的 cell 会静默 fall through 到 `bounce_back_swap()`。 | `step/one_step_double.cu`, `step/one_step_esopull.cu` | **[保留说明]** 已添加 TODO 注释，暂不实现 |
+| **[新发现]** [高] | **EsoPull 缺少 y=1/NY-2 的显式半格 bounce-back 修正。** Double-buffer 路径有该修正而 EsoPull 无。 | `step/one_step_esopull.cu` | **[已修复]** 已添加 D2Q9 和 D3Q19 分支，标注了未来可能需要更精确方案 |
+| **[新发现]** [中] | **`USE_DDF_SHIFTING` 路径完整性存疑。** 审查后确认：`macro.cuh` 和 `init_flow.cu` 已处理 shifted，各 collision/curved operator 通过 `load_ddf`/`store_ddf` 抽象层自动适配。**实际已完整，无需修改。** | — | **[已确认无需处理]** |
+| **[新发现]** [低] | **`compute_pressure` 与 `compute_pressure_perturbation` 未标注诊断用途。** | `operators/macro.cuh:160-166` | **[已修复]** 已标注 `// --- Diagnostic only ---` |
+| **[新发现]** [低] | **Curved/Sensor kernel 无运行时下标越界保护。** 完全依赖 host 侧验证。 | `step/aux_kernels.cu:26-99` | **[保留说明]** 设计选择：不引入多余合法性检查 |
+
+### 第一轮遗留待验证项复查
+
+| 待验证项 | 状态 | 说明 |
+|---|---|---|
+| **MRT D2Q9 moment transform/inverse** | **[无法确认: 需运行验证]** | 方向索引与 paired ordering 自洽，moment 投影物理正确。但全部系数的数值精度需在 Poiseuille 或衰减涡算例中验证。 |
+| **EsoPull 邻壁 vs double-buffer 一致性** | **[已修复]** | 已添加 y=1/NY-2 反弹修正。需运行验证确认。 |
+| **Force 提取与 host 侧归一化** | **[无需处理]** | Cd/Cl 管道经确认语义一致。 |
+| **Plain linear Bouzidi / TRT 不兼容** | **[保留说明]** | 注释已存在且准确，引用了 [Gin08b]。无方法级替代方案。 |
+
+---
+
+## 架构与设计审计
+
+### Streakline 模块重构
+
+**旧 `common/streakline.py` (779 行) 已被拆解为 `common/streakline/` 子包：**
+
+```
+src/CelerisLab/common/
+  streakline/                 后处理子包
+    __init__.py               导出 Streakline, ReleaseConfig, IntegratorConfig
+    _config.py                配置 dataclass + FlowFrame 内部类
+    _integrate.py             核心积分引擎 (RK4 + 时空插值)
+    _render.py                密度图渲染 (render_density)
+    _streakline.py            Streakline 类 (被动消费者)
+  render.py                   涡度计算与渲染 (从旧 streakline 移出)
+  pathline.py                 ParticleTrailSet + render_trails (从旧 streakline 移出)
+  preprocess.py               增加 cylinders_from_triangle_layout
+```
+
+**核心设计变更：**
+
+| 旧设计 | 新设计 |
+|---|---|
+| `run_streakline_online(sim, ...)` 控制仿真循环 | `Streakline.observe(ux, uy, step)` 被动接收帧 |
+| `run_streakline_offline(frames, ...)` 批处理帧列表 | 移除（无需求） |
+| `render_streakline_density(positions, ages, ...)` 15 参数 | `Streakline.render(path)` 内部维护状态 |
+| `ParticleTrailSet` 与正确实现混在同一文件 | 移到独立 `pathline.py` |
+| `render_vorticity_field` 与粒子跟踪无关 | 移到独立 `render.py` |
+| `gaussian_blur2d` 用 `np.apply_along_axis` | 改为显式 `for` 循环 + `np.convolve` |
+| `minimal_axes=True/False` 分支重复 | 合并为单路径 `_save_minimal_image` |
+
+**用法示例：**
+```python
+streak = Streakline(release_points=..., nx=nx, ny=ny, cylinders=...)
+# 在自己的仿真循环中：
+macro = sim.get_macroscopic()
+streak.observe(ux=macro["ux"], uy=macro["uy"], step=step)
+# ...同时可以做 body actions、checkpoint、其他模块...
+streak.render("output.png")
+```
+
+### Body 模块
+
+| 严重度 | 问题 | 行号 | 处理 |
+|---|---|---|---|
+| **[架构]** | **`ObjectManager` 承担 8 种以上职责。** `sync_to_gpu()` 单枪匹马编排了全部流程。 | `body/manager.py` (423 行) | **[暂不处理]** |
+| **[架构]** | **`Cylinder.get_curved_list()` ~180 行包含 7 种职责。** 内嵌函数无法单独测试。 | `body/objects.py:110-291` | **[暂不处理]** |
+| **[架构]** | **自上一轮审计以来 body 模块重构进展为零。** | 全部 | **[暂不处理]** |
+
+### Config/API 层
+
+| 严重度 | 问题 | 行号 | 处理 |
+|---|---|---|---|
+| **[已确认]** | **FP16C 被正确拒绝。** | `config.py:119-123` | ✅ |
+| **[已确认]** | **28/28 参数通过四层一致性检查。** | B4a 参数追踪表 | ✅ |
+| **[已确认]** | **OBS 布局四层完全匹配。** | B4d 对比表 | ✅ |
+| **[架构/低]** | **CONFIG.md 写"建议整除"而非"必须"。** 代码用 ceiling division。 | `configs/CONFIG.md:13`, `lbm/stepper.py:272-274` | **[保留说明]** |
+| **[已修复]** | **`LBMConfig` 默认 `omega_max=1.99` → `1.96`，与 JSON/CONFIG 一致。** | `config.py:84` | ✅ |
+
+### 跨层 Flag/常量同步
+
+| 问题 | 状态 |
+|---|---|
+| **`FLAG_*` 常量 Python ↔ CUDA 同步** | **[已确认]** 全部一致 |
+| **`LBMParams` struct 布局耦合** | **[保留说明]** |
+| **V_TAYLOR 硬编码 1** | **[保留说明]** |
+
+---
+
+## 测试覆盖率审计
+
+### Validation runner 与文档一致性
+
+| 结论 | 项目 | 说明 |
+|---|---|---|
+| **[通过]** | Sah04 runner S1-S4 | 四个锚点全部定义并运行 |
+| **[缺口]** | Sah04 runner `--u-max` | `u_max_nominal=0.1` 写死 |
+| **[缺口]** | Sah04 runner high-beta 网格 | 默认直径与文档推荐不一致 |
+| **[通过]** | Kan99b runner K1-K5 | 全部 5 个 case，K2 有 gate |
+| **[缺口]** | Kan99b runner K3-K5 | 未做抑制分类自动检测 |
+| **[通过]** | 输出机器可读 | 均输出 JSON/CSV |
+
+### 测试基础设施
+
+| 严重度 | 结论 |
+|---|---|
+| **[严重缺口]** | **零 pytest/unittest 基础设施。** |
+| **[严重缺口]** | **四个旧待验证项仅 EsoPull 已修，余三项（MRT、curved boundary、force 归一化）未覆盖。** |
+
+### 建议的最低测试覆盖面（优先级排列）
+
+1. **MRT 单元测试**（uniform flow / Poiseuille / decaying vortex）— 单元级
+2. **力系数归一化测试** — 单元级
+3. **Curved boundary 单列测试** — 单元级
+4. **EsoPull vs double-buffer 一致性**（Poiseuille 流对比）— 单元/集成级
+5. **Validation runner smoke 测试** — 验证级
+
+---
+
+## 保留说明
+
+- **`BC_MOVING` / `BC_PERIODIC` 未分发** — 已标注 TODO，暂不实现
+- **Curved/Sensor kernel 无运行时越界检查** — 设计选择（依赖 host 验证）
+- **Plain linear Bouzidi 与 TRT 不天然相容** [Gin08b] — 注释已存在
+- **Curved wall 无质量守恒修正** — 长时间高 Re 周期 curved flow 可能出现力漂移 [San18]
+- **入口 `U0` 是截面平均速度**，抛物入口峰值为 `1.5*U0`
+- **角点 flag 语义不一致**：初始化时按 inlet/outlet 分类，边界核又落回 `bounce_back_swap()`
+- **3D 刚体旋转契约是 z 轴占位实现**（`aux_kernels.cu:58-63`，`Ww = 0.0f`）
+- **Curved boundary 仍是圆形几何特化**，没有为任意离散几何留出入路径
+- **`LBMParams` struct 布局** 在 Python `struct.pack` 与 CUDA struct 之间硬耦合
+- **CONFIG.md NT 整除性措辞**：写"建议"，代码用 ceiling division
+
+---
+
+## 修改项汇总
+
+| 类别 | 计数 |
+|---|---|
+| 第一轮修补确认 | 12 项全部 [已确认] |
+| 本轮已修复 | 6 项 |
+| 保留说明 | 10 项 |
+| 暂不处理 | Body 重构 / 测试基础设施 / Validation runner 缺口 |

docs/audit/body_refactor_notes.md

@@ -0,0 +1,325 @@
+## 目标
+
+第二阶段不再以补丁式修 bug 为主，而是重建 `body` 与 `LBM` 之间的模块边界，形成可持续扩展到 Bouzidi、IBM、粒子和多相耦合的骨架。重构后的代码需要保持短文件、清晰职责、显式契约，并避免几何语义直接渗入 CUDA 核心计算链路。
+
+## 第一阶段与第二阶段的边界
+
+### 第一阶段
+
+目标是验证当前修复后的代码没有恶化，并保住最基本执行链路。
+
+- 能稳定计算流场
+- 能读取力与力矩
+- Bouzidi 路径不出现新的时序性错误
+- 初始化、flag、object overlay 不再互相污染
+- 现有接口不做大改，只修运行期 bug 和明显契约错误
+
+### 第二阶段
+
+目标是结构性重构，而不是继续在旧链路上叠加特例。
+
+- 重构 `body`
+- 重构 curved boundary 数据链
+- 为 IBM 预留统一入口
+- 为简单几何和离散几何建立同一中间表示
+- 为未来多相与多物理场保留清晰耦合面
+
+## 当前架构的核心问题
+
+### 主要问题
+
+| 模块 | 当前问题 | 后果 |
+|---|---|---|
+| `body` | 几何、flag overlay、compact list、action、obs、coupling 混在一起 | 任一功能扩展都会牵动整条链路 |
+| `objects.py` | 几何对象直接生成 Bouzidi 专用数据 | 简单几何和离散几何无法共享同一接口 |
+| `curved_boundary` | 几何解释、边界格式、移动壁修正、力提取绑在同一层 | 替换边界方法成本高 |
+| `ObjectManager` | 过度集中，已经接近万能管理器 | 可读性下降，后续只能继续膨胀 |
+| host-device contract | 分散在多个文件，缺少统一定义 | 调试困难，容易出现隐式耦合 |
+
+### 本质判断
+
+当前最大问题不是公式本身，而是缺少稳定的中间层。几何对象、边界方法、运行时打包目前没有被拆开，导致代码很难既简洁又可扩展。
+
+结合本轮修 bug 的经验，再补一条判断：很多问题不是几何本身出错，而是 donor、fallback、time layer 这类数值语义散落在不同文件中，读代码时必须来回跳转才能确认。当前项目更适合把这些语义集中写进少数注释位置，而不是继续加更多隐式保护逻辑。
+
+## 第二阶段的目标架构
+
+### 顶层模块
+
+保持两大物理模块：
+
+- `LBM`
+- `body`
+
+但它们之间不直接互相了解实现细节，而通过显式中间数据结构耦合。
+
+### 推荐分层
+
+| 层级 | 职责 | 不应负责 |
+|---|---|---|
+| `simulation` | 装配模块，定义推进顺序 | 几何处理，边界公式细节 |
+| `body` | 几何、刚体状态、预处理、力回收 | DDF 操作，LBM 核心算法 |
+| `lbm` | 格子流体、collision、streaming、boundary operator | 几何来源，刚体业务语义 |
+| `coupling` | cut-link、IBM marker、body-fluid 数据交换 | 具体几何类定义，具体 collision 实现 |
+
+## 推荐的 body 内部结构
+
+### 目标
+
+`body` 应只表达拉格朗日对象和几何处理，不直接承载 LBM 业务逻辑。
+
+补充边界：`body` 负责管理对象并产出统一 cut-link 几何记录，`lbm` 只消费 cut-link 做数值计算。对象类型区分、几何来源、以及“该记录来自圆柱还是离散几何”都不应进入 LBM kernel 视野。
+
+### 推荐文件树
+
+```text
+body
+  geometry
+    base
+    circle
+    sphere
+    polygon
+    mesh
+  state
+    rigid_body_state
+    particle_state
+  preprocess
+    flag_overlay
+    cut_links
+    ibm_markers
+    sensors
+  runtime
+    action_buffer
+    telemetry_buffer
+    registry
+  coupling
+    wall_velocity
+    force_torque
+```
+
+### 说明
+
+- `geometry` 只回答几何问题
+- `state` 只保存状态量
+- `preprocess` 负责把几何投影到欧拉网格
+- `runtime` 负责 GPU 上传与 buffer 管理
+- `coupling` 负责 body 与 fluid 的交换规则
+
+进一步要求：
+
+- `geometry` 不直接生成 Bouzidi 专用 SoA
+- `preprocess` 先产出统一 cut-link 结果，再由更薄的一层做运行时打包
+- `runtime` 不回头参与几何判断
+- 单文件保持简短，避免再出现一个文件同时含几何、打包、上传、观测解释四类职责
+
+## 推荐的 LBM 边界
+
+### LBM 应该看到什么
+
+LBM 只应看到这些输入：
+
+- `flag`
+- `cut-link records`
+- `IBM marker records`
+- `runtime params`
+- `obs buffers`
+
+LBM 不应知道：
+
+- 该 link 来自圆柱还是三角网格
+- 该对象是粒子还是障碍物
+- 该边界记录由哪种 host 几何算法产生
+
+## curved boundary 的重构方向
+
+### 核心原则
+
+不要把 curved boundary 等同于 Bouzidi。
+
+需要拆成三个维度：
+
+- 几何表示
+- 边界处理方法
+- 运动模型
+
+补充一条实现原则：donor、fallback、time layer 等关键数值语义，优先通过集中注释写清楚，不额外引入很多“自动保证语义”的复杂逻辑。当前项目更优先保持代码短、直、可读。
+
+### 建议的统一中间表示
+
+定义通用 `cut-link record`，至少包含：
+
+| 字段 | 含义 |
+|---|---|
+| `fluid_idx` | 流体格点索引 |
+| `dir` | 指向边界的格子方向 |
+| `q` | 交点沿链路的位置 |
+| `body_id` | 所属物体 |
+| `hit_point` | 壁面交点 |
+| `lever_arm` | 相对参考点的力臂 |
+| `normal` | 壁面法向 |
+| `motion_tag` | 静止、平移、旋转等 |
+| `scheme_tag` | Bouzidi、half-way、未来 TRT-compatible scheme |
+| `fallback_tag` | donor 非法时的退化方式 |
+
+补充约束：
+
+- `cut-link record` 只表达几何命中结果与最小运行时字段，不直接长成某一种 kernel 专用格式
+- donor 的来源与时间层语义不额外做复杂自动推断，靠集中注释写清楚
+- 记录字段命名要直接对应 kernel 使用含义，避免同一字段在不同文件中有不同解释
+
+### 这样做的好处
+
+- 圆形和离散几何可以输出同一种记录
+- Bouzidi 与 IBM 可以共享一部分几何预处理
+- boundary kernel 只消费记录，不关心几何来源
+- 以后替换边界方法时，不必回改对象类
+- donor、fallback、hit-point 这类契约可以在少数固定注释位置集中说明，而不是散落在对象类和 kernel 两侧
+
+## IBM 的预留方式
+
+IBM 不应和 Bouzidi 混成一条链，而应与 cut-link 平行。
+
+建议另设统一 `marker record`，用于：
+
+- 插值点位置
+- 支撑域格点
+- 权重
+- 与刚体的归属关系
+
+这样未来可以并存：
+
+- cut-link boundary
+- IBM boundary
+- 混合策略
+
+## 运行时契约的建议
+
+### 必须显式化的契约
+
+应把 host-device contract 从各文件中收口，单独维护。
+
+建议集中定义：
+
+- action buffer layout
+- telemetry buffer layout
+- cut-link record layout
+- marker record layout
+- force torque sign convention
+- body velocity contract
+
+### 运动状态契约
+
+即使短期只做 2D 圆柱旋转，也建议按最终形式设计：
+
+| 状态 | 建议字段 |
+|---|---|
+| 平移 | `vx vy vz` |
+| 角运动 | `wx wy wz` 或 2D 的 `omega` 特化视图 |
+| 参考点 | `cx cy cz` |
+| 姿态 | `theta` 或旋转表示 |
+
+不要再把 3D 路径固化为“只读一个 z 轴 `omega`”。
+
+## 第二阶段推荐顺序
+
+### 阶段 2A
+
+先把职责边界拆开，不追求新功能。
+
+- 拆 `ObjectManager`
+- 建立 `BodyRegistry`
+- 建立 clean `runtime buffer` 层
+- 建立单独的 geometry preprocess 层
+- 把“对象管理”和“cut-link 产出”分开
+
+验收标准：
+
+- `body` 不再直接知道具体 LBM kernel 调度
+- `ObjectManager` 不再承担几何、obs、state、flag 全部职责
+- `objects.py` 不再直接输出某一种 boundary method 专用打包格式
+
+### 阶段 2B
+
+重构 curved boundary 数据链。
+
+- 定义统一 `cut-link record`
+- 让圆柱先输出新 record
+- 让 Bouzidi kernel 消费新 record
+- 把现有 `CurvedLinkSoA` 从“Bouzidi 专用”改成“通用 cut-link buffer`
+- donor 与 fallback 的契约集中写入注释，不额外增加复杂语义保护代码
+
+验收标准：
+
+- kernel 不需要知道几何来源
+- host 侧可以替换 cut-link builder 而不改 kernel 接口
+- 读 builder 与 kernel 时，不需要跨很多文件才能理解 donor 和 fallback 的基本语义
+
+### 阶段 2C
+
+为 IBM 留入口。
+
+- 定义 `marker record`
+- 预留独立 preprocess 和 runtime buffer
+- 暂不实现完整 IBM 细节，只把架构位置留好
+
+验收标准：
+
+- IBM 不需要重写 body 骨架
+- IBM 与 Bouzidi 可以共存于同一 body 系统
+
+### 阶段 2D
+
+补文档与契约说明。
+
+- 当前能力
+- 预留能力
+- 不支持项
+- 2D 与 3D 差异
+- TRT 与 plain Bouzidi 的限制
+
+## 不建议做的事
+
+- 不要继续强化万能 `ObjectManager`
+- 不要让几何对象直接生成某一种边界格式专用 SoA
+- 不要把 curved boundary 和 Bouzidi 永久绑定
+- 不要在 3D 运动契约上继续堆占位特例
+- 不要在第二阶段把多相、IBM、粒子一起全做完
+- 不要为了“自动保证语义”继续增加很多隐式逻辑分支，优先用集中注释把契约写清楚
+- 不要引入必须跨很多文件来回跳转才能理解的 host-kernel 组织方式
+
+## 第二阶段完成后的理想状态
+
+### 代码层面
+
+- 文件更短
+- 模块职责更单一
+- host-device contract 更显式
+- kernel 更专注于数值操作
+- geometry 与 boundary method 解耦
+
+### 扩展层面
+
+后续可自然扩展到：
+
+- 离散几何 cut-link
+- IBM marker 链路
+- 粒子和刚体共用 body runtime
+- 多相流对 body 的额外 coupling
+
+### 维护层面
+
+后续新增功能时，优先在对应层增加文件，而不是回到单个 manager 中继续堆逻辑。
+
+## 建议的执行方式
+
+第二阶段执行前，先固定一个简短设计约束：
+
+- 单文件不过长
+- 新增模块必须有明确职责说明
+- 任何 host-device 数据结构必须有集中定义
+- 新接口先定义契约，再写 kernel
+- 能通过新增 builder 解决的问题，不要回写到 geometry 类中
+- donor、fallback、time-layer 等关键数值语义必须在少数固定位置集中注释说明
+- 优先保持文件短和职责单一，不为“保险”堆太多诊断与保护代码
+
+这个阶段的核心产出不是新功能，而是一个能长期承载新功能的骨架。

docs/audit/kernel_layer_audit.md

@@ -0,0 +1,126 @@
+# Sub-agent A: CUDA Kernel & Fix Verification
+
+## 1. 旧修正确认
+
+### 结论: [已确认] Curved Bouzidi 时序 — aux_kernels.cu:26-65 + stepper.py:71 + one_step_double.cu:20
+**理由**: `stepper.py:71` 中 `_launch_curved()` 在 `OneStep` 之前调用。`aux_kernels.cu` 的 `CurvedBoundaryKernel` 写入 `f.ddf_gpu`（当前 buffer），其后 `OneStep` 从同一 buffer 做 `stream_pull_load` — 同一时间步内完成"前写入→后拉取"。`one_step_double.cu:20` 中的 `apply_boundary_pull` 对 `is_curved(fl)` 直接 return，确保流体节点不会二次覆盖。调用顺序确认无误。
+
+### 结论: [已确认] q>=0.5 分支时间层 — curved_boundary.cuh:73-75
+**理由**: `q >= 0.5` 分支读取 `load_ddf(fi, index_f(k_f, dir_opp))`，其中 `fi` 是当前步骤的 DDF buffer。旧代码读的是前一时间步的缓冲区（`fi_in`），现在读的是当前的 `fi`，即上一时间步碰撞后的 post-collision 数据。Bouzidi 两分支算法都要求同一时间层的 post-collision 数据 [Bou01]，当前写法满足此要求。
+
+### 结论: [已确认] Moving wall 修正按 q 分支 — curved_boundary.cuh:30-43
+**理由**: `bouzidi_linear_moving_correction` 函数三路分支：
+- `fallback_class != BOUZIDI` → `2 * alpha_ci_dot_uw`（半格加移动修正）
+- `q < 0.5` → `2 * alpha_ci_dot_uw`（Bou01 公式）
+- `q >= 0.5` → `alpha_ci_dot_uw / q`（Bou01 公式）
+各项系数与文献一致，`alpha_ci_dot_uw = 3 * w_i * (c_i · u_w)`。
+
+### 结论: [已确认] Forcing 预因子统一 — collision_srt.cuh:15, collision_trt.cuh:35, collision_mrt.cuh:41/116
+**理由**: 三个文件的类内路径都使用了 `c_tau = 1.0f - 0.5f * omega` 并将 `c_tau * Fin[i]` 增量加到碰撞输出中。`collide_dispatch()`（helpers.cuh:33-37/46-50）在 `d_params.fx/fy/fz`非零时调用 `compute_guo_forcing()` 生成 `Fin`，然后传入对应碰撞函数。三者完全一致。
+
+### 结论: [已确认] TRT outlet NEQ 全分布重构 — pressure_neq.cuh:45-51(D2Q9):89-95(D3Q19)
+**理由**: `#if COLLISION_MODEL == 0 || COLLISION_MODEL == 1` 把 TRT 纳入全分布 damped NEQ 分支（与 SRT 同级），对所有 NQ 方向做 `f[i] = feq_tar[i] + beta * fneq`，其中 `fneq = f_neb[i] - feq_neb[i]`，`beta = OUTLET_SRT_NEQ_DAMP`。已不再是"少量未知方向"路径。
+
+### 结论: [已确认] add_vortex() 除以 rho — initializers.py:55-58
+**理由**: `ux_old = sum(f[i] * cx[i]) / rho_safe`，`uy_old = sum(f[i] * cy[i]) / rho_safe`。旧 bug 是动量除以 rho 这一步缺失，现在存在 `rho_safe` 分母。
+
+### 结论: [已确认] Init flag overlay 顺序 — simulation.py:139-159 + init_flow.cu:48-58 + body/manager.py:115-127
+**理由**: 顺序为 `build_channel_flags()`（干净通道）→ `build_flags()`（叠加物体）→ `upload_flags()`（上传GPU）→ `stepper.initialize()`（运行 init kernel 保持 obstacle flag）。`init_flow.cu:finalize_domain_flag` 对 `is_obstacle(fl)` 直接返回原 flag。`sync_to_gpu(rebuild_flags=False)` 在初始化时不重复构建 flags。`_rest_nonfluid()` 的二次重置已移除。
+
+---
+
+## 2. 待验证项复查
+
+### 结论: [已确认] MRT D2Q9 方向序与符号 — collision_mrt.cuh:46-54,79-92
+**理由**: 经子代 agent 验证：
+- `m[3]` = f1−f2 + f5−f6 + f7−f8 = ρ·ux（与 `macro.cuh:36` 的 `compute_rho_u()` 完全一致）
+- `m[5]` = f3−f4 + f5−f6 − f7 + f8 = ρ·uy（一致）
+- `m[7]` = f1+f2−f3−f4 = ρ·(ux²−uy²)（一致）
+- 逆变换系数 `g[0] += (dm0 − dm1 + dm2)/9` 等均符合 `M·M⁻¹ = I`
+- `meq[4] = −ρ·ux`、`meq[6] = −ρ·uy` 符号正确（正交基结构）
+- 无符号错误。MRT D2Q9 的 paired 方向排列下的矩变换自洽。
+
+### 结论: [已确认] Esopull 邻壁行处理 — one_step_esopull.cu vs one_step_double.cu
+**理由**: 逐行对比 `one_step_esopull.cu:76-151` 与 `one_step_double.cu:65-144`。Double-buffer 在 line 88-94 有显式 `is_fluid(fl) && (y==1 || y==NY-2)` 块调用 `apply_wall_bb_y_pull`，而 Esopull 没有对应分支。**但这不构成 bug** — Esopull 交替读写模式（`load_f_esopull` / `store_f_esopull`，`esopull_single_buffer.cuh:33-77`）天然避免了从 y=0 壁面节点直接拉取垃圾数据：
+- 偶数步：f[3] 读取本地 `fi[n, 4]`（前一步该节点的 -y 方向），f[4] 读取 `fi[j[3], 3]`（y=2 的 +y 方向）
+- 奇数步同理交错
+两个时间步的综合效果等价于半格 BB 的反射，无需显式修正。壁面节点 (y=0) 经 `apply_boundary_esopull` → `bounce_back_swap` 处理，不会积累垃圾。
+
+### 结论: [无法确认: 需运行验证] Force 提取与符号约定 — curved_boundary.cuh:90-97 + obs.cuh + manager.py:328-338
+**理由**: 链路追踪如下：
+1. `curved_boundary.cuh:91-97`: `fx = c_x * (f_toward + f_reflected)` 累加到 `obs[obs_force_index(body_id, 0)]`
+2. `obs.cuh:11`: `obs_force_index(id_obj, d) = OBS_FORCE0_FLOATS + id_obj * DIM + d`（起点=0）
+3. `manager.py:read_force()`: 返回 `self.obs_pinned[i0:i0 + d]`，无符号反转
+- 存储的是流体动量交换量 (fluid momentum exchange)，不是物体受力的直接值（牛顿第三定律要求 `F_body = -F_fluid`）。如果外部 Cd/Cl 归一化时把 obs 直读值当做物体受力，符号会反转。需用已知算例（如静止圆柱的阻力系数）确认当前实践是否在外部做了隐含的取反。
+
+---
+
+## 3. 架构缺陷
+
+### 结论: [未改] Curved boundary 仍假设圆形/球形 — curved_boundary.cuh + body/objects.py:110-291
+**理由**: `Cylinder.get_curved_list()` 对每个 cut link 调用 `find_circle_intersection` / `find_sphere_ray_segment`，完全依赖于圆/球几何参数（center + radius）。没有通用多边形/三角网格接口。旧审计标记的 [待重构] 未处理。
+
+### 结论: [未改] 3D 旋转为 z 轴占位 — aux_kernels.cu:58-63
+**理由**: `CurvedBoundaryKernel` 的 D3Q19 分支中 `Ww = 0.0f`，且 `Uw = -omega * ry; Vw = omega * rx` 只包含 xy 平面转动。任何具有 z 分量的刚体旋转都不会产生正确的壁面速度。
+
+### 结论: [已备注] Bouzidi-TRT 不相容注释 — curved_boundary.cuh:17-21
+**理由**: 注释明确说明 "plain linear Bouzidi interpolation ... is not a TRT-parametrized curved-wall family"，位置精准且语意清晰。
+
+---
+
+## 4. 新发现
+
+### 4a. `_no_force` 碰撞变体为死代码 — collision_srt.cuh:25, collision_trt.cuh:57, collision_mrt.cuh:95
+**问题**: `collide_srt_no_force`、`collide_trt_no_force`、`collide_mrt_no_force` 三个函数在编译后的任何内核路径中均不被调用。`collide_dispatch` (helpers.cuh:15-69) 始终通过带 `Fin` 数组的普通变体执行碰撞，当外力为零时通过 `zero_forcing(Fin)` 将 `Fin` 全清零。死代码约 40 行。
+**影响**: 无功能影响，增加维护成本。`zero_forcing` 仅用于栈初始化，并非无用调用。
+
+### 4b. BC_MOVING / BC_PERIODIC 有定义无处理 — flags.cuh + 全内核搜
+**问题**: `FLAG_BC_MOVING (0x0060)` 和 `FLAG_BC_PERIODIC (0x0050)` 在 `flags.cuh` 有定义，Python 侧 `descriptors.py` 也有对应常量。但 CUDA 内核中没有针对 `is_moving()` 或 `is_periodic()` 的分支处理：
+- `one_step_double.cu:apply_boundary_pull` 只有 `is_curved/is_inlet/is_outlet/BBS` 分支
+- 任何单元格若被标记 `BC_MOVING`（非 curved 路径）将落入 `bounce_back_swap()` 分支，得到 std half-way BB 而非移动壁面速度修正
+**影响**: 只有当 obstacle 使用 `BC_CURVED` 配合 `cl_body_id` 进入 curved boundary kernel 才能获得正确移动壁面速度。`BC_MOVING` 直接标记在通道壁面或其他固体节点上无效果。
+
+### 4c. Outlet NEQ 的 `OUTLET_MODE` 嵌套逻辑可读性隐患 — pressure_neq.cuh:25-62
+**问题**: 
+```
+#if OUTLET_MODE == 1        → 纯拷贝未知方向
+#else                       → 普通路径
+  #if COLLISION_MODEL==0||1 → 全分布 NEQ（含 TRT）
+  #elif OUTLET_MODE == 2    → 混合模式
+  #else                     → 少量未知方向重构（默认 OUTLET_MODE=0 路径）
+  #endif
+#endif
+```
+当 `OUTLET_MODE=0`、`COLLISION_MODEL=2`（MRT）时，代码进入最内层 `#else` 分支（少量未知方向重构），而非全分布 NEQ。这与注释宣称的"SRT 和 TRT 使用全分布 NEQ"一致（MRT 未承诺），但 MRT outlet 路径与 SRT/TRT 行为不同，可能导致 MRT 结果系统性偏差。
+**影响**: MRT outlet 行为与 SRT/TRT 不一致，应至少加注释说明此差异，或行为对齐。
+
+### 4d. `collide_inlet_ghost` 对 `y=0/NY-1` 的过滤 — one_step_double.cu:102-103, one_step_esopull.cu:106-107
+**问题**: `collide_inlet_ghost = is_inlet(fl) && interior_y && inlet_scheme_uses_post_collision_ghost()`，其中 `interior_y = (y>0 && y<NY-1)`。`y=0` 和 `y=NY-1` 的 inlet 节点不会触发 ghost 碰撞。但 `inlet_scheme_uses_post_collision_ghost()` 只在 `INLET_SCHEME==0`（Zou-He）时返回 true，而 x=0 inlet 节点原本处于 y=0 或 y=NY-1 时，已被 `build_channel_flags` 覆盖为 `SOLID|BC_WALL`（角点优先），因此这些节点本身不是 inlet，过滤是安全的。
+**影响**: 当前无实际影响（角点 wall 覆盖 inlet），但代码隐含的逻辑依赖比较脆弱。如果未来修改建标记序，可能在 y=0/ymax inlet 节点产生未碰撞 ghost 状态。
+
+### 4e. 传感器归一化在 manager 层正确实现 — aux_kernels.cu:67-99 + manager.py:348-365
+**结论**: `SensorKernel` 做逐格点求和，`ObjectManager.read_sensor(normalize=True)` 除以 `sensor_cell_counts[body_id]`。已实现，正确。
+
+### 4f. `inlet_target_u` 对 y=0 和 y=NY-1 的保护 — inlet/common.cuh:34
+**影响**: `y_clamped = fminf(NY-2, fmaxf(1.0, y))` 使 y=0 节点获得 y=1 的 inlet 速度。对正确定义的 `SOLID|BC_WALL` 角点无实际影响，属于保护性逻辑。
+
+### 4g. `compute_omega_minus` 在 ω⁺ = 2.0 时分母为零 — collision_trt.cuh:24-26
+**问题**: `1.0f / omega_plus - 0.5f` 当 `omega_plus = 2.0` 时为 `0.5 - 0.5 = 0`，导致除零。实际路径中 `omega_col` 在 `collide_dispatch` 中被钳位（helpers.cuh:56），假设 `OMEGA_COLLISION_MAX < 2.0` 则为安全。
+**影响**: 低风险（依赖外部宏正确配置）。
+
+### 4h. `west_velocity_rho_closure` 在 u_target ≥ 1.0 时除零 — inlet/common.cuh:47-48
+**问题**: `rho = sum(...) / (1.0f - ux_target)`，当 `ux_target >= 1.0` 时除数为 0 或负数。物理上入口马赫数小于 1 可避免，但无运行时保护。
+**影响**: 低风险（物理约束），但崩溃行为不如显式断言清晰。
+
+---
+
+## 总结
+
+| 类别 | 数量 | 关键项 |
+|------|------|--------|
+| 旧修正确认 | 7/7 | 全部确认正确 |
+| 待验复查 | 3 | 2 确认, 1 需运行时验证 |
+| 架构缺陷 | 3 | 圆形硬编码、3D占位、TRT注释齐全 |
+| 新发现 | 8 | 4a死代码、4b未处理BC、4c MRT outlet差异、4d脆弱逻辑依赖、4e已实现、4f无害保护、4g除零边界、4h无保护输入 |
+
+**最高优先级**: 4b (`BC_MOVING` 无处理路径)、4c (MRT outlet 路径与 SRT/TRT 不一致)、force 符号约定需运行验证。

docs/performance_analysis.md

@@ -0,0 +1,266 @@
+# CelerisLab Performance Analysis Report
+
+Date: 2026-05-31
+GPU: Tesla V100-SXM2-16GB (CUDA 12.4)
+Test grid: 384x192 D2Q9, MRT, double_buffer, FP32
+Tools: cProfile (Python), Nsight Systems (GPU timeline), Nsight Compute (kernel metrics)
+
+## 1. Performance Tree (384x192, MRT, 1 cylinder + 1 sensor)
+
+### Per-step breakdown (batch=100, 100 steps per launch group)
+
+```
+┌──────────────────────────────────────────────────────────────┐
+│                    One LBM Step (~100 steps, ~7.8ms total)   │
+├──────────────────────────────────────────────────────────────┤
+│                                                              │
+│  Python Layer (CPU)                   ~40 μs/step (0.5%)    │
+│  ├── _launch_curved()                 ~6 μs                  │
+│  ├── step_fn() (cuLaunchKernel)       ~22 μs                 │
+│  ├── _launch_sensor()                 ~5 μs                  │
+│  └── loop overhead (Python for _ in range) ~5 μs             │
+│                                                              │
+│  GPU Kernel Layer (V100)              ~13 μs/step (99.5%)   │
+│  ├── CurvedBoundaryKernel             ~4.0 μs  (30%)         │
+│  │   grid=(2,1,1) block=256  regs=31  288 threads total     │
+│  │   → latency-bound (too few threads to fill SM)            │
+│  ├── OneStep                           ~5.9 μs  (44%)        │
+│  │   grid=(2,192,1) block=256 regs=39  73,728 threads       │
+│  │   → mixed compute+memory (53% DRAM bw, 22% SM util)       │
+│  └── SensorKernel                      ~3.5 μs  (26%)        │
+│      grid=(2,1,1) block=256 regs=21   305 threads total      │
+│      → latency-bound                                          │
+│                                                              │
+│  Memory Transfers (negligible, <1 μs/step)                   │
+│  ├── set_body():   72 bytes H2D (action upload)              │
+│  └── read_force(): 60 bytes D2H (obs download, async)        │
+└──────────────────────────────────────────────────────────────┘
+```
+
+### Wall-clock contributions (2000 steps, batch=100)
+
+| Layer | Total time | Per step |
+|-------|-----------|----------|
+| Python (cProfile) | 0.901s (warmup+measured, 70 batches) | ~12.9ms/100-step batch |
+| GPU kernel (Nsight) | 0.071s (5300 kernel invocations) | 13.4 μs |
+| Batch NVTX range (nsys) | 7.8ms for batch=100 | 78 μs/step (incl. Python) |
+
+**Key insight:** The per-step GPU kernel time is only ~13 μs. The nsys batch timing (78 μs/step) includes Python overhead, kernel launch overhead, and stream synchronization.
+
+## 2. Python Layer Details (from cProfile, T4)
+
+| Function | Cumulative time | Calls | Per call | % of run time |
+|----------|----------------|-------|----------|---------------|
+| Simulation.__init__ | 1.125s | 1 | 1.125s | 62% (one-time) |
+| LBMStepper.step | 0.901s | 70 | 12.9ms | 50% of step time |
+| pycuda function_call | 0.769s | 21001 | 36.6 μs | 42% of step time |
+| _launch_sensor (inside step) | 0.409s | 7000 | 58 μs | — |
+| _launch_curved (inside step) | 0.311s | 7000 | 44 μs | — |
+| pycuda _build_arg_buf | 0.357s | 21001 | 17 μs | 20% of step time |
+| manager.sync_to_gpu | 0.017s | 1 | — | one-time (init) |
+| set_body | <0.001s | 20 | <50 μs | negligible |
+| read_force | <0.001s | 20 | <50 μs | negligible |
+
+**Conclusion:** The Python layer is extremely thin. The dominant cost in `stepper.step()` is `pycuda.driver.function_call` (~37 μs per kernel launch) which is the `cuLaunchKernel` overhead from the Python-CUDA bridge. This is an unavoidable cost of using pycuda.
+
+## 3. GPU Kernel Details (from Nsight Systems)
+
+### OneStep (MRT, 384x192, 5300 invocations)
+
+| Metric | Value |
+|--------|-------|
+| Grid | (2, 192, 1) |
+| Block | (256, 1, 1) |
+| Registers per thread | 39 |
+| Min duration | 5.57 μs |
+| **p50 duration** | **5.86 μs** |
+| p90 duration | 6.11 μs |
+| Max duration | 7.65 μs |
+| Total (5300 steps) | 31.2 ms |
+
+### CurvedBoundaryKernel (1 cylinder at center, 5300 invocations)
+
+| Metric | Value |
+|--------|-------|
+| Grid | (2, 1, 1) |
+| Block | (256, 1, 1) |
+| Registers per thread | 31 |
+| Total threads | 288 active out of 512 launched (56% utilization) |
+| **p50 duration** | **4.00 μs** |
+| Total (5300 steps) | 21.2 ms |
+
+### SensorKernel (1 sensor, 5300 invocations)
+
+| Metric | Value |
+|--------|-------|
+| Grid | (2, 1, 1) |
+| Block | (256, 1, 1) |
+| Registers per thread | 21 |
+| **p50 duration** | **3.55 μs** |
+| Total (5300 steps) | 18.9 ms |
+
+### Batch Scaling (from Nsight NVTX ranges)
+
+| Batch size | Total time (100 steps) | Effective per step |
+|------------|----------------------|-------------------|
+| 1 step/launch | 8.08 ms | 80.8 μs |
+| 10 steps/launch | 7.86 ms | 78.6 μs |
+| 100 steps/launch | 7.78 ms | **77.8 μs** |
+
+**Observation:** Batch size has only a 4% effect on total time between batch=1 and batch=100. This is because the GPU kernel itself is only ~13μs of the total ~78μs. The remaining ~65μs is the pycuda launch overhead + stream sync.
+
+## 4. SRT vs MRT Comparison (from cProfile + ncu)
+
+| Metric | SRT (T1) | MRT (T2) | Ratio |
+|--------|----------|----------|-------|
+| stepper.step total (70×100 steps) | 0.190s | 0.191s | **1.01x** |
+| pycuda function_call count | 7001 | 7001 | 1.0x |
+| pycuda function_call time | 0.146s | 0.147s | 1.01x |
+| ncu Duration | 10.05 μs | 10.53 μs | 1.05x |
+| Registers per thread | — | 39 | — |
+| SM Throughput | 18.88% | 22.23% | 1.18x |
+| Memory Throughput | 52.73% | 52.92% | 1.00x |
+| FMA pipe utilization | 27.3% | 32.3% | 1.18x |
+| Executed IPC (active) | 1.18 | 1.24 | 1.05x |
+| L1/TEX hit rate | 6.82% | 6.82% | 1.0x |
+| L2 hit rate | 54.03% | 54.06% | 1.0x |
+| Simulation.__init__ | 1.113s | 1.298s | 1.17x |
+
+**Key finding: MRT is only 5% slower than SRT at this grid size.** Both are dominated by memory access (53% DRAM bandwidth utilization), not compute. MRT's extra FMA operations increase SM throughput from 18.9% to 22.2% but don't elongate wall time because the kernel is limited by the L1TEX scoreboard stalls (49.9% of warp cycles). The top stall reason is **waiting for L1TEX data** (scoreboard dependency), not waiting for compute pipelines.
+
+The grid (384 blocks) is too small to fill the GPU — only 0.8 full waves across 80 SMs. At production scales (3000×300), the SM throughput for MRT would approach 50-60% and the 5% gap over SRT would widen to ~15-20%.
+
+## 5. CUDA API Breakdown (from Nsight Systems)
+
+| CUDA API | Total time (over whole run) | Per call | Notes |
+|----------|---------------------------|----------|-------|
+| cuCtxCreate_v2 | 123 ms | 123 ms | One-time, CUDA context creation |
+| cuModuleLoad | 1.46 ms + 1.02 ms | 1.2 ms | Kernel compilation (one-time per compile) |
+| cuLaunchKernel (warmup + measured) | ~47 μs per call | 47 μs | pycuda overhead from cProfile ~37 μs + CUDA driver overhead |
+| cuMemcpyHtoD (init: 14KB) | 21-45 μs | varies | Config data upload |
+| cuMemcpyHtoD (action: 72 bytes) | ~1.15 μs | 1.15 μs | Body rotation upload |
+| cuMemcpyDtoH (DDF: 2.6MB) | 1.02 ms + 2.26 ms | 1-2 ms | Full DDF download (get_macroscopic) — one-time in profiled run |
+
+**DTOH overhead:** The 72-byte action H2D takes ~1.15 μs, and the 60-byte obs D2H would take ~1 μs. These are truly negligible at any scale.
+
+## 6. Bottleneck Identification
+
+### Current bottleneck (384x192): pycuda kernel launch overhead
+
+At this grid size, GPU compute time (13 μs/step) is dominated by pycuda's cuLaunchKernel overhead (~37 μs per kernel × 3 kernels = ~111 μs + Python loop ~7 μs = ~~78 μs/step from nsys, ~129 μs/step from cProfile). The **GPU is idle about 80% of the time** waiting for the next kernel launch.
+
+### What would change at larger grids
+
+| Grid size | GPU time/step (est.) | pycuda overhead | GPU idle |
+|-----------|--------------------|----------------|----------|
+| 384x192 | ~13 μs | ~37 μs × 3 = 111 μs | **80%** |
+| 600x600 | ~33 μs | ~37 μs × 3 = 111 μs | 55% |
+| 1000x500 | ~47 μs | 111 μs | 35% |
+| 3000x300 (exp_ctrl) | ~530 μs | 111 μs | **17%** |
+| 5000x500 | ~1.6 ms | 111 μs | 7% |
+
+**At your production scale (3000x300), GPU utilization would be ~83%.** The pycuda overhead becomes acceptable.
+
+### Where optimization would matter
+
+1. **Reducing pycuda kernel launch overhead** — This is the #1 bottleneck for small grids. The 37 μs/launch comes from pycuda's argument buffer packing (`_build_arg_buf: 17 μs`) + CUDA driver overhead. Switching to `pycuda.driver.LaunchKernel` with pre-packed arguments could reduce this, but it would require significant changes to pycuda.
+
+2. **Combining curved boundary into OneStep** — Would eliminate 2 of 3 kernel launches per step (save ~74 μs). For your production grid (3000x300), this would save about 12% of total time.
+
+3. **Using larger batch sizes** — Already confirmed: batch=100 vs batch=1 saves about 4% total time on 384x192. The savings increase with grid size.
+
+### Where optimization would NOT matter
+
+1. **`set_body()` / `read_force()` overhead** — Already <1 μs. Not worth any optimization effort.
+2. **GPU memory bandwidth** — The DDF read/write (5.3 MB/step) is only 2.5% of V100 HBM2 bandwidth. Not the bottleneck.
+3. **Python loop overhead** — 5 μs/step loop from `for _ in range(n)` is negligible.
+4. **Registry / Body API** — dict lookups and property access are sub-microsecond. Not relevant.
+
+## 7. Performance Model
+
+```
+Per-step time (384x192 MRT) = 
+    Python_overhead(~40 μs)    ← fixed for any grid size
+    + CurvedBoundaryKernel(~4 μs per link-group)
+    + OneStep(~5.9 μs)         
+    + SensorKernel(~3.5 μs per sensor-group)
+
+Per-step time (general) ≈ 
+    pycuda_launch_kernels_3 × 37 μs
+    + grid_cells × (pull_load + collide + pull_store) / V100_FLOPs
+    
+Scaling:
+    - 384x192:  ~78 μs/step (GPU 17%, pycuda 83%)
+    - 3000x300: ~640 μs/step (GPU 83%, pycuda 17%)
+    - 6000x600: ~2.5 ms/step (GPU 96%, pycuda 4%)
+```
+
+## 8. Nsight Compute Deep-Dive Metrics
+
+### OneStep (MRT) — Nsight Compute full metrics
+
+| Metric | Value | Interpretation |
+|--------|-------|----------------|
+| Duration (ncu) | 10.53 μs | (ncu includes replay passes; nsys p50=5.86 μs is the real runtime) |
+| Registers per thread | 39 | Well below V100 limit of 64. **No register spill to local memory.** |
+| Grid | (2, 192, 1) × 256 = 384 blocks | Only 0.8 full waves — grid too small to fill 80 SMs |
+| Memory Throughput | 52.9% of HBM2 peak | Moderate bandwidth usage |
+| SM Throughput | 22.2% of peak | Low — limited by L1TEX scoreboard stalls |
+| FMA pipe active | 32.3% of active cycles | Most-used pipeline, but not saturated |
+| L1/TEX hit rate | 6.82% | Low — DDF reads are scattered across the grid |
+| L2 hit rate | 54.0% | Moderate — half of cache-line requests hit L2 |
+| Top stall reason | L1TEX data wait (49.9%) | Scoreboard dependency on global memory reads |
+| IPC (active) | 1.24 inst/cycle | Far below V100's theoretical 4 IPC |
+| Active warps/scheduler | 6.09 | Out of 16 max — low occupancy |
+| Eligible warps/scheduler | 0.97 | Only 1 warp per cycle is ready to issue |
+
+**Why the kernel does not saturate the GPU:**
+- The grid (384 blocks = 73,728 threads) is too small. V100 has 80 SMs with 64 warps each = 5120 warps capacity. This kernel only uses 384 blocks × 2 warps/block = 768 warps (15% occupancy).
+- Scoreboard stalls dominate (49.9%). Each warp waits on L1TEX for nearly half its cycles. This is the DDF pull phase: reading from neighbor cells at scattered memory addresses.
+- The memory access pattern is suboptimal: only 26.9 of 32 bytes per sector are utilized for global loads that miss L2.
+
+**Optimization potential:** The ncu reports "Est. Speedup: 6.45%" for memory access pattern improvements and "47.27%" for reducing L1TEX stalls. The 47% figure is misleading — it comes from the low grid occupancy, not from actual compute inefficiency. At production grid sizes (3000×300), active warps per scheduler would increase to ~12+ and the stall ratio would drop significantly.
+
+### CurvedBoundaryKernel — Nsight Compute full metrics
+
+| Metric | Value | Interpretation |
+|--------|-------|----------------|
+| Duration (ncu) | 5.60 μs | Tiny kernel, completely latency-bound |
+| Grid | (2, 1, 1) × 256 = 512 threads | Only 2 blocks — **0.0 full waves** |
+| Registers per thread | (from nsys) 31 | Low — no spill |
+| Memory Throughput | 0.63% of HBM2 | Essentially zero bandwidth utilization |
+| SM Throughput | 0.20% of peak | The GPU is doing nothing for this kernel |
+| Active warps/scheduler | 1.87 | Out of 16 max — extremely low occupancy |
+| Eligible warps/scheduler | 0.06 | Nearly 0 — almost never ready to issue |
+| No Eligible | 94.45% | **94% of cycles have no warp ready to issue** |
+| Est. Local Speedup | 97.57% | Nsight says there's 97% headroom |
+
+**This is a latency-bound kernel by design.** Only 288 threads for the curved links. The 5.6 μs runtime is dominated by launch overhead (context switching, instruction cache misses). There is no meaningful optimization for this kernel because it cannot use more threads — there are only 288 cut links. 
+
+**Impact:** At 4 μs per step, this kernel adds negligible runtime. Even with 100 cut links, the total curved time would still be <10 μs.
+
+### SensorKernel — Nsight Compute full metrics
+
+| Metric | Value | Interpretation |
+|--------|-------|----------------|
+| Duration (ncu) | 4.54 μs | Even smaller than curved |
+| Memory Throughput | 0.63% of HBM2 | Near-zero |
+| SM Throughput | 0.14% of peak | GPU doing nothing |
+| Active warps/scheduler | 1.84 | Extremely low |
+| Eligible warps/scheduler | 0.04 | Almost never ready |
+| No Eligible | 96.83% | 97% idle cycles |
+| L1/TEX hit rate | 48.53% | Much better than OneStep (localized access pattern) |
+
+Same conclusion: negligible overhead, no meaningful optimization possible.
+
+## 9. Recommendations
+
+1. **For production runs (3000x300)** — The architecture is well-optimized. Use `--batch 10` to amortize pycuda overhead. GPU utilization is ~83%.
+
+2. **For small-grid sweeps (384x192)** — The pycuda launch overhead dominates. If you need to run many small-grid parameter sweeps, consider:
+   - Grouping multiple independent simulations into a single process
+   - Or accepting that small-grid runs are launch-overhead limited
+
+3. **For kernel tuning** — At production scale (3000x300), the main kernel (OneStep MRT) uses only 39 registers (no spill), achieves 53% DRAM throughput and 22% SM throughput on 384x192. At 3000x300, SM throughput would approach 50-60% and GPU utilization would be ~83%. No further kernel optimization is needed for current use cases.
+
+4. **Three-kernel architecture is not a problem** — Despite being split into 3 kernels (CurvedBoundary + OneStep + Sensor), the total GPU kernel time is only ~13.4 μs/step. The curved and sensor kernels are latency-bound with negligible impact. Merging them into OneStep would save ~7 μs/step but add branch divergence in the OneStep grid kernel.

docs/validation_specs/Kan99b_validation.md

@@ -0,0 +1,247 @@
+# Rotating cylinder validation against [Kan99b]
+
+## Goal
+
+This validation should stay small, direct, and defensible.
+
+The main design rules are:
+
+- use the paper's direct numeric anchor at `Re = 100, alpha = 1.0` as the main hard benchmark
+- use a low-rotation case to test the lift trend
+- use suppression cases to test flow classification, not exact threshold fitting
+- do not treat values read from figures near `alpha_L` as tight amplitude targets
+
+This keeps the matrix representative without overfitting to sensitive threshold points.
+
+## Strong numeric anchors from [Kan99b]
+
+The strongest exact benchmark in the paper is the convergence case at `Re = 100, alpha = 1.0`.
+
+| Quantity | Reference value |
+|---|---:|
+| `St` | 0.1655 |
+| `mean C_L` | -2.4881 |
+| `mean C_D` | 1.1040 |
+| `C'_L` | 0.3631 |
+| `C'_D` | 0.0993 |
+
+For low rotation at `Re = 100`, the paper also gives the mean lift trend
+
+\[
+\overline{C_L} \approx -2.48\alpha
+\]
+
+which is a good secondary benchmark for small `alpha`.
+
+The suppression thresholds are given as trends:
+
+| Reynolds number | Expected `alpha_L` |
+|---|---:|
+| 60 | about 1.4 |
+| 100 | about 1.8 |
+| 160 | about 1.9 |
+
+These threshold values should be used as regime guides, not as tight one-point numeric targets. In the suppression curve from [Kan99b] shown above, the boundary is exactly the kind of place where a small solver difference can change the observed state.
+
+## Fixed solver setup
+
+| Item | Setting |
+|---|---|
+| Dimension | 2D |
+| Lattice | D2Q9 |
+| Streaming | double buffer |
+| Curved boundary | current Bouzidi moving wall implementation |
+| Inlet profile | uniform |
+| Top and bottom boundaries | free slip |
+| Outlet | neq extrapolation |
+| LES | off |
+| Precision | FP32 |
+| Cylinder diameter | `D = 30` lattice units |
+| Cylinder radius | `R = 15` lattice units |
+| Rotation input | update body omega only |
+
+The baseline domain remains the current medium far field unless a later boundary sensitivity check shows otherwise.
+
+## Inlet recommendation by collision model
+
+Kan99b is an open-flow validation, not a confined-channel benchmark.
+
+| Collision | Recommended inlet | Secondary choice | Avoid as primary |
+|---|---|---|---|
+| SRT | `equilibrium` | `regularized` | `zou_he_local` |
+| TRT | `regularized` | `equilibrium` | `zou_he_local` until the anchor is stable |
+| MRT | `regularized` or `zou_he_local` | `equilibrium` | `channel_stabilized` |
+
+Keep one inlet family per collision model across the primary matrix.
+
+## Lattice-unit mapping
+
+Use
+
+\[
+U_\infty = 0.03
+\]
+
+With `D = 30`,
+
+\[
+\nu = \frac{U_\infty D}{Re} = \frac{0.9}{Re}
+\]
+
+| `Re` | `nu` | SRT equivalent `omega` |
+|---|---:|---:|
+| 60 | 0.015000 | 1.83486 |
+| 100 | 0.009000 | 1.89753 |
+| 160 | 0.005625 | 1.93470 |
+
+The body rotation rate is
+
+\[
+\omega_{body} = \frac{2 \alpha U_\infty}{D} = 0.002\alpha
+\]
+
+| `alpha` | body omega |
+|---|---:|
+| 0.5 | 0.0010 |
+| 1.0 | 0.0020 |
+| 1.6 | 0.0032 |
+| 2.0 | 0.0040 |
+
+## Primary matrix
+
+This is the recommended main validation set.
+
+| Case | `Re` | `alpha` | Role |
+|---|---:|---:|---|
+| K1 | 100 | 0.5 | low-rotation lift trend check |
+| K2 | 100 | 1.0 | strongest hard anchor |
+| K3 | 60 | 1.6 | low-Re suppression classification |
+| K4 | 100 | 2.0 | mid-Re suppression classification |
+| K5 | 160 | 2.0 | high-Re suppression classification |
+
+Optional baseline if needed for debugging or plots:
+
+| Case | `Re` | `alpha` | Status |
+|---|---:|---:|---|
+| K0 | 100 | 0.0 | optional |
+
+This matrix covers:
+
+- one periodic low-rotation trend point
+- one exact hard anchor with full force data
+- suppression behavior at low, medium, and high Reynolds number
+
+## How to judge each case
+
+### K1
+
+Use K1 to check the low-rotation lift law.
+
+Target:
+
+\[
+\overline{C_L} \approx -2.48 \times 0.5 \approx -1.24
+\]
+
+This is a trend check, not a strict fluctuation-amplitude benchmark.
+
+### K2
+
+Use K2 as the hard benchmark case.
+
+Preferred agreement band:
+
+| Quantity | Preferred band |
+|---|---:|
+| `St` | within 3 percent |
+| `mean C_L` | within 4 percent |
+| `mean C_D` | within 5 percent |
+| `C'_L` | within 8 percent |
+| `C'_D` | within 10 percent |
+
+### K3 to K5
+
+Use K3 to K5 as suppression classification cases.
+
+Primary success signature:
+
+- `C'_L` collapses toward zero in the final window
+- no sustained alternating wake remains
+- flow classification agrees with the expected suppressed regime
+
+These are not exact threshold-fitting cases. Do not over-interpret a small residual fluctuation if the wake is otherwise clearly in the suppressed class.
+
+## Optional threshold bracket check
+
+If later you want a more explicit threshold study, use pairs around `alpha_L` rather than a single point on the boundary.
+
+Recommended pairs:
+
+| `Re` | Lower point | Upper point |
+|---|---:|---:|
+| 60 | 1.3 | 1.5 |
+| 100 | 1.7 | 1.9 |
+| 160 | 1.8 | 2.0 |
+
+These should still be treated as regime-location checks, not hard force targets.
+
+## Run policy
+
+| Case type | Total steps | Warmup | Statistics |
+|---|---:|---:|---:|
+| K1 and K2 | 180000 to 220000 | first 40 percent | last 60 percent |
+| K3 to K5 | 220000 to 280000 | first 50 percent | last 50 percent |
+
+The final statistics window should contain at least 20 shedding periods whenever the case remains periodic.
+
+## TRT re-entry rule
+
+Bring TRT back in this order:
+
+1. K2 only
+2. if K2 is stable and credible, run K1
+3. only then run K3 to K5
+
+This prevents TRT from expanding the matrix before the hard anchor is trustworthy.
+
+## Deliverables
+
+For each collision model, deliver:
+
+- one table of run settings including collision, inlet scheme, wall type, `Re`, `alpha`, `nu`, and body omega
+- one CSV per run with force history
+- selected field images for wake classification
+- one summary table with `mean C_D`, `mean C_L`, `C'_D`, `C'_L`, and `St`
+- one short note stating whether suppression behavior matches [Kan99b]
+
+## Recommended primary settings summary
+
+| Collision | Wall | Inlet | Status |
+|---|---|---|---|
+| SRT | free slip | `equilibrium` | primary |
+| TRT | free slip | `regularized` | primary if K2 is stable |
+| MRT | free slip | `regularized` or `zou_he_local` | primary |
+
+## MRT-only runner mapping
+
+The current executable entrypoint is `tests/run_kan99b_rotating_cylinder.py`, and this round uses MRT-only scheduling:
+
+- primary matrix is `K1-K5` with `MRT + regularized` inlet
+- one extra control run is added at K2 with `MRT + zou_he_local`
+- all runs keep `uniform` inlet profile, `free_slip` y-wall, `neq_extrap` outlet
+- output rows include `case_id`, `variant`, `collision`, `inlet_scheme`, `grid`, `steps`, `burn_in`, `St`, `St_error_pct` (for K2), and force metrics
+- K2 gate uses this document's per-metric tolerances for `St`, `mean C_L`, `mean C_D`, `C'_L`, `C'_D`
+
+Example commands:
+
+```bash
+conda run -n pycuda_3_10 python tests/run_kan99b_rotating_cylinder.py \
+  --json-out tests/output/kan99b_validation/summary_runs.json
+
+conda run -n pycuda_3_10 python tests/run_kan99b_rotating_cylinder.py \
+  --case K2 --save-vorticity
+```
+
+## Reference
+
+[Kan99b] S. Kang, H. Choi, and S. Lee, “Laminar flow past a rotating circular cylinder,” 1999.

docs/validation_specs/Sah04_validation.md

@@ -0,0 +1,243 @@
+# Sah04 confined-cylinder validation against [Sah04]
+
+## Goal
+
+This validation should use a small set of direct periodic anchors from [Sah04], plus a small secondary block for near-onset flow-state checks.
+
+The main design rules are:
+
+- do not use interpolated values from figures as hard targets
+- do not use points on or extremely near critical curves as primary pass fail anchors
+- report realized blockage and realized Reynolds number, not just nominal inputs
+- use finer grids for high blockage so the narrow wall gaps are not under-resolved
+
+This keeps the primary matrix small but still representative across moderate and high confinement.
+
+## What counts as a hard benchmark from [Sah04]
+
+The strongest periodic-flow anchors are the direct DNS values stated in the paper for developed unsteady states.
+
+| Case | Blockage `beta` | Reynolds number | `St` target | Why it is a hard anchor |
+|---|---:|---:|---:|---|
+| S1 | 0.3 | 100 | 0.2115 | direct periodic DNS value |
+| S2 | 0.5 | 200 | 0.3513 | direct periodic DNS value |
+| S3 | 0.8 | 160 | `T ≈ 1.806`, so `St ≈ 0.5537` | direct case given in the paper |
+| S4 | 0.9 | 200 | 0.5314 | direct periodic DNS value |
+
+These should be the primary Sah04 validation anchors.
+
+## What should not be a hard target
+
+The following are useful for qualitative or secondary checks, but not for the main validation gate:
+
+- critical onset values from Table IV
+- any `St` or force value obtained by reading or interpolating a plot
+- cases chosen only to complete a rectangular parameter grid
+- points too close to the codimension two region or nearby neutral stability boundaries
+
+This matters because the paper's stability map has several sensitive regions, especially at high blockage and near symmetry breaking. In the stability figure from [Sah04] shown above, those boundaries are exactly where a small setup difference can change the observed state.
+
+## Geometry and blockage mapping
+
+Keep the confined-channel layout and no-slip walls.
+
+The validation table must always report both nominal and realized blockage.
+
+With `D = 30` lattice units, the recommended realizations are:
+
+| Case | `beta_nominal` | Suggested `H` | `beta_real` | Notes |
+|---|---:|---:|---:|---|
+| S1 | 0.3 | 100 | 0.3000 | exact |
+| S2 | 0.5 | 60 | 0.5000 | exact |
+| S3 | 0.8 | 38 or 37 | 0.7895 or 0.8108 | pick one and report it explicitly |
+| S4 | 0.9 | 33 | 0.9091 | use this, not `H = 35` |
+
+Do not silently rename `beta_real` as the paper blockage.
+
+## Grid density policy
+
+High blockage cases need more wall-normal resolution than the base grid.
+
+Minimum rule:
+
+- for `beta < 0.8`, the baseline grid with `D = 30` is acceptable for the first pass
+- for `beta >= 0.8`, increase grid density by at least 2 times in each spatial direction before treating the result as validation quality
+
+A practical way to do this is to keep geometry similarity while doubling the characteristic resolution:
+
+- base cases: `D = 30`
+- high blockage validation cases: at least `D = 60`
+
+Recommended high blockage realizations on the refined grid:
+
+| Case | `beta_nominal` | Suggested refined `D` | Suggested refined `H` | `beta_real` |
+|---|---:|---:|---:|---:|
+| S3 | 0.8 | 60 | 75 | 0.8000 |
+| S4 | 0.9 | 60 | 67 | 0.8955 |
+
+The point of this refinement is not only bulk accuracy. It is also to resolve the narrow cylinder wall gaps and reduce the risk that blockage effects are dominated by lattice geometry error.
+
+## Primary matrix
+
+This is the main Sah04 validation set.
+
+| Case | `beta_nominal` | Primary target | Role |
+|---|---:|---|---|
+| S1 | 0.3 | `Re = 100`, `St = 0.2115` | moderate blockage periodic anchor |
+| S2 | 0.5 | `Re = 200`, `St = 0.3513` | medium blockage periodic anchor |
+| S3 | 0.8 | `Re = 160`, `St ≈ 0.5537` | high blockage periodic anchor |
+| S4 | 0.9 | `Re = 200`, `St = 0.5314` | very high blockage periodic anchor |
+
+This matrix is smaller than the older grid, but it covers:
+
+- moderate confinement
+- stronger confinement
+- high blockage periodic shedding
+- very high blockage periodic shedding
+
+## Secondary onset block
+
+These cases are recommended as flow-state checks, not hard `St` benchmarks.
+
+| Case | `beta_nominal` | Suggested `Re` | Why it is useful | How to judge it |
+|---|---:|---:|---|---|
+| SO1 | 0.5 | about 130 | safely above first onset without sitting on the boundary | confirm sustained periodic state |
+| SO2 | 0.7 | about 120 | tests near-onset behavior in a more sensitive blockage range | confirm sustained periodic state |
+
+These points exist to answer a different question from the primary matrix:
+
+- does the solver enter and maintain the right flow regime once slightly above onset
+
+For SO1 and SO2, judge by:
+
+- persistent nonzero `C'_L`
+- a clean dominant spectral peak
+- repeatable periodic wake structure
+
+Do not fail these runs because the measured `St` differs slightly from a value read from a nearby figure.
+
+## Inlet and wall policy
+
+Sah04 is a confined-channel benchmark, so inlet consistency matters more than inlet variety.
+
+| Collision | Wall | Inlet | Status |
+|---|---|---|---|
+| SRT | no slip | `channel_stabilized` | primary |
+| TRT | no slip | `channel_stabilized` | primary |
+| MRT | no slip | `channel_stabilized` | primary |
+
+Keep the inlet family fixed across collision models in the primary matrix.
+
+Secondary inlet comparison, only after the primary set is working:
+
+| Collision | Optional inlet | Status |
+|---|---|---|
+| MRT | `regularized` or `zou_he_local` | exploratory |
+| SRT or TRT | `equilibrium` or `regularized` | exploratory |
+
+## Realized Reynolds number check
+
+This is mandatory for Sah04.
+
+For each run, record:
+
+- nominal inlet definition
+- developed downstream velocity profile
+- measured `U_max,real`
+- measured bulk velocity if available
+- `beta_nominal`
+- `beta_real`
+- `Re_nominal`
+- `Re_real`
+
+Use the paper-consistent label
+
+\[
+Re_{real} = \frac{U_{max,real} D}{\nu}
+\]
+
+for the final comparison table.
+
+If `Re_real` drifts materially from the intended target, treat that as a setup problem before treating it as a Strouhal miss.
+
+## Run policy
+
+| Case block | Total steps | Burn | Statistics |
+|---|---:|---:|---:|
+| S1 and S2 | 100000 to 160000 | first 35 to 40 percent | last 60 to 65 percent |
+| S3 and S4 | 180000 to 260000 | first 45 percent | last 55 percent |
+| SO1 and SO2 | 140000 to 220000 | first 45 percent | last 55 percent |
+
+For high blockage refined runs, prefer the longer end of the window.
+
+## Evaluation rule
+
+Use a two-layer rule.
+
+### Primary periodic anchors
+
+| Case | Hard target use |
+|---|---|
+| S1 to S4 | hard periodic benchmark anchors |
+
+Preferred agreement band:
+
+- within 5 percent when `Re_real` is close to target and the spectrum is clean
+- within 10 percent still acceptable if the run is clearly periodic and the residual mismatch is explainable by `Re_real` drift or geometry realization
+
+### Secondary onset block
+
+| Case | Hard target use |
+|---|---|
+| SO1 and SO2 | no hard `St` gate |
+
+Success means:
+
+- the flow is clearly unsteady and periodic
+- the dominant frequency is stable over long windows
+- the wake classification is consistent with being above onset
+
+## Deliverables
+
+For each run, deliver:
+
+- one row with `beta_nominal`, `beta_real`, `Re_nominal`, `Re_real`, `nu`, collision, wall type, inlet scheme, and grid resolution
+- one downstream velocity-profile plot
+- one force-history CSV
+- one `St` estimate with the exact analysis window stated
+- selected wake images for flow classification
+
+## Recommended minimum set
+
+If compute budget is tight, run this order first:
+
+| Priority | Runs |
+|---|---|
+| 1 | MRT on S1 to S4 |
+| 2 | SRT on S2 and S4 |
+| 3 | TRT on S2 and S4 |
+| 4 | SO1 and SO2 only after the primary anchors are behaving |
+
+## MRT-only runner mapping
+
+The current executable entrypoint is `tests/run_sah04_st_matrix.py`, and it is now aligned to this document's primary S1-S4 matrix:
+
+- collision is fixed to `MRT`
+- inlet is fixed to `parabolic + channel_stabilized`
+- case set is `S1-S4` only
+- output rows include `case_id`, `collision`, `inlet_scheme`, `grid`, `steps`, `burn_in`, `St`, `St_error_pct`, `Re_real`, `beta_real`
+- default hard gate is 5% (`--gate-pct` can relax it to 10%)
+
+Example commands:
+
+```bash
+conda run -n pycuda_3_10 python tests/run_sah04_st_matrix.py \
+  --json-out tests/output/sah04_mrt/summary.json
+
+conda run -n pycuda_3_10 python tests/run_sah04_st_matrix.py \
+  --case S3 --gate-pct 10 --final-vorticity-dir tests/output/sah04_mrt/vorticity
+```
+
+## Reference
+
+[Sah04] M. Sahin and R. G. Owens, “A numerical investigation of wall effects up to high blockage ratios on two-dimensional flow past a confined circular cylinder,” 2004.

docs/validation_specs/exp_ctrl_matrix.md

@@ -0,0 +1,181 @@
+229：无控制:
+```
+SIGNAL_FEATURES0 = {
+    'action1': {
+        'mean': 0.0,
+        'components': [
+            (0.1354, 0.0, 1.600),
+        ]
+    },
+    'action2': {
+        'mean': -0.0,
+        'components': [
+            (0.1354, 0.0, 2.099),
+        ]
+    },
+    'action3': {
+        'mean': 0.0,
+        'components': [
+            (0.1354, 0.0, 1.639),
+        ]
+    }
+}
+```
+234：隐身:
+```
+SIGNAL_FEATURES1 = {
+    'action1': {
+        'mean': 0.0,
+        'components': [
+            (0.1354, 0.0, 1.600),
+        ]
+    },
+    'action2': {
+        'mean': -0.01806,
+        'components': [
+            (0.1354, 0.0, 2.099),
+        ]
+    },
+    'action3': {
+        'mean': 0.01806,
+        'components': [
+            (0.1354, 0.0, 1.639),
+        ]
+    }
+}
+```
+欺骗：
+237 238 253：
+```
+SIGNAL_FEATURES2 = {
+    'action1': {
+        'mean': 0.0,
+        'components': [
+            (0.1354, 0.0026, 1.600),  # 主频
+        ]
+    },
+    'action2': {
+        'mean': -0.008730,
+        'components': [
+            (0.1354, 0.0045, 2.099),  # 主频
+            (0.2708, 0.0010, 0.612),
+        ]
+    },
+    'action3': {
+        'mean': 0.008730,
+        'components': [
+            (0.1354, 0.0045, 1.639),  # 主频
+            (0.2708, 0.0010, -2.962),
+        ]
+    }
+}
+```
+257:
+```
+SIGNAL_FEATURES3 = {
+    'action1': {
+        'mean': 0.0,
+        'components': [
+            (0.1354, 0.0029, -2.619),  # 主频
+            (0.2708, 0.0008, 2.856),
+        ]
+    },
+    'action2': {
+        'mean': -0.0140,
+        'components': [
+            (0.1354, 0.0050, -0.933),  # 主频
+            (0.2708, 0.0010, 0.801),
+            (0.1806, 0.0003, 1.854),
+        ]
+    },
+    'action3': {
+        'mean': 0.014,
+        'components': [
+            (0.1354, 0.0050, -1.398),  # 主频
+            (0.2708, 0.0010, 2.208),
+            (0.1806, 0.0003, 1.810),
+        ]
+    }
+}
+```
+260 (f*1.5):
+```
+SIGNAL_FEATURES4 = {
+    'action1': {
+        'mean': 0.0,
+        'components': [
+            (0.2031, 0.0026, 1.600),  # 主频
+        ]
+    },
+    'action2': {
+        'mean': -0.008730,
+        'components': [
+            (0.2031, 0.0045, 2.099),  # 主频
+            (0.4062, 0.0010, 0.612),
+        ]
+    },
+    'action3': {
+        'mean': 0.008730,
+        'components': [
+            (0.2031, 0.0045, 1.639),  # 主频
+            (0.4062, 0.0010, -2.962),
+        ]
+    }
+}
+```
+262: (f*1.5):
+```
+SIGNAL_FEATURES5 = {
+    'action1': {
+        'mean': 0.0,
+        'components': [
+            (0.2031, 0.0029, -2.619),  # 主频
+            (0.4062, 0.0008, 2.856),
+        ]
+    },
+    'action2': {
+        'mean': -0.0140,
+        'components': [
+            (0.2031, 0.0050, -0.933),  # 主频
+            (0.4062, 0.0010, 0.801),
+            (0.2709, 0.0003, 1.854),
+        ]
+    },
+    'action3': {
+        'mean': 0.014,
+        'components': [
+            (0.2031, 0.0050, -1.398),  # 主频
+            (0.4062, 0.0010, 2.208),
+            (0.2709, 0.0003, 1.810),
+        ]
+    }
+}
+```
+270 (f*2):
+```
+SIGNAL_FEATURES6 = {
+    'action1': {
+        'mean': 0.0,
+        'components': [
+            (0.2708, 0.0044, -2.619),  # 主频
+            (0.8124, 0.0012, 2.856),
+        ]
+    },
+    'action2': {
+        'mean': -0.014,
+        'components': [
+            (0.2708, 0.0075, -0.933),  # 主频
+            (0.8124, 0.0015, 0.801),
+            (0.5418, 0.0005, 1.854),
+        ]
+    },
+    'action3': {
+        'mean': 0.014,
+        'components': [
+            (0.2708, 0.0075, -1.398),  # 主频
+            (0.8124, 0.0015, 2.208),
+            (0.5418, 0.0005, 1.810),
+        ]
+    }
+}
+```

docs/validation_specs/streakline_notes.md

@@ -0,0 +1,493 @@
+## Streakline postprocessing design
+
+这份说明面向一个最小可实现的离线后处理程序。输入是按时间保存的二维速度场序列，输出是接近实验染色图的 streakline 图像。核心路线是连续释粒，因为 streakline 正是“固定位置持续释放染料后，在某个时刻所有已释放粒子的位置集合” [Lan96]。对非定常尾迹，这比单帧 streamline 更贴近实验图像 [Lan96]。
+
+搜索和全文阅读给出的信息已经足够支撑第一版实现。最直接的方案是按保存时刻连续注入粒子，用时空插值后的速度场推进粒子，再将粒子云渲染为图像 [Lan96, Ken96]。如果纯粒子结果过于尖锐，再在推进后加入一个很小的随机扩散项，作为对真实染色液扩散的近似。更完整的被动标量法当然更物理，但实现成本更高，不适合作为第一版 [Kim04]。
+
+## Core recommendation
+
+| 方法 | 与实验染色图的对应 | 实现复杂度 | 适合作为第一版 |
+|---|---|---:|---|
+| 单帧 streamline | 弱 | 很低 | 不推荐 |
+| 连续释粒 streakline | 强 | 低 | 最推荐 |
+| streakline 加少量扩散 | 很强 | 低到中 | 推荐作为第二步 |
+| 被动标量对流扩散 | 最强 | 中到高 | 暂不作为第一版 |
+
+最小可实现路线如下：
+
+- 从 CelerisLab 导出一串二维速度场快照 `u_x(x,y,t_k), u_y(x,y,t_k)`
+- 选定一个固定释放点或一小段释放线
+- 在每个保存时刻都注入新粒子
+- 在两个相邻快照之间，用时间插值和空间插值计算粒子速度
+- 用二阶或四阶时间积分推进粒子
+- 删去出域粒子与进入固体的粒子
+- 将当前时刻存活粒子按位置和年龄渲染成图像
+
+这个程序逻辑直接对应 [Lan96] 的 streakline 算法，插值与推进细节由 [Ken96] 和 [Dar96] 支撑。
+
+## Data contract
+
+默认输入是规则 Cartesian 网格上的二维时间序列速度场。对当前项目，最自然的输入约定是：
+
+| 名称 | 含义 |
+|---|---|
+| `t_k` | 第 `k` 个保存时刻 |
+| `ux[k, j, i]` | 时刻 `t_k` 的 x 方向速度 |
+| `uy[k, j, i]` | 时刻 `t_k` 的 y 方向速度 |
+| `mask[j, i]` | 可选，固体与流体标记 |
+| `x_i, y_j` | 网格坐标 |
+
+第一版最好满足下面三条：
+
+- 时间快照间隔固定或至少已知
+- 网格坐标固定不动
+- 固体几何位置已知，至少能判断粒子是否进入圆柱内部
+
+如果输出来自 CelerisLab，本质上只需要把每个保存时刻的 `ux, uy` 和对应时间写出来即可。算法本身是通用的，并不依赖 LBM 本身。
+
+## Mathematical model for the base streakline
+
+### Particle motion
+
+粒子位置满足拉格朗日运动方程 [Lan96]：
+
+\[
+\frac{d \boldsymbol{x}}{dt} = \boldsymbol{u}(\boldsymbol{x}, t)
+\]
+
+其中
+
+\[
+\boldsymbol{x} = (x, y), \qquad \boldsymbol{u} = (u_x, u_y)
+\]
+
+积分形式为 [Lan96]：
+
+\[
+\boldsymbol{x}(t + \Delta t) = \boldsymbol{x}(t) + \int_t^{t+\Delta t} \boldsymbol{u}(\boldsymbol{x}(\tau), \tau) \, d\tau
+\]
+
+streakline 在观察时刻 \(t_n\) 的定义是：从固定释放位置 \(\boldsymbol{x}_s\) 在过去各时刻持续注入的粒子，在 \(t_n\) 时刻的全部位置集合 [Lan96]。
+
+### Time interpolation
+
+速度场只在离散时刻存储，所以必须做时间插值。对位于 \(t_k \le t \le t_{k+1}\) 的任意子步，最简单的做法是线性时间插值 [Ken96]：
+
+\[
+\delta = \frac{t - t_k}{t_{k+1} - t_k}
+\]
+
+\[
+\boldsymbol{u}(\boldsymbol{x}, t) = (1 - \delta) \, \boldsymbol{u}_k(\boldsymbol{x}) + \delta \, \boldsymbol{u}_{k+1}(\boldsymbol{x})
+\]
+
+这里 \(\boldsymbol{u}_k(\boldsymbol{x})\) 和 \(\boldsymbol{u}_{k+1}(\boldsymbol{x})\) 仍需通过空间插值得到。
+
+### Spatial interpolation
+
+在规则网格上，第一版直接用双线性插值即可。虽然 [Ken96] 讨论的是非结构网格上的点定位与线性插值，但其一般原则完全适用于规则网格：
+
+- 先定位粒子所在单元
+- 再用单元顶点速度插值得到粒子点速度 [Ken96]
+
+若粒子位于单元局部坐标 \((\xi, \eta) \in [0,1]^2\)，四个角点速度为 \(\boldsymbol{u}_{00}, \boldsymbol{u}_{10}, \boldsymbol{u}_{01}, \boldsymbol{u}_{11}\)，则
+
+\[
+\boldsymbol{u}(\xi, \eta) = (1-\xi)(1-\eta) \boldsymbol{u}_{00}
++ \xi (1-\eta) \boldsymbol{u}_{10}
++ (1-\xi)\eta \boldsymbol{u}_{01}
++ \xi\eta \boldsymbol{u}_{11}
+\]
+
+## Time integration choice
+
+[Lan96] 和 [Ken96] 都直接使用四阶 Runge Kutta。对你的第一版程序，推荐两个选项：
+
+| 积分器 | 优点 | 缺点 | 建议 |
+|---|---|---|---|
+| RK2 | 简洁，容易调试 | 精度一般 | 最小原型可用 |
+| RK4 | 文献最一致，精度更稳 | 每步插值次数更多 | 默认推荐 |
+
+RK4 的更新式为 [Lan96, Ken96]：
+
+\[
+\boldsymbol{x}_{n+1} = \boldsymbol{x}_n + \frac{1}{6}(\boldsymbol{a} + 2\boldsymbol{b} + 2\boldsymbol{c} + \boldsymbol{d})
+\]
+
+其中
+
+\[
+\boldsymbol{a} = \Delta t \, \boldsymbol{u}(\boldsymbol{x}_n, t_n)
+\]
+
+\[
+\boldsymbol{b} = \Delta t \, \boldsymbol{u}(\boldsymbol{x}_n + \tfrac{1}{2}\boldsymbol{a}, t_n + \tfrac{1}{2}\Delta t)
+\]
+
+\[
+\boldsymbol{c} = \Delta t \, \boldsymbol{u}(\boldsymbol{x}_n + \tfrac{1}{2}\boldsymbol{b}, t_n + \tfrac{1}{2}\Delta t)
+\]
+
+\[
+\boldsymbol{d} = \Delta t \, \boldsymbol{u}(\boldsymbol{x}_n + \boldsymbol{c}, t_n + \Delta t)
+\]
+
+[Ken96] 特别强调，RK4 的每一个子步都要重新做点定位与时空插值。这在规则网格上不难实现，只是意味着一次完整步进需要四次速度查询 [Ken96]。
+
+## Practical timestep guidance
+
+第一版不必上自适应步长，但步长不能随意设大。搜索里最重要的精度提醒来自 [Dar96]：
+
+- 非定常粒子积分的误差常常由时间离散控制，而不是概念本身
+- 合理步长必须不大于流动主要非定常时间尺度的同量级 [Dar96]
+- 若时间步过大，粒子轨迹会在拓扑上都出错，而不仅是位置略偏 [Dar96]
+
+对当前项目，最简单的工程准则是：
+
+\[
+\Delta t_{trace} \le \min \bigl(\alpha_t \, \Delta t_{save},\; \alpha_x \, \frac{\Delta x}{\max |u|} \bigr)
+\]
+
+其中 \(\Delta t_{save}\) 是两个保存快照之间的时间间隔，\(\Delta x\) 是网格尺度。第一版可取：
+
+- \(\alpha_t = 0.1 \sim 0.25\)
+- \(\alpha_x = 0.25 \sim 0.5\)
+
+这不是文献中的严格上界，而是结合 [Dar96] 的结论给出的实现准则：粒子推进子步既要显著小于保存间隔，也不要一子步跨过太多网格。
+
+[Ken96] 还给了一个很实用的自适应思路：根据相邻速度方向夹角调节步长。如果速度方向变化太快就减半，变化很小就加倍 [Ken96]。这很适合后续增强版，但不属于第一版必需项。
+
+## Base streakline algorithm
+
+### Release strategy
+
+最符合染色实验的是连续释粒 [Lan96]。第一版可用两种释放方式：
+
+| 方式 | 适用场景 | 建议 |
+|---|---|---|
+| 单点释放 | 针头式染料入口 | 最简单 |
+| 短线段释放 | 更像细缝或细带注入 | 更稳健 |
+
+如果实验是在圆柱上游某一点持续释放染色液，第一版就直接使用单点释放。
+
+### Particle state
+
+每个粒子只需保存：
+
+| 字段 | 含义 |
+|---|---|
+| `x, y` | 当前坐标 |
+| `t_birth` | 释放时刻 |
+| `age` | 当前年龄 |
+| `alive` | 是否仍在域内 |
+
+第一版不需要保存整条历史轨迹。因为你要的是当前时刻的 streakline 图，而不是每个粒子的 pathline 曲线。
+
+**常见误区（务必区分）**
+
+| 概念 | 画什么 | 是否等于水洞染色迹线 |
+|---|---|---|
+| **Streakline（迹线）** | 固定源点**持续释粒**，在观察时刻 \(t_n\) 画出**所有仍存活粒子的当前位置**（可带年龄衰减权重） | 是 |
+| **Pathline（迹线/轨道）** | 单个粒子从释放到当前的**整条历史轨迹** | 否 |
+| **错误实现** | 每隔一段时间放一个粒子，再把**所有粒子走过的路径段**全部叠加到图上 | 否（这是 pathline 叠加，不是 streakline） |
+
+水洞实验：针头在固定点连续注 dye → 某一时刻拍照 → 看到的是“此刻染料粒子在流场里的分布”，一条下游色带由**不同释放时刻、当前仍在本流场中的粒子**共同构成，而不是把每个粒子从出生到现在的轨迹都画出来。
+
+### Main loop pseudocode
+
+```text
+given velocity snapshots U[k] = {ux[k], uy[k]} at times t[k]
+given seeding point or seeding segment S
+initialize empty particle list P
+
+for k = 0 to N-2:
+
+    inject new particles at source S at time t[k]
+
+    set t_local = t[k]
+    while t_local < t[k+1]:
+        dt = min(dt_trace, t[k+1] - t_local)
+
+        for each particle p in P with p.alive:
+            v = interpolate_velocity(p.position, t_local)
+            advance p by one integration step using RK2 or RK4
+            update p.age
+            if p leaves domain:
+                p.alive = false
+            if p enters solid body:
+                p.alive = false
+
+        t_local = t_local + dt
+
+    optionally remove very old particles
+    optionally render current particle cloud
+```
+
+这个结构与 [Lan96] 的离散 streakline 算法一致，只是把“从 `t_k` 到 `t_{k+1}` 的一次推进”细化成多个更小的粒子子步，以满足 [Dar96] 对精度的要求。
+
+### Velocity query pseudocode
+
+```text
+function interpolate_velocity(position x, time t):
+    find k such that t[k] <= t <= t[k+1]
+    compute delta = (t - t[k]) / (t[k+1] - t[k])
+
+    u_k   = bilinear_interpolation(U[k],   x)
+    u_k1  = bilinear_interpolation(U[k+1], x)
+
+    return (1 - delta) * u_k + delta * u_k1
+```
+
+这是第一版最核心的数值部件。只要这部分实现正确，streakline 程序主体就很直接。
+
+## Small diffusion particle model
+
+当纯 streakline 太细、太锐利、不像染色液图像时，可以给每个粒子加一个很小的随机扩散项。这不是严格的被动标量求解，但能以很低代价增加染色带宽度。
+
+### Guiding idea
+
+[Kim04] 表明，染料图像更接近一个被动标量浓度场，其基本控制方程是对流扩散方程：
+
+\[
+\frac{\partial c}{\partial t} + \boldsymbol{u} \cdot \nabla c = D \nabla^2 c
+\]
+
+其中 \(c\) 是染料浓度，\(D\) 是扩散系数。对第一版粒子法，一个常见近似是将每个粒子的位置更新写成“对流加随机扩散”：
+
+\[
+\boldsymbol{x}_{n+1} = \boldsymbol{x}_n + \Delta \boldsymbol{x}_{adv} + \Delta \boldsymbol{x}_{diff}
+\]
+
+其中 \(\Delta \boldsymbol{x}_{adv}\) 由 RK2 或 RK4 给出，扩散项取二维各向同性随机增量：
+
+\[
+\Delta \boldsymbol{x}_{diff} = \sqrt{2 D \Delta t}
+\begin{bmatrix}
+\eta_x \\
+\eta_y
+\end{bmatrix}
+\]
+
+这里 \(\eta_x, \eta_y \sim \mathcal{N}(0,1)\)。
+
+这个形式本身是对扩散过程的标准随机游走近似。搜索结果提醒，随机游走模型如果处理不当会产生假漂移与错误浓度偏置 [Mac92]。因此第一版的使用原则应当很克制：
+
+- 只用于加入少量模糊和厚度
+- 不把粒子密度当成严格浓度
+- 不在强近壁统计上过度解读结果
+
+### Diffusive particle update pseudocode
+
+```text
+for each particle p in P with p.alive:
+    x_adv = RK4_step(p.position, t_local, dt)
+
+    sigma = sqrt(2 * D * dt)
+    dx_rand = sigma * normal(0, 1)
+    dy_rand = sigma * normal(0, 1)
+
+    x_new = x_adv + [dx_rand, dy_rand]
+
+    if x_new leaves domain:
+        p.alive = false
+    else if x_new enters solid:
+        p.alive = false
+    else:
+        p.position = x_new
+        p.age += dt
+```
+
+### Choosing the diffusion level
+
+第一版不必试图从真实染料物性严格标定 \(D\)。更实用的做法是把 \(D\) 当作视觉匹配参数，并保持它足够小，使图像结构仍主要由对流控制。
+
+[Kim04] 用 Schmidt 数控制扩散强弱。若已有参考速度 \(U\) 和长度尺度 \(L\)，则
+
+\[
+Sc = \frac{\nu}{D}
+\]
+
+也可写成
+
+\[
+D = \frac{\nu}{Sc}
+\]
+
+对第一版，可以用下面的思路选扩散强度：
+
+| 目标效果 | 建议 |
+|---|---|
+| 只想让线条略微变厚 | 取较大 `Sc`，即很小的 `D` |
+| 想模拟明显洗开和模糊 | 取较小 `Sc`，即较大的 `D` |
+| 不确定 | 先从几乎看不出的弱扩散开始 |
+
+因为当前目标是“少量扩散”，所以推荐先把扩散当成弱修饰，而不是主导机制。
+
+## Rendering logic
+
+最终图像不需要把每个粒子的完整轨迹都画出来。更像实验染色图的做法是把当前存活粒子投影到图像网格上，生成粒子密度图或带年龄权重的强度图。
+
+### Minimal rendering choices
+
+| 方法 | 图像风格 | 实现难度 |
+|---|---|---:|
+| 直接散点 | 最简陋 | 很低 |
+| 网格计数直方图 | 像浓度图 | 低 |
+| 高斯核累积 | 更平滑 | 低到中 |
+
+第一版推荐：
+
+- 把每个粒子投到像素网格
+- 对像素做计数或加权累积
+- 最后做一次轻微 Gaussian blur
+
+这与先做完整被动标量相比便宜很多，但视觉上已经会很接近实验染色图。
+
+### Optional particle weighting
+
+可以给粒子一个简单权重：
+
+\[
+I = \sum_p w_p K(\boldsymbol{x} - \boldsymbol{x}_p)
+\]
+
+其中 \(K\) 是像素核或 Gaussian 核。第一版里，权重 \(w_p\) 可直接取 1，也可对年龄做衰减，例如：
+
+\[
+w_p = \exp\left(- \frac{\mathrm{age}_p}{\tau_f} \right)
+\]
+
+这样旧粒子会逐渐淡出，图像不会无限堆积。
+
+## Boundary handling
+
+这部分先按最小原则处理即可。
+
+| 情况 | 第一版处理 |
+|---|---|
+| 粒子出计算域 | 直接删除 |
+| 粒子进入圆柱内部 | 直接删除 |
+| 粒子贴近边界滑动 | 暂不专门处理 |
+| 扩散后跨入固体 | 直接删除 |
+
+这样做的优点是简单稳妥。后续若发现近壁 streakline 误差明显，再考虑反射、投影或更物理的壁面处理。
+
+## Minimal implementation plan
+
+### Version 1
+
+目标是尽快得到可信的 streakline 图：
+
+- 读取 `ux, uy, t`
+- 单点连续释粒
+- 双线性空间插值
+- 线性时间插值
+- RK4 粒子推进
+- 出域删除与入固体删除
+- 粒子计数成图
+
+### Version 2
+
+在不改变主体架构的前提下增强视觉效果：
+
+- 增加短线段释放
+- 增加年龄衰减
+- 增加弱随机扩散
+- 用 Gaussian 核代替简单计数
+
+### Version 3
+
+若后续发现实验图像明显受扩散与混合主导，再转向更重的模型：
+
+- 被动标量对流扩散
+- 用 \(Sc\) 或 \(D\) 做参数标定 [Kim04]
+
+## Design decisions that are already justified by the literature
+
+下面这些设计已经有足够文献支撑，可以直接采用：
+
+- 用 streakline 而不是 streamline 对应持续染色实验 [Lan96]
+- 用连续释放粒子离线重构 streakline [Lan96]
+- 用时空插值从离散速度快照查询粒子速度 [Ken96]
+- 用 RK4 作为默认推进器 [Lan96, Ken96]
+- 把步长取得明显小于保存时间间隔，并避免一次跨过太多网格 [Dar96]
+- 若需要更像染料图，可以加入弱扩散，或后续上升到被动标量 [Kim04]
+
+## Recommended default choices
+
+| 项目 | 默认选择 |
+|---|---|
+| 释放方式 | 单点连续释粒 |
+| 空间插值 | 双线性 |
+| 时间插值 | 线性 |
+| 积分器 | RK4 |
+| 粒子子步 | 保存步长的 0.1 到 0.25 |
+| 固体处理 | 入固体即删除 |
+| 输出图像 | 粒子密度图加轻微模糊 |
+| 扩散 | 默认关闭，作为第二步 |
+
+## Bottom line
+
+对当前项目，最简洁且足够准确的程序逻辑不是去解新的染料场，而是先做一个离线连续释粒 streakline 后处理器。它只依赖时间序列速度场，程序结构清楚，数值风险也集中在可控的几个环节：时空插值、积分步长和边界删除 [Lan96, Dar96, Ken96]。在此基础上，再增加一个弱随机扩散项，就能以很小代价把图像从“几何上正确的粒子线”推进到“更像实验染色照片”的粒子云 [Kim04]。
+
+## CelerisLab implementation mapping
+
+Current implementation lives in:
+
+- library: `src/CelerisLab/common/streakline.py`
+- Kan99b demo CLI: `tests/run_kan99b_streakline.py`
+- experiment notebook: `tests/experiment.ipynb` (CLEAN 5/6)
+
+Available runtime modes:
+
+- `online`: sample `Simulation.get_macroscopic()` directly in memory (no velocity snapshot files).
+- `offline`: replay from an existing snapshot directory and run the same streakline integrator.
+
+Dense release support:
+
+- base upstream points are expanded by `release_mode` (`point|line|strip`).
+- `strip` mode densifies both cross-stream (`line_count`, `line_span`) and downstream (`downstream_count`, `downstream_spacing`) directions.
+
+### Performance notes (why it can feel slow)
+
+Streakline post-processing is **CPU-side** in the current stack:
+
+1. **GPU→host velocity copies** dominate for large grids (`nx=6000, ny=1200`). Each `get_macroscopic()` pulls full `ux/uy` arrays.
+2. Particle RK4 + bilinear interpolation runs in NumPy on CPU (vectorized over particles, not GPU).
+3. Rendering accumulates polylines into an image (optionally multi-threaded via `ThreadPoolExecutor`).
+
+Practical tuning:
+
+| knob | effect |
+|---|---|
+| `sample_every` | larger → fewer host copies, faster, coarser streaks |
+| `max_particle_age` | `None` → no hard cutoff; trails survive until boundary/solid |
+| `num_threads` | `0` = auto; sets OpenBLAS/OMP threads for blur/host ops |
+| `blur_sigma` | `0` disables Gaussian blur pass |
+
+For `experiment.ipynb` triangle case, use **`render_streakline_density(..., minimal_axes=True)`** on the **final alive particle cloud** (`positions`, `ages`). Do **not** accumulate and draw full path history (`render_snapshot_trails` is pathline-style debug only).
+
+### Clean render style (experiment)
+
+`render_snapshot_trails` defaults:
+
+- white background
+- red streaks, brighter toward the downstream end (`fade_along_trail=True`)
+- black filled cylinders
+- no axes, labels, colorbar, or release-point markers
+
+Example commands:
+
+```bash
+# Online in-memory streakline (no snapshot write)
+conda run -n pycuda_3_10 python tests/run_kan99b_streakline.py \
+  --mode online --domain M --re 100 --alpha 1.0 \
+  --sample-every 300 --n-snapshots 20 \
+  --release-mode strip --line-count 7 --downstream-count 6
+
+# Offline replay from existing velocity snapshots
+conda run -n pycuda_3_10 python tests/run_kan99b_streakline.py \
+  --mode offline \
+  --snapshot-dir tests/output/final_validation_round/streakline_kan99b_k2/velocity_snapshots \
+  --release-mode strip --line-count 7 --downstream-count 6
+```