zl_bench/README.md

# zl_bench — FlagGems 算子性能测试工具

针对编译器（FlagTree）改动做算子级 A/B 性能对比的 pytest 插件集 + 驱动脚本。在 FlagGems benchmark 体系之上解决三个问题：

1. **可复现**：固定随机种子、指定 shape，并通过 record/replay 让 A/B 两侧锁定同一套 autotune config，把两次运行之间的差异收敛到"编译器改动"这一个变量（**config 必须靠 replay 锁定，不会自己稳定**，见「为什么需要 replay」）；
2. **可解释**：自动按 shape 收集每次运行实际使用的 ttgir（带可读的变体命名），供 IR 级 diff；
3. **口径统一**：cudagraph 计时消除 launch 开销，小 kernel 的对比不被 CPU 侧噪声淹没；capture 前按 warmup 预算显式预热吸收 autotune/JIT，首轮即稳态、多轮一致。

## 先读这个：口径决定数字的含义

同一个算子、同一台机器、同一份代码，换一条测试口径，测出来的数字可以完全不同。
下表每一行都是本仓踩过的，量级是实测的，不是估计：

| 口径轴 | 改一下的量级 | 由谁固定 | 由谁检查 |
|---|---|---|---|
| autotune config（record vs replay） | router_gemm 差 8pp | `REPLAY_FROM` | `caliber_complaints`：两侧都 record 直接拒绝出报告 |
| timer（cuda_graph vs direct） | replication_pad3d 从 +43% 变成 +124.5% | `ZL_TIMER` / backend 能力 | 同上，`timer` 字段不一致即拒绝 |
| libtuner 持久缓存（共享 vs 隔离） | mm 看起来从 945 个变体降到 33 个——实际是 B 侧继承了 A 侧的缓存 | `ZL_ISOLATE_TUNER_CACHE=1` | 共享缓存上 record 会打 `警告 [caliber]` |
| `--level core` vs `comprehensive` | mm 的 shape 数 5 → 29，能不能覆盖到大 GEMM 全看这个 | `BENCH_LEVEL` | `bench_level` 字段不一致即拒绝 |
| 重复轮数 | 收益落在零附近时，单次测量给不出方向：topk 42 个 shape 有 17 个在两次等效运行间翻了符号 | `ROUNDS` / `ZL_DIRECT_REPEAT` | `direct_repeat` 字段不一致即拒绝 |
| 被测开关本身是否真的接上 | 开关名拼错，两侧跑同一份代码，脚本照样报出"无差异" | 开关名正确 | `ab_repeat.sh`：第 1 轮后比两侧 ttgir 指纹，相同即跳过该算子 |

**结论：一次运行的数字没有意义，只有一对同口径运行的差值有意义。** 每次运行把自己的
口径写进 `runs/<run>/run_identity.json`；`batch_summary.py --ab` / `--ab-rounds` 在出
报告前核对这些字段，不一致就拒绝，因为**口径不同的两次运行会产出一份和有效报告长得
一模一样的表**——没有任何一步会报错，只是数字换了含义。确认两次运行确实可比时，用
`--allow-caliber-mismatch` 显式放行。

文档拦不住人：README 里早就写着"B 侧必须走 REPLAY_FROM"和"只有 record 数据支撑的
结论按未验证处理"，本仓仍然发布过一个 record 对 record 得到的 router_gemm 数字。
所以上表第四列是代码，不是提醒。

单算子精测流程之上另有批量筛查层（`run_batch.py`，见「批量多算子测试」），可对双仓库上千个算子做例行体检，产物结构与单算子完全同构。

## 快速开始

```bash
# 默认算子（fused_marlin_moe_mxfp4，内置 16 组 MoE shape）
bash run_pytest.sh

# 换算子：OP=测试函数名去掉 test_ 前缀，OP_FILE=benchmark 文件名去掉 test_ 前缀/.py 后缀
# （benchmark 文件为 $FLAGGEMS_DIR/benchmark/test_<OP_FILE>.py::test_<OP>）
OP=softmax OP_FILE=softmax SHAPE_FILE=my_shapes.yaml bash run_pytest.sh

# 测 FlagGems-vllm 仓库的算子：换 FLAGGEMS_DIR 即可（插件自动适配其包名 flaggems_vllm）
FLAGGEMS_DIR=/workspace/dev/FlagGems-vllm OP=fused_marlin_moe OP_FILE=fused_marlin_moe bash run_pytest.sh

# 对齐 FlagGems 原生节点（不传 --shape_file，沿用 --level core 的 core_shapes）
ZL_BACKEND=cuda \
TEST_NODE=/workspace/FlagGems/benchmark/test_fused_marlin_moe.py::test_fused_marlin_moe_mxfp4 \
bash run_pytest.sh
```

上游现为两个仓库，benchmark 体系同构、均受支持：主仓 `/workspace/dev/FlagGems`（Python 包名 `flag_gems`，默认；`fused_marlin_moe_mxfp4` 仅此仓有）与 `/workspace/dev/FlagGems-vllm`（包名 `flaggems_vllm`）。两包需各自 `pip install -e <仓库> --no-deps` 安装（`--no-deps` 避免动 FlagTree 的 triton）。

shape yaml 顶层 key 必须是 op 名：

```yaml
softmax:
  shapes:
    - [1024, 1024]
    - [64, 512, 512]
```

上面这条命令是 **record 模式**：现场 autotune、把选中的 config 存档。它用于单次摸底或给 replay 提供基准，**两次 record 的数字不能互相比**——做对比走下面的「A/B 对比测试标准流程」。

跑之前确认目标卡空闲：共享机器上别的任务会把两侧一起等比拖慢，出一份看似自洽实则作废的数据。用 `ZL_BACKEND` 选运行时（默认 `cuda`），再用该后端的可见性变量选卡（CUDA 为 `CUDA_VISIBLE_DEVICES`，也可设 `ZL_DEVICE`）。**测量始终单进程单设备串行**，这是计时可比的前提；多卡仅用于可选的 sweep 加速（见「加速：多卡并行 sweep」）。非 CUDA 后端默认走上游 `do_bench`，不会尝试 CUDA Graph。

脚本内固定了 `--level core --mode kernel`，其余可调项都走环境变量（见下）；需要改动这两个口径本身时直接编辑 `run_pytest.sh` 中的 pytest 命令行。

### 环境变量一览

**测什么**

