Files
zl_bench/README.md
T
Aharrypotter 2fbfe1dd16 bench: make screening a first-class step with the same caliber as the A/B
Six rounds of A/B cost 13-44 minutes per operator on S5000 (rand, 15 shapes:
13.0 min; sort, 25 shapes: 43.4 min). Screening one costs about 0.8 minutes —
21 operators in 15.0 minutes, which knocked 13 of them out of the expensive
queue and cut it from ~7 hours to ~2.7. That ratio is the whole argument for
having a separate first step.

Screening was being done by ad-hoc scripts living in the results tree, and
those got two things wrong that this codebase already knows better than:

Both sides ran `record`. Each swept autotune independently, so the two sides
compiled different configs and produced different IR for reasons that have
nothing to do with the switch — every operator would screen as "changed".
screen.sh pins configs exactly as round 1 of ab_repeat does: A records, B
replays A.

The criterion was the convert_layout count. Equal counts do not mean equal
generated code, and what the A/B measures is whether the generated code
changed. The ttgir fingerprint decides; the cvt counts are still recorded,
because they explain what kind of change it was.

Four verdicts, and three of them are not "the switch does nothing":
IR_CHANGED (measure it), NO_IR_DIFF (skip it), NO_KERNEL (the operator never
reached Triton on this backend), NO_ROWS (it did not run — usually the dtype).
The last two are runnability problems being mistaken for switch results, which
is how `fft` got written off: the screen forced DTYPES=float16 on every
operator and fft only accepts complex64. OPS entries now carry an optional
per-operator dtype and default to the operator's own set.

ir_fingerprint moves to _ir_fingerprint.sh so screening and the A/B cannot
drift apart on the criterion.
2026-09-10 16:49:32 +08:00

42 KiB
Raw Blame History

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,见「批量多算子测试」),可对双仓库上千个算子做例行体检,产物结构与单算子完全同构。

快速开始

# 默认算子(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 名:

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 后仍会落盘已编译部分。

两步:先筛选,再实测

跑一个算子的六轮 A/B 要 1344 分钟(S5000 实测:rand 15 shape 13.0 分钟, sort 25 shape 43.4 分钟)。筛一个算子只要约 0.8 分钟(21 个算子一遍 15.0 分钟)。 **差 2050 倍**,所以先筛后测:

# 第一步:筛掉不值得排队的算子
SWITCH=FLAGTREE_MUSA_RLC_ENHANCE ZL_BACKEND=musa FLAGGEMS_DIR=/path/to/FlagGems   OPS="col2im:col2im topk:topk perf_fft:fft" bash screen.sh

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

screen.sh 对每个算子给四种判定之一:

判定 含义 该怎么办
IR_CHANGED 两侧 ttgir 不同 排六轮 A/B
NO_IR_DIFF 逐字节相同(config 已锁死) 跳过。跑六轮只会测到噪声,并报成一个"无差异"结论
NO_KERNEL 跑出了行,但没编出 .ttgir 该算子在本后端没走 Triton,编译器开关无从谈起。不是"开关无作用"
NO_ROWS 没有测量行 先修可运行性:多半是 dtype 不被该算子接受,或该算子在本后端拿不到 config

三件容易做错的:

  • 筛选也是 A/B,口径一样严。 A 侧 record、B 侧 REPLAY_FROM 它。两侧各自 record 会各自 sweep 出不同 config,IR 自然不同,于是每个算子都判成 IR_CHANGED——筛了等于没筛。
  • 判据是 ttgir 指纹,不是 convert_layout 计数。 cvt 数相同不等于生成代码 相同,而要测的是"生成代码变了没有"。cvt 数照样记,它解释的是变化的性质。
  • 不要全局写死一个 DTYPES。 OPS 元素支持 OP:OP_FILE:DTYPES,留空表示用 该算子自己的全套。fft 只吃 complex64,被一刀切的 DTYPES=float16 判死过一次 ——那是 harness 的问题,被记成了算子的问题。

A/B 对比测试标准流程

# 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),是收益本身落在零附近、单次 测量给不出方向。

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:

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 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 卡并行,再交回单卡串行测量:

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,多卡分片、每卡内部串行,产出逐算子状态表和逐测量行总表。定位是筛查口径——跨卡并行测量存在功耗/散热耦合噪声(幅度可到百分之几),发现可疑算子后回单算子串行流程(前几节)确认,不要直接拿批量数字下精细结论。

# 全量筛查(双仓库全部可收集算子,排除带无条件 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。