Aharrypotter/zl_bench
1
0 Fork 0
forked from zhoulin/zl_bench
Code Pull Requests Activity

Compare commits

base: Aharrypotter/zl_bench:feature/musa-rlc-product-harness
Aharrypotter/zl_bench:main Aharrypotter/zl_bench:feature/musa-benchmark-support Aharrypotter/zl_bench:feature/musa-rlc-product-harness zhoulin/zl_bench:main
..
compare: Aharrypotter/zl_bench:c95474a432725e71cb2e19339060547ca5f90015
Aharrypotter/zl_bench:main Aharrypotter/zl_bench:feature/musa-benchmark-support Aharrypotter/zl_bench:feature/musa-rlc-product-harness zhoulin/zl_bench:main
13 Commits
feature/musa-rlc-product-harness .. c95474a432
Author SHA1 Message Date
AharrypotterandClaude Opus 5 c95474a432 bench: refuse to report a comparison whose two sides have different calibers 
Two runs on different calibers produce a report that looks exactly like a
valid one — no step errors, the numbers just mean something else. This repo
has published three figures that were an artifact of such a difference:
router_gemm -8.2% (record vs record), mm "945 variants -> 33" (shared vs
inherited libtuner cache), replication_pad3d +43% vs +124.5% (timer).

batch_summary --ab / --ab-rounds now reads every participating run's
run_identity.json and compares backend/device/timer/bench_level/shape_mode/
tuner_cache/direct_repeat before parsing a single latency. Mismatch, or a
missing identity, prints !!! [caliber] and exits 4. autotune_mode is judged
separately: record pinning a replay is the normal shape, both sides record
is the exact trap above and is refused by name. --allow-caliber-mismatch
releases it, still printing what it found.

Recording on a *shared* tuner cache warns rather than refuses: the pinned
config depends on what else ran on the machine, which leaves relative
conclusions intact and absolute latencies unquotable. Refusing it would
reject the default every A/B here currently uses.

ab_repeat.sh closes the one axis that had no check: after round 1 it
fingerprints both sides' ttgir (relative path + content, so a changed
variant *set* counts too). Byte-identical means the switch changed no
generated code for this op, and the remaining N-1 rounds would measure
noise and report it as "no difference" — so the op is skipped with
SWITCH_NO_IR_DIFF unless SWITCH_ALLOW_NO_IR_DIFF=1.

README leads with a caliber table: each axis, its measured magnitude, what
pins it, and which code checks it. The fourth column is code because
documentation did not hold — the README already said "the B side must use
REPLAY_FROM" when the record-vs-record number was published.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-10 14:32:49 +08:00
AharrypotterandClaude Opus 5 aa147c938e feat: repeat an A/B over rounds and report whether the verdict holds 
A single pair cannot say whether its own verdict is stable, and on this device
that is not a theoretical gap. Measured on an S5000, topk eliminated 32
convert_layout ops yet 17 of its 42 shapes changed the SIGN of their delta
between two equally valid runs. The cause is not the timer: graph capture and
direct timing agree to 1.00x on that operator, and the verdict moved a median of
1.1pp between them, inside the 1.7pp noise floor a same-timer control
established. The deltas are simply near zero, and one measurement cannot orient
them.

ab_repeat.sh runs an A/B over N rounds in independent processes. Round 1's A
side records the autotune configs and every later run -- including A's own later
rounds -- replays that one set, so configs are pinned across rounds and sides and
only the switch under test and measurement noise remain. Round order alternates
to cancel one-directional drift. An operator with a failed round gets no report,
since a report missing rounds reads exactly like a complete one.

batch_summary --ab-rounds aggregates <root>/r<k>_<A|B>/run.log and separates the
two dispersions, which answer different questions and were previously conflated:
the spread within one side's samples says how noisy the machine was, while the
spread of the per-round deltas says whether the verdict holds. Shapes whose
delta changes sign across rounds are counted and called out, because their
direction is not a conclusion at any timer.

The path is already validated on NVIDIA: six rounds narrowed router_gemm to
-16.6% ~ -16.2% with 0.4% side noise, which is what turned a single pair into a
statement.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-10 13:11:37 +08:00
AharrypotterandClaude Opus 5 9665d36a7c feat: allow the direct timer to repeat and take the median 
do_bench already returns a median over the iterations inside one call, which
suppresses iteration-level jitter. It cannot see drift across a call: clocks
ramping, another process arriving. Direct timing is where that shows, because
it walks the full Python launch path once per iteration, so on small kernels the
number carries launch overhead and CPU jitter that graph capture removes.
ZL_DIRECT_REPEAT=N repeats the call and takes the median over calls, which is
the layer the inner median misses. It defaults to 1, so nothing changes unless
asked.

The knob covers both ways a row ends up on the direct timer: a run that selected
timer=direct outright, and the rows a captured run falls back on, whose marker
now carries repeat=N since those rows were already a different caliber from the
captured ones and can now also be a different sample count. Repetition is
skipped when the caller asks for quantiles or return_mode "all", because the
result is then a sequence and a median across calls is undefined for it.

On MThreads this is not a refinement. A convert-reduction screen over 21
operators puts the largest reductions at sort (65->15), dropout (48->0) and
rand (24->0), and all three are exactly the operators that cannot be
graph-captured there -- RNG seeding and a capture-illegal stream op. The
headline numbers on that device are produced by the noisiest timer, which until
now had no repetition at all.

_cudagraph_plugin imports triton at module scope, so the tests load it against a
stub: the repeat logic is arithmetic and should stay testable without a GPU.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-10 12:56:08 +08:00
AharrypotterandClaude Opus 5 d7cb7efc94 fix: have the campaign worker ask zl_bench for its plugin list 
The worker carried a hand-copied plugin array, taken from zl_bench before
backends/ existed, and it had drifted twice over. It still named
_device_guard_plugin, which is now a shim for _device_setup_plugin. More
consequentially it carried no timing plugin at all, so every measurement this
campaign ever sealed was timer=direct while zl_bench's own driver had been
capturing -- a difference the harness never reported, because the missing plugin
cannot announce itself.

It now calls plugin_plan.py with ZL_BACKEND=musa and uses whatever comes back,
so the caliber follows the backend's declared capabilities instead of a snapshot
of them. The plan is preflighted for non-emptiness, since an empty array would
silently run pytest with no plugins at all: no shape injection, no autotune
record, no IR dump, and a table that still looks well-formed.

Two effects on this campaign. _timing_plugin now enters the list, and since
MUSA declares graph capture the worker will capture where the operator allows
it; operators that cannot are marked per row, as on any other backend. And
_parallel_warmup_plugin joins, inert unless PARALLEL_WARMUP_GPUS is set.

The sealed 32-workload run therefore sits on the far side of a caliber change
and is not comparable with anything produced after this commit. Re-running it is
a campaign decision, not a harness one.

The worker test asserted the literal "-p _musa_benchmark_compat_plugin"; it now
asserts the delegation and that no -p is spelled out locally. What the plan
holds for musa is covered where it belongs, in zl_bench's test_backends.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-10 10:34:05 +08:00
AharrypotterandClaude Opus 5 75405cc555 Merge the S5000 MUSA RLC product campaign 
Brings in campaigns/musa_rlc_product/: the runner, its worker, manifest
validation, the sealed 32-workload manifest, a topology re-accounting tool and
their tests. Developed on a fork branch that never reached main, and it merges
without conflict because it only adds a new directory; the one file it shares
with main, _musa_benchmark_compat_plugin.py, is byte-identical to the copy
already committed here.

The campaign layer sits above zl_bench rather than inside it: freezing a
workload set, classifying outcomes, and producing receipts, where zl_bench turns
one operator into one reproducible measurement. Two of its ideas are worth
naming because they are absent below it. Narrowing what a workload measures --
which correctness variant, which benchmark case -- requires a stated reason that
manifest validation enforces, so an exclusion stays visible in review. And
reaccount_musa_rlc_topology.py re-derives verdicts from sealed artifacts under a
corrected accounting rule without re-running anything, while refusing to promote
a workload on the new count, only to demote it.

Its worker still hardcodes a plugin list copied from a pre-backends zl_bench,
which is why the campaign has been timing with timer=direct. Switching it to
plugin_plan.py is the next change and will alter its caliber, so the sealed run
is not comparable across that boundary.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-10 10:32:33 +08:00
AharrypotterandClaude Opus 5 29af59be9b feat: absorb the MThreads compat plugin and the A/B range report 
Three pieces from the fork's feature/musa-benchmark-support branch that never
reached main, rebased onto the backend abstraction rather than copied verbatim.

_musa_benchmark_compat_plugin drops FlagGems benchmark skip guards written as
`torch.cuda.is_available()` where the code under test dispatches through
flag_gems.device. On MThreads those turn a working operator into "not measured":
test_fft.py skipped entirely, and unskipped on an S5000 it produces 25 rows. The
plugin pins each override to one file and one exact skip reason and only acts
when flag_gems has dispatched to the mthreads vendor; two guards were probed on
device and deliberately left in place, which is why they are named in its
docstring rather than silently absent.

It is wired through a new DeviceBackend.compat_plugins rather than a name test
in the planner, so the policy stays where every other backend difference lives.
Compat plugins are planned first because they rewrite collection-time markers
and nothing may read the item list before them; a test pins that position and
another pins that they do not leak to other backends.

batch_summary gains --ab: the speedup column compares gems against torch, not
one build against another, so an A/B needs the two run.logs paired by operator,
dtype and size detail. Only rows that succeeded on both sides with a positive
latency count, so an incomplete pairing cannot look like a complete range, and
the range is annotated when either side carries a BENCHMARK_DIRECT_NO_CUDAGRAPH
marker, since a fallback row was timed under a different regime.

ab_rlc_test.sh is the same caliber as ab_fold_test.sh with the comparison axis
moved to FLAGTREE_<VENDOR>_RLC_ENHANCE. Both are caliber examples, not
regression scripts: if the switch under test no longer exists the script still
runs every round and prints a full table showing no difference, so confirm the
axis is live before reusing either.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-10 10:31:59 +08:00
AharrypotterandClaude Opus 5 c9f9b7885f feat: check the run's preconditions, and record which tuner cache it used 
Two failure modes that cost this harness real results, both of them silent.

A wrong path does not fail as a wrong path. pytest cannot collect a benchmark
directory that is not there, so benchmark/conftest.py never registers --level
and --mode, and the run dies with "unrecognized arguments: --level core --mode
kernel" -- an error about the flags, several layers away from the missing
directory. run_pytest.sh now checks the FlagGems directory, the benchmark file,
SHAPE_FILE and REPLAY_FROM before starting pytest, prints
"PREFLIGHT_FAILED check=<the named check>" and exits 3, distinct from pytest's
own 1 and 4 so the batch layer can tell a broken environment from a broken
operator. Only the file half of a node is checked; a missing ::test_<fn> already
produces a clear "not found: ..." from pytest and needs no help.

FlagGems keeps libtuner winners in a sqlite DB under $FLAGGEMS_CACHE_DIR
(default ~/.flaggems), shared by every run on the machine. Whichever side of an
A/B runs first therefore hands its winners to the second, which then compiles
only those: measured on router_gemm, the second of two shared-cache runs
compiled 10 variants instead of 234 and finished in 17s instead of 52s. That is
how an mm A/B in this repo produced "945 variants -> 33, 136 converts -> 10" and
read as a compiler result when it was cache inheritance; isolated, both runs
compile 234.

ZL_ISOLATE_TUNER_CACHE=1 gives a run its own cache directory, dropped afterwards
with the Triton one since the chosen configs are already in autotune_records/.
It defaults to 0, because cold tuning is expensive and the standard REPLAY_FROM
A/B pins configs anyway; only record-against-record is exposed. The default is
unchanged, but the caliber is no longer invisible: either way it goes into the
run header and run_identity.json as tuner_cache, which is what was missing when
the mm numbers could not be explained after the fact.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-10 10:31:36 +08:00
AharrypotterandClaude Opus 5 ac818291a6 feat: turn on graph capture for MUSA 
supports_cuda_graph was False for MUSA, so every run there was pinned to
timer=direct. That was correct when it was written -- the capture path named
torch.cuda and could not have worked -- but it is not correct now, and the two
timers are not interchangeable: direct timing walks the full Python launch path
per iteration, so on small kernels it measures mostly launch overhead. Measured
on an S5000, capture against direct: batch_norm 16x16x8x48 0.0758ms -> 0.0089ms,
tanh and gelu at 4096x4096 about 3.3x, GEMMs unchanged at 1.00x. Leaving it off
discards that on every operator that can capture, and it also puts MUSA on a
different caliber from the CUDA numbers the harness is meant to compare against.

torch_musa supplies Stream, Event, synchronize, graph and MUSAGraph, so capture
works through the device interface. Verified on an S5000 over five operators,
one per path:

  tanh      captured, no marker
  dropout   15/15 rows fall back: MUSAGeneratorImpl::current_seed is not
            callable during capture
  rand      same
  perf_sort 25/25 rows fall back: operation not permitted when stream is
            capturing
  unique2   denied up front; rc=0 with 10 rows

MUSA gets its own error wording; none of the CUDA markers match it. "MUSA
error: unknown error" during capture is classified as poisoning rather than
recoverable, because the fallback runs and the process then dies in teardown:
before the deny entry was keyed correctly, unique2 reproduced exactly that,
ending at rc=1 with no rows, but now under an explicit BENCHMARK_CAPTURE_POISONED
marker instead of a bare ExceptionGroup.

capture_deny_ops is keyed by FLAGGEMS_PERF_CURRENT_OP, which holds the pytest
FUNCTION name minus `test_` -- unique2, perf_sort -- not the benchmark file
name. Getting that wrong fails open silently, which is how the first attempt
still hit rc=1 with the entry written as "unique"; a test now pins the key
space. The plugin-plan test is renamed because its assertion that
_cudagraph_plugin is absent from the plan holds on CUDA too and says nothing
about capture support.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-09 17:56:56 +08:00
AharrypotterandClaude Opus 5 6fa5f5d7a2 refactor: let any backend own graph capture and its failure modes 
Two things kept graph-capture timing tied to NVIDIA, and both are mechanism
rather than policy, so neither backend's behaviour changes here.

The capture path named torch.cuda directly for stream, event, synchronize and
the graph type. FlagTree's do_bench already resolves those through
driver.active.get_device_interface(); its do_bench_cudagraph does not. On a
non-CUDA device every capture therefore raised and every row silently took the
fallback while the log still reported capture as enabled. The path now takes
those calls from the interface, with _graph_class picking CUDAGraph or
MUSAGraph. CUDA still calls upstream's do_bench_cudagraph verbatim, so its
numbers do not move; the port only matters where upstream cannot run.

Failure classification moves onto DeviceBackend, because the text it matches is
vendor wording that has no meaning on another runtime, and gains a third
verdict. A kernel refusing capture is ordinary and falls back. A kernel failing
on the device is fatal and re-raises, since a fallback latency there would
report a broken kernel as a measurement. New: the capture ATTEMPT can leave the
context unusable, and falling back then measures inside a dead process -- the
observed shape of that is a fallback marker followed by a teardown crash and
zero rows. Those raise under BENCHMARK_CAPTURE_POISONED instead, naming the
operator as capture-incompatible rather than broken. capture_deny_ops skips the
attempt entirely for operators already known to do this, because for them the
attempt is the damage and rediscovering it per shape costs the run.

CUDA's marker list is carried over unchanged and no poisoning wording is
guessed for it. The scratch-RNG guard around capture is skipped when an
interface exposes no default_generators, since there is then no registration to
protect.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-09 17:56:32 +08:00
AharrypotterandClaude Opus 5 783c0655bc fix: never let the parallel sweep cover only part of the shape set 
The sharded autotune sweep could publish a config set with holes, and a run
using it reports itself as replaying while every uncovered key silently
self-tunes. That is the exact failure the record/replay caliber exists to
prevent, and it was reachable two ways.

Shape lookup was op_name-only, but BLAS operators keep their shapes under the
class name: core_shapes.yaml has BlasBenchmark, not mm. Every BLAS operator
therefore resolved to zero shapes and the plugin disabled itself with a message
that read like a normal skip. It now walks the benchmark class chain the way
FlagGems' set_shapes does. The chain is recovered statically -- the class comes
from the op_name= keyword of the constructor call in the test module, its bases
from that module plus benchmark/base.py -- because the plugin runs at
pytest_configure, before any Benchmark instance exists. Resolved keys are named
in the log, and a total miss lists the candidates it tried.

Two cases can still produce partial coverage, and both now refuse rather than
shard a subset:

- --level comprehensive, where set_shapes() merges set_more_shapes() on top of
  the yaml. That addition is arbitrary per-subclass Python and invisible here.
  Measured on mm: 5 of 29 shapes swept, 90 AUTOTUNE_REPLAY_FALLBACK
  reason=key_missing in the run that consumed the result.
- a shard exiting non-zero, whose shapes were never swept. This used to append a
  WARNING to a success message and publish the merge anyway; observed with 3 of
  5 shards dead. The whole sweep is now discarded, worker logs kept.