| 变量 | 默认 | 说明 |
|-----|------|------|
| `OP` | `fused_marlin_moe_mxfp4` | 测试函数名（`test_` 之后的部分）；设了 `TEST_NODE` 时由节点覆盖 |
| `OP_FILE` | `fused_marlin_moe` | benchmark 文件名（`test_` 与 `.py` 之间的部分）；设了 `TEST_NODE` 时由节点覆盖 |
| `TEST_NODE` | 空 | FlagGems pytest 节点（`…/benchmark/test_<op_file>.py::test_<op>`）。设了且未指定 `SHAPE_FILE` 时默认不传 `--shape_file` |
| `SHAPE_FILE` | 空 | shape yaml 路径；空且无 `TEST_NODE` 时用脚本内置 yaml；`-`/`upstream`/`none` 表示沿用上游在 `BENCH_LEVEL` 下的原生 shape 集 |
| `SHAPE_MODE` | 空 | `inline` / `file` / `upstream`；通常不用手设，`TEST_NODE` 会把空 `SHAPE_FILE` 推成 `upstream` |
| `FLAGGEMS_DIR` | `/workspace/dev/FlagGems` | FlagGems 仓库路径。测试文件缺失时回退 `/workspace/FlagGems`；`TEST_NODE` 也可直接带出仓库根 |
| `ZL_BACKEND` | `cuda` | 设备后端：`cuda` / `musa` / `rocm` / `ascend`（别名 `nvidia`/`amd`/`npu` 等；未知名直接失败） |
| `ZL_DEVICE` | 空 | 父进程看到的设备 id，写入 `run_identity.json`；批量层绑卡后也会设 |

**怎么测**

| 变量 | 默认 | 说明 |
|-----|------|------|
| `REPLAY_FROM` | 空（record 模式） | 指向某次历史 run 目录，replay 其 autotune 选择。做 A/B 时 B 侧必须设（见「A/B 对比测试标准流程」） |
| `USE_FLAGTUNE` | `0` | `0` 走普通 autotune（跳过搜索、快速验证）；`1` 用 FlagTune 扩展调优空间（首跑全量搜索、慢）。A/B 两侧须取同值 |
| `BENCH_LEVEL` | `core` | 透传上游 `--level`。`comprehensive` 相对 `core` 做两件事：① 在 `set_shapes()` 里并入子类 `set_more_shapes()`（`mm` 因此从 5 个 shape 变 29 个，含 N/K 到 152064 的大 K 与 `models_shapes/*.yaml` 的模型 shape）；② 在 `get_input_iter()` 里对每个 shape 再产出一次变体（BLAS 是转置侧），用例数翻倍。注意 ① 只在 `SHAPE_MODE=upstream` 下可见——传了 `SHAPE_FILE` 时 `_shape_inject_plugin` 会在上游合并之后再覆盖一次、把 shape 集拉回 yaml 内容（BLAS 那类 arity 不一致的情形除外，插件会退让）；② 不经插件，任何 shape 来源下都生效。sweep 成本同步放大，shape 多时配合 `PARALLEL_WARMUP_GPUS`。A/B 两侧须取同值 |
| `ZL_ISOLATE_TUNER_CACHE` | `0` | FlagGems 的 libtuner winner 存在 `$FLAGGEMS_CACHE_DIR/config_cache` 的 sqlite 里（默认 `~/.flaggems`），跨 run 长存。`1` 给本次 run 一个私有缓存目录，两侧各自冷调优（README「各自最优」口径），跑完随 Triton 缓存一起删——选中的 config 已在 `autotune_records/` 里，DB 本身无留存价值。共享缓存的风险是单向的：先跑的一侧写入 winner，后跑的一侧直接命中、只编译那几个胜出变体。实测 `router_gemm` 共享模式下第二次运行编译变体从 234 掉到 10、耗时 52s→17s，隔离模式下两次都是 234 个变体。走 `REPLAY_FROM` 的标准 A/B 不受影响（config 本就锁死），record 对 record 才会中招。取值无论如何都写进 `run_identity.json` 的 `tuner_cache` |
| `ZL_DIRECT_REPEAT` | `1` | direct 计时的重复次数，取中位。只作用于 direct 路径：整体 `timer=direct` 的运行，以及图捕获运行里逐行回退的那些点（标记会带上 `repeat=N`）；捕获成功的行不受影响。`do_bench` 自己已在调用内取中位，压住的是迭代间抖动；跨调用重复补的是它看不见的那层——时钟爬坡、邻居进程进来这类漂移。调用方要 `quantiles` 或 `return_mode="all"` 时返回序列、跨调用中位无定义，自动不重复。在 MThreads 上这不是可选项：cvt 消除量最大的 `sort`/`dropout`/`rand` 恰好都无法捕获，那批结论正是由这条路径产出的。取值写进 `run_identity.json` 的 `direct_repeat` |
| `DTYPES` | 空（上游默认 dtype 扫描） | 空格分隔的 dtype 白名单（如 `"bfloat16 float16"`），逐个转为上游 `--dtypes`。算子不支持时上游报 `can't be supported by this op` |
| `PARALLEL_WARMUP_GPUS` | 空（串行） | 设为 `N`（≥2）时把 autotune sweep 分片到 N 卡并行，测量仍单卡串行；不设或 `<2` 则完全串行（见「加速：多卡并行 sweep」） |
| `ZL_ENABLE_CUDAGRAPH` | `auto` | `auto`/`1` 在**声明了 `supports_cuda_graph` 的后端**上走 Graph 计时（目前 cuda 与 musa），其余后端即使设 `1` 也拒绝并保持 `timer=direct`。注意 `timer=cuda_graph` 只表示**尝试过**捕获，不表示有任何一行捕获成功——S5000 上 `dropout`/`rand`/`masked_select` 是逐 shape 全部回退（每行 `BENCHMARK_DIRECT_NO_CUDAGRAPH`，零毒化、不丢行），整张表实际是 direct 计时。这种情况下设 `0` 走的是同一条路径，只是记录更诚实 |
| `ZL_ENABLE_AUTOTUNE` | `1` | `0` 关闭 autotune record/replay 插件 |
| `ZL_ENABLE_IR` | `1` | `0` 关闭 TTGIR / 编译元数据插件 |
| `ZL_ENABLE_MM_CLUSTER` | `0` | `1` 且后端支持时加载 Hopper mm cluster 规避 |
| `ZL_EXTRA_PLUGINS` | 空 | 空格分隔、追加到 `-p` 列表的插件名 |

**输出**

| 变量 | 默认 | 说明 |
|-----|------|------|
| `FLAGGEMS_PERF_COLOR` | 空（按 tty 自动判断） | `always`/`never` 强制开/关终端颜色；`run.log` 始终为去色纯文本 |

## 输出目录结构

每次运行产出 `runs/<op>_<时间戳>/`：

```
runs/softmax_20260716-031752/
├── run.log                  # 完整日志（含 SUCCESS 行的 latency/speedup 表；已去 ANSI 色的纯文本）
├── run_identity.json        # 本次运行的口径：provider/backend/device/timer/vendor/
│                          # bench_level/shape_mode/tuner_cache/direct_repeat/autotune_mode
│                          # （shell 写 stub，插件合并写入；对比工具据此判可比性）
├── shapes.yaml              # 本次实际使用的 shape（存档）
├── autotune_records/        # Triton autotune 选中的 config（供 replay）
│   ├── softmax.json
│   └── meta.json            # provider/backend；跨后端 replay 会被拒绝
└── ttgir/                   # 按 shape 分组的 IR 落盘
    ├── 1024x1024/
    │   └── softmax_kernel_inner/
    │       ├── w4s3__bf16.ttgir
    │       ├── w4s3__fp16.ttgir
    │       └── w4s3__fp32.ttgir
    ├── 4096x4096/...
    ├── index.tsv            # 每个编译变体的完整参数、启动次数、cache hash
    └── naming.md            # 文件名缩写图例
```

