Covers the three screen cells and why replay locks them to one variable, how to
freeze and run a campaign, what the runner refuses to start on, the six
classifications, and the three scoping rules that decide whether a verdict
means anything.
Also states what a terminal campaign does not prove, since the manifest fixes
production_default_policy to unchanged-off and a passing run is not an argument
for flipping it.
A sealed run already holds every TTGIR dump and its launch index, so a change
in how convert topology is counted can be applied to it directly instead of
spending accelerator time.
The tool re-reads each cell's dump, recomputes the scoped count, re-runs the
campaign classifier against the sealed correctness verdicts, and writes a
derived receipt and matrix outside the run root. It refuses to write inside a
sealed run and rejects a directory with no campaign receipt.
A derived verdict can only demote a topology claim. Promoting one to G or C
needs reversed-order timing, which a sealed run holds only for tasks that
already showed a reduction, so a reduction appearing for the first time under
scoping is reported as X for missing round-b evidence rather than guessed.
Run against the sealed 32-workload S5000 campaign, a floor of 1 changes exactly
one task: gelu_backward moves from C to B, and the only excluded files in the
whole campaign are its five once-launched randn_kernel dumps. A floor of 0
reproduces the sealed counts with no task changed, which is the equivalence
that makes the scoped number auditable rather than merely different.
A frozen 32-workload campaign that measures how FlagTree's RemoveLayoutConversions
pass behaves on a Moore Threads S5000, through FlagGems public dispatch and this
repository's existing plugins.
The manifest is the contract: source commits, the single physical device, the
three screen cells off/mask0/candidate, the native-event timer, and per-workload
dtypes, timeouts and canary membership. prepare validates it against
ops/inventory_main.csv and freezes an immutable task matrix; the runner executes
the matrix serially, and the worker drives one operator cell through pytest with
_ir_meta_plugin, _autotune_record_plugin and the rest.
Three properties are worth calling out because they decide whether a verdict
means anything.
Convert topology is scoped to the measured operator. A benchmark call builds its
inputs once and then runs the operator many times under the timer, so counting
ttg.convert_layout over every executed kernel credits the candidate with
converts removed from input construction that no timer can observe. Each dumped
TTGIR file is joined to its launch count from index.tsv and only variants
launched more than measured_launch_floor contribute; the unscoped total and the
excluded files stay in the receipt, and a cell that cannot be joined to its
index is classified X rather than trusted.
Correctness is scoped to the variant a workload measures. Running a whole test
file lets an unrelated variant decide the verdict, so a workload may pin node
ids or a -k selector, never both, and any narrowing must state its reason.
Targets are resolved before the campaign starts. The inventory can lag the
pinned FlagGems commit, and such a task previously failed only inside the
worker's preflight, which exited silently and left an empty log. The runner now
resolves every benchmark and correctness target against the real tree, a
workload may name a target the inventory does not know about if it says why,
and each preflight check names itself on failure.
This repository's drivers are shell scripts executed on Linux accelerator
hosts, and its consumers record the SHA-256 of the files they ran. With Git's
default autocrlf on Windows a checkout rewrites those files with CRLF, which
breaks the shell drivers outright and changes every digest a receipt refers to.
Several upstream FlagGems tests gate themselves on torch.cuda.is_available()
even though the code under test runs on whatever accelerator FlagGems
dispatches to. On MThreads those guards skip the operator instead of
exercising it, so a working path reports as "not measured".
The plugin drops such a guard only when it is pinned to one file and one exact
skip reason, only when FlagGems has dispatched to the MThreads vendor, and only
after the unskipped test was observed to pass on an S5000. It walks the node
chain rather than the item's own markers, because a module-level pytestmark
does not appear on the item.
Two guards were probed on device and deliberately left in place, which is why
they are named in the module docstring rather than silently absent.
test_grid_sample.py builds its inputs through a helper whose device argument
defaults to the literal "cuda": unskipped, all 90 collected cases fail in
aten::empty on a CUDA backend that is not present. test_cutlass_scaled_mm.py
gates on vLLM as well as CUDA, so removing its guard would raise on the missing
dependency instead of measuring anything.
run_batch.py drives run_pytest.sh per operator across GPUs (stable-hash
sharding, per-op subprocess isolation, process-group timeouts, retry with
deterministic-failure cutoff, two-level dtype fallback, .complete resume,
per-op REPLAY_FROM). batch_summary.py aggregates run.log tables into
summary.csv. ops/ holds the curated assets: dual-repo inventories rebuilt
via AST scan + pytest collect verification, shape sets migrated from the
old regression harness and merged with upstream core_shapes class-name
keys (upstream's set_shapes falls back op_name -> MRO class name ->
1-D DEFAULT_SHAPES, so replacing the shape file without class keys
crashes the BLAS family), and a dismiss list where all 76 entries carry
verified reasons. Validated end to end: 1036-op full screen with zero
failures.
Also fix two cudagraph plugin regressions: newer torch appends "enable
device-side assertions" to every CUDA error, so the loose fatal-error
marker disabled the documented do_bench fallback entirely; and an aborted
graph capture can leave the default CUDA RNG generator stuck in capturing
state, poisoning every later torch.randn - captures now run under a
throwaway RNG state. run_pytest.sh gains an optional DTYPES passthrough.
Upstream moved from single /workspace/FlagGems-dev to /workspace/dev/FlagGems
(flag_gems) plus /workspace/dev/FlagGems-vllm (flaggems_vllm). Patch LibTuner
per actually-imported package at collection finish instead of hardcoding
flag_gems, and update the default FLAGGEMS_DIR.
Also stop dropping pre_hook configs for libtuner kernels: hopper mm's TMA
configs carry a pre_hook even with USE_FLAGTUNE=0, so record silently skipped
them and replay fell back. Upstream LibTuner.run now re-attaches the hook by
kwargs match on cache read (same path as its own ConfigCache round-trip), so
pre_hook-less injection is safe there; plain triton Autotuner keeps the drop.
For shape-heavy ops the autotune sweep dominates wall time, not the
measurement. _parallel_warmup_plugin shards the sweep across N GPUs and
hands the merged configs to the normal single-GPU serial measurement, so
timing stays comparable. Activated by PARALLEL_WARMUP_GPUS=N (silent
no-op when unset, so it can stay in the plugin list); output layout,
REPLAY_FROM usage and the latency table are unchanged.
Measurement is deliberately not parallelized: N processes saturating one
box couple through the power/thermal budget, so per-card latency gets
dragged by its neighbours by an amount that does not reproduce. Config
keys whose winner differs across shards are counted and reported as a
WARNING -- that count is how much to trust the run.
Fixes found while auditing:
- Shards no longer overwrite CUDA_VISIBLE_DEVICES with a bare shard
index, which a caller who had selected idle cards (e.g. 6,7) would see
re-interpreted as absolute ids 0,1 -- silently benchmarking on the busy
cards they were avoiding.
- Interrupting the sweep no longer leaves N subprocesses holding GPUs;
children are terminated and reaped before the exception propagates
(BaseException, since KeyboardInterrupt is not an Exception).
- run_pytest.sh claims its output dir with a bare mkdir and retreats to
a -2/-3 suffix on collision. The second-resolution timestamp meant two
concurrent runs shared one directory and overwrote each other.
- Replay eviction failures now emit a distinct compile_*_no_evict marker.
LibTuner.cache is a sqlite-backed ConfigCache with no __delitem__, so
the bad config could not be dropped and the retry re-read it, while the
log still claimed a clean fallback to live autotune.
Also trims comment density in both shell scripts and reworks the README:
promotes the parallel and cudagraph sections out from under the A/B flow,
groups the env table by purpose, documents that two record-mode runs are
not comparable, and marks ab_fold_test.sh as a caliber example whose
switch upstream has already removed.
LibTuner.cache is a view onto flag_gems' persistent sqlite config DB
(~/.flaggems/config_cache/TunedConfig_*.db), which survives across runs.
Both record and replay assumed a cold, in-process cache:
- record captured the chosen config by diffing cache keys before/after
run(). On a warm DB the key is already present, LibTuner.run takes the
cached branch without writing it again, so the diff was always empty
and nothing was recorded for any libtuner kernel. Only @triton.autotune
kernels (in-process cache) made it into the json -- e.g. a
fused_marlin_moe_mxfp4 run recorded moe_sum_kernel alone, missing both
MXFP4 GEMMs.
- replay only injected when the key was absent from the cache, so a warm
DB skipped injection entirely: the run reported "replaying N entries"
while actually self-tuning.
Record now reads back this call's own self.cache[key] after run();
replay overwrites unconditionally. Verified on fused_marlin_moe_mxfp4:
recorded entries 1 -> 6 (both GEMMs present), replay injects all 6 with
zero AUTOTUNE_REPLAY_FALLBACK and reproduces latency.
README: note that "fresh tune per side" requires dropping the sqlite DB
(not merely omitting REPLAY_FROM), and how to verify record coverage.
Plain autotune is the common case (fast, same-config A/B); FlagTune's
exhaustive first-run search is now opt-in via USE_FLAGTUNE=1.
ab_fold_test.sh automates the FLAGGEMS_MXFP4_FOLDSCALE=0/1 record-replay
comparison over the four DeepSeek-V4-Flash trace shape sets, which are
now tracked under runs/shape/.
- _cudagraph_plugin: warm up autotune/JIT explicitly before graph capture
(internal 5-iter warmup is too short), tag fallback markers with the
failing phase; first-run latency no longer jitters run-to-run.
- _device_guard_plugin (new): set GEMS_VENDOR via torch probe before
importing flag_gems, avoiding its timeout-less nvidia-smi subprocess
probe that can hang import in fork-broken environments.
- _pretty_report_plugin (new) + _term_style (new): fold inputs identical
across all result rows into a legend line, color status/plugin
tags/markers on the live terminal; run.log is ANSI-stripped and keeps
upstream SUCCESS/column wording for grep compatibility.
- run_pytest.sh: make USE_FLAGTUNE overridable, group all knobs into a
config section with Chinese comments, add start/end banners.
- README: document the warmup semantics, new plugins/env vars, and the
A/B rule that both sides must use the same USE_FLAGTUNE.
Aharrypotter
zhoulin
