Piping `telemetry.py --run` into anything that stops reading first — head, or
the campaign shutting the sampler down — raised BrokenPipeError and printed a
traceback. Under ab_repeat that traceback lands in telemetry.err, where it
reads as "telemetry is broken" rather than "telemetry ended normally", and
telemetry.err is exactly the file someone checks when the clock record looks
short.
Swallowing the exception alone is not enough: the interpreter flushes stdout
again on exit and emits a second "Exception ignored" to stderr, so stdout is
pointed at devnull the way CPython documents.
Both the _timing_plugin docstring and the README env table still described the
behaviour from before MUSA declared supports_cuda_graph: "other backends do
not attempt CUDA Graph at all", "非 CUDA 即使设 1 也拒绝". Capture selection has
been driven by the backend capability, not by "is this CUDA", since the MUSA
fallback policy landed. Reading the old text is most likely why the S5000
campaigns were launched with ZL_ENABLE_CUDAGRAPH=0 without anyone checking
what auto would do.
Also states what timer=cuda_graph does and does not mean, since the two are
easy to conflate: it says capture was attempted. Measured on S5000 with auto
and configs pinned, dropout/rand/masked_select select the graph path and then
fall back on every single shape — BENCHMARK_DIRECT_NO_CUDAGRAPH on all 15
rows, zero poisoned, zero rows lost — so the whole table is direct-timed under
a cuda_graph label. ZL_ENABLE_CUDAGRAPH=0 selects the same path and records it
plainly, which is why the existing campaign data stands as measured.
A 6-round A/B here reports side noise at the 0.01% level. Clock throttling and
thermal drift move latency by far more than that and leave no trace in
run.log, so the report comes out looking exactly as clean as a run on a stable
machine. The NV run that reproduced PR763 sampled nvidia-smi every 5s
alongside the campaign; the MUSA campaigns had no equivalent, and the parallel
plan makes it worse by heating the whole box.
telemetry.py samples through the backend, so cuda and musa emit the same
columns (index, sm_clock_mhz, mem_clock_mhz, temp_c, power_w, util_pct) and
two machines' campaigns can be read side by side. mthreads-gmi offers no CSV
query mode, so the MUSA path parses --json and strips the units glued onto
each value ("36C", "1888MHz", "477.68W"); its "Power Draw " key carries a
trailing space, matched by stripping rather than by literal.
The device comes from the backend's own visibility variable when not pinned. A
list of devices resolves to no -i at all, sampling every one, rather than
silently picking the first. Dropping the placeholder drops its flag too: a
bare trailing -i fails the sampler instead of widening it.
Telemetry is evidence, not a gate. A missing or unpermitted sampler exits
quietly with its reason in telemetry.err and the campaign runs unchanged.
The three markers the README requires be checked — AUTOTUNE_REPLAY_FALLBACK,
BENCHMARK_DIRECT_NO_CUDAGRAPH, BENCHMARK_CAPTURE_POISONED — were only
reachable by grepping run.log afterwards, which happens when someone
remembers to. The NV run that reproduced PR763 wrote a per-cell status row
instead; a 6-round campaign here had no equivalent, so a missing cell, a
stalled cell, or a marker count going from 0 to nonzero was invisible until
the batch ended.
run_cell now writes round/op/side/switch/rc/elapsed_s/success/failed plus the
three marker counts to <root>/status.tsv as each cell completes, and wraps the
cell in CELL_TIMEOUT (default 2400s, matching the NV run) so one hung cell
cannot hold the batch. rc 124/137 is called out by name.
grep -c prints 0 and exits 1 when it finds nothing, so each count swallows its
status; without that the bookkeeping itself would mark a clean cell failed.
triton normally comes from a PYTHONPATH pointing at a compiler build and
flag_gems from a source checkout, so a shell that reaches every correct path
can still be the wrong Python. run_pytest.sh checked that the FlagGems
directory and benchmark file exist and then handed the run to pytest, where
the real problem surfaced per cell as `pytest exit 4` from a conftest
ImportError.
A multi-round campaign burns every cell that way before anyone reads a log:
today's noise-calibration run lost 12 measurement cells and 160 load runs to
`ModuleNotFoundError: No module named 'triton'`, each failing in under a
second and each looking, from the outside, like a cell that ran.
The check names the interpreter it used, since which Python is on PATH is the
whole question. Its stderr redirect stays on the child — putting `2>/dev/null`
on the preflight call swallows preflight's own `>&2` message and turns a named
refusal into a bare exit 3.
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>
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>
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>
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>
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>
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>
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>
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>
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>
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>
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>
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>
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>
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