ttgir 关键设计：

- **只落盘"实际使用"的变体**：autotune sweep 中测过但落选的 config 不拷贝（只在 index.tsv 里留 `sweep loser` 记录）。mm 这类 sweep 上千个变体的算子，最终只留真正被选中的几个文件。
- **命名 = 组内有区分度的参数**：同一 (shape, kernel) 组内取值相同的 constexpr 不进文件名；多词参数缩写为首字母（`BLOCK_SIZE_M` → `BSM`，图例见 naming.md）；同名冲突依次用 dtype（`__fp16`）、参数对齐特化（`__EMdiv16`）、hash 前缀消歧。
- **崩溃也能拿到 IR**：dump 挂在进程 atexit 上，CUDA crash 后仍会落盘已编译部分。

## 多卡跑一批：campaign.sh

`ab_repeat.sh` 是单卡串行的。一批算子要并行时用 `campaign.sh` 编排——它只分活，
测量仍然全部由 `ab_repeat.sh` 做，那个脚本一行不动。

```bash
SWITCH=FLAGTREE_MUSA_RLC_ENHANCE DEVICES=0,1,2,3,4,5,6,7   SCREEN=runs/ab_rlc_XXX/screen.tsv ZL_BACKEND=musa   FLAGGEMS_DIR=/path/to/FlagGems DRY_RUN=1 bash campaign.sh
```

`DRY_RUN=1` 先看方案，确认了再去掉它起跑。

- **只取筛选表里 `IR_CHANGED` 的算子**——第一步的判定在这里生效
- **按实测耗时贪心分片**（`partition.py`），不是 hash 取模。S5000 上 16 个算子实测：
  hash 分片关键路径 120 分钟、有一张卡空转；贪心 70 分钟。**下界由最重的单个算子
  决定**（`topk` 一个就 70.6 分钟），加卡不解决，分片只保证别的卡不空等
- **每卡独立 `AB_ROOT`**——八个 `ab_repeat.sh` 同一秒启动，默认的时间戳目录名会撞
- **强制 `ZL_ISOLATE_TUNER_CACHE=1`**。8 个进程写同一个 sqlite 是要避开的事。代价是
  pin 的 config 与共享缓存批次不同，**两批的绝对延迟不可互引**——口径守卫会自己拒绝
- `EXTRA_OPS` 追加筛选表里没有的算子（如 benchmark 刚修好的），放到**负载最小**的卡上
- 跑完合并各卡 `status.tsv`（首列是卡号），并统计非零退出的格与三个标记的总数

设备可见性变量名从 backends 层取，不写死 `MUSA_VISIBLE_DEVICES`。

## 两步：先筛选，再实测

跑一个算子的六轮 A/B 要 13~44 分钟（S5000 实测：`rand` 15 shape 13.0 分钟，
`sort` 25 shape 43.4 分钟）。用 `ab_rlc_test.sh` 筛一对只要约 0.8 分钟。
**差 20~50 倍**，所以先筛后测：

```bash
# 第一步：一对定去留（A 侧 record，B 侧 REPLAY_FROM 它，config 锁死）
OPS="col2im:col2im topk:topk perf_fft:fft" VENDOR=MUSA bash ab_rlc_test.sh

# 第二步：只把 IR_CHANGED 的送进来
SWITCH=FLAGTREE_MUSA_RLC_ENHANCE ROUNDS=6 OPS="topk:topk" bash ab_repeat.sh
```

`ab_rlc_test.sh` 除了原有的 speedup 区间，还给每个算子写一行判定到
`<root>/screen.tsv`：

| 判定 | 含义 | 该怎么办 |
|---|---|---|
| `IR_CHANGED` | 两侧 ttgir 不同 | 排六轮 A/B |
| `NO_IR_DIFF` | 逐字节相同（config 已锁死） | 跳过。跑六轮只会测到噪声，并报成一个"无差异"结论 |
| `NO_TRITON` | 跑出了行，但没编出 `.ttgir` | 本次运行没有任何东西经过 Triton。算子本身可能完全正常——多半跑的是 aten 或厂商库的 kernel——只是被测的编译器 pass 没有作用对象。**既不是"开关无作用"，也不是"没有 kernel"**。`run_pytest.sh` 每个 run 用独立 `TRITON_CACHE_DIR`，所以"缓存命中没重新编译"这个解释已被排除 |
| `NO_ROWS` | 没有测量行 | 先修可运行性：多半是 dtype 不被该算子接受，或它在本后端拿不到 config |

两件容易做错的：

- **判据是 ttgir 指纹，不是 `convert_layout` 计数。** cvt 数相同不等于生成代码
  相同，而要测的是"生成代码变了没有"。两个方向都会骗人：`addmm` 的 cvt 是 15→15
  没变，15 个 ttgir 文件却实质全不同；`tanh`/`softmax` 的 cvt 是 0→0，生成代码也
  真的没动。cvt 数照样记，它解释的是变化的性质。
- **指纹先抹掉开关自己盖的注解。** FlagTree 的 `_apply_musa_rlc_policy(mod)` 在开关
  打开时把整套 RLC 策略旋钮盖成 module 属性（`ttg.rlc-*`），跟这个 pass 后面有没有
  真的改动什么无关。不抹的话判据有个常量偏移——**凡是能编出 kernel 的算子都必然判成
  "变了"，筛选器谁都不排除**。24 个算子实测：抹之前 23 个 `IR_CHANGED`，抹之后 7 个
  落回 `NO_IR_DIFF`。零对照抓不到这个，因为两侧同值时属性也相同。
- **不要全局写死一个 `DTYPES`。** `OPS` 元素支持 `OP:OP_FILE:DTYPES`，留空表示用
  该算子自己的全套。`fft` 只吃 complex64，被一刀切的 `DTYPES=float16` 判死过一次
  ——那是 harness 的问题，被记成了算子的问题。

同一份 `_ir_fingerprint.sh` 也被 `ab_repeat.sh` 用于第 1 轮后的对比轴有效性检查，
两处判据不会漂开。

判定还管住报告：`NO_ROWS` 与 `NO_TRITON` 不出 speedup 区间（前者没有可比的测量，
后者与被测开关无关）；`NO_IR_DIFF` 照出，但明说那是**本机对该算子的噪声底**，不是
开关的收益——它离零多远，就是该算子单次测量的可信下限。判定刚说完这一对不可比，
下一行就给个 `+0.3%`，正是这套工具一路在防的"看起来像结论的表"。

`screen.tsv` 末七列是口径（`backend device timer bench_level shape_mode
tuner_cache direct_repeat`），由 `python batch_summary.py --caliber <run_dir>` 取——
`CALIBER_FIELDS` 只有那一处定义，筛选表不重列一遍。两份不同 `--level` 的
`screen.tsv` 放一起比时，这几列是唯一能拦住的东西。

