fix(gemm): stop stamping self-reported speedups as validated, enable tuned-config hit logging - #1222
Merged
Merged
Conversation
![github-advanced-security[bot]](https://avatars.githubusercontent.com/in/57789?s=60&v=4){kind=small}
![github-code-quality[bot]](https://avatars.githubusercontent.com/u/9919?s=60&v=4){kind=small}
CI E2E report — ✅ Succeeded
|
Contributor
Author
8-GPU MI355X verificationVerified on two 8-card boxes ( The injected flag produces the evidence the gate needs
KEEP is now refused when the artifact provably never appliedThe e2e loop already computed reachability and then credited the gain anyway.
A related finding from the forge side reinforces why the merge check matters: Paired A/B is wired, and the ledger now says which it got
What is no longer optional is labelling: MoE stage detectionA serving log showing both 1-stage and 2-stage dispatch stays tunable; only an 168 tests pass, 10 of them new. |
6 tasks
…ed-config hits Three changes on the measurement-trust side of GEMM tuning. _promote_gemm_tuning_keep derived tuned_tput = baseline * speedup from an externally reported number and wrote the result straight into cumulative_gain_validated. The comment above it already said the figure was "inferred from a micro-benchmark's speedup, never measured end to end" and stamped it anyway. On this fleet three rounds of one unchanged configuration span 58%, so a lone reported speedup > 1.0 is within what drift alone produces. The entry is still promoted -- the artifact is real -- but the gain is now recorded as reported, not validated, unless the result carries an end-to-end number we measured, in which case that number is used directly instead of being reconstructed from a ratio. The e2e candidate filter accepted status "ok" plus candidate-or-improved_shapes. improved_shapes cannot exceed 0 for tuners that never time an untuned baseline (TunableOp, the candidate-CSV fallback, hipblaslt-only bf16), and those now report unverified_shapes; partial_output is likewise a real artifact whose rows are deployable. Both are accepted so the forced-e2e path is not undone here. AITER_LOG_TUNED_CONFIG is set for every serving run. aiter logs tuned-config misses unconditionally but hits only behind this flag, and a scan of 60 production server logs found it set in none of them -- which leaves "0 hits" and "hit logging was off" indistinguishable, and makes the GEMM shape list fall back to config.json, which covers 0.4% of real lookups. The bypass backend launches its own server and forwards almost nothing, so it carries the flag explicitly. Adds orchestrator/measurement/convergence.py: discard the warm-up round, then require the remaining rounds to agree before a throughput counts as steady. A controlled repeat of five identical passes spanned 117.6%, and 3.9% once the cold round was dropped, so the fix is to discard it rather than to loosen the threshold. Not yet wired into the measurement rounds. Co-authored-by: Cursor <cursoragent@cursor.com>
…ing it The double-run discards round 1 by design, which leaves exactly one usable measurement -- and one measurement cannot be shown to be steady. The whole gain ledger is then graded against that number with a 3% KEEP threshold, while a real session produced 14,202 -> 19,374 -> 22,425 tok/s from one unchanged configuration. So the convergence verdict is now recorded on the result (it reads insufficient_rounds, which is the honest answer for two rounds where the first is known-cold; establishing convergence needs a third), together with the cold-start delta. A warning is raised only when that delta is large enough to suggest round 2 had not settled either. Deliberately does not fail the baseline: halting here would stall the session, and the point is to make the anchor's provenance visible to the gates that grade against it. Co-authored-by: Cursor <cursoragent@cursor.com>
Two judges, both deliberately unwired for now: they are the pieces the KEEP gate is missing, and landing them separately keeps the behavioural change reviewable. paired.py -- convergence establishes that each side is steady; it does not establish that both sides saw the same machine. Measuring all of A then all of B leaves drift between the blocks indistinguishable from the effect. One resident server held ~137 req/s for three identical passes and then fell to 117 when a neighbour landed, a 16% swing owing nothing to the configuration. So pairs are interleaved and judged on the median paired difference, and pairs that disagree in sign report inconclusive rather than letting an average invent a winner. apply_verification.py -- artifact exists, env var set, throughput up, and the tuning can still have done nothing: the keys may be unreachable, or the merge step may never have picked the table up and the server loaded its bundled default. Neither is detectable by choosing a better tuner. The trap this had to avoid: aiter logs a miss unconditionally but a hit only under AITER_LOG_TUNED_CONFIG. Reading "no hit lines" as "zero hits" would revert every arm that ran without the flag -- in a scan of 60 production logs, all of them. "Cannot tell" is therefore its own verdict, and only "not merged" and "zero hit" block a KEEP. Co-authored-by: Cursor <cursoragent@cursor.com>
The e2e loop already computed whether the tuned table was reachable, then logged the answer and credited the gain anyway. A run whose artifact the runtime never read measures drift, not tuning, so both positive findings now block the promotion: the shape keys never resolving (coverage) and the table never reaching the server's merge list (apply verdict). "Cannot tell" still does not block -- aiter only logs a hit under AITER_LOG_TUNED_CONFIG=1, and a scan of 60 production logs found it set in none of them, so reading a missing hit line as a failed apply would revert every arm. Also wires the paired A/B judge, which until now no caller used. base_tput and new_tput are measured at different times, so their difference carries whatever the machine did in between -- 16% on one controlled repeat on this fleet, 58% across three rounds of one unchanged configuration. Interleaving separates the two but costs two extra rounds per pair, so it is opt-in via HYPERLOOM_GEMM_PAIRED_PAIRS; what is no longer optional is saying which of the two produced the number, so a block comparison cannot arrive downstream as an unqualified "validated". Verified on an 8-GPU MI355X box: a real vLLM run with the injected flag emitted 83 tuned-config hits against 122 misses and the parser called it "served" -- the first non-empty apply verdict available, since the flag was previously set nowhere. 164 tests pass, 10 of them new. Co-authored-by: Cursor <cursoragent@cursor.com>
Every serving run now materializes AITER_LOG_TUNED_CONFIG=1, which is the whole point of that change -- aiter logs a tuned-config miss unconditionally but the matching hit only under the flag, and both GEMM shape discovery and apply verification read hits. The three golden snapshots still pinned the YAML from before it, so all four CI shards failed on them alone. Its position differs between the snapshots because a caller passing extra_server_args seeds EXTRA_*_ARGS earlier in the dict, so the flag lands after it in the two baseline snapshots and before it in the profile one. Co-authored-by: Cursor <cursoragent@cursor.com>
…ed-gain assertion Two leftovers from earlier commits on this branch, both caught by CI shards 3 and 4: * `measurement/__init__.py` and its tests' `__init__.py` were written through a shell that neither interpreted the escapes nor omitted the BOM, so the module docstring arrived as one line of literal backtick-n behind a U+FEFF. The LLM architecture guard walks every source file and cannot parse one, so it failed the whole shard rather than the file. * `test_kernel_entry_auto_runs_gemm_tuning_for_fp8_sglang` still asserted that a self-reported micro speedup lands in `cumulative_gain_validated`. That is the behaviour 563eb1d deliberately removed: the 1.28x is the tuner's own figure, and reconstructing tuned_tput from it is not a measurement we made. The reported gain is still recorded; only the validated stamp waits for an end-to-end number of our own. Co-authored-by: Cursor <cursoragent@cursor.com>
`reports/trace/gemm_tuning.jsonl` exists to answer one question -- did this tuner run -- and recorded only `tuner`, `best_micro_speedup` and `kept`, none of which can separate a tuner that crashed from one that searched and found nothing. The `result.json` beside it held `status`, `elapsed_s`, `error` and `error_class`, and all four were dropped on the way in. The envelope reported no `error_class` either, even when a tuner had named one. Replaying the campaign's own artefacts under `/shared_nfs/hyperloom-claw` through the writer, before and after: 349 result.json, 378 tuner runs, of which 43 ended `failed` or `empty_output`. The old row made 0 of those 43 visible; the new row makes all 43 visible, and 28 envelopes gain an `error_class` a tuner had already supplied. What appears is diagnostic rather than noise -- e.g. a `sglang_dense_bf16` row that now reads `status=failed`, `error_class=validation_error`, `intermediate_size not set in model config` instead of three nulls. This is the blind spot that hid #1211: 82 runs rejected by argparse in ~11s and recorded as a clean `no_improvement`, indistinguishable in the trace from the six that really tuned for ~1600s. It also matters going forward, because the tuner now emits finer outcomes than it used to (`partial_output`, `unsupported_argument`, dropped-as-inaccurate rows) and this row is where those distinctions were being thrown away again. A clean run stays compact: everything added is dropped when null, so a successful row gains only `status` and `elapsed_s`. `tuner`, `best_micro_speedup` and `kept` are still emitted even when null -- `kept` is null on every row observed so far, and an absent key would be indistinguishable from `false`. `error` is truncated to 400 characters; it is a pointer, and the full text stays in the run's own result.json and tune.log. Closes #1210. Co-authored-by: Cursor <cursoragent@cursor.com>
Review findings on the KEEP gate. Both defects pointed the same way: the gate
looked complete and could not actually fire, or would fire on everything.
**`zero_hit` was dead code.** `BLOCKING_VERDICTS` lists it, but the parser
answers `inconclusive_no_hit_logging` for every `hits == 0 and misses > 0`,
which is the exact condition `verify_applied` tests before returning `zero_hit`.
The branch was unreachable, so one of the two blocking verdicts did nothing.
That matters more now than it would have before: this same PR makes
`AITER_LOG_TUNED_CONFIG=1` the default for every serving run, and with the flag
on, "0 hits and N misses" is a genuine zero -- precisely what the gate exists to
catch -- yet it was being waved through as "cannot tell". The caller knows
whether the flag was set, so it now says, and unknown still means inconclusive.
**The arrival check compared the wrong names, and would have reverted
everything.** It required our artifact's basename in aiter's merge line, but by
then the env value is `merged_<candidate>.csv`, a name Hyperloom invented, while
aiter's merge line lists canonical table names. Worse, on a real MI355X run the
comparison never ran at all: setting `AITER_CONFIG_*` -- what every candidate
does -- makes aiter skip the merge step entirely, print no merge line, and name
our file directly in the lookup lines. So the only branch that can revert every
candidate is the one the live verification never reached, and the unit test
constructed a log where the basenames happened to line up, pinning the
assumption instead of testing it.
Arrival is now judged on the tables the lookups actually named, which is where
the evidence is in both modes, and accepts either the candidate's filename or
the canonical table name the runtime resolves it under. It only blocks when none
of the deployed artifacts appear anywhere the runtime looked.
**The paired confirmation served its B leg differently from the KEEP it was
confirming.** `fmoe_ck` only takes effect under `--moe-runner-backend aiter`;
the confirmation payload carried no server args, so B ran without the CK path
active, never read the tuned table, and measured the same thing as A. A real
gain would come back `within_noise` and be labelled unconfirmed.
**The apply-verification tests never ran in CI.** `importorskip
("forge_gemm_tune")` skipped the module, and forge is not a Hyperloom
dependency, so the gate's whole decision surface had no automated coverage --
9 tests, all skipped. They are parametrised over a stand-in parser now, so the
verdict logic is tested everywhere, and over the real parser as well wherever
forge is installed. Both agree on every case: 42 pass without forge, 55 with.
Also: a missing `server.log` is logged rather than silently returning None (the
one way this check can fail quietly), and `has_e2e_evidence` is documented as a
forward-looking contract with no producer today.
541 tests pass.
Co-authored-by: Cursor <cursoragent@cursor.com>
Routing a vLLM run to the aiter tuner family is a substring scan: an aiter fused-MoE marker anywhere sends the whole run that way, and `vllm_moe_triton` then never runs. A run can dispatch both -- aiter CK over part of the token range, vLLM's Triton path over the rest -- and forge's own parser records that case as `impl="mixed"`. Whichever way the single flag falls, the range the other backend served is left untuned. It is the same shape as the defect already fixed one level down, where a single 1-stage sighting disabled CK tuning for the tokens 2-stage was serving. Reported rather than acted on: changing this routing changes which tuners run for every aiter-served vLLM model, which is a larger step than it looks and does not belong in the same change as the checks around it. Forge already adds `fmoe_ck` from the same evidence, so what this names is the Triton half of a mixed run -- the part still going untuned, and now visible instead of implicit. Best-effort throughout: no forge installed, an unreadable log or a parse failure all leave routing exactly as it was. 522 tests pass. Co-authored-by: Cursor <cursoragent@cursor.com>
…rift The branch was cut before #1188, which deleted _promote_gemm_tuning_keep -- the synthetic-throughput promoter this branch had been carefully teaching not to stamp self-reported speedups as validated. Rebasing brings that deletion back, so the two hunks guarding it and the two tests covering it are gone; #1188 solved the same problem the stronger way, by measuring every candidate end to end instead of declining to certify a ratio. What survives is the layer above it: an end-to-end number can be measured correctly and still be measuring the wrong thing, because the tuned table may never have been read. The apply-verification test module skipped on the top-level forge package while production imports forge_gemm_tune.evidence. A box with the package but not that submodule passed the skip, got a None parser, scored every verdict 'unknown', and failed eleven cases for reasons unrelated to what they test. Now 44 pass and 14 skip without forge, 58 pass with it. The env-var-to-table map existed twice in this file, and a name that drifts makes the apply check compare our deployed file against the wrong table, conclude the artifact never arrived, and revert a candidate that was fine. One constant now, with a test pinning the third copy in KernelForge wherever it is importable. AITER_LOG_TUNED_CONFIG and HYPERLOOM_GEMM_PAIRED_PAIRS are documented. The first changes the environment of every serving run, so an operator wanting to turn it off has to be able to find out it exists. Co-authored-by: Cursor <cursoragent@cursor.com>
chennyiiis
force-pushed
the
feature/chenyi/gemm-measurement-trust
branch
from
August 19, 2026 11:28
fc59174 to
28c3dec
Compare
| # ``_enqueue_internal_stack_rebench``. Doubles as the placeholder that reserves | ||
| # ``geak_pending`` before the task row exists, so the phase guard already sees a | ||
| # pending revalidation while the enqueue is in flight. | ||
| _GEAK_REVALIDATE_IDEMPOTENCY_KEY = "geak-revalidate" |
| # ``_enqueue_internal_stack_rebench``. Doubles as the placeholder that reserves | ||
| # ``geak_pending`` before the task row exists, so the phase guard already sees a | ||
| # pending revalidation while the enqueue is in flight. | ||
| _GEAK_REVALIDATE_IDEMPOTENCY_KEY = "geak-revalidate" |
7 tasks
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What this is
#1188made every GEMM tuning candidate measured end to end instead of promotedfrom a self-reported ratio. This branch is the layer above that: an end-to-end
number can be measured correctly and still be measuring the wrong thing.
Three ways that happens, all observed:
runtime, the server loads its own bundled default. Our CSV is still on disk,
so nothing looks wrong — the throughput difference is measuring drift.
AITER_LOG_TUNED_CONFIGset in none of them. Without hit lines, shapediscovery falls back to
config.json, which covered 0.4% of the keys theruntime actually asks for.
unchanged configuration spanned 58% on this fleet; one controlled repeat
moved 16%.
What it changes
measurement/apply_verification.py, wired into the KEEP decisionAITER_LOG_TUNED_CONFIG=1injected for every serving run (operator value wins)partial_outputandunverified_shapesresults now reach the e2e validator_confirm_gemm_gain_paired+measurement/paired.py, opt-inmeasurement/convergence.pyThe one behaviour change that can reject a KEEP
_gemm_apply_verdictcan now block a KEEP that would previously have beenaccepted. It blocks on two positive findings only:
not_merged— the artifact is absent from the runtime's merge list, so theserver is serving its bundled default;
zero_hit— the artifact was merged, hit logging was on, and every lookupmissed it.
Everything else (
inconclusive_no_hit_logging,no_lookups,unknown) doesnot block. "Cannot tell" is never treated as "failed", and a box without
the forge parser degrades to
unknown/non-blocking.Environment variables
Both are documented in
docs/reference/environment-variables.md:AITER_LOG_TUNED_CONFIG— default1, injected per serving run. An operatorvalue wins; set
0to turn hit logging off, at the cost of the demand listand the apply verdict going inconclusive.
HYPERLOOM_GEMM_PAIRED_PAIRS— default0(off). Interleavedbaseline/tuned pairs before a gain is reported as confirmed. When off, the
gain is still promoted but labelled an unpaired block comparison.
Rebase note
This branch was cut before
#1188and originally carried a fix to_promote_gemm_tuning_keep, teaching it not to stamp self-reported speedups asvalidated.
#1188deleted that function outright and solved the same problemthe stronger way. After rebasing, those hunks and their two tests are gone, and
one commit was dropped as fully obsolete.
cumulative_gain_validatedisunaffected by this branch:
#1188already made it come from a realmeasurement.
Test plan
orchestrator/measurement— 44 passed / 14 skipped without KernelForge,58 passed / 0 skipped with it (both checked on a Linux box)
test_decision_framework.py,test_coordinator_gemm_promote_units.py,test_kernel_request_handlers_units.py— 521 passed, 14 skippeddemand derived from a real serving log (122 misses),
vllm_moe_triton9/10 shapes improved at 1.783x,
fmoe_ck10/10 at 1.097xinference_optimizersuite on a Linux box (cannot run on Windows:fcntl)