Both fall back to the normal serial sweep, which costs wall time rather than the
same-config premise.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-09 16:40:12 +08:00
AharrypotterandClaude Opus 5 92dae82c7f feat: expose the upstream benchmark level as BENCH_LEVEL 
run_pytest.sh hardcoded --level core, so the only shapes ever reachable were
each operator's core_shapes entry. FlagGems merges a subclass's
set_more_shapes() on top of that at --level comprehensive, and for BLAS that
addition is where the large GEMMs live: mm goes from 5 shapes to 29, taking in
four large-K tuples plus models_shapes/*.yaml, with N/K up to 152064. The
comprehensive path also yields each shape a second time (transposed, for BLAS),
so the case count doubles. None of that was reachable without editing the driver.

BENCH_LEVEL defaults to core, so existing runs are unchanged, and rejects
anything other than core/comprehensive rather than passing it through to pytest.
The chosen level is echoed in the run header, written into the archived
shapes.yaml, and recorded in run_identity.json, because an A/B whose two sides
used different levels is not comparable and is otherwise invisible afterwards.

Sweep cost scales with the shape count, so comprehensive on a BLAS operator
wants PARALLEL_WARMUP_GPUS.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-09 16:39:45 +08:00
AharrypotterandClaude Opus 5 c8ea586d0f feat: make the harness backend-agnostic 
Everything about which accelerator a run targets was previously spelled CUDA:
device discovery went through nvidia-smi, pinning through CUDA_VISIBLE_DEVICES,
the vendor guard only recognised NVIDIA, and the -p plugin list was hardcoded in
run_pytest.sh with no way to drop the CUDA-only members. Running the same
operator corpus on another vendor's device therefore meant editing the drivers.

The selection now hangs off one ZL_BACKEND variable (cuda/musa/rocm/ascend, with
the usual aliases):

- backends/ owns per-backend device enumeration, visibility env keys and
  capability flags, so nothing above it names a vendor;
- plugin_plan.py derives the -p list from those capabilities instead of a
  hardcoded array, which is what lets a non-CUDA run omit the graph-capture and
  cluster-fix plugins rather than watch them fail;
- _timing_plugin selects the timer from the same capabilities: CUDA keeps
  graph capture with its per-shape fallback, other backends stay on upstream
  do_bench and record timer=direct, because capture-then-fallback on a device
  without graph support only produces noise;
- _device_setup_plugin replaces the NVIDIA-only vendor guard;
- providers/flaggems.py resolves the pytest node, so a caller may pass TEST_NODE
  directly instead of the OP/OP_FILE pair, and SHAPE_MODE=upstream can defer to
  FlagGems' own core_shapes;
- _run_identity.py writes run_identity.json next to each run, recording backend,
  device, timer, autotune mode and the resolved test node. Two runs whose
  calibers differ are otherwise indistinguishable after the fact.

run_batch gains --backend and --devices (--gpus kept as an alias) and carries
backend/device/timer into ops_status.csv. test_backends.py covers the registry,
alias resolution, visibility env and plugin planning.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-09 16:39:14 +08:00
AharrypotterandClaude Opus 5 35a928a516 chore: pin shipped text formats to LF 
Everything in this repository runs on Linux accelerator hosts, and campaign
receipts record the SHA-256 of the harness files that produced them. With Git's
default autocrlf on Windows a checkout rewrites the shell drivers with CRLF,
which breaks them outright on the target host and changes every recorded digest.

The two inventory CSVs are renormalized here because their content is unchanged;
the remaining files normalize as part of the commits that actually modify them.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
 2026-09-09 16:38:33 +08:00
28 changed files with 2894 additions and 256 deletions
Expand all files Collapse all files

No files matched your search

README.md
+153 -20
View File
@@ -6,6 +6,30 @@
2. **可解释**：自动按 shape 收集每次运行实际使用的 ttgir（带可读的变体命名），供 IR 级 diff； 2. **可解释**：自动按 shape 收集每次运行实际使用的 ttgir（带可读的变体命名），供 IR 级 diff；
3. **口径统一**：cudagraph 计时消除 launch 开销，小 kernel 的对比不被 CPU 侧噪声淹没；capture 前按 warmup 预算显式预热吸收 autotune/JIT，首轮即稳态、多轮一致。 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`，见「批量多算子测试」），可对双仓库上千个算子做例行体检，产物结构与单算子完全同构。 单算子精测流程之上另有批量筛查层（`run_batch.py`，见「批量多算子测试」），可对双仓库上千个算子做例行体检，产物结构与单算子完全同构。
## 快速开始 ## 快速开始
@@ -20,6 +44,11 @@ OP=softmax OP_FILE=softmax SHAPE_FILE=my_shapes.yaml bash run_pytest.sh
# 测 FlagGems-vllm 仓库的算子：换 FLAGGEMS_DIR 即可（插件自动适配其包名 flaggems_vllm） # 测 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_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）。 上游现为两个仓库，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）。
@@ -35,7 +64,7 @@ softmax:
上面这条命令是 **record 模式**：现场 autotune、把选中的 config 存档。它用于单次摸底或给 replay 提供基准，**两次 record 的数字不能互相比**——做对比走下面的「A/B 对比测试标准流程」。 上面这条命令是 **record 模式**：现场 autotune、把选中的 config 存档。它用于单次摸底或给 replay 提供基准，**两次 record 的数字不能互相比**——做对比走下面的「A/B 对比测试标准流程」。
跑之前先 `nvidia-smi` 确认目标卡空闲：共享机器上别的任务会把两侧一起等比拖慢，出一份看似自洽实则作废的数据。多卡机器用 `CUDA_VISIBLE_DEVICES=<空闲卡号>` 选卡。**测量始终单进程单卡串行**，这是计时可比的前提；多卡仅用于可选的 sweep 加速（见「加速：多卡并行 sweep」）。 跑之前确认目标卡空闲：共享机器上别的任务会把两侧一起等比拖慢，出一份看似自洽实则作废的数据。用 `ZL_BACKEND` 选运行时（默认 `cuda`），再用该后端的可见性变量选卡（CUDA 为 `CUDA_VISIBLE_DEVICES`，也可设 `ZL_DEVICE`）。**测量始终单进程单设备串行**，这是计时可比的前提；多卡仅用于可选的 sweep 加速（见「加速：多卡并行 sweep」）。非 CUDA 后端默认走上游 `do_bench`，不会尝试 CUDA Graph。
脚本内固定了 `--level core --mode kernel`，其余可调项都走环境变量（见下）；需要改动这两个口径本身时直接编辑 `run_pytest.sh` 中的 pytest 命令行。 脚本内固定了 `--level core --mode kernel`，其余可调项都走环境变量（见下）；需要改动这两个口径本身时直接编辑 `run_pytest.sh` 中的 pytest 命令行。
@@ -45,10 +74,14 @@ softmax:
| 变量 | 默认 | 说明 | | 变量 | 默认 | 说明 |
|-----|------|------| |-----|------|------|
| `OP` | `fused_marlin_moe_mxfp4` | 测试函数名（`test_` 之后的部分） | | `OP` | `fused_marlin_moe_mxfp4` | 测试函数名（`test_` 之后的部分）；设了 `TEST_NODE` 时由节点覆盖 |
| `OP_FILE` | `fused_marlin_moe` | benchmark 文件名（`test_` 与 `.py` 之间的部分） | | `OP_FILE` | `fused_marlin_moe` | benchmark 文件名（`test_` 与 `.py` 之间的部分）；设了 `TEST_NODE` 时由节点覆盖 |
| `SHAPE_FILE` | 空（用脚本内置 yaml） | shape yaml 路径 | | `TEST_NODE` | 空 | FlagGems pytest 节点（`…/benchmark/test_<op_file>.py::test_<op>`）。设了且未指定 `SHAPE_FILE` 时默认不传 `--shape_file` |
| `FLAGGEMS_DIR` | `/workspace/dev/FlagGems` | FlagGems 仓库路径（测 vllm 仓库时指向 `/workspace/dev/FlagGems-vllm`） | | `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`；批量层绑卡后也会设 |
**怎么测** **怎么测**
@@ -56,8 +89,16 @@ softmax:
|-----|------|------| |-----|------|------|
| `REPLAY_FROM` | 空（record 模式） | 指向某次历史 run 目录，replay 其 autotune 选择。做 A/B 时 B 侧必须设（见「A/B 对比测试标准流程」） | | `REPLAY_FROM` | 空（record 模式） | 指向某次历史 run 目录，replay 其 autotune 选择。做 A/B 时 B 侧必须设（见「A/B 对比测试标准流程」） |
| `USE_FLAGTUNE` | `0` | `0` 走普通 autotune（跳过搜索、快速验证）；`1` 用 FlagTune 扩展调优空间（首跑全量搜索、慢）。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` | | `DTYPES` | 空（上游默认 dtype 扫描） | 空格分隔的 dtype 白名单（如 `"bfloat16 float16"`），逐个转为上游 `--dtypes`。算子不支持时上游报 `can't be supported by this op` |
| `PARALLEL_WARMUP_GPUS` | 空（串行） | 设为 `N`（≥2）时把 autotune sweep 分片到 N 卡并行，测量仍单卡串行；不设或 `<2` 则完全串行（见「加速：多卡并行 sweep」） | | `PARALLEL_WARMUP_GPUS` | 空（串行） | 设为 `N`（≥2）时把 autotune sweep 分片到 N 卡并行，测量仍单卡串行；不设或 `<2` 则完全串行（见「加速：多卡并行 sweep」） |
| `ZL_ENABLE_CUDAGRAPH` | `auto` | `auto`/`1` 仅在 CUDA 上走 Graph 计时；非 CUDA 即使设 `1` 也拒绝并保持 `timer=direct` |
| `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` 列表的插件名 |
**输出** **输出**
@@ -72,9 +113,13 @@ softmax:
``` ```
runs/softmax_20260716-031752/ runs/softmax_20260716-031752/
├── run.log # 完整日志（含 SUCCESS 行的 latency/speedup 表；已去 ANSI 色的纯文本） ├── 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（存档） ├── shapes.yaml # 本次实际使用的 shape（存档）
├── autotune_records/ # Triton autotune 选中的 config（供 replay） ├── autotune_records/ # Triton autotune 选中的 config（供 replay）
│ └── softmax.json │ ├── softmax.json
│ └── meta.json # provider/backend；跨后端 replay 会被拒绝
└── ttgir/ # 按 shape 分组的 IR 落盘 └── ttgir/ # 按 shape 分组的 IR 落盘
├── 1024x1024/ ├── 1024x1024/
│ └── softmax_kernel_inner/ │ └── softmax_kernel_inner/
@@ -131,7 +176,78 @@ libtuner 的持久缓存何时失效：FlagGems kernel 源码改动、tune_confi
同一层缓存也决定了 `_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 那几个。 同一层缓存也决定了 `_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 那几个。
### ab_fold_test.sh：A/B 测试参考示例 ### 前置检查
`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: ...`。
### 口径核对：出报告前的拒绝
`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 轮。
第 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 轮： `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 轮：
@@ -140,6 +256,8 @@ bash ab_fold_test.sh
# 汇总 log：runs/ab_fold_<时间戳>.log（已去色）；快速对账：grep -E '^#####|FAILED' <log> # 汇总 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 集、逐对成组——对比其它开关 / 编译器改动 / 算子，改循环变量与环境变量即可，口径部分不用动。 可复用的是**口径**——同 config replay、cudagraph 计时、真实 trace shape 集、逐对成组——对比其它开关 / 编译器改动 / 算子，改循环变量与环境变量即可，口径部分不用动。
但**它是口径示例，不是可长期直接执行的回归脚本**：对比轴是 FlagGems 侧的算子开关，会随上游演进被改名、改语义或移除，本仓库不跟随同步更新（该脚本引用的开关目前就已被上游移除）。失效后果是**静默的**：脚本照样跑完全部轮次、照样输出完整对比表，只是两侧执行同一份代码，得出"无差异"的假结论。所以套用前先确认对比轴仍有效——`grep -rn '<开关名>' $FLAGGEMS_DIR/src` 要有命中，且两侧 latency 表确有差异。 但**它是口径示例，不是可长期直接执行的回归脚本**：对比轴是 FlagGems 侧的算子开关，会随上游演进被改名、改语义或移除，本仓库不跟随同步更新（该脚本引用的开关目前就已被上游移除）。失效后果是**静默的**：脚本照样跑完全部轮次、照样输出完整对比表，只是两侧执行同一份代码，得出"无差异"的假结论。所以套用前先确认对比轴仍有效——`grep -rn '<开关名>' $FLAGGEMS_DIR/src` 要有命中，且两侧 latency 表确有差异。
@@ -152,7 +270,7 @@ shape 多时墙上时间几乎全在 autotune sweep，而不在测量——被
PARALLEL_WARMUP_GPUS=8 SHAPE_FILE=shape/DeepSeek-V4-Flash-p32768d1024.yaml bash run_pytest.sh PARALLEL_WARMUP_GPUS=8 SHAPE_FILE=shape/DeepSeek-V4-Flash-p32768d1024.yaml bash run_pytest.sh
``` ```
加速比取决于 sweep 占比：shape 多、config 空间大的算子收益最明显，可达数倍。不设该变量（或设 <2）时插件静默 no-op，常驻 `PLUGINS` 数组即可。**产物结构、`REPLAY_FROM` 用法、latency 表格式全不变**——合并后的 config 就写进本次 run 的 `autotune_records/<op>.json`，分片临时目录跑完即删。 加速比取决于 sweep 占比：shape 多、config 空间大的算子收益最明显，可达数倍。不设该变量（或设 <2）时插件静默 no-op。**产物结构、`REPLAY_FROM` 用法、latency 表格式全不变**——合并后的 config 就写进本次 run 的 `autotune_records/<op>.json`，分片临时目录跑完即删。
**测量本身不并行**：N 个进程压满同机多卡会通过功耗墙/散热耦合，单卡 latency 被邻居拖慢且不可复现。并行只用于"决定哪个 config 胜出"，测量仍是单进程单卡、与不开插件走同一条路径。 **测量本身不并行**：N 个进程压满同机多卡会通过功耗墙/散热耦合，单卡 latency 被邻居拖慢且不可复现。并行只用于"决定哪个 config 胜出"，测量仍是单进程单卡、与不开插件走同一条路径。
@@ -162,7 +280,17 @@ PARALLEL_WARMUP_GPUS=8 SHAPE_FILE=shape/DeepSeek-V4-Flash-p32768d1024.yaml bash
WARNING 1 config key(s) got different winners across shards — parallel interference reached the sweep WARNING 1 config key(s) got different winners across shards — parallel interference reached the sweep
``` ```
**冲突数就是这次加速的可信度指标**：0 可放心用；偏多说明 config 选择已被污染，出正式结论前不设该变量重跑一遍。其余行为（`REPLAY_FROM` 已设或可见卡不足 2 张时跳过、分片失败回退串行、轮转分片而非按块切、用子进程而非 xdist 的原因）见插件 docstring。分片绑卡遵守继承来的 `CUDA_VISIBLE_DEVICES`——用它选过卡时，分片只会落在你选的那几张上。 **冲突数就是这次加速的可信度指标**：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_batch.py）
@@ -170,7 +298,8 @@ WARNING 1 config key(s) got different winners across shards — parallel interfe
```bash ```bash
# 全量筛查（双仓库全部可收集算子，排除带无条件 skip 标记的与 ops/dismiss.txt 里的） # 全量筛查（双仓库全部可收集算子，排除带无条件 skip 标记的与 ops/dismiss.txt 里的）
python run_batch.py --all --gpus 0,1,2,3,4,5,6,7 --dtypes bfloat16 --op-timeout 1800 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 仓加前缀；同名测试函数用 @文件名 消歧 # 指定子集：默认 main 仓，vllm 仓加前缀；同名测试函数用 @文件名 消歧
python run_batch.py --ops softmax vllm:fused_marlin_moe nextafter_@nextafter_ python run_batch.py --ops softmax vllm:fused_marlin_moe nextafter_@nextafter_
@@ -195,7 +324,7 @@ python run_batch.py --ops-file my_ops.txt --replay-root runs/batch_xxx
**执行语义**：每算子独立子进程（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` 据此复用。 **执行语义**：每算子独立子进程（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、autotune_records/、ttgir/……），所以任何一个算子都可以事后单独 `REPLAY_FROM` 复测。批量层额外产出 `ops_status.csv`（每算子一行：状态/耗时/行数/备注，边跑边原子更新）与 `summary.csv`（每测量行一行：dtype、latency、speedup、Size Detail，来源是 run.log 的结果表，`batch_summary.py` 也可单独对旧批次重跑）。备注列聚合 `no_cudagraph_rows`/`replay_fallback` 计数，解读口径时先看这列。 **产物**：`<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 的预热、回退与精度 ## 计时口径：cudagraph 的预热、回退与精度
@@ -208,27 +337,29 @@ python run_batch.py --ops-file my_ops.txt --replay-root runs/batch_xxx
回退本身不影响 A/B 公平性（两侧同一算子回退行为一致），但**回退口径的测量误差更大**：do_bench 每次迭代都走完整的 Python → launch 路径，kernel 越小，launch 开销和 CPU 侧抖动在数字里占比越高——亚毫秒级 kernel 上两种口径可差 2 倍以上，且行间波动更明显。解读这类算子的结果时： 回退本身不影响 A/B 公平性（两侧同一算子回退行为一致），但**回退口径的测量误差更大**：do_bench 每次迭代都走完整的 Python → launch 路径，kernel 越小，launch 开销和 CPU 侧抖动在数字里占比越高——亚毫秒级 kernel 上两种口径可差 2 倍以上，且行间波动更明显。解读这类算子的结果时：
- 小 shape 行的绝对值和小幅（<10%）差异不要过度解读，优先看大 shape 行； - 小 shape 行的绝对值和小幅（<10%）差异不要过度解读，优先看大 shape 行；
- 需要更高置信度时，同一配置多跑几次取中位数，或对该算子直接注释掉 `_cudagraph_plugin` 统一用 do_bench 口径（消除同表混两种口径的问题）。 - 需要更高置信度时，同一配置多跑几次取中位数，或设 `ZL_ENABLE_CUDAGRAPH=0` 统一用 `timer=direct`（消除同表混两种口径的问题）。
回退是按测量点发生的，同一份 latency 表里可能混有两种口径的行；若不确定，先 `grep BENCHMARK_DIRECT_NO_CUDAGRAPH run.log` 确认哪些行是回退口径再下结论。 回退是按测量点发生的，同一份 latency 表里可能混有两种口径的行；若不确定，先 `grep BENCHMARK_DIRECT_NO_CUDAGRAPH run.log` 确认哪些行是回退口径再下结论。
## 插件说明 ## 插件说明
脚本通过 `-p` 加载以下插件（`run_pytest.sh` 的 `PLUGINS` 数组，可按需注释）： `-p` 列表由 `plugin_plan.py` 按 `ZL_BACKEND` 生成，不再在 `run_pytest.sh` 里写死。`_cudagraph_plugin` **不**出现在默认列表上：只有 CUDA 且 `_timing_plugin` 选中 Graph 时才会 import 它。
| 插件 | 作用 | 何时关闭 | | 插件 | 作用 | 何时关闭 |
|-----|------|---------| |-----|------|---------|
| `_device_guard_plugin` | 在 import flag_gems 之前直接用 torch 探测到 NVIDIA 卡就设 `GEMS_VENDOR=nvidia`，跳过 flag_gems 启动时 `nvidia-smi` 子进程探测（该探测在部分 fork 环境下会挂在 `wait4` 上导致 import 卡死）。副作用：import 期即初始化 CUDA context，故与 pytest-xdist 不兼容（本框架不用 xdist） | 一般无需关：已设 `GEMS_VENDOR`/`FLAGGEMS_VENDOR` 等 env 时自动跳过，非 NVIDIA 卡上自动 no-op | | `_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」） | 不设该变量即关闭 | | `_parallel_warmup_plugin` | 由 `PARALLEL_WARMUP_GPUS=N` 激活（未设则静默 no-op）：autotune sweep 分片到 N 卡并行，测量仍单卡串行、产物结构不变（见「加速：多卡并行 sweep」） | 不设该变量即关闭 |
| `_seed_plugin` | 固定 random/numpy/torch 种子，数据相关算子（sort/topk 等）输入逐字节一致 | 不关 | | `_seed_plugin` | 固定 random/numpy/torch 种子，数据相关算子（sort/topk 等）输入逐字节一致 | 不关 |
| `_shape_inject_plugin` | 让 shape yaml 覆盖子类硬编码的 `set_shapes()` | 不关 | | `_shape_inject_plugin` | 让 shape yaml 覆盖子类硬编码的 `set_shapes()` | 不关 |
| `_shape_iter_inject_plugin` | 覆盖在 `get_input_iter` 里硬编码 shape 的类（conv/pool 等） | 不关 | | `_shape_iter_inject_plugin` | 覆盖在 `get_input_iter` 里硬编码 shape 的类（conv/pool 等） | 不关 |
| `_bespoke_shape_plugin` | 覆盖特殊输入构造的算子（upsample/flash_mla/cutlass 等，按类名注册） | 测这些算子之外可关 | | `_bespoke_shape_plugin` | 覆盖特殊输入构造的算子（upsample/flash_mla/cutlass 等，按类名注册） | 测这些算子之外可关 |
| `_autotune_record_plugin` | record/replay Triton autotune 选择（由 RECORD/REPLAY 环境变量二选一激活） | 不关 | | `_autotune_record_plugin` | record/replay Triton autotune 选择；`autotune_records/meta.json` 记录 backend，跨后端 replay 拒绝 | `ZL_ENABLE_AUTOTUNE=0` |
| `_cudagraph_plugin` | `do_bench` → `do_bench_cudagraph`（kernel 纯耗时；内部先做跨流同步，修过一个间歇性 illegal instruction）。**capture 前按 warmup 时间预算显式预热**（先一次丢弃跑吸收 autotune/JIT 编译，再稳态预热），消除首轮抖动。无法 graph capture 的 kernel 自动回退并打 `BENCHMARK_DIRECT_NO_CUDAGRAPH` 标记 | 需要与他人的普通 do_bench 数据对齐时注释掉 | | `_timing_plugin` | 按后端选择计时路径：CUDA 可走 Graph（内部调用 `_cudagraph_plugin`），其余后端一开始就是上游 `do_bench`（`timer=direct`） | `ZL_ENABLE_CUDAGRAPH=0` 强制 direct |
| `_pretty_report_plugin` | 结果表整理 + 着色：全表相同的输入折叠成表头下一行图例，每行 Size Detail 只留随行变化的部分（shape 保持 `torch.Size([...])` 原样，MoE 类算子单行从 ~400 字符缩到一屏内）；SUCCESS 绿 / FAILED 红。列名与 `SUCCESS` 字样保持上游原文，`run.log` 的 grep/解析不受影响 | 需要与上游原始表格逐字对齐时注释掉 | | `_cudagraph_plugin` | 仅由 `_timing_plugin` 在 CUDA 上调用。`do_bench` → `do_bench_cudagraph`，capture 前按 warmup 预算显式预热；无法 capture 时回退并打 `BENCHMARK_DIRECT_NO_CUDAGRAPH` | 不要直接 `-p` 加到非 CUDA 运行 |
| `_ir_meta_plugin` | 编译期记录每个变体的 constexpr/签名/特化；launch 钩子统计每个 (kernel, shape) 的实际使用；退出时按上述结构 dump ttgir | 不关 | | `_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` 声明，其他后端不加载 |
| `_mm_cluster_fix_plugin` | Hopper fp16 mm cluster kernel 越界崩溃的运行时规避 | 默认注释；测 fp16 mm 崩溃时打开 | | `_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` |
## 常见问题 ## 常见问题
@@ -257,5 +388,7 @@ shape 多时这部分会主导墙上时间，可设 `PARALLEL_WARMUP_GPUS=N` 并
## 依赖假设 ## 依赖假设
- FlagGems benchmark 体系（`benchmark/base.py` 的 `Benchmark` 类、conftest 的 `--shape_file/--level/--mode` 选项）；主仓与 FlagGems-vllm 仓库同构，`_autotune_record_plugin` 会按实际被 import 的包（`flag_gems` / `flaggems_vllm`）挂 LibTuner 补丁； - FlagGems benchmark 体系（`benchmark/base.py` 的 `Benchmark` 类、conftest 的 `--shape_file/--level/--mode` 选项）；主仓与 FlagGems-vllm 仓库同构，`_autotune_record_plugin` 会按实际被 import 的包（`flag_gems` / `flaggems_vllm`）挂 LibTuner 补丁；
- Triton 需支持 `knobs.compilation.listener`、`kernel_load_end_hook`、`launch_enter_hook`（当前 FlagTree 的 triton 3.6 满足）； - 设备后端由 `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 输出的注册日志是否还正常打印。 - 插件通过 monkeypatch 挂钩上游内部结构，FlagGems/Triton 大版本升级后若行为异常，优先检查各插件 pytest_configure 输出的注册日志是否还正常打印。
- 无 GPU 的适配器自检：`python test_backends.py`。
_autotune_record_plugin.py
+71 -1
View File
@@ -214,6 +214,57 @@ def _replay_run(original):
return runner return runner
def _write_record_meta(record_dir: str) -> None:
"""Sidecar next to <op>.json. Replay refuses a different backend.
The config JSON itself stays `{kernel: {key: cfg}}` so existing REPLAY_FROM
directories keep working.
"""
try:
from backends import get_backend
backend_name = get_backend().name
except Exception:
backend_name = os.environ.get("ZL_BACKEND", "")
meta = {
"schema_version": 1,
"provider": "flaggems",
"backend": backend_name,
"op": _op_name(),
}
path = Path(record_dir) / "meta.json"
try:
path.parent.mkdir(parents=True, exist_ok=True)
path.write_text(json.dumps(meta, indent=2, ensure_ascii=False) + "\n")
except Exception as exc:
from _term_style import tag
print(f"{tag('[autotune-record-plugin]')} meta.json write failed: {exc}",
file=sys.stderr, flush=True)
def _replay_backend_ok(replay_dir: str) -> bool:
path = Path(replay_dir) / "meta.json"
if not path.is_file():
return True
try:
meta = json.loads(path.read_text())
except Exception:
return True
if not isinstance(meta, dict):
return True
recorded = str(meta.get("backend") or "").strip()
if not recorded:
return True
try:
from backends import get_backend
current = get_backend().name
except Exception:
current = os.environ.get("ZL_BACKEND", "").strip()
if recorded != current:
_emit_marker(f"backend_mismatch recorded={recorded} current={current}")
return False
return True
def _load_replay_dir(replay_dir: str) -> int: def _load_replay_dir(replay_dir: str) -> int:
path = _record_path(replay_dir) path = _record_path(replay_dir)
if not path.is_file(): if not path.is_file():
@@ -256,29 +307,48 @@ def _dump_record(record_dir: str) -> None:
def pytest_configure(config): def pytest_configure(config):
global _wrap global _wrap
from _term_style import tag from _term_style import tag
from _run_identity import update_identity
record_dir = os.environ.get(_RECORD_DIR_ENV, "").strip() record_dir = os.environ.get(_RECORD_DIR_ENV, "").strip()
replay_dir = os.environ.get(_REPLAY_DIR_ENV, "").strip() replay_dir = os.environ.get(_REPLAY_DIR_ENV, "").strip()
if record_dir and replay_dir: if record_dir and replay_dir:
# Mutually exclusive: recording while replaying is pointless. Surface loudly. # Mutually exclusive: recording while replaying is pointless. Surface loudly.
print(f"{tag('[autotune-record-plugin]')} error: both RECORD_DIR and " print(f"{tag('[autotune-record-plugin]')} error: both RECORD_DIR and "
"REPLAY_DIR set; ignoring both (no-op)", file=sys.stderr, flush=True) "REPLAY_DIR set; ignoring both (no-op)", file=sys.stderr, flush=True)
update_identity(autotune_mode="uncontrolled")
return return
if not record_dir and not replay_dir: if not record_dir and not replay_dir:
update_identity(autotune_mode="uncontrolled")
return
try:
import triton.runtime.autotuner as _autotuner
except Exception as exc:
print(f"{tag('[autotune-record-plugin]')} Triton Autotuner unavailable "
f"({exc}); autotune_mode=uncontrolled", file=sys.stderr, flush=True)
update_identity(autotune_mode="uncontrolled")
return
if replay_dir and not _replay_backend_ok(replay_dir):
print(f"{tag('[autotune-record-plugin]')} refusing replay across backends",
file=sys.stderr, flush=True)
update_identity(autotune_mode="replay_fallback")
return return
import triton.runtime.autotuner as _autotuner
_wrap = _record_run if record_dir else _replay_run _wrap = _record_run if record_dir else _replay_run
_autotuner.Autotuner.run = _wrap(_autotuner.Autotuner.run) _autotuner.Autotuner.run = _wrap(_autotuner.Autotuner.run)
if record_dir: if record_dir:
# atexit (not sessionfinish): persist whatever was recorded even if an op # atexit (not sessionfinish): persist whatever was recorded even if an op
# crash kills the session; a later replay run falls back for missing keys. # crash kills the session; a later replay run falls back for missing keys.
_write_record_meta(record_dir)
atexit.register(_dump_record, record_dir) atexit.register(_dump_record, record_dir)
update_identity(autotune_mode="record")
print(f"{tag('[autotune-record-plugin]')} recording autotune configs to " print(f"{tag('[autotune-record-plugin]')} recording autotune configs to "
f"{_record_path(record_dir)} (Autotuner)", f"{_record_path(record_dir)} (Autotuner)",
file=sys.stderr, flush=True) file=sys.stderr, flush=True)
else: else:
loaded = _load_replay_dir(replay_dir) loaded = _load_replay_dir(replay_dir)
update_identity(autotune_mode="replay")
print(f"{tag('[autotune-record-plugin]')} replaying {loaded} recorded entries from " print(f"{tag('[autotune-record-plugin]')} replaying {loaded} recorded entries from "
f"{_record_path(replay_dir)} (Autotuner)", f"{_record_path(replay_dir)} (Autotuner)",
file=sys.stderr, flush=True) file=sys.stderr, flush=True)
_cudagraph_plugin.py
+224 -67
View File
@@ -1,16 +1,25 @@
"""pytest plugin: swap `triton.testing.do_bench` for cudagraph-based timing. """pytest plugin: swap `triton.testing.do_bench` for graph-capture timing.
Loaded via `-p _cudagraph_plugin` in run_pytest.sh's PLUGINS list (comment that Invoked by `_timing_plugin` when the active backend declares
line out to time with plain do_bench). Monkey-patches do_bench → `supports_cuda_graph`. Do not put this module on the default `-p` list; the
timing plugin owns the decision. Monkey-patches do_bench →
do_bench_cudagraph (kernel-only latency) at session configure, after Triton is do_bench_cudagraph (kernel-only latency) at session configure, after Triton is
loaded but before any benchmark calls it. Kernels that can't be graph-captured loaded but before any benchmark calls it. Kernels that can't be graph-captured
fall back to plain do_bench and print a BENCHMARK_DIRECT_NO_CUDAGRAPH marker fall back to plain do_bench and print a BENCHMARK_DIRECT_NO_CUDAGRAPH marker
into the log. Before capturing, fn is explicitly warmed up for the caller's into the log. Before capturing, fn is explicitly warmed up for the caller's
warmup time budget (do_bench_cudagraph's own 5-iter warmup is too short to warmup time budget (do_bench_cudagraph's own 5-iter warmup is too short to
settle autotune/JIT for the MoE kernels, which makes the first captured settle autotune/JIT for the MoE kernels, which makes the first captured
measurement unstable run-to-run). The capture itself runs under a scratch CUDA measurement unstable run-to-run). The capture itself runs under a scratch
RNG state so that an aborted capture cannot poison the process-wide generator RNG state so that an aborted capture cannot poison the process-wide generator
(rationale: _capture_with_scratch_rng). (rationale: _capture_with_scratch_rng).
Stream/event/graph calls come from Triton's active device interface, not from
torch.cuda. FlagTree's do_bench is already written that way; its
do_bench_cudagraph is not, so on a non-CUDA backend every capture raised and
every row silently took the fallback while the log still claimed capture was on.
CUDA still calls upstream's do_bench_cudagraph verbatim, so its numbers do not
move. Which failures are recoverable is likewise a backend question, because the
error wording is vendor text; see DeviceBackend.classify_capture_error.
""" """
from __future__ import annotations from __future__ import annotations
@@ -24,31 +33,74 @@ import triton.testing as _tt
_ORIGINAL_DO_BENCH = _tt.do_bench _ORIGINAL_DO_BENCH = _tt.do_bench
def _is_fatal_cuda_error(exc: Exception) -> bool: def direct_repeat() -> int:
"""True for CUDA context errors (illegal address, OOM, device-side assert): """ZL_DIRECT_REPEAT, clamped to >= 1. 1 keeps today's single call."""
the context is poisoned, so re-raise rather than fall back and record a bogus try:
latency as success. return max(1, int(os.environ.get("ZL_DIRECT_REPEAT", "1")))
except ValueError:
Markers must match the error text itself, not torch's generic footer: newer return 1
torch appends "Compile with `TORCH_USE_CUDA_DSA` to enable device-side
assertions." to EVERY AcceleratorError, so a loose "device-side assert"
marker turned all capture-incompatible ops into fatal re-raises and killed
the documented do_bench fallback path entirely."""
msg = f"{type(exc).__name__}: {exc}".lower()
markers = (
"illegal memory access",
"cudaerrorillegaladdress",
"out of memory",
"device-side assert triggered",
"an illegal instruction",
"misaligned address",
"uncorrectable ecc",
)
return any(m in msg for m in markers)
def _warmup_before_capture(fn, warmup_ms, grad_to_none): def run_direct(fn, warmup, rep, grad_to_none, quantiles, return_mode):
"""Warm fn up before cudagraph capture. """Upstream do_bench, optionally repeated with the median taken.
Direct timing walks the full Python launch path once per iteration, so on
small kernels it carries launch overhead and CPU jitter that graph capture
does not. do_bench already returns a median over the iterations inside one
call, which suppresses iteration-level jitter but not drift across a call:
clocks ramping, another process arriving. Repeating the call and taking the
median over calls is the part the inner median cannot see.
It matters most where it is least optional. On MThreads the operators with
the largest RLC convert reductions -- sort, dropout, rand -- are exactly the
ones that cannot be graph-captured, so the headline numbers there are
produced by this path.
Repetition is skipped when the caller asks for quantiles or return_mode
"all", because the result is then a sequence and a median across calls is
not defined for it; the single call is returned unchanged.
"""
n = direct_repeat()
if n == 1 or quantiles is not None or return_mode == "all":
return _ORIGINAL_DO_BENCH(
fn, warmup=warmup, rep=rep, grad_to_none=grad_to_none,
quantiles=quantiles, return_mode=return_mode)
import statistics
samples = [
_ORIGINAL_DO_BENCH(fn, warmup=warmup, rep=rep, grad_to_none=grad_to_none,
quantiles=None, return_mode=return_mode)
for _ in range(n)
]
return statistics.median(samples)
def _device_interface():
"""Triton's active device interface: torch.cuda on NVIDIA, torch_musa on MThreads."""
from triton.runtime import driver
return driver.active.get_device_interface()
def _is_cuda(di) -> bool:
import torch
return di is torch.cuda
def _graph_class(di):
"""The interface's graph type: CUDAGraph, MUSAGraph, or whatever it names."""
for name in ("CUDAGraph", "MUSAGraph"):
graph = getattr(di, name, None)
if graph is not None:
return graph
for name in dir(di):
if name.endswith("Graph"):
return getattr(di, name)
raise RuntimeError("device interface exposes no graph type")
def _warmup_before_capture(di, fn, warmup_ms, grad_to_none):
"""Warm fn up before graph capture.
do_bench_cudagraph only warms up 5 iterations internally — too few for the do_bench_cudagraph only warms up 5 iterations internally — too few for the
MoE kernels, where the first calls trigger Triton autotune config MoE kernels, where the first calls trigger Triton autotune config
@@ -66,8 +118,6 @@ def _warmup_before_capture(fn, warmup_ms, grad_to_none):
""" """
import time import time
import torch
def _run(): def _run():
if grad_to_none is not None: if grad_to_none is not None:
for x in grad_to_none: for x in grad_to_none:
@@ -76,23 +126,23 @@ def _warmup_before_capture(fn, warmup_ms, grad_to_none):
# First call absorbs autotune/JIT compilation (can take seconds); discard it. # First call absorbs autotune/JIT compilation (can take seconds); discard it.
_run() _run()
torch.cuda.synchronize() di.synchronize()
# Size the warmup loop from one timed steady-state run. dt_ms includes # Size the warmup loop from one timed steady-state run. dt_ms includes
# launch overhead, so n errs low for tiny kernels; warmup needn't be exact. # launch overhead, so n errs low for tiny kernels; warmup needn't be exact.
t0 = time.perf_counter() t0 = time.perf_counter()
_run() _run()
torch.cuda.synchronize() di.synchronize()
dt_ms = (time.perf_counter() - t0) * 1000.0 dt_ms = (time.perf_counter() - t0) * 1000.0
budget = warmup_ms if (warmup_ms and warmup_ms > 0) else 25.0 budget = warmup_ms if (warmup_ms and warmup_ms > 0) else 25.0
n = int(budget / dt_ms) if dt_ms > 0 else 25 n = int(budget / dt_ms) if dt_ms > 0 else 25
n = max(5, min(n, 200)) n = max(5, min(n, 200))
for _ in range(n): for _ in range(n):
_run() _run()
torch.cuda.synchronize() di.synchronize()
def _capture_with_scratch_rng(call): def _capture_with_scratch_rng(di, call):
"""Run a graph capture under a throwaway CUDA RNG state. """Run a graph capture under a throwaway device RNG state.
torch.cuda.graph's capture_begin unconditionally registers the default CUDA torch.cuda.graph's capture_begin unconditionally registers the default CUDA
RNG generator and sets its capturing flag; when a capture aborts midway the RNG generator and sets its capturing flag; when a capture aborts midway the
@@ -104,12 +154,20 @@ def _capture_with_scratch_rng(call):
before capturing and restore the original state object afterwards, success before capturing and restore the original state object afterwards, success
or failure. The real generator never takes part in a capture, so poisoning or failure. The real generator never takes part in a capture, so poisoning
can only land on the discarded scratch state, and _seed_plugin determinism can only land on the discarded scratch state, and _seed_plugin determinism
is left untouched.""" is left untouched.
An interface that exposes no default_generators has no such registration to
protect against (torch_musa is one), so the capture simply runs."""
import torch import torch
dev = torch.cuda.current_device()
gen = torch.cuda.default_generators[dev] generators = getattr(di, "default_generators", None)
if not generators:
return call()
gen = generators[di.current_device()]
if not hasattr(gen, "graphsafe_get_state"):
return call()
saved = gen.graphsafe_get_state() saved = gen.graphsafe_get_state()
scratch = torch.Generator(device="cuda") scratch = torch.Generator(device=torch.device(gen.device))
scratch.manual_seed(gen.initial_seed()) scratch.manual_seed(gen.initial_seed())
gen.graphsafe_set_state(scratch) gen.graphsafe_set_state(scratch)
try: try:
@@ -118,24 +176,111 @@ def _capture_with_scratch_rng(call):
gen.graphsafe_set_state(saved) gen.graphsafe_set_state(saved)
def _do_bench_graph(di, fn, rep, grad_to_none, quantiles, return_mode, n_retries=10):
"""do_bench_cudagraph, taken from the device interface instead of torch.cuda.
A step-for-step port of upstream's implementation: estimate one iteration,
unroll that many calls into a captured graph so host launch overhead drops
out, then time independent replays. Upstream is used verbatim on CUDA; this
exists because upstream names torch.cuda directly and so cannot run anywhere
else.
Replay timings go to _summarize_statistics as a plain list of floats, which
is what both do_bench and upstream's do_bench_cudagraph pass: this fork's
helper runs Python's statistics module over them, so wrapping them in a
tensor makes "mean" raise and makes a quantile request return tensors that
later fail to serialize.
"""
with di.stream(di.Stream()):
fn()
if grad_to_none is not None:
for x in grad_to_none:
x.detach_()
x.requires_grad_(True)
x.grad = None
start_event = di.Event(enable_timing=True)
end_event = di.Event(enable_timing=True)
start_event.record()
for _ in range(5):
fn()
end_event.record()
di.synchronize()
estimate_ms = start_event.elapsed_time(end_event) / 5
n_repeat = 1000 if estimate_ms == 0 else max(1, int(rep / estimate_ms))
graph = _graph_class(di)()
with di.graph(graph):
for _ in range(n_repeat):
if grad_to_none is not None:
for x in grad_to_none:
x.grad = None
fn()
di.synchronize()
ret = []
for _ in range(n_retries):
start_event = di.Event(enable_timing=True)
end_event = di.Event(enable_timing=True)
start_event.record()
graph.replay()
end_event.record()
di.synchronize()
ret.append(start_event.elapsed_time(end_event) / n_repeat)
del graph
return _tt._summarize_statistics(ret, quantiles, return_mode)
def _patched_do_bench(fn, warmup=25, rep=100, grad_to_none=None, def _patched_do_bench(fn, warmup=25, rep=100, grad_to_none=None,
quantiles=None, return_mode="mean"): quantiles=None, return_mode="mean"):
"""Drop-in replacement for triton.testing.do_bench using cudagraph capture.""" """Drop-in replacement for triton.testing.do_bench using graph capture."""
op = os.environ.get("FLAGGEMS_PERF_CURRENT_OP", "<unknown>") op = os.environ.get("FLAGGEMS_PERF_CURRENT_OP", "<unknown>")
phase = "sync" # which step rejected cudagraph, for the fallback marker
def _direct(note):
from _term_style import YELLOW, paint
n = direct_repeat()
# The repeat count belongs in the marker: these rows are already a
# different caliber from the captured ones, and now they can also be a
# different number of samples.
rpt = "" if n == 1 else f" repeat={n}"
print(paint(f"BENCHMARK_DIRECT_NO_CUDAGRAPH op={op} {note}{rpt}", YELLOW),
flush=True)
return run_direct(fn, warmup, rep, grad_to_none, quantiles, return_mode)
backend = None
try: try:
# do_bench_cudagraph runs fn on a fresh side stream without syncing from backends import get_backend
# with the default stream first; input tensors produced just before backend = get_backend()
# (quantize/topk in the input iter) may still be in flight, and racing except Exception:
# on them intermittently kills the context (illegal instruction from a pass
# device trap on garbage indices). Sync before switching streams.
import torch # Known context-poisoning operators never get a capture attempt: the point
torch.cuda.synchronize() # of the deny list is that for these the attempt itself is the damage, so
# discovering it again per shape would cost the run, not just the precision.
if backend is not None and backend.denies_capture(op):
return _direct("phase=denylist reason=KnownPoisoningOp "
f"detail='{backend.name} denies capture for this op'")
phase = "sync" # which step rejected capture, for the marker
try:
di = _device_interface()
# The graph path runs fn on a fresh side stream without syncing with the
# default stream first; input tensors produced just before (quantize/topk
# in the input iter) may still be in flight, and racing on them
# intermittently kills the context (illegal instruction from a device
# trap on garbage indices). Sync before switching streams.
di.synchronize()
phase = "warmup" phase = "warmup"
# Settle autotune/JIT before capture (rationale: _warmup_before_capture). # Settle autotune/JIT before capture (rationale: _warmup_before_capture).
_warmup_before_capture(fn, warmup, grad_to_none) _warmup_before_capture(di, fn, warmup, grad_to_none)
phase = "capture" phase = "capture"
return _capture_with_scratch_rng(lambda: _tt.do_bench_cudagraph( if _is_cuda(di):
return _capture_with_scratch_rng(di, lambda: _tt.do_bench_cudagraph(
fn,
rep=rep,
grad_to_none=grad_to_none,
quantiles=quantiles,
return_mode=return_mode,
))
return _capture_with_scratch_rng(di, lambda: _do_bench_graph(
di,
fn, fn,
rep=rep, rep=rep,
grad_to_none=grad_to_none, grad_to_none=grad_to_none,
@@ -143,30 +288,42 @@ def _patched_do_bench(fn, warmup=25, rep=100, grad_to_none=None,
return_mode=return_mode, return_mode=return_mode,
)) ))
except Exception as exc: except Exception as exc:
# Re-raise genuine device failures; anything else falls back to plain verdict = (backend.classify_capture_error(exc)
# do_bench, with the failing phase (sync/warmup/capture) in the marker. if backend is not None else "recoverable")
if _is_fatal_cuda_error(exc): if verdict == "fatal":
# The kernel failed on the device. A fallback latency here would
# report a broken kernel as a measurement.
raise
if verdict == "poisoned":
# The attempt, not the kernel, killed the context. Falling back now
# would measure inside a dead process -- observed as a fallback
# marker followed by a teardown crash and zero rows. Surface it
# under its own marker and stop, so the operator can be re-run with
# ZL_ENABLE_CUDAGRAPH=0 and added to the backend's deny list.
from _term_style import RED, paint
detail = " ".join(str(exc).split())[:160]
print(paint(f"BENCHMARK_CAPTURE_POISONED op={op} phase={phase} "
f"reason={type(exc).__name__} detail={detail!r}", RED),
flush=True)
raise raise
reason = type(exc).__name__
detail = " ".join(str(exc).split())[:160] detail = " ".join(str(exc).split())[:160]
from _term_style import YELLOW, paint return _direct(f"phase={phase} reason={type(exc).__name__} "
print( f"detail={detail!r}")
paint(f"BENCHMARK_DIRECT_NO_CUDAGRAPH op={op} phase={phase} "
f"reason={reason} detail={detail!r}", YELLOW),
flush=True,
)
return _ORIGINAL_DO_BENCH(
fn,
warmup=warmup,
rep=rep,
grad_to_none=grad_to_none,
quantiles=quantiles,
return_mode=return_mode,
)
def pytest_configure(config): def pytest_configure(config):
from _term_style import tag from _term_style import tag
try:
from backends import get_backend
if not get_backend().supports_cuda_graph:
print(f"{tag('[cudagraph-plugin]')} skipped: backend "
f"{get_backend().name!r} does not support CUDA Graph",
file=sys.stderr, flush=True)
return
except Exception as exc:
print(f"{tag('[cudagraph-plugin]')} backend probe failed ({exc}); "
"not patching do_bench", file=sys.stderr, flush=True)
return
_tt.do_bench = _patched_do_bench _tt.do_bench = _patched_do_bench
print(f"{tag('[cudagraph-plugin]')} triton.testing.do_bench → do_bench_cudagraph", print(f"{tag('[cudagraph-plugin]')} triton.testing.do_bench → do_bench_cudagraph",
file=sys.stderr, flush=True) file=sys.stderr, flush=True)
_device_guard_plugin.py
+5 -43
View File
@@ -1,45 +1,7 @@
"""Guard against FlagGems' device-detection hang (benchmark-side fix). """Backward-compatible alias of `_device_setup_plugin`.
FlagGems' `runtime.backend.device_finder` falls back to Kept so an old `-p _device_guard_plugin` line still sets GEMS_VENDOR. New
`subprocess.run(nvidia-smi)` **without a timeout**. Under this conda env's plugin lists should load `_device_setup_plugin` instead.
forked/inconsistent subprocess (`_posixsubprocess` symbol mismatch), that
child can hang indefinitely, leaving `import flag_gems` stuck in `wait4`
(seen as run_pytest hanging with no output, GPU at 0%).
This plugin runs at import time — before any test module imports flag_gems —
detects the vendor via torch (no subprocess), and sets GEMS_VENDOR so
device_finder takes its env fast-path (`_get_vendor_from_env`) and never
reaches the hang-prone subprocess probe. No change to the FlagGems repo.
Side effect: get_device_properties initializes the CUDA context very early in
the pytest process. Fine for the current single-process runs; revisit if
fork-based parallelism (e.g. pytest-xdist) is ever introduced.
""" """
import os from _device_setup_plugin import * # noqa: F401,F403
import sys from _device_setup_plugin import pytest_configure # noqa: F401
_VENDOR_ENV_KEYS = ("GEMS_VENDOR", "FLAGGEMS_VENDOR", "GEMS_BACKEND", "FLAGGEMS_BACKEND")
def _guard_device_vendor():
# Respect an explicit choice if the user already set one.
if any(k in os.environ for k in _VENDOR_ENV_KEYS):
return
try:
import torch
if torch.cuda.is_available():
name = torch.cuda.get_device_properties(0).name.upper()
if "NVIDIA" in name:
os.environ["GEMS_VENDOR"] = "nvidia"
from _term_style import tag
print(f"{tag('[device-guard-plugin]')} set GEMS_VENDOR=nvidia "
"(skip flag_gems nvidia-smi subprocess probe, avoids import hang)",
file=sys.stderr, flush=True)
except Exception as e: # pragma: no cover
print(f"[device-guard-plugin] torch vendor probe failed ({e}); "
"leaving detection to flag_gems",
file=sys.stderr, flush=True)
_guard_device_vendor()
_device_setup_plugin.py
+69
View File
@@ -0,0 +1,69 @@
"""Set FlagGems vendor + record backend identity before flag_gems is imported.
Replaces the NVIDIA-only `_device_guard_plugin`. Setting GEMS_VENDOR from the
zl_bench backend name skips FlagGems' hang-prone nvidia-smi subprocess probe
on CUDA, and maps musa/rocm/ascend onto the matching FlagGems vendor.
Import-time side effect is intentional: FlagGems reads GEMS_VENDOR while the
benchmark module is being collected. Device name / arch are filled later in
pytest_configure so import does not initialize a device context.
"""
from __future__ import annotations
import os
import sys
from backends import get_backend
from providers.flaggems import apply_vendor_env
from _run_identity import update_identity
def _setup_at_import() -> None:
try:
backend = get_backend()
except Exception as exc:
print(f"[device-setup-plugin] cannot resolve backend ({exc}); "
"leaving GEMS_VENDOR to flag_gems",
file=sys.stderr, flush=True)
return
vendor = apply_vendor_env(backend)
os.environ["ZL_BACKEND"] = backend.name
device = os.environ.get("ZL_DEVICE", "").strip()
if not device:
for key in backend.visibility_env_keys:
raw = os.environ.get(key, "").strip()
if raw:
device = raw.split(",")[0].strip()
break
update_identity(
backend=backend.name,
device=device,
flaggems_vendor=vendor or "",
visibility_env=backend.visibility_env_keys[0] if backend.visibility_env_keys else "",
)
if vendor:
from _term_style import tag
print(f"{tag('[device-setup-plugin]')} backend={backend.name} "
f"GEMS_VENDOR={vendor}",
file=sys.stderr, flush=True)
def pytest_configure(config):
try:
backend = get_backend()
ident = backend.device_identity()
update_identity(
backend=backend.name,
device_name=ident.get("device_name", ""),
device_arch=ident.get("device_arch", ""),
timer_api=ident.get("timer_api", ""),
flaggems_vendor=ident.get("flaggems_vendor", ""),
)
except Exception as exc:
from _term_style import tag
print(f"{tag('[device-setup-plugin]')} identity probe failed ({exc})",
file=sys.stderr, flush=True)
_setup_at_import()
_ir_meta_plugin.py
+21 -5
View File
@@ -402,17 +402,33 @@ def _dump_ttgir(cache_dir: str, dump_dir: str) -> None:
def pytest_configure(config): def pytest_configure(config):
global _dump_enabled global _dump_enabled
import triton from _term_style import tag
import triton.runtime.autotuner as _autotuner try:
import triton
import triton.runtime.autotuner as _autotuner
except Exception as exc:
print(f"{tag('[ir-meta-plugin]')} Triton knobs unavailable ({exc}); "
"IR collection disabled", file=sys.stderr, flush=True)
return
if not hasattr(triton, "knobs"):
print(f"{tag('[ir-meta-plugin]')} triton.knobs missing; "
"IR collection disabled", file=sys.stderr, flush=True)
return
prev = triton.knobs.compilation.listener try:
compilation = triton.knobs.compilation
prev = compilation.listener
except Exception as exc:
print(f"{tag('[ir-meta-plugin]')} compilation listener unavailable "
f"({exc}); IR collection disabled", file=sys.stderr, flush=True)
return
if prev is None: if prev is None:
triton.knobs.compilation.listener = _compile_listener compilation.listener = _compile_listener
else: else:
def chained(**kwargs): def chained(**kwargs):
prev(**kwargs) prev(**kwargs)
_compile_listener(**kwargs) _compile_listener(**kwargs)
triton.knobs.compilation.listener = chained compilation.listener = chained
dump_dir = os.environ.get(_DUMP_DIR_ENV, "").strip() dump_dir = os.environ.get(_DUMP_DIR_ENV, "").strip()
if dump_dir: if dump_dir:
_mm_cluster_fix_plugin.py
+12
View File
@@ -105,6 +105,18 @@ def _warmup_load_backend():
def pytest_configure(config): def pytest_configure(config):
from _term_style import tag from _term_style import tag
try:
from backends import get_backend
if not get_backend().supports_mm_cluster_fix:
print(f"{tag('[mm-cluster-fix-plugin]')} skipped: backend "
f"{get_backend().name!r} is not the CUDA Hopper path",
file=sys.stderr, flush=True)
return
except Exception as exc:
print(f"{tag('[mm-cluster-fix-plugin]')} backend probe failed ({exc}); "
"leaving cluster_remote_mm_scenario unpatched",
file=sys.stderr, flush=True)
return
_patch_loaded() _patch_loaded()
_warmup_load_backend() _warmup_load_backend()
_patch_loaded() _patch_loaded()
_parallel_warmup_plugin.py
+171 -30
View File
@@ -33,7 +33,7 @@ import sys
import tempfile import tempfile
import time import time
from pathlib import Path from pathlib import Path
from typing import Any, Dict, List, Tuple from typing import Any, Dict, List, Optional, Tuple
import pytest import pytest
@@ -60,8 +60,8 @@ def _requested_gpus() -> int:
if want < 2: if want < 2:
return 0 return 0
try: try:
import torch from backends import get_backend
have = torch.cuda.device_count() have = get_backend().probe_device_count()
except Exception: except Exception:
return 0 return 0
return max(0, min(want, have)) return max(0, min(want, have))
@@ -70,30 +70,132 @@ def _requested_gpus() -> int:
def _visible_devices() -> List[str]: def _visible_devices() -> List[str]:
"""The device ids a shard may be pinned to, in parent-visible order. """The device ids a shard may be pinned to, in parent-visible order.
Must respect an inherited CUDA_VISIBLE_DEVICES: writing a bare shard index Must respect inherited visibility env for the *active* backend. Writing a
into the child would re-interpret it as an absolute id, so a caller who bare shard index into the child would re-interpret it as an absolute id,
picked idle cards (CUDA_VISIBLE_DEVICES=5,6) would silently get cards 0,1 — so a caller who picked idle cards (CUDA_VISIBLE_DEVICES=5,6) would silently
exactly the busy-GPU case the caller was avoiding. get cards 0,1 — exactly the busy-GPU case the caller was avoiding.
""" """
raw = os.environ.get("CUDA_VISIBLE_DEVICES", "").strip() from backends import get_backend
if raw: backend = get_backend()
return [d.strip() for d in raw.split(",") if d.strip()] for key in backend.visibility_env_keys:
raw = os.environ.get(key, "").strip()
if raw:
return [d.strip() for d in raw.split(",") if d.strip()]
try: try:
import torch n = backend.probe_device_count()
return [str(i) for i in range(torch.cuda.device_count())] return [str(i) for i in range(n)] if n else []
except Exception: except Exception:
return [] return []
def _read_shape_yaml(path: str, op: str) -> Tuple[List[Any], Dict[str, Any]]: def _attr_name(node: Any) -> str:
"""Shapes for this op, plus the sibling keys to carry into each shard.""" """Last component of a Name/Attribute node: `base.BlasBenchmark` -> `BlasBenchmark`."""
import ast
if isinstance(node, ast.Name):
return node.id
if isinstance(node, ast.Attribute):
return node.attr
return ""
def _class_bases(py_path: Path) -> Dict[str, List[str]]:
"""{class name: [base class names]} parsed statically from one module."""
import ast
try:
tree = ast.parse(py_path.read_text(errors="replace"))
except Exception:
return {}
return {
node.name: [n for n in (_attr_name(b) for b in node.bases) if n]
for node in ast.walk(tree)
if isinstance(node, ast.ClassDef)
}
def _class_chain(test_py: Path, op: str) -> List[str]:
"""Benchmark class names to try for `op`, nearest ancestor first.
FlagGems' `Benchmark.set_shapes` looks up `op_name` and, failing that, walks
`type(self).__mro__` trying each class name as a yaml key -- which is how
BLAS ops reach the shared `BlasBenchmark` entry in core_shapes.yaml. This
plugin runs at pytest_configure, before any Benchmark exists, so the real
MRO is unavailable; the class is recovered statically instead, from the
`op_name=` keyword of the constructor call inside the test module, then its
bases from that module plus benchmark/base.py.
Breadth-first over the inheritance graph, not a true C3 linearisation: the
benchmark classes are single-inheritance in practice, and a miss only costs
the sharding (see the caller's skip path), never correctness.
"""
import ast
try:
tree = ast.parse(test_py.read_text(errors="replace"))
except Exception:
return []
start = ""
for node in ast.walk(tree):
if not isinstance(node, ast.Call):
continue
for kw in node.keywords:
if (kw.arg == "op_name" and isinstance(kw.value, ast.Constant)
and kw.value.value == op):
start = _attr_name(node.func)
break
if start:
break
if not start:
return []
graph = _class_bases(test_py)
for name, bases in _class_bases(test_py.parent / "base.py").items():
graph.setdefault(name, bases)
chain: List[str] = []
seen: set = set()
queue = [start]
while queue:
name = queue.pop(0)
if not name or name in seen:
continue
seen.add(name)
chain.append(name)
queue.extend(graph.get(name, []))
return chain
def _test_module_path(argv: List[str]) -> Optional[Path]:
"""The benchmark module under test, from pytest's own argv."""
for a in argv:
if a.startswith("-"):
continue
p = Path(a.split("::", 1)[0])
if p.suffix == ".py" and p.is_file():
return p
return None
def _read_shape_yaml(path: str, op: str, test_py: Optional[Path] = None
) -> Tuple[List[Any], Dict[str, Any], str]:
"""Shapes for this op, the sibling keys to carry into each shard, and the
yaml key they came from ("" when nothing matched).
Tries `op` first, then the benchmark class chain, mirroring FlagGems. Without
the class fallback every BLAS op (mm/addmm/bmm/baddbmm, whose shapes live
under `BlasBenchmark`) resolved to zero shapes and sharding silently turned
itself off.
"""
import yaml import yaml
with open(path) as f: with open(path) as f:
doc = yaml.safe_load(f) or {} doc = yaml.safe_load(f) or {}
entry = doc.get(op) or {} keys = [op]
shapes = entry.get("shapes") or [] if not isinstance(doc.get(op), dict) and test_py is not None:
extra = {k: v for k, v in entry.items() if k != "shapes"} keys += _class_chain(test_py, op)
return list(shapes), extra for key in keys:
entry = doc.get(key)
if isinstance(entry, dict) and entry.get("shapes"):
extra = {k: v for k, v in entry.items() if k != "shapes"}
return list(entry["shapes"]), extra, key
return [], {}, ""
def _shard(shapes: List[Any], n: int) -> List[List[Any]]: def _shard(shapes: List[Any], n: int) -> List[List[Any]]:
@@ -111,8 +213,8 @@ def _spawn(shards: List[List[Any]], op: str, extra: Dict[str, Any],
) -> Tuple[List[Path], List[int]]: ) -> Tuple[List[Path], List[int]]:
"""One subprocess per shard, each pinned to its own GPU. """One subprocess per shard, each pinned to its own GPU.
Subprocesses rather than pytest-xdist: _device_guard_plugin initializes a Subprocesses rather than pytest-xdist: `_device_setup_plugin` may initialize
CUDA context at import time, which does not survive fork-based parallelism. a device context at import time, which does not survive fork-based parallelism.
""" """
import yaml import yaml
procs: List[Tuple[int, subprocess.Popen, Any]] = [] procs: List[Tuple[int, subprocess.Popen, Any]] = []
@@ -143,7 +245,8 @@ def _spawn(shards: List[List[Any]], op: str, extra: Dict[str, Any],
child_argv += ["--shape_file", str(w / "shapes.yaml")] child_argv += ["--shape_file", str(w / "shapes.yaml")]
env = dict(os.environ) env = dict(os.environ)
env["CUDA_VISIBLE_DEVICES"] = devices[i] from backends import get_backend
env = get_backend().apply_visibility(env, devices[i])
env[_WORKER_ENV] = "1" # stops recursion env[_WORKER_ENV] = "1" # stops recursion
env[_RECORD_DIR_ENV] = str(rec) # shard records its own picks env[_RECORD_DIR_ENV] = str(rec) # shard records its own picks
env.pop(_REPLAY_DIR_ENV, None) # a shard must sweep env.pop(_REPLAY_DIR_ENV, None) # a shard must sweep
@@ -222,6 +325,23 @@ def pytest_configure(config):
return return
op = _op_name() op = _op_name()
# Only --level core is shardable. At comprehensive, FlagGems' set_shapes()
# merges set_more_shapes() on top of whatever the yaml holds, and that
# addition is arbitrary per-subclass Python (hardcoded tuples plus
# models_shapes/*.yaml) which this plugin cannot see. Sharding the yaml list
# alone would sweep a strict subset -- measured on mm: 5 of 29 shapes, 90
# AUTOTUNE_REPLAY_FALLBACK reason=key_missing -- and every uncovered key
# then self-tunes while the run reports itself as replaying. Refusing costs
# wall time; sharding a subset costs the same-config premise.
level = str(getattr(config.option, "level", "") or "").strip().lower()
if level and level != "core":
print(f"{tag('[parallel-warmup-plugin]')} --level {level} merges shapes "
"this plugin cannot enumerate (set_more_shapes); sharding a subset "
"would break the same-config premise, so sweeping serially instead",
flush=True)
return
shape_file = getattr(config.option, "shape_file", "") or "" shape_file = getattr(config.option, "shape_file", "") or ""
if not shape_file or not Path(shape_file).is_file(): if not shape_file or not Path(shape_file).is_file():
print(f"{tag('[parallel-warmup-plugin]')} no --shape_file; " print(f"{tag('[parallel-warmup-plugin]')} no --shape_file; "
@@ -232,15 +352,24 @@ def pytest_configure(config):
"already pinned, nothing to sweep, skipping", flush=True) "already pinned, nothing to sweep, skipping", flush=True)
return return
test_py = _test_module_path(list(config.invocation_params.args))
try: try:
shapes, extra = _read_shape_yaml(shape_file, op) shapes, extra, key = _read_shape_yaml(shape_file, op, test_py)
except Exception as exc: except Exception as exc:
print(f"{tag('[parallel-warmup-plugin]')} cannot read shapes " print(f"{tag('[parallel-warmup-plugin]')} cannot read shapes "
f"({type(exc).__name__}: {exc}); skipping", flush=True) f"({type(exc).__name__}: {exc}); skipping", flush=True)
return return
if not key:
# Distinct from "found a short list": nothing in the yaml describes this
# op at all, so the shard files would be built from an empty set.
print(f"{tag('[parallel-warmup-plugin]')} no yaml entry for {op!r} "
f"(tried: {', '.join([op] + _class_chain(test_py, op)) if test_py else op}); "
"skipping", flush=True)
return
if len(shapes) < 2: if len(shapes) < 2:
print(f"{tag('[parallel-warmup-plugin]')} only {len(shapes)} shape(s); " print(f"{tag('[parallel-warmup-plugin]')} only {len(shapes)} shape(s) "
"sharding would not pay off, skipping", flush=True) f"under yaml key {key!r}; sharding would not pay off, skipping",
flush=True)
return return
devices = _visible_devices() devices = _visible_devices()
@@ -251,8 +380,9 @@ def pytest_configure(config):
shards = _shard(shapes, min(gpus, len(shapes), len(devices))) shards = _shard(shapes, min(gpus, len(shapes), len(devices)))
scratch = Path(tempfile.mkdtemp(prefix=f"warmup_{op}_")) scratch = Path(tempfile.mkdtemp(prefix=f"warmup_{op}_"))
print(f"{tag('[parallel-warmup-plugin]')} sweeping {len(shapes)} shapes on " print(f"{tag('[parallel-warmup-plugin]')} sweeping {len(shapes)} shapes "
f"{len(shards)} GPUs (shards: {[len(s) for s in shards]}); this phase's " f"(yaml key {key!r}) on {len(shards)} GPUs "
f"(shards: {[len(s) for s in shards]}); this phase's "
"latency is discarded, only configs are kept", flush=True) "latency is discarded, only configs are kept", flush=True)
keep_scratch = False keep_scratch = False
@@ -260,9 +390,23 @@ def pytest_configure(config):
try: try:
recs, bad = _spawn(shards, op, extra, scratch, recs, bad = _spawn(shards, op, extra, scratch,
list(config.invocation_params.args), devices) list(config.invocation_params.args), devices)
merged, total, conflicts = _merge(recs, op)
elapsed = time.time() - t0 elapsed = time.time() - t0
if bad:
# A dead shard means its shapes were never swept, so publishing the
# merge would hand the measurement a config set with holes; every
# hole then silently self-tunes under a run that claims to replay.
# Discard the whole sweep rather than publish a partial one.
keep_scratch = True
print(f"{tag('[parallel-warmup-plugin]')} shard(s) {bad} exited "
f"non-zero after {elapsed:.0f}s; their shapes were never swept, "
"so the merged configs would have holes. Discarding this sweep "
f"and falling back to normal serial autotune (worker logs kept "
f"in {scratch})", flush=True)
return
merged, total, conflicts = _merge(recs, op)
if not total: if not total:
keep_scratch = True keep_scratch = True
print(f"{tag('[parallel-warmup-plugin]')} no configs recovered in " print(f"{tag('[parallel-warmup-plugin]')} no configs recovered in "
@@ -287,9 +431,6 @@ def pytest_configure(config):
msg = (f"{tag('[parallel-warmup-plugin]')} sweep done in {elapsed:.0f}s: " msg = (f"{tag('[parallel-warmup-plugin]')} sweep done in {elapsed:.0f}s: "
f"{len(merged)} kernels / {total} configs -> {out / f'{op}.json'}; " f"{len(merged)} kernels / {total} configs -> {out / f'{op}.json'}; "
"measurement continues serially on one GPU") "measurement continues serially on one GPU")
if bad:
msg += f" (WARNING shard(s) {bad} exited non-zero; missing keys will "
msg += "fall back to live autotune)"
print(msg, flush=True) print(msg, flush=True)
if conflicts: if conflicts:
print(f"{tag('[parallel-warmup-plugin]')} WARNING {conflicts} config " print(f"{tag('[parallel-warmup-plugin]')} WARNING {conflicts} config "
_run_identity.py
+96
View File
@@ -0,0 +1,96 @@
"""Merge-update `$OUT_DIR/run_identity.json` so plugins can record backend facts.
Shell writes a stub before pytest; plugins add device_name / timer / vendor
without clobbering earlier keys. Missing OUT_DIR is a silent no-op so unit
tests and ad-hoc pytest invocations do not fail.
"""
from __future__ import annotations
import argparse
import json
import os
from pathlib import Path
from typing import Any, Dict, Optional
IDENTITY_NAME = "run_identity.json"
def identity_path(out_dir: Optional[str] = None) -> Optional[Path]:
root = out_dir or os.environ.get("OUT_DIR", "").strip()
if not root:
return None
return Path(root) / IDENTITY_NAME
def load_identity(out_dir: Optional[str] = None) -> Dict[str, Any]:
path = identity_path(out_dir)
if path is None or not path.is_file():
return {}
try:
data = json.loads(path.read_text(encoding="utf-8"))
except Exception:
return {}
return data if isinstance(data, dict) else {}
def update_identity(out_dir: Optional[str] = None, **fields: Any) -> Dict[str, Any]:
path = identity_path(out_dir)
if path is None:
return dict(fields)
path.parent.mkdir(parents=True, exist_ok=True)
current = load_identity(out_dir)
for key, value in fields.items():
if value is None:
continue
current[key] = value
tmp = path.with_name(path.name + ".tmp")
tmp.write_text(json.dumps(current, indent=2, ensure_ascii=False) + "\n",
encoding="utf-8")
os.replace(tmp, path)
return current
def init_identity(out_dir: str, **fields: Any) -> Dict[str, Any]:
payload = {"schema_version": 1, "provider": "flaggems"}
payload.update(fields)
os.environ["OUT_DIR"] = str(out_dir)
return update_identity(out_dir, **payload)
def main(argv=None) -> int:
p = argparse.ArgumentParser()
p.add_argument("out_dir")
p.add_argument("--backend", default=os.environ.get("ZL_BACKEND", "cuda"))
p.add_argument("--op", default=os.environ.get("OP", ""))
p.add_argument("--op-file", default=os.environ.get("OP_FILE", ""))
p.add_argument("--flaggems-dir", default=os.environ.get("FLAGGEMS_DIR", ""))
p.add_argument("--autotune-mode", default="")
p.add_argument("--device", default=os.environ.get("ZL_DEVICE", ""))
p.add_argument("--test-file", default=os.environ.get("TEST_FILE", ""))
p.add_argument("--shape-mode", default=os.environ.get("SHAPE_MODE", ""))
p.add_argument("--bench-level", default=os.environ.get("BENCH_LEVEL", "core"))
# "isolated" or "shared(<path>)". Recorded because a shared libtuner cache
# lets whichever side runs first hand its winners to the other, and after
# the fact the two runs are otherwise indistinguishable.
p.add_argument("--tuner-cache", default="")
args = p.parse_args(argv)
init_identity(
args.out_dir,
backend=args.backend,
op=args.op,
op_file=args.op_file,
flaggems_dir=args.flaggems_dir,
autotune_mode=args.autotune_mode,
device=args.device,
test_file=args.test_file,
shape_mode=args.shape_mode,
bench_level=args.bench_level,
tuner_cache=args.tuner_cache,
)
return 0
if __name__ == "__main__":
raise SystemExit(main())
_timing_plugin.py
+65
View File
@@ -0,0 +1,65 @@
"""Select the timing path from backend capabilities, before any shape runs.
CUDA keeps the existing `_cudagraph_plugin` behaviour (graph capture, with
per-shape fallback to plain do_bench). Other backends do not attempt CUDA
Graph at all: they leave `triton.testing.do_bench` in place and record
`timer=direct`. Trying capture-then-fallback on a non-CUDA device would
produce FATAL_DEVICE noise and mixed-caliber rows.
"""
from __future__ import annotations
import os
import sys
from backends import get_backend
from _run_identity import update_identity
def _graph_requested(backend) -> bool:
raw = os.environ.get("ZL_ENABLE_CUDAGRAPH", "auto").strip().lower()
if raw in ("0", "false", "no", "off", "never", "direct"):
return False
if raw in ("1", "true", "yes", "on", "always"):
return bool(backend.supports_cuda_graph)
return bool(backend.supports_cuda_graph)
def pytest_configure(config):
from _term_style import tag
backend = get_backend()
raw = os.environ.get("ZL_ENABLE_CUDAGRAPH", "auto").strip().lower()
force_on = raw in ("1", "true", "yes", "on", "always")
if force_on and not backend.supports_cuda_graph:
print(f"{tag('[timing-plugin]')} ZL_ENABLE_CUDAGRAPH=1 ignored on "
f"backend={backend.name} (no CUDA Graph); using timer=direct",
file=sys.stderr, flush=True)
if _graph_requested(backend):
import _cudagraph_plugin as _cg
_cg.pytest_configure(config)
# Recorded on this branch too: the fallback rows a captured run produces
# are timed by the direct path and take the same repeat count.
update_identity(timer="cuda_graph", graph_capture="attempted",
direct_repeat=_cg.direct_repeat())
print(f"{tag('[timing-plugin]')} timer=cuda_graph backend={backend.name}",
file=sys.stderr, flush=True)
return
# A wholly-direct run gets the same repetition the fallback rows get, so the
# "put everything on one timer" advice does not silently cost the samples.
# Patching only when repeat > 1 keeps the default path untouched: upstream's
# do_bench stays in place exactly as before.
import _cudagraph_plugin as _cg
repeat = _cg.direct_repeat()
if repeat > 1:
def _repeated(fn, warmup=25, rep=100, grad_to_none=None,
quantiles=None, return_mode="mean"):
return _cg.run_direct(fn, warmup, rep, grad_to_none,
quantiles, return_mode)
import triton.testing as _tt
_tt.do_bench = _repeated
update_identity(timer="direct", graph_capture="disabled",
direct_repeat=repeat)
print(f"{tag('[timing-plugin]')} timer=direct backend={backend.name} "
f"(upstream do_bench; graph capture not selected)"
+ (f"; median of {repeat} calls" if repeat > 1 else ""),
file=sys.stderr, flush=True)
ab_repeat.sh
+129
View File
@@ -0,0 +1,129 @@
#!/bin/bash
# 多轮 A/B：同一对比在 N 轮独立进程里各跑一次，逐 shape 取轮次中位。
#
# 为什么需要它：单对观测无法说明自己稳不稳。S5000 实测 topk 消除了 32 个
# convert_layout，42 个 shape 里却有 17 个在两次同等有效的运行之间改变了 delta
# 的符号——那不是计时口径的问题（图捕获与 direct 的延迟比是 1.00x），是收益本身
# 落在零附近、单次测量给不出方向。多轮之后才谈得上"这个判定成立"。
#
# 口径：第 1 轮 A 侧 record autotune config，此后每一次运行（含 A 侧自己的后续
# 轮次）都 REPLAY_FROM 那一份。config 因此跨轮跨侧全部锁死，剩下的变量只有被测
# 开关与测量噪声。逐轮交替先后顺序，抵消单向漂移。
#
# 用法：
# SWITCH=FLAGTREE_MUSA_RLC_ENHANCE OPS="topk masked_select" ROUNDS=6 \
# ZL_BACKEND=musa FLAGGEMS_DIR=/path/to/FlagGems bash ab_repeat.sh
#
# OPS 元素为 OP 或 OP:OP_FILE（省略时 OP_FILE=OP）。
# SWITCH 是被测环境变量，A 侧设 0、B 侧设 1。
#
# 对比轴有效性由脚本自己查：第 1 轮 A/B 跑完后比两侧的 ttgir 指纹，完全一致说明
# 这个开关在这个算子上没改变任何生成代码——继续跑满 N 轮只会得到一个"无差异"的假
# 结论（同 ab_fold_test.sh 的告诫）。这类算子默认跳过；开关确实只影响运行时而不影响
# IR 时，用 SWITCH_ALLOW_NO_IR_DIFF=1 放行。
#
# 产物：runs/ab_repeat_<开关>_<时间戳>/<op>/r<k>_<A|B>/，跑完对每个算子调用
# batch_summary.py --ab-rounds 出报告。
set -uo pipefail
cd "$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
SWITCH_ALLOW_NO_IR_DIFF="${SWITCH_ALLOW_NO_IR_DIFF:-0}"
# 一个 run 生成代码的指纹：ttgir 相对路径 + 内容。路径进指纹是因为"变体集合变了"
# 与"某个变体内容变了"同样算改变，而只 cat 内容会把前者洗掉。
ir_fingerprint() {
local dir="$1/ttgir"
[[ -d "$dir" ]] || return 1
local out
out=$( { find "$dir" -type f | LC_ALL=C sort | while IFS= read -r f; do
printf '%s
' "${f#"$dir"/}"; cat "$f"
done; } | sha256sum | cut -d' ' -f1 )
[[ -n "$out" ]] && printf '%s' "$out"
}
: "${SWITCH:?SWITCH is required, e.g. FLAGTREE_MUSA_RLC_ENHANCE}"
OPS="${OPS:-}"
: "${OPS:?OPS is required, e.g. \"topk masked_select\"}"
ROUNDS="${ROUNDS:-6}"
if [[ -t 1 && -z "${NO_COLOR:-}" ]]; then
export FLAGGEMS_PERF_COLOR=always
fi
ROOT="runs/ab_repeat_${SWITCH}_$(date +%Y%m%d-%H%M%S)"
mkdir -p "$ROOT"
LOG="$ROOT/summary.log"
{
echo ">>> switch=$SWITCH rounds=$ROUNDS ops=$OPS"
for entry in $OPS; do
op="${entry%%:*}"
op_file="${entry#*:}"
[[ "$op_file" == "$entry" ]] && op_file="$op"
out="$ROOT/$op"
mkdir -p "$out"
# 第 1 轮 A 侧 record，其余全部回放它，config 跨轮跨侧一致
echo "########## $op | round 1 | A ($SWITCH=0, record) ##########"
if ! OP="$op" OP_FILE="$op_file" OUT_DIR="$out/r1_A" \
env "$SWITCH=0" bash run_pytest.sh; then
echo "!!! FAILED $op round 1 A; skipping this operator"
continue
fi
REC="$out/r1_A"
ok=1; reason="有失败轮次"
for ((k = 1; k <= ROUNDS; k++)); do
# 奇数轮先 A 后 B，偶数轮反过来
if (( k % 2 == 1 )); then order="A B"; else order="B A"; fi
for side in $order; do
[[ $k -eq 1 && $side == A ]] && continue # 已经跑过
val=$([[ $side == A ]] && echo 0 || echo 1)
echo "########## $op | round $k | $side ($SWITCH=$val, replay) ##########"
if ! OP="$op" OP_FILE="$op_file" OUT_DIR="$out/r${k}_${side}" \
REPLAY_FROM="$REC" env "$SWITCH=$val" bash run_pytest.sh; then
echo "!!! FAILED $op round $k $side"
ok=0
fi
done
if (( k == 1 && ok )); then
fa=$(ir_fingerprint "$out/r1_A") || fa=""
fb=$(ir_fingerprint "$out/r1_B") || fb=""
if [[ -z "$fa" || -z "$fb" ]]; then
echo "??? $op: 没有 ttgir 产物，跳过对比轴有效性检查"
elif [[ "$fa" == "$fb" ]]; then
echo "!!! SWITCH_NO_IR_DIFF op=$op switch=$SWITCH"
echo "!!! 两侧 ttgir 逐字节相同：这个开关没有改变 $op 的生成代码。"
echo "!!! 继续跑 $ROUNDS 轮只会测出噪声，并报成一个'无差异'结论。"
echo "!!! 先确认开关名拼写与它在被测源码里的读取点；确属只影响运行时的"
echo "!!! 开关，用 SWITCH_ALLOW_NO_IR_DIFF=1 重跑。"
if [[ "$SWITCH_ALLOW_NO_IR_DIFF" != "1" ]]; then
ok=0; reason="开关未改变该算子的生成代码"
break
fi
else
echo ">>> 对比轴有效：$op 两侧 ttgir 不同（A=${fa:0:12} B=${fb:0:12}）"
fi
fi
done
echo "########## $op | report ##########"
if (( ok )); then
python batch_summary.py --ab-rounds "$out" --label "$op" \
|| echo "!!! report failed: $op"
else
# 缺轮的算子照样出报告会让人误以为是完整的 N 轮
echo "!!! $op $reason，不出报告；产物在 $out"
fi
done
echo "########## ALL DONE ##########"
} 2>&1 | tee "$LOG"
sed -i -E $'s/\x1b\\[[0-9;]*[A-Za-z]//g' "$LOG"
echo ">>> 总 log: $(pwd)/$LOG"
echo ">>> 逐算子重出报告: python batch_summary.py --ab-rounds $ROOT/<op> --label <op>"
echo ">>> 解读要点: 逐轮区间宽说明判定不稳；单侧噪声大说明机器不稳；两者含义不同"
ab_rlc_test.sh
+68
View File
@@ -0,0 +1,68 @@
#!/bin/bash
# A/B 对比：FlagTree RemoveLayoutConversions 增强关 vs 开，逐算子两轮。
#
# 口径：每个算子先跑关（record autotune config），再用 REPLAY_FROM 回放同一套
# config 跑开，两侧只差 FLAGTREE_<VENDOR>_RLC_ENHANCE 这一个开关。产物沿用
# run_pytest.sh 的目录结构，所以每个算子恰好留下两份 run.log：
#
# runs/<op>_<时间戳>/run.log 关（before）
# runs/<op>_<时间戳>/run.log 开（after）
#
# 跑完用 batch_summary.py --ab <before> <after> 得到该算子的 speedup 区间。
#
# 用法：
# OPS="softmax gelu_backward" bash ab_rlc_test.sh
# OPS="sort:perf_sort topk" bash ab_rlc_test.sh # 冒号分隔 OP:OP_FILE
#
# 注意：RLC 开关名带厂商前缀，用 VENDOR 指定（默认 MUSA）。套用前先确认开关有效：
# grep 该开关名于 FlagTree 源码应有命中，否则两侧跑的是同一份代码。
set -uo pipefail
cd "$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
VENDOR="${VENDOR:-MUSA}"
ENHANCE_VAR="FLAGTREE_${VENDOR}_RLC_ENHANCE"
# 空格分隔的算子表。元素为 OP 或 OP:OP_FILE（省略时 OP_FILE=OP）。
OPS="${OPS:-softmax}"
if [[ -t 1 && -z "${NO_COLOR:-}" ]]; then
export FLAGGEMS_PERF_COLOR=always
fi
LOG="runs/ab_rlc_$(date +%Y%m%d-%H%M%S).log"
mkdir -p runs
{
echo ">>> vendor=$VENDOR switch=$ENHANCE_VAR ops=$OPS"
for entry in $OPS; do
op="${entry%%:*}"
op_file="${entry#*:}"
[[ "$op_file" == "$entry" ]] && op_file="$op"
echo "########## $op | $ENHANCE_VAR=0 before (record) ##########"
if ! OP="$op" OP_FILE="$op_file" env "$ENHANCE_VAR=0" bash run_pytest.sh; then
echo "!!! FAILED before: $op"
continue
fi
BEFORE=$(ls -dt "runs/${op}_"* | head -1)
echo "########## $op | $ENHANCE_VAR=1 after (replay $BEFORE) ##########"
if ! OP="$op" OP_FILE="$op_file" REPLAY_FROM="$BEFORE" env "$ENHANCE_VAR=1" bash run_pytest.sh; then
echo "!!! FAILED after: $op"
continue
fi
AFTER=$(ls -dt "runs/${op}_"* | head -1)
echo "########## $op | summary ##########"
python batch_summary.py --ab "$BEFORE/run.log" "$AFTER/run.log" --label "$op" \
|| echo "!!! summary failed: $op"
done
echo "########## ALL DONE ##########"
} 2>&1 | tee "$LOG"
# 去掉 ANSI 转义以便 grep/diff（终端输出保留颜色）
sed -i -E $'s/\x1b\\[[0-9;]*[A-Za-z]//g' "$LOG"
echo ">>> 总 log: $(pwd)/$LOG"
echo ">>> 每算子两份 run.log 在 runs/<op>_<时间戳>/ 下，按时间先后即 before / after"
backends/__init__.py
+76
View File
@@ -0,0 +1,76 @@
"""Registry of zl_bench device backends.
Aliases map common vendor spellings onto the canonical backend name. Unknown
names fail closed rather than defaulting to CUDA: a silent CUDA fallback would
load the cudagraph plugin on the wrong device.
"""
from __future__ import annotations
import os
from typing import Dict, Optional
from .ascend import AscendBackend
from .base import DeviceBackend, parse_device_list
from .cuda import CudaBackend
from .musa import MusaBackend
from .rocm import RocmBackend
ALIASES: Dict[str, str] = {
"nvidia": "cuda",
"nv": "cuda",
"amd": "rocm",
"hip": "rocm",
"npu": "ascend",
"huawei": "ascend",
}
_INSTANCES: Dict[str, DeviceBackend] = {
"cuda": CudaBackend(),
"musa": MusaBackend(),
"rocm": RocmBackend(),
"ascend": AscendBackend(),
}
class UnknownBackendError(ValueError):
pass
def known_backends() -> tuple[str, ...]:
return tuple(sorted(_INSTANCES))
def normalize_backend_name(name: str) -> str:
key = (name or "").strip().lower()
if not key:
return "cuda"
return ALIASES.get(key, key)
def get_backend(name: Optional[str] = None) -> DeviceBackend:
raw = name if name is not None else os.environ.get("ZL_BACKEND", "cuda")
key = normalize_backend_name(raw)
backend = _INSTANCES.get(key)
if backend is None:
raise UnknownBackendError(
f"unknown backend {raw!r}; known: {', '.join(known_backends())} "
f"(aliases: {', '.join(sorted(ALIASES))})"
)
return backend
def enumerate_devices(backend: Optional[DeviceBackend] = None,
explicit: Optional[str] = None) -> list[str]:
return (backend or get_backend()).enumerate_devices(explicit)
__all__ = [
"DeviceBackend",
"UnknownBackendError",
"enumerate_devices",
"get_backend",
"known_backends",
"normalize_backend_name",
"parse_device_list",
]
backends/ascend.py
+17
View File
@@ -0,0 +1,17 @@
from __future__ import annotations
from .base import DeviceBackend
class AscendBackend(DeviceBackend):
name = "ascend"
visibility_env_keys = (
"ASCEND_RT_VISIBLE_DEVICES",
"ASCEND_VISIBLE_DEVICES",
)
torch_namespace = "npu"
flaggems_vendor = "ascend"
supports_cuda_graph = False
supports_mm_cluster_fix = False
supports_triton_hooks = True
list_command = None
backends/base.py
+193
View File
@@ -0,0 +1,193 @@
"""Device backend contract for zl_bench screening runs.
This is the orchestration-layer adapter: device enumeration, process visibility
env, and capability flags that decide which pytest plugins are safe to load.
It does not implement kernel timing receipts or compiler-artifact identity;
those stay with the caller (FlagGems pytest for screening, a separate harness
for qualification).
"""
from __future__ import annotations
import os
import subprocess
from typing import Dict, Iterable, List, Optional, Sequence, Tuple
# Every vendor visibility variable zl_bench knows about. apply_visibility()
# clears the whole set before writing the active backend's keys so a parent
# CUDA_VISIBLE_DEVICES cannot leak onto a MUSA/Ascend child.
ALL_VISIBILITY_KEYS: Tuple[str, ...] = (
"CUDA_VISIBLE_DEVICES",
"MUSA_VISIBLE_DEVICES",
"HIP_VISIBLE_DEVICES",
"ROCR_VISIBLE_DEVICES",
"ASCEND_RT_VISIBLE_DEVICES",
"ASCEND_VISIBLE_DEVICES",
)
def parse_device_list(raw: Optional[str]) -> List[str]:
if not raw:
return []
return [tok.strip() for tok in raw.split(",") if tok.strip()]
class DeviceBackend:
"""One hardware runtime (CUDA, MUSA, ROCm, Ascend, ...).
Subclasses only fill class attributes and optional probe hooks. Callers
should not branch on ``backend.name`` for visibility or plugin policy;
they should read these attributes / methods.
"""
name: str = ""
# First key is the canonical one written into run_identity.json.
visibility_env_keys: Tuple[str, ...] = ()
# torch.<namespace> used for optional device_count / identity probes.
torch_namespace: str = ""
# FlagGems GEMS_VENDOR value. Lives on the backend object because zl_bench
# still only drives FlagGems, but the mapping is owned by providers.flaggems
# and copied here so DeviceBackend.prepare_environment() can stay generic.
flaggems_vendor: str = ""
supports_cuda_graph: bool = False
supports_mm_cluster_fix: bool = False
supports_triton_hooks: bool = True
# --- graph-capture failure classification -----------------------------
# A kernel that refuses capture is normal and falls back. A device error is
# not, and the two are told apart by vendor error text, so the wording lives
# on the backend rather than in the timing plugin.
#
# fatal : the kernel itself failed on the device. Re-raise; recording a
# fallback latency here would report a broken kernel as a result.
# poisoning: the capture ATTEMPT left the context unusable. Falling back runs
# the measurement inside a dead process, which is how an operator
# ends up emitting a fallback marker and then still dying at
# teardown with zero rows. Re-raise too, but under its own marker,
# because the operator is capture-incompatible rather than broken:
# the actionable response is to re-run it with ZL_ENABLE_CUDAGRAPH=0.
fatal_error_markers: Tuple[str, ...] = ()
poisoning_error_markers: Tuple[str, ...] = ()
# Operators whose capture attempt is known to poison the context on this
# backend. Skipped up front so one of them cannot cost a whole run.
capture_deny_ops: Tuple[str, ...] = ()
def classify_capture_error(self, exc: BaseException) -> str:
"""-> "fatal" | "poisoned" | "recoverable".
Matched against the exception text itself, never a generic footer: torch
appends "Compile with `TORCH_USE_CUDA_DSA` to enable device-side
assertions." to every AcceleratorError, so a loose "device-side assert"
marker would classify every capture refusal as fatal and disable the
documented fallback path outright.
"""
msg = f"{type(exc).__name__}: {exc}".lower()
if any(m in msg for m in self.fatal_error_markers):
return "fatal"
if any(m in msg for m in self.poisoning_error_markers):
return "poisoned"
return "recoverable"
def denies_capture(self, op: str) -> bool:
return bool(op) and op in self.capture_deny_ops
# Plugins that patch around upstream assumptions this backend violates, e.g.
# FlagGems benchmarks that gate themselves on torch.cuda.is_available() even
# though the operator under test runs on whatever device FlagGems dispatched
# to. Loaded first so collection-time overrides land before anything reads
# the item list. Each such plugin must scope itself to its own vendor.
compat_plugins: Tuple[str, ...] = ()
# Optional listing command; (argv, line_prefix). prefix="" means count
# non-empty lines. Must be a real vendor CLI; do not invent flags.
list_command: Optional[Tuple[Sequence[str], str]] = None
def visibility_env(self, device: str) -> Dict[str, str]:
return {key: device for key in self.visibility_env_keys}
def apply_visibility(self, env: Dict[str, str], device: str) -> Dict[str, str]:
out = dict(env)
for key in ALL_VISIBILITY_KEYS:
out.pop(key, None)
out.update(self.visibility_env(device))
return out
def enumerate_devices(self, explicit: Optional[str] = None) -> List[str]:
listed = parse_device_list(explicit)
if listed:
return listed
for key in self.visibility_env_keys:
listed = parse_device_list(os.environ.get(key, ""))
if listed:
return listed
n = self.probe_device_count()
if n > 0:
return [str(i) for i in range(n)]
return ["0"]
def probe_device_count(self) -> int:
n = self._probe_list_command()
if n > 0:
return n
return self._probe_torch_count()
def _probe_list_command(self) -> int:
if not self.list_command:
return 0
argv, prefix = self.list_command
try:
proc = subprocess.run(
list(argv), text=True, capture_output=True, timeout=30,
)
except Exception:
return 0
if proc.returncode != 0:
return 0
lines = [ln for ln in proc.stdout.splitlines() if ln.strip()]
if prefix:
lines = [ln for ln in lines if ln.startswith(prefix)]
return len(lines)
def _probe_torch_count(self) -> int:
if not self.torch_namespace:
return 0
try:
import torch
api = getattr(torch, self.torch_namespace, None)
if api is None or not hasattr(api, "device_count"):
return 0
return int(api.device_count())
except Exception:
return 0
def device_identity(self) -> Dict[str, str]:
"""Best-effort identity for screening logs. Empty strings are allowed."""
info = {
"backend": self.name,
"device_name": "",
"device_arch": "",
"timer_api": f"torch.{self.torch_namespace}" if self.torch_namespace else "",
"flaggems_vendor": self.flaggems_vendor,
}
if not self.torch_namespace:
return info
try:
import torch
api = getattr(torch, self.torch_namespace, None)
if api is None:
return info
count = int(api.device_count()) if hasattr(api, "device_count") else 0
if count <= 0:
return info
index = int(api.current_device()) if hasattr(api, "current_device") else 0
if hasattr(api, "get_device_name"):
info["device_name"] = str(api.get_device_name(index))
if hasattr(api, "get_device_capability"):
cap = api.get_device_capability(index)
info["device_arch"] = f"{cap[0]}.{cap[1]}"
elif hasattr(api, "get_device_properties"):
props = api.get_device_properties(index)
arch = getattr(props, "name", "") or getattr(props, "arch", "")
info["device_arch"] = str(arch)
except Exception:
return info
return info
backends/cuda.py
+28
View File
@@ -0,0 +1,28 @@
from __future__ import annotations
from .base import DeviceBackend
class CudaBackend(DeviceBackend):
name = "cuda"
visibility_env_keys = ("CUDA_VISIBLE_DEVICES",)
torch_namespace = "cuda"
flaggems_vendor = "nvidia"
supports_cuda_graph = True
supports_mm_cluster_fix = True
supports_triton_hooks = True
list_command = (("nvidia-smi", "--list-gpus"), "GPU ")
# Unchanged from when this list lived in _cudagraph_plugin; CUDA numbers do
# not move because of the classification refactor.
fatal_error_markers = (
"illegal memory access",
"cudaerrorillegaladdress",
"out of memory",
"device-side assert triggered",
"an illegal instruction",
"misaligned address",
"uncorrectable ecc",
)
# No CUDA wording is known to poison the context without also matching one
# of the fatal markers above, so nothing is listed here rather than guessed.
poisoning_error_markers = ()
backends/musa.py
+52
View File
@@ -0,0 +1,52 @@
from __future__ import annotations
from .base import DeviceBackend
class MusaBackend(DeviceBackend):
name = "musa"
visibility_env_keys = ("MUSA_VISIBLE_DEVICES",)
torch_namespace = "musa"
flaggems_vendor = "musa"
# torch_musa supplies Stream, Event, synchronize, graph and MUSAGraph, so
# capture works once the timing path takes them from Triton's device
# interface instead of torch.cuda. Measured on an S5000 over 15 operators:
# 11 captured, 3 fell back cleanly with a marker (dropout/rand on the RNG
# seed, sort on a capture-illegal stream op), 1 poisoned the context (see
# capture_deny_ops). Capture is worth having because it removes launch
# overhead the direct timer cannot: batch_norm 16x16x8x48 went 0.0758ms ->
# 0.0089ms, tanh/gelu 4096x4096 about 3.3x, while GEMMs stayed at 1.00x.
supports_cuda_graph = True
supports_mm_cluster_fix = False
supports_triton_hooks = True
# MUSA reports its own wording; none of the CUDA markers match it.
fatal_error_markers = (
"musa error: an illegal memory access",
"out of memory",
"device-side assert triggered",
"misaligned address",
"uncorrectable ecc",
)
# "unknown error" during capture is not recoverable on this runtime: the
# fallback runs, and the process then dies in teardown with zero rows. Treat
# it as poisoning so the operator is reported as capture-incompatible rather
# than silently losing its measurement.
poisoning_error_markers = ("musa error: unknown error",)
# Keyed by FLAGGEMS_PERF_CURRENT_OP, which is the pytest FUNCTION name minus
# `test_` -- `unique2`, `perf_sort` -- not the benchmark file name. Getting
# that wrong is silent: the entry simply never matches.
#
# Observed on an S5000: capturing test_unique2 takes two recoverable
# fallbacks and then raises "MUSA error: unknown error", ending the run at
# rc=1 with no measured rows. Skipping the attempt costs that operator's
# timing precision instead of its measurement. Revisit if a newer torch_musa
# changes the behaviour.
capture_deny_ops = ("unique2",)
# Several FlagGems benchmarks skip on `not torch.cuda.is_available()` while
# the code under test dispatches through flag_gems.device. On MThreads that
# turns a working operator into "not measured"; the plugin drops those
# guards, pinned to one file and one exact skip reason each.
compat_plugins = ("_musa_benchmark_compat_plugin",)
# Device count comes from torch.musa or MUSA_VISIBLE_DEVICES. Do not guess
# a vendor CLI here; confirm the real listing tool before adding it.
list_command = None
backends/rocm.py
+16
View File
@@ -0,0 +1,16 @@
from __future__ import annotations
from .base import DeviceBackend
class RocmBackend(DeviceBackend):
name = "rocm"
# HIP runtimes honor both; write both so a parent CUDA/HIP mix cannot
# leave the child looking at the wrong cards.
visibility_env_keys = ("HIP_VISIBLE_DEVICES", "ROCR_VISIBLE_DEVICES")
torch_namespace = "cuda" # ROCm PyTorch still exposes torch.cuda
flaggems_vendor = "amd"
supports_cuda_graph = False
supports_mm_cluster_fix = False
supports_triton_hooks = True
list_command = None
batch_summary.py
+337 -4
View File
@@ -16,8 +16,10 @@ from __future__ import annotations
import argparse import argparse
import csv import csv
import itertools
import os import os
import re import re
import statistics
import sys import sys
from pathlib import Path from pathlib import Path
from typing import Dict, List, Optional, Tuple from typing import Dict, List, Optional, Tuple
@@ -39,12 +41,31 @@ DIRECT_NO_CUDAGRAPH_RE = re.compile(r"BENCHMARK_DIRECT_NO_CUDAGRAPH\b")
REPLAY_FALLBACK_RE = re.compile(r"AUTOTUNE_REPLAY_FALLBACK\b") REPLAY_FALLBACK_RE = re.compile(r"AUTOTUNE_REPLAY_FALLBACK\b")
SUMMARY_FIELDS = [ SUMMARY_FIELDS = [
"repo", "op", "fg_op_name", "dtype", "case_id", "row_status", "repo", "op", "backend", "device", "timer",
"fg_op_name", "dtype", "case_id", "row_status",
"torch_lat", "gems_lat", "speedup", "extra_metrics", "size_detail", "torch_lat", "gems_lat", "speedup", "extra_metrics", "size_detail",
"op_status", "notes", "op_status", "notes",
] ]
def _identity_from_status(status_row: Dict[str, str], op_dir: Path) -> Dict[str, str]:
"""Backend/device/timer from ops_status.csv, falling back to run_identity.json."""
backend = (status_row.get("backend") or "").strip()
device = (status_row.get("device") or status_row.get("gpu") or "").strip()
timer = (status_row.get("timer") or "").strip()
ident_path = op_dir / "run_identity.json"
if ident_path.is_file() and not (backend and device and timer):
try:
from _run_identity import load_identity
ident = load_identity(str(op_dir))
backend = backend or str(ident.get("backend") or "")
device = device or str(ident.get("device") or ident.get("gpu") or "")
timer = timer or str(ident.get("timer") or "")
except Exception:
pass
return {"backend": backend, "device": device, "timer": timer}
def parse_run_log(text: str) -> List[Dict[str, str]]: def parse_run_log(text: str) -> List[Dict[str, str]]:
"""把 run.log 里的全部结果表解析为行列表（跨多个 dtype 表）。""" """把 run.log 里的全部结果表解析为行列表（跨多个 dtype 表）。"""
rows: List[Dict[str, str]] = [] rows: List[Dict[str, str]] = []
@@ -133,10 +154,11 @@ def summarize(batch_dir: Path, out_path: Optional[Path] = None) -> Path:
text = log.read_text(errors="replace") text = log.read_text(errors="replace")
status_row = status_by_dir.get(f"{repo}/{op_dir.name}", {}) status_row = status_by_dir.get(f"{repo}/{op_dir.name}", {})
op = status_row.get("op") or op_dir.name op = status_row.get("op") or op_dir.name
ident = _identity_from_status(status_row, op_dir)
table_rows = parse_run_log(text) table_rows = parse_run_log(text)
if not table_rows: if not table_rows:
all_rows.append({ all_rows.append({
"repo": repo, "op": op, "fg_op_name": "", "dtype": "", "repo": repo, "op": op, **ident, "fg_op_name": "", "dtype": "",
"case_id": "", "row_status": "", "torch_lat": "", "case_id": "", "row_status": "", "torch_lat": "",
"gems_lat": "", "speedup": "", "extra_metrics": "", "gems_lat": "", "speedup": "", "extra_metrics": "",
"size_detail": "", "size_detail": "",
@@ -150,7 +172,8 @@ def summarize(batch_dir: Path, out_path: Optional[Path] = None) -> Path:
idx = case_counter.get(key, 0) idx = case_counter.get(key, 0)
case_counter[key] = idx + 1 case_counter[key] = idx + 1
all_rows.append({ all_rows.append({
"repo": repo, "op": op, "case_id": f"case_{idx:03d}", "repo": repo, "op": op, **ident,
"case_id": f"case_{idx:03d}",
"op_status": status_row.get("status", ""), "op_status": status_row.get("status", ""),
"notes": status_row.get("notes", ""), "notes": status_row.get("notes", ""),
**row, **row,
@@ -165,12 +188,322 @@ def summarize(batch_dir: Path, out_path: Optional[Path] = None) -> Path:
return out_path return out_path
def ab_rows(before_text: str, after_text: str) -> List[Dict[str, str]]:
"""把两次 run.log 的测量行按 (op, dtype, size_detail) 配对。
speedup 列是 gems 对 torch，不是前后对比，所以这里用 gems_lat 自己算
before/after - 1。只有两侧都 SUCCESS 且延迟为正的行才计入；对不上的行
单独返回，避免用不完整的配对报一个看起来完整的区间。
"""
def index(text: str) -> Dict[Tuple[str, str, str], float]:
out: Dict[Tuple[str, str, str], float] = {}
for row in parse_run_log(text):
if row.get("row_status") != "SUCCESS":
continue
try:
lat = float(row.get("gems_lat", ""))
except ValueError:
continue
if lat > 0:
out[(row.get("fg_op_name", ""), row.get("dtype", ""),
row.get("size_detail", ""))] = lat
return out
before, after = index(before_text), index(after_text)
paired: List[Dict[str, str]] = []
for key in sorted(before.keys() & after.keys()):
op, dtype, size = key
b, a = before[key], after[key]
paired.append({
"fg_op_name": op, "dtype": dtype, "size_detail": size,
"before_lat": f"{b:.6f}", "after_lat": f"{a:.6f}",
"speedup_percent": f"{(b / a - 1.0) * 100.0:.2f}",
})
return paired
def report_ab(before: Path, after: Path, label: str = "",
strict: bool = True) -> List[Dict[str, str]]:
"""打印一个算子的前后对比：逐行 speedup 与整体区间。
先核对两侧口径；不一致时默认拒绝出报告，因为不可比的两次运行会产出
一份和有效报告长得一模一样的表。
"""
if not _report_caliber([("before", before.parent), ("after", after.parent)],
label or "ab", strict):
raise CaliberMismatch()
before_text = before.read_text(encoding="utf-8", errors="replace")
after_text = after.read_text(encoding="utf-8", errors="replace")
rows = ab_rows(before_text, after_text)
name = label or (rows[0]["fg_op_name"] if rows else "?")
if not rows:
print(f"[ab] {name}: 没有可配对的测量行"
f"（before={before} after={after}）")
return rows
values = [float(r["speedup_percent"]) for r in rows]
for r in rows:
print(f" {r['dtype']:<16} {float(r['speedup_percent']):+7.2f}% "
f"{r['before_lat']} -> {r['after_lat']} {r['size_detail']}")
fallback = ""
if DIRECT_NO_CUDAGRAPH_RE.search(before_text) or \
DIRECT_NO_CUDAGRAPH_RE.search(after_text):
# 回退行用的是 do_bench 口径，与 cudagraph 行不同表可比性
fallback = " [含 BENCHMARK_DIRECT_NO_CUDAGRAPH 回退行]"
print(f"[ab] {name}: {len(rows)} 行配对，speedup "
f"{min(values):+.2f}% ~ {max(values):+.2f}%{fallback}")
return rows
# Fields whose disagreement makes two runs incomparable. Measured magnitudes,
# so the list is not precautionary: an unpinned autotune config moved
# router_gemm by 8pp; the timer moved replication_pad3d from +43% to +124.5%;
# a shared tuner cache made mm look like 945 compiled variants had become 33;
# --level core sees 5 mm shapes where comprehensive sees 29.
CALIBER_FIELDS = (
"backend", "device", "timer", "bench_level", "shape_mode",
"tuner_cache", "direct_repeat",
)
def caliber_of(run_dir: Path) -> Optional[Dict[str, str]]:
"""Caliber fields from a run's run_identity.json, or None if unreadable."""
if not (run_dir / "run_identity.json").is_file():
return None
try:
from _run_identity import load_identity
ident = load_identity(str(run_dir))
except Exception:
return None
return {k: str(ident.get(k, "")) for k in CALIBER_FIELDS} | {
"autotune_mode": str(ident.get("autotune_mode", ""))
}
def caliber_complaints(runs: List[Tuple[str, Path]]) -> List[str]:
"""Everything about this set of runs that makes their comparison unsound.
Reported rather than silently tolerated because the failure is invisible in
the numbers: two runs on different calibers produce a report that looks
exactly like a valid one. This repository has already published three
figures that were an artifact of one such difference.
"""
out: List[str] = []
cals = [(name, caliber_of(d)) for name, d in runs]
missing = [n for n, c in cals if c is None]
if missing:
out.append(f"口径不可核实：{', '.join(missing)} 缺 run_identity.json"
"（早于该字段的产物）")
known = [(n, c) for n, c in cals if c is not None]
if len(known) < 2:
return out
for field in CALIBER_FIELDS:
seen = {c[field] for _, c in known}
if len(seen) > 1:
vals = ", ".join(f"{n}={c[field] or '<空>'}" for n, c in known)
out.append(f"{field} 不一致：{vals}")
# A/B needs one side to pin the configs the other replays. Two record runs
# each sweep independently, and the winner can differ; the README puts that
# at the 10% level and one-directional per shape, which is enough to swamp
# what is being measured.
modes = [(n, c["autotune_mode"]) for n, c in known]
if all(m.startswith("record") for _, m in modes):
out.append("两侧都是 record 模式：config 未锁定，各自 sweep 可能选出不同 "
"winner。按 README，只有 record 数据支撑的结论按未验证处理")
return out
def caliber_hazards(runs: List[Tuple[str, Path]]) -> List[str]:
"""Not mismatches — shared state that makes a matched caliber less solid.
Warned about rather than refused, because refusing would reject the
default configuration every A/B in this repo currently uses.
"""
cals = [(n, c) for n, c in ((n, caliber_of(d)) for n, d in runs) if c]
recording = [n for n, c in cals if c["autotune_mode"].startswith("record")]
shared = [n for n, c in cals if c["tuner_cache"].startswith("shared")]
if recording and shared:
# 一次 record 跑在共享 libtuner cache 上，sweep 的起点是别的运行留下的
# 结果。曾把 mm 的 945 个变体读成 33 个，误报为编译期收益。A/B 两侧都
# replay 同一份 config 时结论仍成立，绝对值不成立。
return [f"{', '.join(recording)} 在共享 tuner cache 上 record："
"被 pin 的 config 取决于此前跑过什么。相对结论不受影响，"
"绝对延迟/变体数不可跨机器引用（ZL_ISOLATE_TUNER_CACHE=1 可消除）"]
return []
class CaliberMismatch(SystemExit):
"""Strict-mode refusal. Exit 4 so a driver script can tell it from a crash."""
def __init__(self) -> None:
super().__init__(4)
def _report_caliber(runs: List[Tuple[str, Path]], label: str, strict: bool) -> bool:
"""Print caliber problems. Returns False when strict mode should stop."""
for h in caliber_hazards(runs):
print(f"警告 [caliber] {label}: {h}")
problems = caliber_complaints(runs)
if not problems:
return True
head = "!!!" if strict else "警告"
for p in problems:
print(f"{head} [caliber] {label}: {p}")
if strict:
print(f"!!! [caliber] {label}: 拒绝出报告。确认口径确实可比后加 "
"--allow-caliber-mismatch")
return False
return True
def _lat_index(text: str) -> Dict[Tuple[str, str, int], float]:
"""(dtype, size_detail, occurrence) -> gems latency, SUCCESS rows only.
Occurrence disambiguates a shape a benchmark measures twice in one table,
which BLAS does at --level comprehensive.
"""
out: Dict[Tuple[str, str, int], float] = {}
seen: Dict[Tuple[str, str], int] = {}
for row in parse_run_log(text):
if row.get("row_status") != "SUCCESS":
continue
try:
lat = float(row.get("gems_lat", ""))
except ValueError:
continue
if lat <= 0:
continue
base = (row.get("dtype", ""), row.get("size_detail", ""))
i = seen.get(base, 0)
seen[base] = i + 1
out[base + (i,)] = lat
return out
def ab_round_rows(root: Path) -> Tuple[List[Dict[str, str]], List[float], int]:
"""Aggregate an N-round A/B laid out as <root>/r<k>_<A|B>/run.log.
Returns the per-shape rows, the per-round deltas, and the round count.
A single pair cannot say whether its own verdict is stable, which is the
question that matters once a delta lands near zero: measured on an S5000,
topk eliminated 32 converts yet 17 of its 42 shapes changed the SIGN of
their delta between two equally valid runs. Two dispersions answer different
halves of that and both are reported:
side spread max/min-1 within one side's samples -> how noisy the machine
was while measuring
round spread min~max of the per-round deltas -> whether the VERDICT holds,
which is what a single pair cannot show
"""
rounds: List[Tuple[Dict, Dict]] = []
for k in itertools.count(1):
a, b = root / f"r{k}_A" / "run.log", root / f"r{k}_B" / "run.log"
if not (a.is_file() and b.is_file()):
break
rounds.append((
_lat_index(a.read_text(encoding="utf-8", errors="replace")),
_lat_index(b.read_text(encoding="utf-8", errors="replace")),
))
if not rounds:
return [], [], 0
keys = set(rounds[0][0]) & set(rounds[0][1])
for a, b in rounds[1:]:
keys &= set(a) & set(b)
per_round = [
statistics.median((a[k] / b[k] - 1.0) * 100.0 for k in keys)
for a, b in rounds
] if keys else []
rows: List[Dict[str, str]] = []
for key in sorted(keys):
av = [a[key] for a, _ in rounds]
bv = [b[key] for _, b in rounds]
am, bm = statistics.median(av), statistics.median(bv)
deltas = [(a / b - 1.0) * 100.0 for a, b in zip(av, bv)]
rows.append({
"dtype": key[0], "size_detail": key[1],
"before_lat": f"{am:.6f}", "after_lat": f"{bm:.6f}",
"speedup_percent": f"{(am / bm - 1.0) * 100.0:.2f}",
"delta_min": f"{min(deltas):.2f}", "delta_max": f"{max(deltas):.2f}",
"sign_stable": "yes" if all(d > 0 for d in deltas)
or all(d <= 0 for d in deltas) else "NO",
"before_spread": f"{(max(av) / min(av) - 1.0) * 100.0:.2f}",
"after_spread": f"{(max(bv) / min(bv) - 1.0) * 100.0:.2f}",
})
return rows, per_round, len(rounds)
def report_ab_rounds(root: Path, label: str = "",
strict: bool = True) -> List[Dict[str, str]]:
"""Print the multi-round A/B: per-shape medians plus both dispersions."""
name = label or root.name
cells = sorted(d for d in root.glob("r*_[AB]") if (d / "run.log").is_file())
if cells and not _report_caliber([(d.name, d) for d in cells], name, strict):
raise CaliberMismatch()
rows, per_round, n = ab_round_rows(root)
if n < 2:
print(f"[ab-rounds] {name}: 需要至少 2 轮 <root>/r<k>_<A|B>/run.log"
f"（找到 {n} 轮）")
return rows
if not rows:
print(f"[ab-rounds] {name}: {n} 轮，但没有全轮次都成功的配对行")
return rows
for r in rows:
flag = "" if r["sign_stable"] == "yes" else " 符号不稳"
print(f" {r['dtype']:<16} {float(r['speedup_percent']):+7.2f}% "
f"[{float(r['delta_min']):+.2f}%~{float(r['delta_max']):+.2f}%] "
f"{r['before_lat']} -> {r['after_lat']} "
f"{r['size_detail']}{flag}")
med = [float(r["speedup_percent"]) for r in rows]
spread = [float(r["before_spread"]) for r in rows] + \
[float(r["after_spread"]) for r in rows]
unstable = sum(1 for r in rows if r["sign_stable"] != "yes")
print(f"[ab-rounds] {name}: {n} 轮 × {len(rows)} shape，"
f"speedup 中位 {statistics.median(med):+.2f}%，"
f"逐轮 {min(per_round):+.2f}% ~ {max(per_round):+.2f}%，"
f"单侧噪声中位 {statistics.median(spread):.2f}%")
if unstable:
# The verdict, not the machine, is what these rows fail to pin down.
print(f"[ab-rounds] {name}: {unstable}/{len(rows)} 个 shape 的 delta "
"在轮次间改变符号——这些行的方向不可结论，与计时口径无关")
return rows
def main() -> None: def main() -> None:
p = argparse.ArgumentParser(description=__doc__.splitlines()[0]) p = argparse.ArgumentParser(description=__doc__.splitlines()[0])
p.add_argument("batch_dir", type=Path, help="run_batch.py 的产物目录") p.add_argument("batch_dir", type=Path, nargs="?",
help="run_batch.py 的产物目录")
p.add_argument("-o", "--output", type=Path, default=None, p.add_argument("-o", "--output", type=Path, default=None,
help="输出 CSV（默认 <batch_dir>/summary.csv）") help="输出 CSV（默认 <batch_dir>/summary.csv）")
p.add_argument("--ab", nargs=2, type=Path, metavar=("BEFORE", "AFTER"),
help="两份 run.log，按 shape 配对输出该算子的 speedup 区间")
p.add_argument("--ab-rounds", type=Path, default=None, metavar="ROOT",
help="多轮 A/B 目录，含 r<k>_A/run.log 与 r<k>_B/run.log；"
"逐 shape 取轮次中位，并报逐轮区间与符号稳定性")
p.add_argument("--label", default="", help="--ab / --ab-rounds 报告里的算子名")
p.add_argument("--allow-caliber-mismatch", action="store_true",
help="口径不一致时仍出报告（默认拒绝）。只在确认两次运行"
"确实可比时使用")
args = p.parse_args() args = p.parse_args()
if args.ab_rounds:
if not args.ab_rounds.is_dir():
raise SystemExit(f"目录不存在: {args.ab_rounds}")
report_ab_rounds(args.ab_rounds, args.label,
strict=not args.allow_caliber_mismatch)
return
if args.ab:
for path in args.ab:
if not path.is_file():
raise SystemExit(f"run.log 不存在: {path}")
report_ab(args.ab[0], args.ab[1], args.label,
strict=not args.allow_caliber_mismatch)
return
if args.batch_dir is None:
raise SystemExit("需要 batch_dir，或用 --ab BEFORE AFTER")
if not args.batch_dir.is_dir(): if not args.batch_dir.is_dir():
raise SystemExit(f"目录不存在: {args.batch_dir}") raise SystemExit(f"目录不存在: {args.batch_dir}")
out = summarize(args.batch_dir, args.output) out = summarize(args.batch_dir, args.output)
campaigns/musa_rlc_product/run_pytest_musa_rlc_product.sh
+10 -11
View File
@@ -53,17 +53,16 @@ export FLAGGEMS_PERF_CURRENT_OP="$OP"
export FLAGGEMS_PERF_COLOR=never export FLAGGEMS_PERF_COLOR=never
export NO_COLOR=1 export NO_COLOR=1
plugins=( # The plugin list comes from zl_bench, not from here. Hand-copying it is how
-p _musa_benchmark_compat_plugin # this campaign drifted: the copy was taken before backends/ existed, so it
-p _device_guard_plugin # still named _device_guard_plugin (now _device_setup_plugin) and, more
-p _seed_plugin # importantly, carried no timing plugin at all -- which is why every sealed
-p _shape_inject_plugin # measurement is timer=direct while zl_bench's own driver had been capturing.
-p _shape_iter_inject_plugin # Asking the planner means the caliber tracks whatever the backend declares.
-p _bespoke_shape_plugin export ZL_BACKEND=musa
-p _autotune_record_plugin mapfile -t plugins < <("$PYTHON_BIN" "$ZL_BENCH_DIR/plugin_plan.py" \
-p _pretty_report_plugin --backend "$ZL_BACKEND" --pytest-args | tr ' ' '\n')
-p _ir_meta_plugin preflight "plugin plan is non-empty" test "${#plugins[@]}" -gt 0
)
dtype_args=() dtype_args=()
for dtype in ${DTYPES:-float16}; do for dtype in ${DTYPES:-float16}; do
campaigns/musa_rlc_product/tests/test_musa_rlc_product.py
+7 -1
View File
@@ -66,7 +66,13 @@ class MusaRlcProductTests(unittest.TestCase):
self.assertIn( self.assertIn(
'FLAGGEMS_PERF_TTGIR_DUMP_DIR="$OUT_DIR/ttgir"', worker 'FLAGGEMS_PERF_TTGIR_DUMP_DIR="$OUT_DIR/ttgir"', worker
) )
self.assertIn("-p _musa_benchmark_compat_plugin", worker) # The plugin list is asked for, not written here: a hand-copied list is
# what let this worker keep _device_guard_plugin and no timing plugin
# long after zl_bench had moved on. What the plan contains for musa --
# the compat plugin included -- is asserted in zl_bench's test_backends.
self.assertIn("plugin_plan.py", worker)
self.assertIn('--backend "$ZL_BACKEND"', worker)
self.assertNotIn("-p _", worker)
self.assertIn( self.assertIn(
'PYTHONPATH="$worker_dir:$ZL_BENCH_DIR:$FLAGGEMS_DIR/src:$FLAGTREE_SITE"', 'PYTHONPATH="$worker_dir:$ZL_BENCH_DIR:$FLAGGEMS_DIR/src:$FLAGTREE_SITE"',
worker, worker,
ops/gen_inventory.py
+2
View File
@@ -41,6 +41,8 @@ import yaml
OPS_DIR = Path(__file__).resolve().parent OPS_DIR = Path(__file__).resolve().parent
# Inventory is FlagGems-logical (repo/op/op_file), not per-hardware.
# One case is reused across backends; do not duplicate CSVs per vendor.
DEFAULT_REPOS = [ DEFAULT_REPOS = [
("main", Path("/workspace/dev/FlagGems")), ("main", Path("/workspace/dev/FlagGems")),
("vllm", Path("/workspace/dev/FlagGems-vllm")), ("vllm", Path("/workspace/dev/FlagGems-vllm")),
plugin_plan.py
+109
View File
@@ -0,0 +1,109 @@
"""Decide which pytest plugins a zl_bench run should load.
CUDA-specific plugins (cudagraph patch, Hopper mm cluster workaround) are not
placed on the default list; `_timing_plugin` enables graph capture only when
the backend declares `supports_cuda_graph`. IR / autotune plugins stay on for
Triton-compatible compilers and can be turned off with env flags.
"""
from __future__ import annotations
import argparse
import os
from typing import List, Optional, Sequence
from backends import get_backend
CORE_PLUGINS: tuple[str, ...] = (
"_device_setup_plugin",
"_parallel_warmup_plugin",
"_seed_plugin",
"_shape_inject_plugin",
"_shape_iter_inject_plugin",
"_bespoke_shape_plugin",
)
ALWAYS_AFTER: tuple[str, ...] = (
"_pretty_report_plugin",
)
def _env_flag(name: str, default: str = "1") -> bool:
raw = os.environ.get(name, default).strip().lower()
return raw not in ("0", "false", "no", "off", "")
def plan_plugins(
backend_name: Optional[str] = None,
extra: Optional[Sequence[str]] = None,
enable_autotune: Optional[bool] = None,
enable_ir: Optional[bool] = None,
enable_mm_cluster: Optional[bool] = None,
) -> List[str]:
backend = get_backend(backend_name)
# Compat first: these override collection-time markers, so they have to be
# in place before anything else inspects the collected items.
plugins: List[str] = list(backend.compat_plugins) + list(CORE_PLUGINS)
autotune = (
backend.supports_triton_hooks
if enable_autotune is None
else enable_autotune
)
if enable_autotune is None:
autotune = autotune and _env_flag("ZL_ENABLE_AUTOTUNE", "1")
if autotune:
plugins.append("_autotune_record_plugin")
plugins.append("_timing_plugin")
plugins.extend(ALWAYS_AFTER)
ir = backend.supports_triton_hooks if enable_ir is None else enable_ir
if enable_ir is None:
ir = ir and _env_flag("ZL_ENABLE_IR", "1")
if ir:
plugins.append("_ir_meta_plugin")
mm = enable_mm_cluster
if mm is None:
mm = backend.supports_mm_cluster_fix and _env_flag("ZL_ENABLE_MM_CLUSTER", "0")
elif mm:
mm = backend.supports_mm_cluster_fix
if mm:
plugins.append("_mm_cluster_fix_plugin")
if extra is None:
extras = [p for p in os.environ.get("ZL_EXTRA_PLUGINS", "").split() if p]
else:
extras = [p for p in extra if p]
for name in extras:
if name not in plugins:
plugins.append(name)
return plugins
def pytest_args(plugins: Optional[Sequence[str]] = None, **kwargs) -> List[str]:
names = list(plugins) if plugins is not None else plan_plugins(**kwargs)
args: List[str] = []
for name in names:
args.extend(("-p", name))
return args
def main(argv: Optional[Sequence[str]] = None) -> int:
p = argparse.ArgumentParser(description="Print zl_bench pytest -p args")
p.add_argument("--backend", default=None)
p.add_argument("--pytest-args", action="store_true",
help="print `-p name` tokens on one line")
p.add_argument("--one-per-line", action="store_true")
args = p.parse_args(argv)
names = plan_plugins(backend_name=args.backend)
if args.one_per_line:
print("\n".join(names))
else:
print(" ".join(pytest_args(names)))
return 0
if __name__ == "__main__":
raise SystemExit(main())
providers/__init__.py
+1
View File
@@ -0,0 +1 @@
"""Benchmark providers. Phase 1 only implements FlagGems."""
providers/flaggems.py
+219
View File
@@ -0,0 +1,219 @@
"""FlagGems-specific mapping kept out of the generic scheduler.
GEMS_VENDOR values are FlagGems' vendor names, not zl_bench backend names.
cuda → nvidia is the important one: FlagGems' device_finder otherwise shells
out to nvidia-smi without a timeout.
"""
from __future__ import annotations
import argparse
import os
import shlex
import sys
from pathlib import Path
from typing import Dict, Optional
_ROOT = Path(__file__).resolve().parent.parent
if str(_ROOT) not in sys.path:
sys.path.insert(0, str(_ROOT))
from backends.base import DeviceBackend
# Live checkouts used by the NV FlagGems product path. Order is the fallback
# sequence when the configured root does not contain the requested test file.
KNOWN_FLAGGEMS_ROOTS = (
"/workspace/FlagGems",
"/workspace/dev/FlagGems",
)
VENDOR_ENV_KEYS = (
"GEMS_VENDOR", "FLAGGEMS_VENDOR", "GEMS_BACKEND", "FLAGGEMS_BACKEND",
)
VENDOR_BY_BACKEND: Dict[str, str] = {
"cuda": "nvidia",
"musa": "musa",
"rocm": "amd",
"ascend": "ascend",
}
def flaggems_vendor(backend: DeviceBackend) -> str:
return backend.flaggems_vendor or VENDOR_BY_BACKEND[backend.name]
def apply_vendor_env(backend: DeviceBackend) -> Optional[str]:
"""Set GEMS_VENDOR unless the user already chose one. Returns the value."""
for key in VENDOR_ENV_KEYS:
if key in os.environ and os.environ[key].strip():
return os.environ[key].strip()
vendor = flaggems_vendor(backend)
if vendor:
os.environ["GEMS_VENDOR"] = vendor
return vendor or None
def make_test_node(flaggems_dir: str, op_file: str, op: str) -> str:
root = flaggems_dir.replace("\\", "/").rstrip("/")
return f"{root}/benchmark/test_{op_file}.py::test_{op}"
# Back-compat alias. Do not call this from emit_run_env: the parameter of the
# same name would shadow it and raise TypeError.
test_node = make_test_node
def parse_test_node(node: str) -> Dict[str, Optional[str]]:
"""Parse a FlagGems pytest node into op / op_file / optional repo root.
Example:
/workspace/FlagGems/benchmark/test_fused_marlin_moe.py::test_fused_marlin_moe_mxfp4
"""
raw = (node or "").strip()
if "::" not in raw:
raise ValueError(f"pytest node must contain '::': {node!r}")
file_part, func = raw.rsplit("::", 1)
if not func.startswith("test_"):
raise ValueError(f"test function must start with 'test_': {func!r}")
stem = Path(file_part.replace("\\", "/").rsplit("/", 1)[-1]).stem
if not stem.startswith("test_"):
raise ValueError(f"benchmark filename must start with 'test_': {file_part!r}")
posix_file = file_part.replace("\\", "/")
parent = posix_file.rsplit("/", 1)[0] if "/" in posix_file else ""
parent_name = parent.rsplit("/", 1)[-1] if parent else ""
flaggems_dir = (
parent.rsplit("/", 1)[0] if parent_name == "benchmark" and "/" in parent else None
)
return {
"test_file": f"{file_part}::{func}",
"file": file_part,
"op": func[len("test_"):],
"op_file": stem[len("test_"):],
"flaggems_dir": flaggems_dir,
}
def _test_py(flaggems_dir: str, op_file: str) -> Path:
return Path(flaggems_dir) / "benchmark" / f"test_{op_file}.py"
def resolve_flaggems_dir(configured: str, op_file: str) -> str:
"""Return the first root that actually contains benchmark/test_<op_file>.py."""
candidates = []
for item in (configured, *KNOWN_FLAGGEMS_ROOTS):
if item and item not in candidates:
candidates.append(item)
for root in candidates:
if _test_py(root, op_file).is_file():
return root
return configured
def resolve_existing_test_file(test_file: str, op_file: str) -> str:
"""If the node path is missing, retry the same relative path under known roots."""
if "::" not in test_file:
return test_file
file_part, func = test_file.rsplit("::", 1)
path = Path(file_part)
if path.is_file():
return test_file
for root in KNOWN_FLAGGEMS_ROOTS:
candidate = _test_py(root, op_file)
if candidate.is_file():
return f"{candidate.as_posix()}::{func}"
return test_file
def emit_run_env(
test_node: str = "",
op: str = "",
op_file: str = "",
flaggems_dir: str = "",
shape_file: str = "",
shape_mode: str = "",
) -> Dict[str, str]:
"""Resolve OP / OP_FILE / TEST_FILE / SHAPE_MODE for run_pytest.sh."""
parsed = parse_test_node(test_node) if test_node.strip() else None
if parsed:
op = parsed["op"] or op
op_file = parsed["op_file"] or op_file
if parsed["flaggems_dir"]:
flaggems_dir = parsed["flaggems_dir"]
test_file = parsed["test_file"] or ""
else:
flaggems_dir = resolve_flaggems_dir(flaggems_dir, op_file)
test_file = make_test_node(flaggems_dir, op_file, op)
flaggems_dir = resolve_flaggems_dir(flaggems_dir, op_file)
test_file = resolve_existing_test_file(test_file, op_file)
shape_file = shape_file.strip()
shape_mode = shape_mode.strip().lower()
if shape_file in ("-", "upstream", "none"):
shape_file = ""
shape_mode = "upstream"
if not shape_mode:
shape_mode = "upstream" if (test_node.strip() and not shape_file) else "inline"
if shape_file:
shape_mode = "file"
pass_shape_file = "1"
elif shape_mode == "upstream":
pass_shape_file = "0"
else:
pass_shape_file = "1"
shape_mode = "inline"
return {
"OP": op,
"OP_FILE": op_file,
"FLAGGEMS_DIR": flaggems_dir,
"TEST_FILE": test_file,
"SHAPE_FILE": shape_file,
"SHAPE_MODE": shape_mode,
"PASS_SHAPE_FILE": pass_shape_file,
}
def _print_shell(env: Dict[str, str]) -> None:
for key, value in env.items():
print(f"{key}={shlex.quote(str(value))}")
def main(argv: Optional[list] = None) -> int:
p = argparse.ArgumentParser(description="FlagGems pytest node helpers")
p.add_argument("--parse-node", default="")
p.add_argument("--emit-shell", action="store_true")
p.add_argument("--op", default="")
p.add_argument("--op-file", default="")
p.add_argument("--flaggems-dir", default="")
p.add_argument("--shape-file", default="")
p.add_argument("--shape-mode", default="")
p.add_argument("--shell", action="store_true",
help="with --parse-node, print OP/OP_FILE/FLAGGEMS_DIR/TEST_FILE")
args = p.parse_args(argv)
if args.emit_shell:
_print_shell(emit_run_env(
test_node=args.parse_node,
op=args.op,
op_file=args.op_file,
flaggems_dir=args.flaggems_dir,
shape_file=args.shape_file,
shape_mode=args.shape_mode,
))
return 0
if args.parse_node:
parsed = parse_test_node(args.parse_node)
if args.shell:
_print_shell({k: v or "" for k, v in parsed.items()
if k in ("op", "op_file", "flaggems_dir", "test_file")})
return 0
for key, value in parsed.items():
print(f"{key}={value or ''}")
return 0
p.error("pass --emit-shell or --parse-node")
return 2
if __name__ == "__main__":
raise SystemExit(main())
run_batch.py
+84 -47
View File
@@ -1,13 +1,16 @@
#!/usr/bin/env python3 #!/usr/bin/env python3
"""多算子批量测试驱动：按算子清单逐个调 run_pytest.sh，多卡分片、串行测量。 """多算子批量测试驱动：按算子清单逐个调 run_pytest.sh，多卡分片、串行测量。
每个算子一个独立子进程（CUDA crash 只废单个算子）、独立产物目录 每个算子一个独立子进程（设备 crash 只废单个算子）、独立产物目录
<batch>/<repo>/<op>/（与单算子跑 run_pytest.sh 的产物结构完全一致，含 <batch>/<repo>/<op>/（与单算子跑 run_pytest.sh 的产物结构完全一致，含
run.log / shapes.yaml / autotune_records / ttgir）。批量层只做编排： run.log / shapes.yaml / autotune_records / ttgir / run_identity.json）。
批量层只做编排：
- 分片：按算子名稳定哈希到 GPU（每卡内部串行，测量口径与单算子一致； - 分片：按算子名稳定哈希到设备（每卡内部串行，测量口径与单算子一致；
跨卡并行测量存在功耗耦合噪声，本模式定位为筛查口径，可疑算子回 跨卡并行测量存在功耗耦合噪声，本模式定位为筛查口径，可疑算子回
单算子串行流程确认）； 单算子串行流程确认）；
- 后端：--backend / ZL_BACKEND 选择 cuda/musa/rocm/ascend，设备可见性
由 backends/ 写入，不再写死 CUDA_VISIBLE_DEVICES；
- 超时：--op-timeout 杀死挂死算子（记 TIMEOUT，不重试，继续跑）； - 超时：--op-timeout 杀死挂死算子（记 TIMEOUT，不重试，继续跑）；
- 重试：pytest rc=1（用例失败/flaky）与信号杀（rc>=128）最多重试 2 次， - 重试：pytest rc=1（用例失败/flaky）与信号杀（rc>=128）最多重试 2 次，
连续两次同 rc 视为确定性失败提前止损；--dtypes 下两级降级——遇上游 连续两次同 rc 视为确定性失败提前止损；--dtypes 下两级降级——遇上游
@@ -23,7 +26,7 @@ run.log / shapes.yaml / autotune_records / ttgir）。批量层只做编排：
示例： 示例：
# 全量筛查（两仓库全部可收集算子，排除 marked_skip 与 dismiss） # 全量筛查（两仓库全部可收集算子，排除 marked_skip 与 dismiss）
python run_batch.py --all --gpus 0,1,2,3,4,5,6,7 --op-timeout 1800 python run_batch.py --all --devices 0,1,2,3,4,5,6,7 --backend cuda --op-timeout 1800
# 指定算子子集（repo 前缀默认 main；重名用 @文件名 消歧） # 指定算子子集（repo 前缀默认 main；重名用 @文件名 消歧）
python run_batch.py --ops softmax vllm:fused_marlin_moe nextafter_@nextafter_ python run_batch.py --ops softmax vllm:fused_marlin_moe nextafter_@nextafter_
@@ -50,6 +53,8 @@ from pathlib import Path
from typing import Dict, List, Optional, Set, Tuple from typing import Dict, List, Optional, Set, Tuple
import batch_summary import batch_summary
from backends import UnknownBackendError, get_backend
from _run_identity import load_identity
ZL_BENCH = Path(__file__).resolve().parent ZL_BENCH = Path(__file__).resolve().parent
OPS_DIR = ZL_BENCH / "ops" OPS_DIR = ZL_BENCH / "ops"
@@ -62,7 +67,9 @@ REPO_DIRS = {
} }
STATUS_FIELDS = [ STATUS_FIELDS = [
"repo", "op", "op_file", "run_name", "gpu", "status", "return_code", "repo", "op", "op_file", "run_name",
"backend", "device", "gpu", "timer",
"status", "return_code",
"elapsed_ms", "rows_total", "rows_success", "notes", "out_dir", "elapsed_ms", "rows_total", "rows_success", "notes", "out_dir",
] ]
@@ -194,48 +201,44 @@ def resolve_selection(args: argparse.Namespace,
return cases return cases
# ---------------- GPU 分片 ---------------- # ---------------- 设备分片 ----------------
def visible_gpus(arg: Optional[str]) -> List[str]: def resolve_device_arg(devices: Optional[str], gpus: Optional[str]) -> Optional[str]:
if arg: """Prefer --devices; --gpus is an alias. Passing both is an error."""
return [g.strip() for g in arg.split(",") if g.strip()] if devices is not None and gpus is not None:
env = os.environ.get("CUDA_VISIBLE_DEVICES", "").strip() raise ValueError("do not pass both --devices and --gpus")
if env: return devices if devices is not None else gpus
return [g.strip() for g in env.split(",") if g.strip()]
try:
out = subprocess.run(["nvidia-smi", "--list-gpus"], text=True,
capture_output=True, timeout=30).stdout
n = sum(1 for line in out.splitlines() if line.startswith("GPU "))
if n:
return [str(i) for i in range(n)]
except Exception:
pass
return ["0"]
def partition(cases: List[Case], gpus: List[str]) -> Dict[str, List[Case]]: def visible_devices(backend, explicit: Optional[str] = None) -> List[str]:
return list(backend.enumerate_devices(explicit))
def partition(cases: List[Case], devices: List[str]) -> Dict[str, List[Case]]:
"""按 (repo, run_name) 稳定哈希分片：与选择顺序无关，resume 时同卡复用。""" """按 (repo, run_name) 稳定哈希分片：与选择顺序无关，resume 时同卡复用。"""
parts: Dict[str, List[Case]] = {g: [] for g in gpus} parts: Dict[str, List[Case]] = {d: [] for d in devices}
for c in cases: for c in cases:
digest = hashlib.sha256(f"{c.repo}/{c.run_name}".encode()).digest() digest = hashlib.sha256(f"{c.repo}/{c.run_name}".encode()).digest()
parts[gpus[int.from_bytes(digest[:8], "big") % len(gpus)]].append(c) parts[devices[int.from_bytes(digest[:8], "big") % len(devices)]].append(c)
return parts return parts
# ---------------- 单算子执行 ---------------- # ---------------- 单算子执行 ----------------
def _spawn(case: Case, out_dir: Path, gpu: str, args: argparse.Namespace, def _spawn(case: Case, out_dir: Path, device: str, args: argparse.Namespace,
dtypes: str) -> subprocess.Popen: dtypes: str) -> subprocess.Popen:
env = os.environ.copy() env = os.environ.copy()
env.pop("PARALLEL_WARMUP_GPUS", None) # 批量层已分片，算子内不再并行 env.pop("PARALLEL_WARMUP_GPUS", None) # 批量层已分片，算子内不再并行
env.pop("REPLAY_FROM", None) env.pop("REPLAY_FROM", None)
env.pop("OUT_DIR", None) env.pop("OUT_DIR", None)
env = args.backend_obj.apply_visibility(env, device)
env.update({ env.update({
"OP": case.op, "OP": case.op,
"OP_FILE": case.op_file, "OP_FILE": case.op_file,
"FLAGGEMS_DIR": REPO_DIRS[case.repo], "FLAGGEMS_DIR": REPO_DIRS[case.repo],
"OUT_DIR": str(out_dir), "OUT_DIR": str(out_dir),
"CUDA_VISIBLE_DEVICES": gpu, "ZL_BACKEND": args.backend,
"ZL_DEVICE": device,
"FLAGGEMS_PERF_COLOR": "never", "FLAGGEMS_PERF_COLOR": "never",
"NO_COLOR": "1", "NO_COLOR": "1",
"USE_FLAGTUNE": str(args.use_flagtune), "USE_FLAGTUNE": str(args.use_flagtune),
@@ -271,13 +274,14 @@ def _wait(proc: subprocess.Popen, timeout: int) -> Tuple[int, bool]:
return 124, True return 124, True
def run_case(case: Case, batch_dir: Path, gpu: str, def run_case(case: Case, batch_dir: Path, device: str,
args: argparse.Namespace) -> Dict[str, str]: args: argparse.Namespace) -> Dict[str, str]:
out_dir = batch_dir / case.repo / case.run_name out_dir = batch_dir / case.repo / case.run_name
marker = out_dir / ".complete" marker = out_dir / ".complete"
if args.resume and marker.is_file() and (out_dir / "run.log").is_file(): if args.resume and marker.is_file() and (out_dir / "run.log").is_file():
row = _classify(case, out_dir, gpu, 0, 0, "", timed_out=False) row = _classify(case, out_dir, device, 0, 0, "", timed_out=False,
args=args)
row["notes"] = _join("resumed", row["notes"]) row["notes"] = _join("resumed", row["notes"])
return row return row
@@ -295,7 +299,7 @@ def run_case(case: Case, batch_dir: Path, gpu: str,
log = out_dir / "run.log" log = out_dir / "run.log"
if log.is_file(): if log.is_file():
log.rename(out_dir / f"run.attempt{attempt - 1}.log") log.rename(out_dir / f"run.attempt{attempt - 1}.log")
rc, timed_out = _wait(_spawn(case, out_dir, gpu, args, dtypes), rc, timed_out = _wait(_spawn(case, out_dir, device, args, dtypes),
args.op_timeout) args.op_timeout)
rcs.append(rc) rcs.append(rc)
if timed_out: if timed_out:
@@ -315,20 +319,22 @@ def run_case(case: Case, batch_dir: Path, gpu: str,
# 覆盖 baseline 对受限 dtype 编译失败（torch jiterator 的 special 函数族等） # 覆盖 baseline 对受限 dtype 编译失败（torch jiterator 的 special 函数族等）
# 但报错文案不是上游 "can't be supported by this op" 的场景。 # 但报错文案不是上游 "can't be supported by this op" 的场景。
if dtypes and not timed_out: if dtypes and not timed_out:
probe = _classify(case, out_dir, gpu, rcs[-1], 0, "", False) probe = _classify(case, out_dir, device, rcs[-1], 0, "", False,
args=args)
if probe["status"] == "FAIL" and probe["rows_success"] == "0": if probe["status"] == "FAIL" and probe["rows_success"] == "0":
log = out_dir / "run.log" log = out_dir / "run.log"
if log.is_file(): if log.is_file():
log.rename(out_dir / f"run.attempt{len(rcs) - 1}.log") log.rename(out_dir / f"run.attempt{len(rcs) - 1}.log")
notes = _join(notes, "dtype_rescue") notes = _join(notes, "dtype_rescue")
rc, timed_out = _wait(_spawn(case, out_dir, gpu, args, ""), rc, timed_out = _wait(_spawn(case, out_dir, device, args, ""),
args.op_timeout) args.op_timeout)
rcs.append(rc) rcs.append(rc)
elapsed_ms = int((time.time() - t0) * 1000) elapsed_ms = int((time.time() - t0) * 1000)
if len(rcs) > 1: if len(rcs) > 1:
notes = _join(notes, "retry_rc=" + "->".join(map(str, rcs))) notes = _join(notes, "retry_rc=" + "->".join(map(str, rcs)))
row = _classify(case, out_dir, gpu, rcs[-1], elapsed_ms, notes, timed_out) row = _classify(case, out_dir, device, rcs[-1], elapsed_ms, notes, timed_out,
args=args)
if row["status"] in ("PASS", "SKIP"): if row["status"] in ("PASS", "SKIP"):
tmp = marker.with_name(".complete.tmp") tmp = marker.with_name(".complete.tmp")
tmp.write_text(row["status"] + "\n") tmp.write_text(row["status"] + "\n")
@@ -345,8 +351,21 @@ def _join(*notes: str) -> str:
return ";".join(n for n in notes if n) return ";".join(n for n in notes if n)
def _classify(case: Case, out_dir: Path, gpu: str, rc: int, elapsed_ms: int, def _identity_fields(out_dir: Path, device: str, backend: str) -> Dict[str, str]:
notes: str, timed_out: bool) -> Dict[str, str]: ident = load_identity(str(out_dir))
backend_v = str(ident.get("backend") or backend or "")
device_v = str(ident.get("device") or device or "")
timer_v = str(ident.get("timer") or "")
return {
"backend": backend_v,
"device": device_v,
"gpu": device_v,
"timer": timer_v,
}
def _classify(case: Case, out_dir: Path, device: str, rc: int, elapsed_ms: int,
notes: str, timed_out: bool, args: argparse.Namespace) -> Dict[str, str]:
text = _read_log(out_dir) text = _read_log(out_dir)
rows = batch_summary.parse_run_log(text) rows = batch_summary.parse_run_log(text)
if timed_out: if timed_out:
@@ -356,12 +375,13 @@ def _classify(case: Case, out_dir: Path, gpu: str, rc: int, elapsed_ms: int,
n_skipped = batch_summary.pytest_skipped_count(text) n_skipped = batch_summary.pytest_skipped_count(text)
if status == "PASS" and n_skipped: if status == "PASS" and n_skipped:
notes = _join(notes, f"partial_skip={n_skipped}") notes = _join(notes, f"partial_skip={n_skipped}")
ident = _identity_fields(out_dir, device, args.backend)
return { return {
"repo": case.repo, "repo": case.repo,
"op": case.op, "op": case.op,
"op_file": case.op_file, "op_file": case.op_file,
"run_name": case.run_name, "run_name": case.run_name,
"gpu": gpu, **ident,
"status": status, "status": status,
"return_code": str(rc), "return_code": str(rc),
"elapsed_ms": str(elapsed_ms), "elapsed_ms": str(elapsed_ms),
@@ -393,8 +413,12 @@ def parse_args() -> argparse.Namespace:
help="--all 时不剔除带无条件 skip 标记的算子") help="--all 时不剔除带无条件 skip 标记的算子")
run = p.add_argument_group("执行") run = p.add_argument_group("执行")
run.add_argument("--backend", default=os.environ.get("ZL_BACKEND", "cuda"),
help="设备后端：cuda/musa/rocm/ascend（默认 ZL_BACKEND 或 cuda）")
run.add_argument("--devices", default=None,
help="逗号分隔设备列表（默认该后端可见设备；优先于 --gpus）")
run.add_argument("--gpus", default=None, run.add_argument("--gpus", default=None,
help="逗号分隔 GPU 列表（默认 CUDA_VISIBLE_DEVICES 或全部）") help="--devices 的别名（两者同时给出则报错）")
run.add_argument("--batch-dir", type=Path, default=None, run.add_argument("--batch-dir", type=Path, default=None,
help="产物目录（默认 runs/batch_<时间戳>；resume 时必须指定）") help="产物目录（默认 runs/batch_<时间戳>；resume 时必须指定）")
run.add_argument("--op-timeout", type=int, default=1800, run.add_argument("--op-timeout", type=int, default=1800,
@@ -429,19 +453,30 @@ def parse_args() -> argparse.Namespace:
def main() -> None: def main() -> None:
args = parse_args() args = parse_args()
try:
backend = get_backend(args.backend)
except UnknownBackendError as exc:
raise SystemExit(f"{LOG} {exc}") from exc
args.backend_obj = backend
args.backend = backend.name
try:
explicit = resolve_device_arg(args.devices, args.gpus)
except ValueError as exc:
raise SystemExit(f"{LOG} {exc}") from exc
devices = visible_devices(backend, explicit)
inv = load_inventory() inv = load_inventory()
cases = resolve_selection(args, inv) cases = resolve_selection(args, inv)
if not cases: if not cases:
raise SystemExit(f"{LOG} 选择结果为空") raise SystemExit(f"{LOG} 选择结果为空")
gpus = visible_gpus(args.gpus)
batch_dir = (args.batch_dir or batch_dir = (args.batch_dir or
ZL_BENCH / "runs" / f"batch_{datetime.now():%Y%m%d-%H%M%S}") ZL_BENCH / "runs" / f"batch_{datetime.now():%Y%m%d-%H%M%S}")
batch_dir = batch_dir.resolve() batch_dir = batch_dir.resolve()
batch_dir.mkdir(parents=True, exist_ok=True) batch_dir.mkdir(parents=True, exist_ok=True)
parts = {g: cs for g, cs in partition(cases, gpus).items() if cs} parts = {d: cs for d, cs in partition(cases, devices).items() if cs}
print(f"{LOG} {len(cases)} 个算子 -> {len(parts)} 张卡 " print(f"{LOG} backend={args.backend} {len(cases)} 个算子 -> {len(parts)} 台设备 "
f"(shards: {[len(cs) for cs in parts.values()]}) out={batch_dir}") f"(shards: {[len(cs) for cs in parts.values()]}) out={batch_dir}")
if args.replay_root: if args.replay_root:
print(f"{LOG} replay 模式: REPLAY_FROM={args.replay_root}/<repo>/<op>") print(f"{LOG} replay 模式: REPLAY_FROM={args.replay_root}/<repo>/<op>")
@@ -461,17 +496,19 @@ def main() -> None:
writer.writerows(rows) writer.writerows(rows)
os.replace(tmp, status_csv) os.replace(tmp, status_csv)
def run_partition(gpu: str, part: List[Case]) -> None: def run_partition(device: str, part: List[Case]) -> None:
for case in part: for case in part:
try: try:
row = run_case(case, batch_dir, gpu, args) row = run_case(case, batch_dir, device, args)
except Exception as exc: # 编排层异常不拖垮整卡分片 except Exception as exc: # 编排层异常不拖垮整卡分片
out_dir = batch_dir / case.repo / case.run_name
ident = _identity_fields(out_dir, device, args.backend)
row = { row = {
"repo": case.repo, "op": case.op, "op_file": case.op_file, "repo": case.repo, "op": case.op, "op_file": case.op_file,
"run_name": case.run_name, "gpu": gpu, "status": "FAIL", "run_name": case.run_name, **ident, "status": "FAIL",
"return_code": "EXC", "elapsed_ms": "0", "rows_total": "0", "return_code": "EXC", "elapsed_ms": "0", "rows_total": "0",
"rows_success": "0", "notes": f"harness_exc={exc}", "rows_success": "0", "notes": f"harness_exc={exc}",
"out_dir": str(batch_dir / case.repo / case.run_name), "out_dir": str(out_dir),
} }
with lock: with lock:
done[0] += 1 done[0] += 1
@@ -480,16 +517,16 @@ def main() -> None:
print(f" ({done[0]:>4}/{len(cases)}) {row['status']:<7} " print(f" ({done[0]:>4}/{len(cases)}) {row['status']:<7} "
f"{row['elapsed_ms']:>8}ms rows={row['rows_success']}/" f"{row['elapsed_ms']:>8}ms rows={row['rows_success']}/"
f"{row['rows_total']:<4} {case.repo}:{case.run_name:<{wname}} " f"{row['rows_total']:<4} {case.repo}:{case.run_name:<{wname}} "
f"[gpu{gpu}]" f"[device{device}]"
+ (f" {row['notes']}" if row["notes"] else ""), + (f" {row['notes']}" if row["notes"] else ""),
flush=True) flush=True)
if len(parts) == 1: if len(parts) == 1:
gpu, part = next(iter(parts.items())) device, part = next(iter(parts.items()))
run_partition(gpu, part) run_partition(device, part)
else: else:
with ThreadPoolExecutor(max_workers=len(parts)) as pool: with ThreadPoolExecutor(max_workers=len(parts)) as pool:
futures = [pool.submit(run_partition, g, p) for g, p in parts.items()] futures = [pool.submit(run_partition, d, p) for d, p in parts.items()]
for f in futures: for f in futures:
f.result() f.result()
run_pytest.sh
+114 -27
View File
@@ -8,15 +8,21 @@ set -euo pipefail
# 每项均可用同名环境变量覆盖，详见 README「环境变量一览」。 # 每项均可用同名环境变量覆盖，详见 README「环境变量一览」。
# 被测算子，对应 $FLAGGEMS_DIR/benchmark/test_<OP_FILE>.py::test_<OP> # 被测算子，对应 $FLAGGEMS_DIR/benchmark/test_<OP_FILE>.py::test_<OP>
# 也可直接传 FlagGems pytest 节点，例如：
# TEST_NODE=/workspace/FlagGems/benchmark/test_fused_marlin_moe.py::test_fused_marlin_moe_mxfp4
OP="${OP:-fused_marlin_moe_mxfp4}" OP="${OP:-fused_marlin_moe_mxfp4}"
OP_FILE="${OP_FILE:-fused_marlin_moe}" OP_FILE="${OP_FILE:-fused_marlin_moe}"
# 上游已拆分为两个仓库：主仓 /workspace/dev/FlagGems（包名 flag_gems）与 # 上游已拆分为两个仓库：主仓 /workspace/dev/FlagGems（包名 flag_gems）与
# /workspace/dev/FlagGems-vllm（包名 flaggems_vllm）。测 vllm 仓库算子时覆盖本变量即可。 # /workspace/dev/FlagGems-vllm（包名 flaggems_vllm）。测 vllm 仓库算子时覆盖本变量即可。
# NV 产品路径也常见 /workspace/FlagGems；TEST_NODE 或缺失测试文件时会回退到该路径。
FLAGGEMS_DIR="${FLAGGEMS_DIR:-/workspace/dev/FlagGems}" FLAGGEMS_DIR="${FLAGGEMS_DIR:-/workspace/dev/FlagGems}"
TEST_NODE="${TEST_NODE:-${ZL_TEST_NODE:-}}"
# shape 来源：非空则用该 yaml，否则用 INLINE_YAML。顶层 key 必须是 op 名。 # shape 来源：非空则用该 yaml；SHAPE_MODE=upstream 则不传 --shape_file，沿用 FlagGems
# 在 $BENCH_LEVEL 下的原生 shape 集。设了 TEST_NODE 且未指定 SHAPE_FILE 时默认 upstream。
SHAPE_FILE="${SHAPE_FILE:-}" SHAPE_FILE="${SHAPE_FILE:-}"
SHAPE_MODE="${SHAPE_MODE:-}"
read -r -d '' INLINE_YAML <<'YAML' || true read -r -d '' INLINE_YAML <<'YAML' || true
fused_marlin_moe_mxfp4: fused_marlin_moe_mxfp4:
shapes: shapes:
@@ -53,25 +59,67 @@ REPLAY_FROM="${REPLAY_FROM:-}"
# always/never 强制开关终端颜色，空则按 tty 判断（run.log 始终去色） # always/never 强制开关终端颜色，空则按 tty 判断（run.log 始终去色）
FLAGGEMS_PERF_COLOR="${FLAGGEMS_PERF_COLOR:-}" FLAGGEMS_PERF_COLOR="${FLAGGEMS_PERF_COLOR:-}"
# 各插件作用见 README「插件说明」；注释掉某行即停用该插件 # 设备后端：cuda / musa / rocm / ascend（别名见 backends/__init__.py）。
PLUGINS=( # 插件列表由 plugin_plan.py 按后端能力生成，不再在此硬编码 -p。
-p _device_guard_plugin # 非 CUDA 默认走上游 do_bench，不会加载 _cudagraph_plugin。
-p _parallel_warmup_plugin ZL_BACKEND="${ZL_BACKEND:-cuda}"
-p _seed_plugin
-p _shape_inject_plugin # 上游 --level：core（默认）只用各算子的 core_shapes；comprehensive 额外并入
-p _shape_iter_inject_plugin # 子类的 set_more_shapes()，并对同一 shape 再产出一次变体（BLAS 是转置侧）。
-p _bespoke_shape_plugin # 差距可以很大：mm 在 core 下是 5 个 shape，comprehensive 下并入 4 个大 K
# -p _mm_cluster_fix_plugin # （N/K 到 152064）与 models_shapes/*.yaml 的模型 shape，共 29 个，且用例翻倍。
-p _autotune_record_plugin # 代价同步放大——sweep 是墙上时间的大头，shape 多时配合 PARALLEL_WARMUP_GPUS。
-p _cudagraph_plugin BENCH_LEVEL="${BENCH_LEVEL:-core}"
-p _pretty_report_plugin case "$BENCH_LEVEL" in
-p _ir_meta_plugin core|comprehensive) ;;
) *) echo "!!! BENCH_LEVEL 只能是 core 或 comprehensive（当前：$BENCH_LEVEL）" >&2; exit 2 ;;
esac
# FlagGems 的 libtuner winner 存在 $FLAGGEMS_CACHE_DIR/config_cache 的 sqlite 里，
# 默认 ~/.flaggems，跨 run 长存。1=本次 run 用独立缓存目录，两侧各自冷调优
# （README「各自最优」口径）；0=沿用共享缓存（默认，快）。
#
# 共享缓存的风险是单向的：先跑的一侧把 winner 写进去，后跑的一侧直接命中，
# 于是它只编译那几个胜出变体。实测 mm 上后跑侧的 TTGIR 变体从 945 塌到 33、
# cvt 从 136 变 10，看起来像被测改动的效果，其实只是继承了对侧的 config。
# 走 REPLAY_FROM 的标准 A/B 不受此影响（config 本来就锁死），record 对 record
# 才会中招。无论取哪个值都会写进 run_identity.json，口径不再隐形。
ZL_ISOLATE_TUNER_CACHE="${ZL_ISOLATE_TUNER_CACHE:-0}"
# ============================== 执行逻辑 ============================== # ============================== 执行逻辑 ==============================
# 前置检查。静默失败最贵：路径不存在时 pytest 收集不到 benchmark 目录，注册
# --level/--mode 的 conftest 就没加载，报错变成 "unrecognized arguments"，指向
# 完全错误的方向。每项检查自报名字，失败即停。
preflight() {
local what="$1"; shift
if ! "$@"; then
echo "!!! PREFLIGHT_FAILED check=$what" >&2
exit 3
fi
}
SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
TEST_FILE="$FLAGGEMS_DIR/benchmark/test_${OP_FILE}.py::test_${OP}" export PYTHONPATH="$SCRIPT_DIR${PYTHONPATH:+:$PYTHONPATH}"
eval "$(python "$SCRIPT_DIR/providers/flaggems.py" --emit-shell \
--parse-node "$TEST_NODE" \
--op "$OP" \
--op-file "$OP_FILE" \
--flaggems-dir "$FLAGGEMS_DIR" \
--shape-file "$SHAPE_FILE" \
--shape-mode "$SHAPE_MODE")"
: "${TEST_FILE:=$FLAGGEMS_DIR/benchmark/test_${OP_FILE}.py::test_${OP}}"
preflight "FlagGems directory exists: $FLAGGEMS_DIR" test -d "$FLAGGEMS_DIR"
# The node's file half; the ::test_<fn> half is left to pytest, which reports a
# missing function clearly ("not found: ...") while a missing FILE does not.
preflight "benchmark file exists: ${TEST_FILE%%::*}" test -f "${TEST_FILE%%::*}"
if [[ -n "$SHAPE_FILE" ]]; then
preflight "shape file exists: $SHAPE_FILE" test -f "$SHAPE_FILE"
fi
if [[ -n "$REPLAY_FROM" ]]; then
preflight "replay run directory exists: $REPLAY_FROM" test -d "$REPLAY_FROM"
fi
# 产物目录，可用 OUT_DIR 指定。时间戳只到秒，并发启动会撞名，故用不带 -p 的 # 产物目录，可用 OUT_DIR 指定。时间戳只到秒，并发启动会撞名，故用不带 -p 的
# mkdir 抢占（已存在即失败），撞上就退让到 -2、-3…… # mkdir 抢占（已存在即失败），撞上就退让到 -2、-3……
@@ -104,9 +152,12 @@ else
PYTEST_COLOR=() PYTEST_COLOR=()
fi fi
export PYTHONPATH="$SCRIPT_DIR${PYTHONPATH:+:$PYTHONPATH}"
export FLAGGEMS_PERF_CURRENT_OP="$OP" export FLAGGEMS_PERF_CURRENT_OP="$OP"
export PYTHONUNBUFFERED=1 # 实时输出不缓冲 export PYTHONUNBUFFERED=1 # 实时输出不缓冲
export ZL_BACKEND
# 按 ZL_BACKEND / ZL_ENABLE_* 生成 -p 列表（见 plugin_plan.py）
read -r -a PLUGINS <<< "$(python "$SCRIPT_DIR/plugin_plan.py" --backend "$ZL_BACKEND" --pytest-args)"
# DTYPES 非空时逐个转为上游 --dtypes 选项（action=append，每个 dtype 一次） # DTYPES 非空时逐个转为上游 --dtypes 选项（action=append，每个 dtype 一次）
DTYPE_ARGS=() DTYPE_ARGS=()
@@ -126,17 +177,52 @@ else
MODE_DESC=record MODE_DESC=record
fi fi
# 未指定 shape 文件时，把 INLINE_YAML 落到临时文件供 pytest 读取 # 未指定 shape 文件时：inline 用 INLINE_YAML；upstream 不传 --shape_file。
if [[ -z "$SHAPE_FILE" ]]; then SHAPE_ARGS=()
SHAPE_FILE="$(mktemp --suffix=.yaml)" if [[ "${PASS_SHAPE_FILE:-1}" == "0" ]]; then
printf '%s\n' "$INLINE_YAML" > "$SHAPE_FILE" if [[ -f "$FLAGGEMS_DIR/benchmark/core_shapes.yaml" ]]; then
trap 'rm -f "$SHAPE_FILE"' EXIT cp -f "$FLAGGEMS_DIR/benchmark/core_shapes.yaml" "$OUT_DIR/shapes.yaml"
else
printf '%s\n' "source: upstream" "level: $BENCH_LEVEL" > "$OUT_DIR/shapes.yaml"
fi
else
if [[ -z "$SHAPE_FILE" ]]; then
SHAPE_FILE="$(mktemp --suffix=.yaml)"
printf '%s\n' "$INLINE_YAML" > "$SHAPE_FILE"
trap 'rm -f "$SHAPE_FILE"' EXIT
fi
cp -f "$SHAPE_FILE" "$OUT_DIR/shapes.yaml" # 存档本次实际使用的 shape
SHAPE_ARGS=(--shape_file "$SHAPE_FILE")
fi fi
cp -f "$SHAPE_FILE" "$OUT_DIR/shapes.yaml" # 存档本次实际使用的 shape
# libtuner 缓存：隔离时用本 run 的私有目录，跑完随 Triton 缓存一起删（选中的
# config 已经在 autotune_records/ 里，DB 本身没有留存价值）。
TUNER_CACHE_ARGS=()
if [[ "$ZL_ISOLATE_TUNER_CACHE" == "1" ]]; then
TUNER_CACHE_DIR="$OUT_DIR/.flaggems_cache"
rm -rf "$TUNER_CACHE_DIR"; mkdir -p "$TUNER_CACHE_DIR"
TUNER_CACHE_ARGS=(FLAGGEMS_CACHE_DIR="$TUNER_CACHE_DIR")
TUNER_DESC=isolated
else
TUNER_DESC="shared(${FLAGGEMS_CACHE_DIR:-~/.flaggems})"
fi
python "$SCRIPT_DIR/_run_identity.py" "$OUT_DIR" \
--backend "$ZL_BACKEND" \
--op "$OP" \
--op-file "$OP_FILE" \
--flaggems-dir "$FLAGGEMS_DIR" \
--autotune-mode "$MODE_DESC" \
--device "${ZL_DEVICE:-}" \
--test-file "$TEST_FILE" \
--shape-mode "$SHAPE_MODE" \
--bench-level "$BENCH_LEVEL" \
--tuner-cache "$TUNER_DESC"
status=0 status=0
{ {
echo "${C_BOLD}>>> op=$OP mode=$MODE_DESC USE_FLAGTUNE=$USE_FLAGTUNE${C_RESET}" echo "${C_BOLD}>>> op=$OP backend=$ZL_BACKEND mode=$MODE_DESC USE_FLAGTUNE=$USE_FLAGTUNE shape=$SHAPE_MODE level=$BENCH_LEVEL tuner=$TUNER_DESC${C_RESET}"
echo "${C_DIM}>>> node=$TEST_FILE${C_RESET}"
echo "${C_DIM}>>> out=$OUT_DIR${C_RESET}" echo "${C_DIM}>>> out=$OUT_DIR${C_RESET}"
# 每次用独立的 Triton 缓存目录，跑完即删：保证编译过程可复现，且 ttgir 落盘 # 每次用独立的 Triton 缓存目录，跑完即删：保证编译过程可复现，且 ttgir 落盘
@@ -146,13 +232,14 @@ status=0
status=0 status=0
TRITON_CACHE_DIR="$CACHE_DIR" \ TRITON_CACHE_DIR="$CACHE_DIR" \
FLAGGEMS_PERF_TTGIR_DUMP_DIR="$OUT_DIR/ttgir" \ FLAGGEMS_PERF_TTGIR_DUMP_DIR="$OUT_DIR/ttgir" \
env "$AUTOTUNE_ENV" \ env "$AUTOTUNE_ENV" ${TUNER_CACHE_ARGS[@]+"${TUNER_CACHE_ARGS[@]}"} \
USE_FLAGTUNE=$USE_FLAGTUNE python -u -m pytest -s "$TEST_FILE" \ USE_FLAGTUNE=$USE_FLAGTUNE python -u -m pytest -s "$TEST_FILE" \
"${PLUGINS[@]}" "${PYTEST_COLOR[@]}" \ "${PLUGINS[@]}" "${PYTEST_COLOR[@]}" \
--shape_file "$SHAPE_FILE" ${DTYPE_ARGS[@]+"${DTYPE_ARGS[@]}"} \ ${SHAPE_ARGS[@]+"${SHAPE_ARGS[@]}"} ${DTYPE_ARGS[@]+"${DTYPE_ARGS[@]}"} \
--level core --mode kernel || status=$? --level "$BENCH_LEVEL" --mode kernel || status=$?
rm -rf "$CACHE_DIR" rm -rf "$CACHE_DIR"
[[ -n "${TUNER_CACHE_DIR:-}" ]] && rm -rf "$TUNER_CACHE_DIR"
if (( status == 0 )); then if (( status == 0 )); then
echo "${C_GREEN}>>> done. outputs in $OUT_DIR${C_RESET}" echo "${C_GREEN}>>> done. outputs in $OUT_DIR${C_RESET}"
test_backends.py
+545
View File
@@ -0,0 +1,545 @@
#!/usr/bin/env python3
"""No-GPU unit tests for zl_bench backend adapters and plugin policy.
FlagGems pytest invocations pass an explicit node, so this file is not
collected during operator runs. Run it with `python test_backends.py`.
"""
from __future__ import annotations
import json
import os
import sys
import tempfile
import types
import unittest
from pathlib import Path
ROOT = Path(__file__).resolve().parent
if str(ROOT) not in sys.path:
sys.path.insert(0, str(ROOT))
from backends import ( # noqa: E402
UnknownBackendError,
get_backend,
normalize_backend_name,
)
from backends.base import ALL_VISIBILITY_KEYS # noqa: E402
from batch_summary import SUMMARY_FIELDS # noqa: E402
from plugin_plan import plan_plugins # noqa: E402
from providers.flaggems import ( # noqa: E402
VENDOR_BY_BACKEND,
emit_run_env,
flaggems_vendor,
parse_test_node,
)
from run_batch import STATUS_FIELDS, resolve_device_arg # noqa: E402
from _run_identity import init_identity, load_identity, update_identity # noqa: E402
class BackendResolveTests(unittest.TestCase):
def test_empty_and_alias_resolve_to_cuda(self):
self.assertEqual(normalize_backend_name(""), "cuda")
self.assertEqual(get_backend("nvidia").name, "cuda")
self.assertEqual(get_backend("nv").name, "cuda")
self.assertEqual(get_backend("").name, "cuda")
def test_other_aliases(self):
self.assertEqual(get_backend("amd").name, "rocm")
self.assertEqual(get_backend("hip").name, "rocm")
self.assertEqual(get_backend("npu").name, "ascend")
self.assertEqual(get_backend("huawei").name, "ascend")
def test_unknown_backend_fails_closed(self):
with self.assertRaises(UnknownBackendError):
get_backend("tpu")
with self.assertRaises(UnknownBackendError):
get_backend("not-a-backend")
class VisibilityTests(unittest.TestCase):
def test_musa_drops_leftover_cuda(self):
env = {"CUDA_VISIBLE_DEVICES": "7", "PATH": "/bin"}
out = get_backend("musa").apply_visibility(env, "1")
self.assertNotIn("CUDA_VISIBLE_DEVICES", out)
self.assertEqual(out["MUSA_VISIBLE_DEVICES"], "1")
self.assertEqual(out["PATH"], "/bin")
for key in ALL_VISIBILITY_KEYS:
if key != "MUSA_VISIBLE_DEVICES":
self.assertNotIn(key, out)
def test_rocm_writes_hip_and_rocr(self):
out = get_backend("rocm").apply_visibility(
{"CUDA_VISIBLE_DEVICES": "0"}, "3")
self.assertEqual(out["HIP_VISIBLE_DEVICES"], "3")
self.assertEqual(out["ROCR_VISIBLE_DEVICES"], "3")
self.assertNotIn("CUDA_VISIBLE_DEVICES", out)
def test_ascend_writes_both_keys(self):
out = get_backend("ascend").apply_visibility({}, "2")
self.assertEqual(out["ASCEND_RT_VISIBLE_DEVICES"], "2")
self.assertEqual(out["ASCEND_VISIBLE_DEVICES"], "2")
self.assertNotIn("CUDA_VISIBLE_DEVICES", out)
def test_cuda_only_sets_cuda(self):
out = get_backend("cuda").apply_visibility(
{"MUSA_VISIBLE_DEVICES": "9"}, "0")
self.assertEqual(out["CUDA_VISIBLE_DEVICES"], "0")
self.assertNotIn("MUSA_VISIBLE_DEVICES", out)
class FlagGemsVendorTests(unittest.TestCase):
def test_vendor_map(self):
self.assertEqual(VENDOR_BY_BACKEND["cuda"], "nvidia")
self.assertEqual(VENDOR_BY_BACKEND["rocm"], "amd")
self.assertEqual(flaggems_vendor(get_backend("cuda")), "nvidia")
self.assertEqual(flaggems_vendor(get_backend("rocm")), "amd")
self.assertEqual(flaggems_vendor(get_backend("musa")), "musa")
self.assertEqual(flaggems_vendor(get_backend("ascend")), "ascend")
class FlagGemsNodeTests(unittest.TestCase):
_NODE = (
"/workspace/FlagGems/benchmark/test_fused_marlin_moe.py"
"::test_fused_marlin_moe_mxfp4"
)
def test_parse_user_command_node(self):
parsed = parse_test_node(self._NODE)
self.assertEqual(parsed["op"], "fused_marlin_moe_mxfp4")
self.assertEqual(parsed["op_file"], "fused_marlin_moe")
self.assertEqual(parsed["flaggems_dir"], "/workspace/FlagGems")
self.assertTrue(parsed["test_file"].endswith(
"test_fused_marlin_moe.py::test_fused_marlin_moe_mxfp4"))
def test_test_node_defaults_to_upstream_shapes(self):
env = emit_run_env(
test_node=self._NODE,
op="softmax",
op_file="softmax",
flaggems_dir="/workspace/dev/FlagGems",
)
self.assertEqual(env["OP"], "fused_marlin_moe_mxfp4")
self.assertEqual(env["OP_FILE"], "fused_marlin_moe")
self.assertEqual(env["FLAGGEMS_DIR"], "/workspace/FlagGems")
self.assertEqual(env["SHAPE_MODE"], "upstream")
self.assertEqual(env["PASS_SHAPE_FILE"], "0")
self.assertEqual(env["TEST_FILE"], self._NODE)
def test_plain_run_keeps_inline_shapes(self):
env = emit_run_env(
op="fused_marlin_moe_mxfp4",
op_file="fused_marlin_moe",
flaggems_dir="/workspace/dev/FlagGems",
)
self.assertEqual(env["SHAPE_MODE"], "inline")
self.assertEqual(env["PASS_SHAPE_FILE"], "1")
self.assertIn("test_fused_marlin_moe.py::test_fused_marlin_moe_mxfp4",
env["TEST_FILE"])
def test_explicit_shape_file_wins(self):
env = emit_run_env(
test_node=self._NODE,
shape_file="/tmp/my_shapes.yaml",
)
self.assertEqual(env["SHAPE_MODE"], "file")
self.assertEqual(env["PASS_SHAPE_FILE"], "1")
def test_resolve_uses_existing_local_root(self):
with tempfile.TemporaryDirectory() as tmp:
root = Path(tmp) / "FlagGems"
py = root / "benchmark" / "test_fused_marlin_moe.py"
py.parent.mkdir(parents=True)
py.write_text("# stub\n", encoding="utf-8")
env = emit_run_env(
op="fused_marlin_moe_mxfp4",
op_file="fused_marlin_moe",
flaggems_dir=str(root),
)
self.assertEqual(env["FLAGGEMS_DIR"], str(root))
self.assertTrue(Path(env["TEST_FILE"].split("::", 1)[0]).is_file())
class PluginPlanTests(unittest.TestCase):
def _plan(self, backend, **kwargs):
kwargs.setdefault("extra", [])
kwargs.setdefault("enable_autotune", True)
kwargs.setdefault("enable_ir", True)
kwargs.setdefault("enable_mm_cluster", False)
return plan_plugins(backend, **kwargs)
def test_cuda_does_not_put_cudagraph_on_the_list(self):
names = self._plan("cuda")
self.assertIn("_timing_plugin", names)
self.assertIn("_device_setup_plugin", names)
self.assertNotIn("_cudagraph_plugin", names)
self.assertNotIn("_mm_cluster_fix_plugin", names)
def test_musa_loads_its_compat_plugin_first(self):
names = self._plan("musa")
self.assertIn("_musa_benchmark_compat_plugin", names)
# It rewrites collection-time markers, so nothing may inspect the item
# list before it runs.
self.assertEqual(names[0], "_musa_benchmark_compat_plugin")
def test_compat_plugins_do_not_leak_to_other_backends(self):
for name in ("cuda", "rocm", "ascend"):
self.assertNotIn("_musa_benchmark_compat_plugin", self._plan(name), name)
def test_musa_same_core_without_cluster_fix(self):
# _cudagraph_plugin is absent from the plan on every backend: the timing
# plugin loads it, so its absence here says nothing about capture support.
names = self._plan("musa")
self.assertIn("_timing_plugin", names)
self.assertNotIn("_cudagraph_plugin", names)
self.assertNotIn("_mm_cluster_fix_plugin", names)
def test_mm_cluster_ignored_on_non_cuda(self):
names = self._plan("musa", enable_mm_cluster=True)
self.assertNotIn("_mm_cluster_fix_plugin", names)
names = self._plan("cuda", enable_mm_cluster=True)
self.assertIn("_mm_cluster_fix_plugin", names)
def test_extra_plugins_isolated(self):
names = self._plan("cuda", extra=["_cudagraph_plugin"])
self.assertIn("_cudagraph_plugin", names)
names = self._plan("cuda", extra=[])
self.assertNotIn("_cudagraph_plugin", names)
def test_extra_none_reads_env(self):
old = os.environ.get("ZL_EXTRA_PLUGINS")
os.environ["ZL_EXTRA_PLUGINS"] = "_cudagraph_plugin"
try:
names = plan_plugins(
"cuda", extra=None, enable_autotune=True,
enable_ir=True, enable_mm_cluster=False)
self.assertIn("_cudagraph_plugin", names)
finally:
if old is None:
os.environ.pop("ZL_EXTRA_PLUGINS", None)
else:
os.environ["ZL_EXTRA_PLUGINS"] = old
class CaptureClassificationTests(unittest.TestCase):
"""Graph-capture failures split three ways; the wording is vendor text."""
def test_cuda_and_musa_both_capture(self):
self.assertTrue(get_backend("cuda").supports_cuda_graph)
# torch_musa supplies Stream/Event/graph/MUSAGraph; measured on S5000.
self.assertTrue(get_backend("musa").supports_cuda_graph)
def test_recoverable_is_the_default(self):
# A kernel refusing capture is ordinary and must reach the fallback.
for name in ("cuda", "musa"):
v = get_backend(name).classify_capture_error(
RuntimeError("operation not permitted when stream is capturing"))
self.assertEqual(v, "recoverable", name)
def test_cuda_fatal_markers(self):
b = get_backend("cuda")
self.assertEqual(
b.classify_capture_error(RuntimeError(
"CUDA error: an illegal memory access was encountered")), "fatal")
self.assertEqual(
b.classify_capture_error(RuntimeError("CUDA out of memory")), "fatal")
def test_torch_dsa_footer_does_not_make_everything_fatal(self):
# torch appends this to every AcceleratorError; matching it loosely
# would classify each capture refusal as fatal and kill the fallback.
exc = RuntimeError("Cannot call current_seed during capture. Compile with "
"`TORCH_USE_CUDA_DSA` to enable device-side assertions.")
self.assertEqual(get_backend("cuda").classify_capture_error(exc),
"recoverable")
def test_musa_unknown_error_is_poisoning_not_recoverable(self):
b = get_backend("musa")
exc = RuntimeError("MUSA error: unknown error")
self.assertEqual(b.classify_capture_error(exc), "poisoned")
# The same text carries no meaning for CUDA, so it must not leak across.
self.assertEqual(get_backend("cuda").classify_capture_error(exc),
"recoverable")
def test_musa_rng_seed_failure_stays_recoverable(self):
exc = RuntimeError("Cannot call MUSAGeneratorImpl::current_seed during "
"MUSA graph capture")
self.assertEqual(get_backend("musa").classify_capture_error(exc),
"recoverable")
def test_deny_list_is_per_backend(self):
self.assertTrue(get_backend("musa").denies_capture("unique2"))
self.assertFalse(get_backend("musa").denies_capture("softmax"))
self.assertFalse(get_backend("cuda").denies_capture("unique2"))
self.assertFalse(get_backend("musa").denies_capture(""))
def test_deny_list_is_keyed_by_pytest_function_name(self):
# FLAGGEMS_PERF_CURRENT_OP holds the pytest function name minus `test_`,
# so a deny entry written as the benchmark FILE name never matches and
# fails open. Observed once: "unique" listed, op=unique2 at runtime.
musa = get_backend("musa")
self.assertTrue(musa.denies_capture("unique2"))
self.assertFalse(musa.denies_capture("unique"))
class DirectRepeatTests(unittest.TestCase):
"""Repetition on the direct timer.
_cudagraph_plugin imports triton at module scope, so it is loaded here
against a stub: the repeat logic is plain arithmetic and should stay
testable on a machine with no GPU and no Triton.
"""
@classmethod
def setUpClass(cls):
cls._saved = {k: sys.modules.get(k) for k in ("triton", "triton.testing")}
triton = types.ModuleType("triton")
testing = types.ModuleType("triton.testing")
testing.do_bench = lambda *a, **k: 0.0
testing.do_bench_cudagraph = lambda *a, **k: 0.0
testing._summarize_statistics = lambda *a, **k: 0.0
triton.testing = testing
sys.modules["triton"] = triton
sys.modules["triton.testing"] = testing
import _cudagraph_plugin # noqa: E402
cls.cg = _cudagraph_plugin
@classmethod
def tearDownClass(cls):
for name, mod in cls._saved.items():
if mod is None:
sys.modules.pop(name, None)
else:
sys.modules[name] = mod
sys.modules.pop("_cudagraph_plugin", None)
def _run(self, value, **kwargs):
"""-> (result, how many times do_bench was called)"""
seq, calls = [3.0, 1.0, 2.0, 9.0, 5.0], []
self.cg._ORIGINAL_DO_BENCH = (
lambda fn, **kw: (calls.append(1), seq[(len(calls) - 1) % len(seq)])[1])
old = os.environ.get("ZL_DIRECT_REPEAT")
os.environ["ZL_DIRECT_REPEAT"] = value
try:
out = self.cg.run_direct(lambda: None, 25, 100, None,
kwargs.get("quantiles"),
kwargs.get("return_mode", "mean"))
finally:
if old is None:
os.environ.pop("ZL_DIRECT_REPEAT", None)
else:
os.environ["ZL_DIRECT_REPEAT"] = old
return out, len(calls)
def test_default_is_one_call(self):
self.assertEqual(self._run("1"), (3.0, 1))
def test_unset_is_one_call(self):
old = os.environ.pop("ZL_DIRECT_REPEAT", None)
try:
self.assertEqual(self.cg.direct_repeat(), 1)
finally:
if old is not None:
os.environ["ZL_DIRECT_REPEAT"] = old
def test_median_over_calls(self):
# 3, 1, 2, 9, 5 -> 3
self.assertEqual(self._run("5"), (3.0, 5))
def test_garbage_and_zero_fall_back_to_one(self):
self.assertEqual(self._run("abc"), (3.0, 1))
self.assertEqual(self._run("0"), (3.0, 1))
self.assertEqual(self._run("-4"), (3.0, 1))
def test_sequence_results_are_not_repeated(self):
# quantiles / return_mode="all" make do_bench return a sequence, and a
# median across calls is undefined for it; take the single call.
self.assertEqual(self._run("5", quantiles=[0.5])[1], 1)
self.assertEqual(self._run("5", return_mode="all")[1], 1)
class DeviceArgTests(unittest.TestCase):
def test_devices_preferred(self):
self.assertEqual(resolve_device_arg("0,1", None), "0,1")
self.assertEqual(resolve_device_arg(None, "2"), "2")
self.assertIsNone(resolve_device_arg(None, None))
def test_both_set_is_error(self):
with self.assertRaises(ValueError):
resolve_device_arg("0", "1")
class CsvFieldTests(unittest.TestCase):
def test_status_fields(self):
for name in ("backend", "device", "gpu", "timer"):
self.assertIn(name, STATUS_FIELDS)
def test_summary_fields(self):
for name in ("backend", "device", "timer"):
self.assertIn(name, SUMMARY_FIELDS)
self.assertLess(SUMMARY_FIELDS.index("backend"),
SUMMARY_FIELDS.index("fg_op_name"))
class IdentityMergeTests(unittest.TestCase):
def test_init_then_update(self):
old_out = os.environ.get("OUT_DIR")
try:
with tempfile.TemporaryDirectory() as tmp:
init_identity(tmp, backend="cuda", device="0", op="softmax")
update_identity(tmp, timer="direct")
ident = load_identity(tmp)
self.assertEqual(ident["schema_version"], 1)
self.assertEqual(ident["provider"], "flaggems")
self.assertEqual(ident["backend"], "cuda")
self.assertEqual(ident["device"], "0")
self.assertEqual(ident["timer"], "direct")
self.assertEqual(ident["op"], "softmax")
finally:
if old_out is None:
os.environ.pop("OUT_DIR", None)
else:
os.environ["OUT_DIR"] = old_out
class AutotuneReplayBackendTests(unittest.TestCase):
def setUp(self):
self._old = os.environ.get("ZL_BACKEND")
os.environ["ZL_BACKEND"] = "cuda"
def tearDown(self):
if self._old is None:
os.environ.pop("ZL_BACKEND", None)
else:
os.environ["ZL_BACKEND"] = self._old
def test_missing_meta_still_replays(self):
from _autotune_record_plugin import _replay_backend_ok
with tempfile.TemporaryDirectory() as tmp:
self.assertTrue(_replay_backend_ok(tmp))
def test_matching_backend_ok(self):
from _autotune_record_plugin import _replay_backend_ok
with tempfile.TemporaryDirectory() as tmp:
Path(tmp, "meta.json").write_text(
json.dumps({"backend": "cuda"}), encoding="utf-8")
self.assertTrue(_replay_backend_ok(tmp))
def test_mismatch_refuses_replay(self):
from _autotune_record_plugin import _replay_backend_ok
with tempfile.TemporaryDirectory() as tmp:
Path(tmp, "meta.json").write_text(
json.dumps({"backend": "musa"}), encoding="utf-8")
self.assertFalse(_replay_backend_ok(tmp))
class CaliberGuardTests(unittest.TestCase):
"""两次运行口径不同产出的报告，和有效报告长得一模一样。
这些用例钉住的是本仓已经踩过的三次：record 对 record、共享 tuner cache、
timer 不同。它们都不会让任何一步报错，只会让数字变成另一个含义。
"""
BASE = {
"backend": "musa", "device": "S5000", "timer": "cuda_graph",
"bench_level": "core", "shape_mode": "default",
"tuner_cache": "isolated", "direct_repeat": 5,
"autotune_mode": "replay",
}
def _run(self, root, name, **over):
d = Path(root, name)
d.mkdir(parents=True)
if over.get("_no_identity"):
return d
ident = dict(self.BASE)
ident.update({k: v for k, v in over.items() if not k.startswith("_")})
Path(d, "run_identity.json").write_text(
json.dumps(ident), encoding="utf-8")
return d
def test_same_caliber_has_no_complaints(self):
from batch_summary import caliber_complaints
with tempfile.TemporaryDirectory() as tmp:
runs = [("A", self._run(tmp, "a")), ("B", self._run(tmp, "b"))]
self.assertEqual(caliber_complaints(runs), [])
def test_record_side_pinning_the_other_is_normal(self):
"""A/B 的正常形态：一侧 record 定 config，另一侧 replay 它。"""
from batch_summary import caliber_complaints
with tempfile.TemporaryDirectory() as tmp:
runs = [("r1_A", self._run(tmp, "a", autotune_mode="record")),
("r1_B", self._run(tmp, "b"))]
self.assertEqual(caliber_complaints(runs), [])
def test_record_against_record_is_flagged(self):
from batch_summary import caliber_complaints
with tempfile.TemporaryDirectory() as tmp:
runs = [("A", self._run(tmp, "a", autotune_mode="record")),
("B", self._run(tmp, "b", autotune_mode="record"))]
self.assertTrue(any("record" in c for c in caliber_complaints(runs)))
def test_each_caliber_field_is_checked(self):
from batch_summary import CALIBER_FIELDS, caliber_complaints
others = {"backend": "cuda", "device": "H20", "timer": "direct",
"bench_level": "comprehensive", "shape_mode": "yaml",
"tuner_cache": "shared", "direct_repeat": 1}
self.assertEqual(set(others), set(CALIBER_FIELDS))
for field, other in others.items():
with tempfile.TemporaryDirectory() as tmp:
runs = [("A", self._run(tmp, "a")),
("B", self._run(tmp, "b", **{field: other}))]
got = caliber_complaints(runs)
self.assertTrue(any(c.startswith(field) for c in got),
f"{field} 不一致时没被拦下: {got}")
def test_missing_identity_is_flagged_not_assumed_equal(self):
from batch_summary import caliber_complaints
with tempfile.TemporaryDirectory() as tmp:
runs = [("A", self._run(tmp, "a")),
("B", self._run(tmp, "b", _no_identity=True))]
self.assertTrue(any("run_identity" in c
for c in caliber_complaints(runs)))
def test_shared_cache_record_warns_but_does_not_block(self):
from batch_summary import caliber_complaints, caliber_hazards
with tempfile.TemporaryDirectory() as tmp:
runs = [("r1_A", self._run(tmp, "a", autotune_mode="record",
tuner_cache="shared(~/.flaggems)")),
("r1_B", self._run(tmp, "b",
tuner_cache="shared(~/.flaggems)"))]
self.assertEqual(caliber_complaints(runs), [])
self.assertTrue(caliber_hazards(runs))
def test_isolated_cache_has_no_hazard(self):
from batch_summary import caliber_hazards
with tempfile.TemporaryDirectory() as tmp:
runs = [("A", self._run(tmp, "a", autotune_mode="record")),
("B", self._run(tmp, "b"))]
self.assertEqual(caliber_hazards(runs), [])
def test_strict_refuses_and_permissive_continues(self):
from batch_summary import _report_caliber
with tempfile.TemporaryDirectory() as tmp:
runs = [("A", self._run(tmp, "a")),
("B", self._run(tmp, "b", timer="direct"))]
self.assertFalse(_report_caliber(runs, "op", strict=True))
self.assertTrue(_report_caliber(runs, "op", strict=False))
LOG = ("Operator: op (dtype=torch.float16, mode=kernel, level=core)\n"
"SUCCESS 1.000 2.000 0.500 shape\n")
def test_report_ab_refuses_mismatched_pair(self):
from batch_summary import CaliberMismatch, report_ab
with tempfile.TemporaryDirectory() as tmp:
a = self._run(tmp, "a")
b = self._run(tmp, "b", timer="direct")
for d in (a, b):
Path(d, "run.log").write_text(self.LOG, encoding="utf-8")
with self.assertRaises(CaliberMismatch):
report_ab(a / "run.log", b / "run.log", "op")
# 放行开关在，但必须是显式的
self.assertEqual(
len(report_ab(a / "run.log", b / "run.log", "op", strict=False)), 1)
if __name__ == "__main__":
unittest.main()