每对有 `CELL_TIMEOUT`（默认 2400 秒）兜底，超时打 `!!! CELL_TIMEOUT` 并按耗时问题
报告，不混进 dtype 那条解释里。`elapsed_before/after` 也进表：两步流程的全部依据
就是"一对 0.8 分钟对六轮 13~44 分钟"这个比例，那这一步该记下自己那一半。

## 算子跑不起来时的排查清单

`NO_ROWS` / `NO_TRITON` 说的是"这一格没测出东西"，不是"开关无作用"。按下面的顺序查，
每一条都踩过：

| # | 症状 | 查什么 | 实例 |
|---|---|---|---|
| 1 | `PREFLIGHT_FAILED check=interpreter imports triton + flag_gems` | 用的是哪个 python。triton 通常来自 PYTHONPATH 指向的编译器构建 | 一次噪声标定因此废掉 12 格 + 160 次负载 run |
| 2 | `PREFLIGHT_FAILED check=benchmark file exists` | `OPS` 元素是 `OP:OP_FILE`，写反就找不到文件 | `test_perf_sort` 在 `test_sort.py` 里，要写 `perf_sort:sort` |
| 3 | `ValueError: Given dtype(s) '...' can't be supported by this op` | 该算子只吃特定 dtype，别全局写死 `DTYPES` | `fft` 只吃 complex64，被一刀切的 fp16 判成"跑不出行"，而它是 cvt 消除量最大的算子 |
| 4 | `dynamic_func() missing ... BLOCK_M, BLOCK_N ...` 之类 JIT 深处的报错 | **benchmark 是不是导入了别家厂商的实现**（见下） | `router_gemm` 拿 NVIDIA Hopper 的 kernel 跑在 MUSA 上 |
| 5 | 跑出行但 `kernels=0` | 该算子在本后端没走 Triton，走的是 aten 或厂商库 | — |

### 第 4 条：import 成功不等于这个后端支持

FlagGems 的厂商实现放在 `flag_gems/runtime/backend/_<vendor>/ops/`。benchmark 里常见
这种可用性判据：

```python
try:
    from flag_gems.runtime.backend._nvidia.hopper.ops.mm import router_gemm
    ROUTER_GEMM_AVAILABLE = True          # ← import 成功就置 True
except Exception:
    ROUTER_GEMM_AVAILABLE = False
```

**`import` 只证明那个 `.py` 在磁盘上，它不碰设备，所以在任何后端上都成功。** gate 于是
形同虚设：MUSA 上照样置 True，benchmark 拿 NVIDIA Hopper 的 kernel 跑在摩尔线程卡上，
最后死在 JIT 深处的 `dynamic_func() missing BLOCK_M, BLOCK_N, BLOCK_K, SPLIT_K`——
**那个报错完全指不向真正的原因**。而 `_mthreads/ops/mm.py` 下明明有一份能用的实现，
benchmark 从头到尾没导入过。

正确写法是先判 vendor（同一个文件里 Cambricon 那支就是对的）：

```python
if base.vendor_name == "nvidia":
    try:
        from flag_gems.runtime.backend._nvidia.hopper.ops.mm import router_gemm
        ROUTER_GEMM_AVAILABLE = True
    except Exception:
        pass
```

`_device_setup_plugin` 会在 session 结束时扫 `sys.modules`，载入过别家厂商实现时打：

```
FOREIGN_VENDOR_MODULE active=mthreads loaded=nvidia
```

**它不拦，只报**——载入别家模块本身不一定有问题（有些是共享工具），但一个算子结果可疑
时，这行是第一个该看的地方。

## A/B 对比测试标准流程

```bash
# A 侧（基线编译器）：正常跑，自动 record autotune 选择
bash run_pytest.sh    # -> runs/<op>_<ts_A>/

# B 侧（改动后编译器）：replay A 侧的 config，保证两侧同 config
REPLAY_FROM=$PWD/runs/<op>_<ts_A> bash run_pytest.sh
```

B 侧必须走 `REPLAY_FROM`。不能靠「改动前后各跑一次」来比——两次 record 会各自重新 sweep、可能选出不同 config，这个差异足以盖过被测改动本身（成因见「为什么需要 replay」）。

A/B 两侧的 `USE_FLAGTUNE` 必须取同值：该开关决定调优空间，两侧不一致时 replay 会大量 fallback，libtuner 持久缓存也各自独立命中，"同 config"前提不再成立。

对比 `run.log` 的 latency 表看性能差异；diff 两侧 `ttgir/<shape>/<kernel>/` 下的同名文件看 IR 差异。

带 `pre_hook` 的 config（如 hopper mm 的 TMA configs）：libtuner kernel 正常 record/replay——注入的 config 会由上游按 kwargs 匹配自动接回 pre_hook（与其自身 ConfigCache 的 DB 回读同一条路径）；普通 `@triton.autotune` kernel 无此恢复机制，这类 config 不落 record，replay 时表现为下述 `key_missing` fallback。

replay 的兜底行为：B 侧遇到记录中没有的 key、或记录的 config 在新编译器下编译失败时，自动回退到现场 autotune 并在 `run.log` 打 `AUTOTUNE_REPLAY_FALLBACK reason=...` 标记——出现该标记的测量点不再满足"同 config"前提，解读时注意。标记以 `_no_evict` 结尾时更弱一层：坏 config 存在 libtuner 的 sqlite 缓存里、没有 `__delitem__` 可摘除，重试会再读到同一个 config——这类测量点按未验证处理。另外 replay 模式的 run 目录不产生 `autotune_records/`，后续 run 的 `REPLAY_FROM` 应始终指向最初 record 的那次 A 侧目录，不要链式指向 replay 产物。

### 为什么需要 replay（以及它管不到什么）

本项目涉及两层调优机制，对 A/B 的影响不同：

| 机制 | 结果存储 | 编译器改动后 | A/B 风险 |
|-----|---------|------------|----------|
| Triton `@triton.autotune` | 仅进程内存 | 每次进程重新 sweep | 计时噪声可能让 A/B 选中**不同 config** → 用 REPLAY_FROM 固定（影响比预期大，见下） |
| FlagGems `@libtuner`（含 FlagTune 扩展空间） | `~/.flaggems/config_cache/*.db`（sqlite，跨进程持久） | **不失效**（表名只含 kernel 源码与 config 空间的 hash），A/B 自动命中同一 winner | 反向风险：B 侧沿用 A 侧选的旧 winner，测的是"旧 config 下的编译器差异"而非"各自最优" |

新版 libtuner 在同一 db 里还持久化了 **BenchmarkCache**（sweep 中每个 config 的实测 latency）：db 已热时 record 模式的 sweep 也不重测，直接按历史延迟选 winner——"record 每次现场 sweep"仅在冷 db 下严格成立。下文"删 db 换各自最优口径"的操作会同时清掉 winner 与延迟两层，仍然有效。

**这不是理论风险，量级足以吞掉被测优化本身。** 同一份 `shapes.yaml`、同一口径、相隔十几分钟的两次 record，平均 speedup 可以差出 10% 量级，且逐 shape 单向偏移（不是随机噪声）。diff 两侧 `autotune_records/*.json` 能看到差异往往不是微调而是换挡——`num_stages`、`BLOCK_SIZE_*`、`num_warps` 整档跳变。怀疑遇到这种情况时，先 diff 两侧的 config 再看 latency。

所以 record 模式的数字只用来给 replay 提供 config 基准。**若某次结论只有 record 数据支撑，按未验证处理、重跑补 replay。**

libtuner 的持久缓存何时失效：FlagGems kernel 源码改动、tune_configs.yaml / expand yaml / `USE_FLAGTUNE` 开关变化、Triton 大版本或 GPU 型号变化。该 db 跨 run 长存，**本脚本只清理自己的 `TRITON_CACHE_DIR`（Triton 编译产物），不碰它**——清编译缓存不等于重新调优，config 命中后照样沿用旧 winner。如果需要"各自最优"口径（让两侧各自重新 sweep），**删掉 `~/.flaggems/config_cache/TunedConfig_*.db` 或设 `FLAGGEMS_DB_URL` 指向一次性文件**——注意"各自 fresh tune"不等于"不设 `REPLAY_FROM`"：缓存已热时两侧都会直接命中同一 winner，看似独立调优实则同 config。**两种口径都合理，报告结论时注明用的哪种。**

同一层缓存也决定了 `_autotune_record_plugin` 的实现方式：record 不能靠"cache 新增了哪个 key"来判断本次选中的 config（缓存一热就走 cached 分支、不写新 key，键集差集恒为空），改为 `run()` 之后直接读回本次调用的 `self.cache[key]`；replay 同理不能加"key 不在 cache 里才注入"的前置条件，否则注入被跳过、该 run 表面在 replay 实际在自调优。想确认 record 真的覆盖到目标 kernel，查 `runs/<run>/autotune_records/<op>.json` 里有无对应 kernel 条目——漏记时该文件照样生成，只是少了 libtuner 那几个。

### 前置检查

`run_pytest.sh` 在启动 pytest 之前校验 FlagGems 目录、benchmark 文件、`SHAPE_FILE`
与 `REPLAY_FROM`，失败时打 `!!! PREFLIGHT_FAILED check=<具名检查>` 并以 3 退出。

这些检查存在的理由是**静默失败会指向错误的方向**：路径不存在时 pytest 收集不到
benchmark 目录，注册 `--level`/`--mode` 的 `conftest.py` 就没被加载，报错变成
`unrecognized arguments: --level core --mode kernel`，看起来像参数或版本问题。

只校验节点的文件那一半；`::test_<fn>` 交给 pytest，函数不存在时它自己会清楚地报
`not found: ...`。

还校验**解释器**——`python -c "import triton, flag_gems"`。triton 通常来自 PYTHONPATH
指向的编译器构建产物、flag_gems 来自源码检出，所以一个能摸到全部正确路径的 shell
仍然可能是错的那个 Python。不查的话，失败以每格一次 `pytest exit 4`（conftest 的
ImportError）出现，多轮 campaign 会在有人读日志之前把每一格都烧掉——这不是假想，
2026-09-10 的噪声标定实验就这么废了 12 格加 160 次负载 run。报错里带上
`$(command -v python)`，因为要看的就是"用的到底是哪个解释器"。

### 口径核对：出报告前的拒绝

`batch_summary.py --ab` / `--ab-rounds` 在解析任何 latency 之前，先读参与对比的每个
run 目录的 `run_identity.json`，逐字段核对 `backend / device / timer / bench_level /
shape_mode / tuner_cache / direct_repeat`。任何一项不一致，或任何一侧缺 identity，
就打 `!!! [caliber] ...` 并以 **4** 退出，不出报告。

```
!!! [caliber] sort: timer 不一致：before=cuda_graph, after=direct
!!! [caliber] sort: 拒绝出报告。确认口径确实可比后加 --allow-caliber-mismatch
```

单独判 `autotune_mode`：一侧 record 定 config、另一侧 replay 它是正常形态；**两侧都
record 直接拒绝**——那正是本仓发布过错误 router_gemm 数字的那次形态。

不拒绝但会警告的一类：在**共享** libtuner 缓存上 record。被 pin 的 config 取决于此前
在这台机器上跑过什么，相对结论仍成立，绝对延迟与变体数不可跨机器引用。

`--allow-caliber-mismatch` 放行，仍会把问题打出来。用它之前先能说清为什么这两次运行
可比；跨口径引用别人的数字，比自己重跑一遍贵得多。

### 多轮重复：ab_repeat.sh 与 --ab-rounds

单对观测无法说明自己稳不稳。S5000 实测 `topk` 消除了 32 个 `convert_layout`，
42 个 shape 里却有 17 个在两次同等有效的运行之间**改变了 delta 的符号**——那不是
计时口径问题（图捕获与 direct 的延迟比是 1.00x），是收益本身落在零附近、单次
测量给不出方向。

```bash
SWITCH=FLAGTREE_MUSA_RLC_ENHANCE OPS="topk masked_select" ROUNDS=6   ZL_BACKEND=musa FLAGGEMS_DIR=/path/to/FlagGems bash ab_repeat.sh
```

口径：第 1 轮 A 侧 record config，此后每一次运行（含 A 侧的后续轮次）都
`REPLAY_FROM` 那一份，config 因此跨轮跨侧全部锁死；逐轮交替先后顺序抵消单向漂移。
有轮次失败的算子不出报告——缺轮的报告会让人误以为是完整的 N 轮。

每跑完一格往 `<root>/status.tsv` 追一行：`round op side switch rc elapsed_s success
failed replay_fallback no_cudagraph capture_poisoned`。README 要求必查的三个标记就在
这张表里，跑的时候就能看见它们从 0 变成非 0，不用等结束再 grep——事后 grep 只在有人
想起来的时候才发生。每格另有 `CELL_TIMEOUT`（默认 2400 秒）兜底，超时记 `rc=124` 并打
`!!! CELL_TIMEOUT`，不会让一格卡死拖住整批。

全程还采一份 `<root>/telemetry.tsv`：

```
timestamp            index  sm_clock_mhz  mem_clock_mhz  temp_c  power_w  util_pct
2026-09-10T07:41:12  7      1888          2500           36      477.68   87
```

列在 cuda 与 musa 上完全一致（`nvidia-smi --query-gpu` 与 `mthreads-gmi -q --json`
各自解析后归一），所以两台机器的 campaign 可以并排读。理由很直接：六轮 A/B 声称的是
0.01% 量级的单侧噪声，**降频和热漂移能移动比这大得多的延迟，而且在 run.log 里不留
任何痕迹**——报告照样干净。`TELEMETRY_INTERVAL`（默认 5 秒）设 0 关闭；采样器缺失或
没权限时 `telemetry.py` 安静退出并把原因写进 `telemetry.err`，不影响测试本身。

第 1 轮 A/B 跑完后，脚本比两侧 `ttgir/` 的指纹（相对路径 + 内容）。**逐字节相同就
跳过这个算子**，并打 `!!! SWITCH_NO_IR_DIFF op=... switch=...`：这个开关没改变该算子的
生成代码，剩下 N-1 轮测的全是噪声，最后还会报成一个"无差异"的结论。先查开关名拼写
与它在被测源码里的读取点；确属只影响运行时、不影响 IR 的开关，用
`SWITCH_ALLOW_NO_IR_DIFF=1` 放行。

报告分开给两种离散度，它们回答不同的问题：

| 指标 | 含义 |
|---|---|
| 逐轮区间 | N 轮各自 delta 的 min~max，说明**判定**稳不稳 |
| 单侧噪声 | 同一侧 N 次延迟的极差，说明**机器**吵不吵 |
| 符号不稳行数 | delta 在轮次间改变符号的 shape 数，这些行的方向不可结论 |

已有产物可以单独重出报告，目录形态是 `<root>/r<k>_<A|B>/run.log`：

```bash
python batch_summary.py --ab-rounds runs/ab_repeat_XXX/topk --label topk
```

NV 上验证过这条路径：`router_gemm` 六轮把逐轮区间收窄到 -16.6% ~ -16.2%、单侧
噪声 0.4%，把单对观测升级成了结论。

### ab_fold_test.sh / ab_rlc_test.sh：A/B 测试参考示例

`ab_fold_test.sh` 是上述 A/B 流程的现成参考：以 FlagGems `FLAGGEMS_MXFP4_FOLDSCALE`（fold_scale 优化）开关为对比对象，对 4 个 DeepSeek-V4-Flash 真实 trace shape 集（`shape/DeepSeek-V4-Flash-p{1024,4096,32768,65536}d1024.yaml`，随仓库提供）各跑一对 FOLD=0（基线，record）→ FOLD=1（replay 同 config），共 8 轮：

```bash
bash ab_fold_test.sh
# 汇总 log：runs/ab_fold_<时间戳>.log（已去色）；快速对账：grep -E '^#####|FAILED' <log>
```

`ab_rlc_test.sh` 是同一口径的另一个实例，对比轴换成 FlagTree 的 `FLAGTREE_<VENDOR>_RLC_ENHANCE`（`VENDOR` 默认 `MUSA`），算子表由 `OPS` 指定，跑完调 `batch_summary.py --ab <before> <after>` 出该算子的 speedup 区间。

可复用的是**口径**——同 config replay、cudagraph 计时、真实 trace shape 集、逐对成组——对比其它开关 / 编译器改动 / 算子，改循环变量与环境变量即可，口径部分不用动。

但**它是口径示例，不是可长期直接执行的回归脚本**：对比轴是 FlagGems 侧的算子开关，会随上游演进被改名、改语义或移除，本仓库不跟随同步更新（该脚本引用的开关目前就已被上游移除）。失效后果是**静默的**：脚本照样跑完全部轮次、照样输出完整对比表，只是两侧执行同一份代码，得出"无差异"的假结论。所以套用前先确认对比轴仍有效——`grep -rn '<开关名>' $FLAGGEMS_DIR/src` 要有命中，且两侧 latency 表确有差异。

## 加速：多卡并行 sweep（可选）

shape 多时墙上时间几乎全在 autotune sweep，而不在测量——被测 kernel 本身往往只有毫秒量级，绝大部分编译变体是 sweep 里测完就丢的 loser（`ttgir/index.tsv` 的 `sweep loser` 行）。`_parallel_warmup_plugin` 把 sweep 分片到 N 卡并行，再交回单卡串行测量：

```bash
PARALLEL_WARMUP_GPUS=8 SHAPE_FILE=shape/DeepSeek-V4-Flash-p32768d1024.yaml bash run_pytest.sh
```

加速比取决于 sweep 占比：shape 多、config 空间大的算子收益最明显，可达数倍。不设该变量（或设 <2）时插件静默 no-op。**产物结构、`REPLAY_FROM` 用法、latency 表格式全不变**——合并后的 config 就写进本次 run 的 `autotune_records/<op>.json`，分片临时目录跑完即删。

**测量本身不并行**：N 个进程压满同机多卡会通过功耗墙/散热耦合，单卡 latency 被邻居拖慢且不可复现。并行只用于"决定哪个 config 胜出"，测量仍是单进程单卡、与不开插件走同一条路径。

因此要留意日志里的冲突警告——同一 config key 在不同分片选出了不同 winner，说明互扰已经影响到 sweep 结果：

```
WARNING 1 config key(s) got different winners across shards — parallel interference reached the sweep
```

**冲突数就是这次加速的可信度指标**：0 可放心用；偏多说明 config 选择已被污染，出正式结论前不设该变量重跑一遍。

插件会在三种情况下**主动放弃分片、退回串行 sweep**，日志各有独立措辞，都不是静默的：

- **`BENCH_LEVEL` 不是 `core`**。`comprehensive` 会让上游在 `set_shapes()` 里并入 `set_more_shapes()`，那部分是各子类自己的 Python（硬编码元组加 `models_shapes/*.yaml`），插件只看得到 yaml 里的 shape 列表。按子集分片会让未覆盖的 key 全部现场自调优，而这一轮却自称在 replay——`mm` 上实测是 29 个 shape 只扫到 5 个、90 次 `AUTOTUNE_REPLAY_FALLBACK reason=key_missing`。
- **任一分片非零退出**。死掉的分片的 shape 从未被扫过，合并出来的 config 集是带洞的，publish 之后每个洞都会静默自调优。因此整轮 sweep 直接作废，worker 日志保留在临时目录里备查。
- yaml 里查不到该算子的 shape（见下）、`REPLAY_FROM` 已设、或可见卡不足 2 张。

yaml key 的查找与上游对齐：先按 `op_name`，查不到再沿 Benchmark 类链逐级尝试类名（`mm` 因此命中 `BlasBenchmark`）。类链是静态解析测试模块里 `op_name=` 那次构造调用得到的，不 import FlagGems。日志会打出命中的 key；一个都没命中时会列出尝试过的候选。

其余行为（轮转分片而非按块切、用子进程而非 xdist 的原因）见插件 docstring。分片绑卡遵守当前后端的可见性变量（CUDA 为 `CUDA_VISIBLE_DEVICES`）——用它选过卡时，分片只会落在你选的那几张上。

## 批量多算子测试（run_batch.py）

单算子流程之上的批量筛查层：按清单逐算子调 `run_pytest.sh`，多卡分片、每卡内部串行，产出逐算子状态表和逐测量行总表。**定位是筛查口径**——跨卡并行测量存在功耗/散热耦合噪声（幅度可到百分之几），发现可疑算子后回单算子串行流程（前几节）确认，不要直接拿批量数字下精细结论。

```bash
# 全量筛查（双仓库全部可收集算子，排除带无条件 skip 标记的与 ops/dismiss.txt 里的）
python run_batch.py --all --devices 0,1,2,3,4,5,6,7 --backend cuda --dtypes bfloat16 --op-timeout 1800
# --gpus 仍是 --devices 的别名；两者同时给出会报错

# 指定子集：默认 main 仓，vllm 仓加前缀；同名测试函数用 @文件名 消歧
python run_batch.py --ops softmax vllm:fused_marlin_moe nextafter_@nextafter_

# 断点续跑（PASS/SKIP 的复用，FAIL/TIMEOUT 的重跑）
python run_batch.py --all --batch-dir runs/batch_xxx --resume

# 锁 config 复测：逐算子 REPLAY_FROM <root>/<repo>/<op>，语义同单算子 REPLAY_FROM
python run_batch.py --ops-file my_ops.txt --replay-root runs/batch_xxx
```

**算子资产（`ops/`）**：

| 文件 | 内容 | 维护方式 |
|-----|------|---------|
| `inventory_main.csv` / `inventory_vllm.csv` | 双仓库全部 benchmark 测试函数（op、op_file、无条件 skip 标记、pytest collect 核实结果、op_name 提取） | 上游更新后重跑 `python ops/gen_inventory.py --verify-collect` |
| `shapes_single.yaml` | 每算子单 shape 的筛查集（537 个精选 op_name 键 + 上游 core_shapes 的类名键底座——上游 shape 回退链是 op_name → MRO 类名 → 基类默认，缺类名键会让 BLAS 族跌到 1 维默认值崩溃） | 手工增补；上游变动后原地刷新：`python ops/gen_inventory.py --migrate-shapes ops/shapes_single.yaml --shapes-out ops/shapes_single.yaml`（未匹配键落 `.unmatched.yaml` 供复核） |
| `shapes_multi.yaml` | 每算子多 shape 的深查集（结构同上） | 同上 |
| `dismiss.txt` | 批量排除清单（`[repo:]op` 每行一个），只放本仓复核确认的失败项并注明原因/日期 | 批量跑出 FAIL 并确认原因后手工添加 |

无 shape 条目的算子用上游默认 shape，照常可测。`--shape-file ops/shapes_multi.yaml` 切换深查集。

**执行语义**：每算子独立子进程（CUDA crash 只废单个算子）；`--op-timeout` 对进程组 SIGTERM→SIGKILL，记 `TIMEOUT` 不重试、批次继续；pytest rc=1 与信号杀最多重试 2 次，连续同 rc 视为确定性失败提前止损；`--dtypes` 有两级降级——遇上游 `can't be supported by this op` 即去掉限制重跑（备注 `dtype_fallback`），其余失败且无成功行时也去掉限制最后救一次（备注 `dtype_rescue`，覆盖 torch baseline 对受限 dtype 编译失败的场景）。runtime skipif 与上游 `xfail` 都判 SKIP（上游声明的无信号状态）；`.complete` 标记只给 PASS/SKIP，`--resume` 据此复用。

**产物**：`<batch>/<repo>/<op>/` 与单算子 run 目录结构完全一致（run.log、run_identity.json、autotune_records/、ttgir/……），所以任何一个算子都可以事后单独 `REPLAY_FROM` 复测。批量层额外产出 `ops_status.csv`（每算子一行：backend/device/timer、状态/耗时/行数/备注，边跑边原子更新；`gpu` 列与 `device` 同值）与 `summary.csv`（每测量行一行：backend/device/timer、dtype、latency、speedup、Size Detail，来源是 run.log 的结果表，`batch_summary.py` 也可单独对旧批次重跑）。备注列聚合 `no_cudagraph_rows`/`replay_fallback` 计数，解读口径时先看这列。不同后端、不同 timer 的行不要当成同口径。

## 计时口径：cudagraph 的预热、回退与精度

**capture 前的显式 warmup**：`do_bench_cudagraph` 自带的内部预热只有 5 次迭代，对首跑要触发 autotune 编译（尤其含 FlagTune 扩展空间）、libtuner 选择、lazy JIT 的算子远远不够——这些一次性开销若漏进被捕获的 graph 或第一个计时迭代，测出的 latency 会 run-to-run 抖动（M=1 多 kernel 路径最明显，实测首轮可低到稳态的 ~1/3）。`_cudagraph_plugin` 因此在 capture 前显式预热：先跑一次并丢弃（吸收 autotune/JIT 编译），再按调用方传入的 warmup 时间预算（`Config.warm_up`）循环稳态预热，然后才 capture。预热次数按稳态单次耗时换算，并夹在 5–200 次之间——亚毫秒 kernel 的实际预热时长因此低于名义预算（1000ms），实测足够；若换新算子仍见首轮抖动，优先调大 `_warmup_before_capture` 里的次数上限。这样首轮即稳态、多轮一致（实测同一 M=1 shape 两轮 speedup 差 <0.1%）。

部分算子本身不支持 CUDA graph capture（测量函数内含 host 同步、动态显存分配、autograd backward、不合法的流操作等），这类算子会自动回退到普通 do_bench 计时，`run.log` 中打 `BENCHMARK_DIRECT_NO_CUDAGRAPH` 标记（含失败阶段与具体原因；`phase=warmup` 表示预热阶段就失败了，并非 graph capture 被拒）。

回退路径的两个防御措施（都吃过亏）：（1）致命错误判定按错误原文匹配而非宽泛子串——新版 torch 给每个 CUDA 错误都追加 "enable device-side assertions" 提示语，宽松匹配会把所有 capture 失败误判为致命错误、令回退路径整体失效；（2）capture 在一次性 RNG state 替身下执行——`torch.cuda.graph` 会无条件注册默认 CUDA RNG 生成器，capture 中途失败可能让生成器卡在 capturing 态，此后进程内任何 `torch.randn` 都抛 "Offset increment outside graph capture"（毒化显形时机不定，事后修复不可靠），替身隔离让真身永不参与 capture，成败都原样归位。

回退本身不影响 A/B 公平性（两侧同一算子回退行为一致），但**回退口径的测量误差更大**：do_bench 每次迭代都走完整的 Python → launch 路径，kernel 越小，launch 开销和 CPU 侧抖动在数字里占比越高——亚毫秒级 kernel 上两种口径可差 2 倍以上，且行间波动更明显。解读这类算子的结果时：

- 小 shape 行的绝对值和小幅（<10%）差异不要过度解读，优先看大 shape 行；
- 需要更高置信度时，同一配置多跑几次取中位数，或设 `ZL_ENABLE_CUDAGRAPH=0` 统一用 `timer=direct`（消除同表混两种口径的问题）。

回退是按测量点发生的，同一份 latency 表里可能混有两种口径的行；若不确定，先 `grep BENCHMARK_DIRECT_NO_CUDAGRAPH run.log` 确认哪些行是回退口径再下结论。

## 插件说明

`-p` 列表由 `plugin_plan.py` 按 `ZL_BACKEND` 生成，不再在 `run_pytest.sh` 里写死。`_cudagraph_plugin` **不**出现在默认列表上：只有 CUDA 且 `_timing_plugin` 选中 Graph 时才会 import 它。

| 插件 | 作用 | 何时关闭 |
|-----|------|---------|
| `_device_setup_plugin` | 按后端设 `GEMS_VENDOR`（cuda→nvidia），跳过 flag_gems 启动时可能挂死的 `nvidia-smi` 探测，并写入 `run_identity.json`。旧名 `_device_guard_plugin` 仍是兼容 shim | 已设 `GEMS_VENDOR` 时不覆盖 |
| `_parallel_warmup_plugin` | 由 `PARALLEL_WARMUP_GPUS=N` 激活（未设则静默 no-op）：autotune sweep 分片到 N 卡并行，测量仍单卡串行、产物结构不变（见「加速：多卡并行 sweep」） | 不设该变量即关闭 |
| `_seed_plugin` | 固定 random/numpy/torch 种子，数据相关算子（sort/topk 等）输入逐字节一致 | 不关 |
| `_shape_inject_plugin` | 让 shape yaml 覆盖子类硬编码的 `set_shapes()` | 不关 |
| `_shape_iter_inject_plugin` | 覆盖在 `get_input_iter` 里硬编码 shape 的类（conv/pool 等） | 不关 |
| `_bespoke_shape_plugin` | 覆盖特殊输入构造的算子（upsample/flash_mla/cutlass 等，按类名注册） | 测这些算子之外可关 |
| `_autotune_record_plugin` | record/replay Triton autotune 选择；`autotune_records/meta.json` 记录 backend，跨后端 replay 拒绝 | `ZL_ENABLE_AUTOTUNE=0` |
| `_timing_plugin` | 按后端的 `supports_cuda_graph` 选计时路径：声明了的走 Graph（内部调用 `_cudagraph_plugin`，逐 shape 回退），没声明的一开始就是上游 `do_bench`（`timer=direct`） | `ZL_ENABLE_CUDAGRAPH=0` 强制 direct |
| `_cudagraph_plugin` | 仅由 `_timing_plugin` 在声明了 `supports_cuda_graph` 的后端上调用。`do_bench` → `do_bench_cudagraph`，capture 前按 warmup 预算显式预热；无法 capture 时回退并打 `BENCHMARK_DIRECT_NO_CUDAGRAPH` | 不要直接 `-p` 加到非 CUDA 运行 |
| `_musa_benchmark_compat_plugin` | 剥掉 FlagGems 上游若干 benchmark 上过时的 `torch.cuda.is_available()` 跳过门——被测代码本就经 `flag_gems.device` 分发，在 MThreads 上那些门把可用算子变成「未测量」。每条覆盖都钉死到一个文件与一条确切的 skip 理由，且仅在 `flag_gems.vendor_name == "mthreads"` 时生效 | 由 `backends/musa.py` 的 `compat_plugins` 声明，其他后端不加载 |
| `_pretty_report_plugin` | 结果表整理 + 着色：全表相同的输入折叠成表头下一行图例，每行 Size Detail 只留随行变化的部分（shape 保持 `torch.Size([...])` 原样，MoE 类算子单行从 ~400 字符缩到一屏内）；SUCCESS 绿 / FAILED 红。列名与 `SUCCESS` 字样保持上游原文，`run.log` 的 grep/解析不受影响 | 需要与上游原始表格逐字对齐时去掉 |
| `_ir_meta_plugin` | 编译期记录每个变体的 constexpr/签名/特化；launch 钩子统计每个 (kernel, shape) 的实际使用；退出时按上述结构 dump ttgir。Triton knobs 缺失时 fail-soft | `ZL_ENABLE_IR=0` |
| `_mm_cluster_fix_plugin` | Hopper fp16 mm cluster kernel 越界崩溃的运行时规避 | 默认关；`ZL_ENABLE_MM_CLUSTER=1` 或 `ZL_EXTRA_PLUGINS` |

## 常见问题

**Q: latency 和别人跑的差很多？**

先检查 GPU 是否被共占（`nvidia-smi`）；再确认对方是否开了 cudagraph——M=1 这类小 shape 下 launch 开销占比大，两种口径可差 2 倍以上，大 shape 基本一致。

**Q: 同一份 shape、什么都没改，两次跑出来的 speedup 不一样？**

正常，且幅度可能不小——record 模式每次进程都重新 sweep Triton autotune，计时噪声会让不同 run 选中不同 config。想让两次可比，B 侧必须 `REPLAY_FROM` A 侧的 record 目录；diff 两侧 `autotune_records/*.json` 可确认 config 是否真的一致。实测幅度与成因见「为什么需要 replay」。

**Q: 第一次跑某算子特别慢？**

两个来源：（1）开了 `USE_FLAGTUNE=1` 时，首跑要在 FlagTune 扩展空间做全量搜索（fused_marlin_moe_mxfp4 单进程可达数分钟甚至十几分钟，期间 GPU 满载、`run.log` 停在测试名不动属正常，不是卡死）；（2）libtuner 算子首跑的全量 sweep（mm 约 50 分钟）。winner 持久化到 `~/.flaggems/config_cache/` 后同 shape 秒级命中（但该缓存会因源码/开关/GPU 变化失效，失效后又需重搜）。默认 `USE_FLAGTUNE=0` 走普通 autotune，单 shape 通常 10 秒级出结果。

shape 多时这部分会主导墙上时间，可设 `PARALLEL_WARMUP_GPUS=N` 并行做 sweep，见「加速：多卡并行 sweep」。

**Q: `--warmup/--iter` 要设吗？**

不用。cudagraph 计时路径下，传入的 warmup 时间预算会被用来在 capture 前显式预热（先吸收 autotune/JIT 编译再稳态预热，稳定首轮，见「计时口径：cudagraph 的预热、回退与精度」），iter 默认 100ms 预算按 kernel 耗时自适应换算次数。

**Q: ttgir 目录里某个 shape 少了文件？**

看 `index.tsv` 的 `launches` 列——没在该 shape 下真正启动过的变体不落盘。`run.log` 里的 `BENCHMARK_DIRECT_NO_CUDAGRAPH` / `AUTOTUNE_REPLAY_FALLBACK` 标记可解释异常回退。

## 依赖假设

- FlagGems benchmark 体系（`benchmark/base.py` 的 `Benchmark` 类、conftest 的 `--shape_file/--level/--mode` 选项）；主仓与 FlagGems-vllm 仓库同构，`_autotune_record_plugin` 会按实际被 import 的包（`flag_gems` / `flaggems_vllm`）挂 LibTuner 补丁；
- 设备后端由 `ZL_BACKEND` / `--backend` 选择（`cuda` / `musa` / `rocm` / `ascend`）。inventory 仍是 FlagGems 逻辑算子清单，同一 case 可在多个后端上跑，不要按硬件复制 CSV；
- Triton 需支持 `knobs.compilation.listener`、`kernel_load_end_hook`、`launch_enter_hook`（当前 FlagTree 的 triton 3.6 满足）。缺 hook 时 IR / autotune 插件 fail-soft，不会把进程打挂；
- 插件通过 monkeypatch 挂钩上游内部结构，FlagGems/Triton 大版本升级后若行为异常，优先检查各插件 pytest_configure 输出的注册日志是否还正常打印。
- 无 GPU 的适配器自检：`python test_backends.py`。