fix(breakdown): credit the GEAK kernels the collector is currently dropping - #1209
fix(breakdown): credit the GEAK kernels the collector is currently dropping#1209iraj465 wants to merge 10 commits into
Conversation
collect_geak() only reconstructs accepted_kernels from kernel_journey.json
when status == "ok". status is derived from throughput_speedup, which GEAK
reports on the run's headline basis -- frequently "cold". A run can therefore
be stamped no_gain on the cold basis while alignment_metrics.hot_geak_speedup
records a large measured hot win and the journey holds genuine KEEP rows.
Those rows were dropped, and the session reported kernels_optimized: 0.
Widen the gate to ("ok", "no_gain"). error/timeout stay excluded: those runs
never produced a trustworthy workflow return.
Verified against all 64 GEAK result.json files in the hyperloom-claw campaign:
sessions with kernels_optimized > 0 goes 12 -> 13, kernels credited 23 -> 26,
one session changes, zero regressions. The recovered session is
GLM-5.2-MXFP4/20260814T163244Z -- e2e_gain_pct 29.994, hot_geak_speedup 2.5722,
a 31 KB final_patch.diff on disk, previously credited to nothing.
Note: counting accepted_heads alongside accepted_kernels is NOT the fix. In 7
of the 8 campaign runs carrying heads, the head is already the same kernel as
an accepted_kernels entry, so summing would double-count.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…last ``_geak_accepted_kernels_from_journey`` resolved a single journey file from ``kernel_journey_path`` / ``eval_dir``. That pointer names the LAST e2e cycle. A run that keeps kernels in cycle 0 and then opens a cycle 1 that keeps nothing back-fills nothing, and the run is reported as zero kernels optimized. Read the pointer plus every sibling ``e2e_cycle*/kernel_journey.json``, de-duplicating on ``kernel_id`` with the pointer cycle winning. No kernel in the 64 campaign result files is KEEP in one cycle and rejected in a later one, so the union cannot over-credit. Two campaign runs are affected: Qwen3-14B-FP8/20260816T050457Z cycle0 3 KEEP, cycle1 0 -> 0 credited gemma-4-26B-A4B-it/20260816T112750Z cycle0 2 KEEP, cycle1 0 -> 0 credited Replayed over all 64 GEAK result.json in the campaign, combined with the no_gain gate change in the previous commit: sessions with kernels_optimized > 0 12 -> 15 kernels credited 23 -> 31 regressions 0
GEAK's journey records one acceptance twice when the profiler resolves a dispatched candidate to a library symbol: once under the candidate id with the measurement (``gpu_pct``), once under the resolved symbol with ``gpu_pct: null``. Both rows repeat the same ``e2e_gain_pct``. Back-filling both doubled the kernel. Group accepted rows by rounded gain -- the twin is sometimes the rounded copy (10.38337292749906 vs 10.383) -- and fold a group that holds both a measured and an unmeasured row into the measured row, recording the dropped ids under ``aliases``. Groups that are all-measured or all-unmeasured are left intact, so two real kernels of equal gain both survive and a shape-split pair of unmeasured kernels is not mistaken for an alias. Also extract journey resolution into ``_journey_paths``. Replayed over all 64 campaign result files, with the two earlier commits: sessions with kernels_optimized > 0 12 -> 15 kernels credited 23 -> 23 (+3 recovered, -8 duplicates) alias twins collapsed 8 regressions 0 Affected sessions carry an alias twin in every case: GLM-5.2-MXFP4/20260814T163244Z c0_triton <- dsa_sparse_attn_prefill_main_kernel Kimi-K3/20260816T122327Z decode_attention_grouped_mla MiniMax-M3-MXFP8/20260801T023259Z c0_flydsl <- mxfp8_linear_kernel Qwen3-14B-FP8/20260814T163051Z c1_ck <- ck_gemm_a8w8_blockscale_bpreshuffle Qwen3.5-122B-A10B-FP8/20260801T034513Z, 20260805T080211Z gemma-4-26B-A4B-it/20260814T155153Z, 20260816T112750Z
CI E2E report — ✅ Succeeded
|
Cross-workload view: what GEAK reported, what Hyperloom creditedAsked for during review: how the numbers in this PR break down per workload, Measured on 17 Aug 2026 over Definitions, so the three columns are not read as the same thing:
1. GEAK-reported e2e improvement, by workload
26 runs report a gain above zero, but half of those are below 1% — well inside 2. Hyperloom-reported e2e improvement, by workload
Caveat, stated rather than hidden: The two columns diverge because they answer different questions. Hyperloom's 3. GEAK kernel claims, and what Hyperloom credits
Sessions credited with at least one kernel: 12 → 15. Credited rows stay at Read the per-workload rows rather than only the total. Qwen3-14B-FP8 goes 2 real 4. Three worked runs
That run is not an isolated quirk of one model: 24 of 66 runs report on the 5. One thing this PR does NOT fixReconciling the tables above leaves a gap of exactly 3 kernels — 26 distinct
Whether a non-empty producer list should be reconciled against a richer journey Reproduce# per-workload GEAK verdicts
python3 - <<'PY'
import glob, json
from collections import defaultdict
per = defaultdict(lambda: [0, 0])
for rp in sorted(glob.glob('/shared_nfs/hyperloom-claw/*/*/geak/result.json')):
r = json.load(open(rp)); m = rp.split('/')[3]
per[m][0] += 1
if (r.get("throughput_speedup") or 0) > 1.0: per[m][1] += 1
for m, (runs, up) in sorted(per.items(), key=lambda kv: -kv[1][0]):
print(f"{m:32} {runs:3} {up:3}")
PYFor the credited columns, run the collector on each branch: |
Follow-up: does Hyperloom's e2e gain come from GEAK's kernels?The earlier comment counted GEAK-reported gains and Hyperloom-reported gains side Scope: the 66 GEAK runs across 15 workloads in The three counts
Per workload
Why the last column is 4 and not 0 — and then 0 againFour runs pass a naive test: GEAK positive, Hyperloom positive, within 3x.
Read as a coincidence check, that looks like agreement. It is not. Open each Every one of the four is a server argument. None is a kernel. Each single So the answer to "both reported a gain of the same order after taking in The same result from Hyperloom's own kernel report
There is no gain field on these rows at all — only And from the patch traceAcross the 62 GEAK sessions with a readable
29 {"change": "integrate_patch", "component": "orchestration",
"gain_pct": 8.702391612340799, "kind": "other",
"operation_kind": "other", "outcome": "KEEP",
"task_id": "a3962a52e1de4d2a9d687240e19446b9"}
The
That is GEMM-tuning plumbing, not a GEAK kernel. The sessions barely overlapPut GEAK's kernel claims and Hyperloom's kept patches side by side per session:
The two records are describing different work. The worked case
GEAK's journey holds 3 Hyperloom's
GEAK claimed three kernels. Hyperloom's gain contains none of them. What this does and does not sayIt does not say GEAK's kernels are worthless. It says the CI cannot That is the honest reading of the dashboard question. The three defects filed
None of the three closes the gap this comment measures. Fixing them makes the Reproducepython - <<'PY'
import json, glob, os
sess = sorted({os.path.dirname(os.path.dirname(p))
for p in glob.glob('/shared_nfs/hyperloom-claw/*/*/geak/result.json')})
from collections import Counter
cat = Counter()
for s in sess:
try: d = json.load(open(f"{s}/reports/kernel_optimization_summary.json"))
except Exception: continue
for r in d.get("by_kernel") or []:
cat[r.get("category")] += 1
print(len(sess), "GEAK sessions;", dict(cat))
PY
# 66 GEAK sessions; {'UNATTEMPTED': 357, 'ATTEMPTED_REJECTED': 14, 'IN_FLIGHT': 4} |
GEAK's dashboard row is ~0 because of plumbing, not kernel quality. Measured
over /shared_nfs/hyperloom-claw (66 sessions with a result): GEAK is credited
12 times, +1.2% to +68.7%, and every credit is filed as a server flag. It has
13 accepted, parity-checked kernel results and Hyperloom credits 0 of them to
a kernel. Four separate places drop the identity, the baseline or the evidence.
Change 1 - attribution: the GEAK family is resolved from provenance
The revalidation dispatches as kind="explore" with provenance
"geak_revalidate", so _action_family saw "explore" and all 12 wins landed in
the explore family - exactly what the comment above _ACTION_FAMILY_TABLE says
the geak bucket exists to prevent. The provenance names come from
_CONFIG_REPLAY_PROVENANCE in the executor that stamps them, so there is one
list in the tree. A new phase bucket splits the gain by what was running:
config / kernel / joint. A joint row is reported whole under joint - the
stack rebench cannot divide it, and an invented split is worse than none.
Change 2 - dead-run recovery reads the file that exists
_geak_reconstruct_from_disk recovered accepted kernels from
kernel_journey.json, which is written last: 0 of 30 killed runs have one,
while 10 have per-candidate overlay/*/integrate_result.json. Those are now a
second source, same admission test (gate == accepted, positive
e2e_delta_pct), tagged accepted_kernels_source=integrate_result_backfill and
validated: False. Recovers _mxfp8_linear_kernel at +40.626% and nothing else.
Change 4 - one adoption ledger
state.kernel_integrate_attempts is what by_kernel, kernel_lifecycle.adopted,
attribution and the timeline all read; GEAK wrote only optimization_stack +
geak_pending. The promotion now writes both from the same measurement.
_record_geak_candidate also stops dropping a result whose status is no_gain
when it carries an accepted kernel with a positive same-config delta: that
status is a verdict on GEAK's headline basis, not on its kernels.
Acceptances are read from BOTH lanes - accepted_heads and accepted_kernels
differ only by which queue proposed them, and 8 of the 11 sessions with an
acceptance carry it in accepted_heads alone. kind == "env" entries are
excluded: they select an existing library, so they are config gain, not an
authored kernel. Alias twins collapse on (op_kind, e2e_delta_pct), keeping
the kernel symbol over the candidate tag.
Change 5 - prove the kernel ran before crediting it
canonical_fingerprint covers (args, envs) only, so a revalidation that
silently dropped the overlay still passed its own identity assertion. Of the
62 results carrying a final_overlay, 24 name a directory that does not exist
and 29 hold no sitecustomize.py. The dispatcher now requires a loadable
overlay, carries the overlay's manifest digest beside the config hash,
re-checks both after the run, and names the accepted kernels on the grid
entry. An overlay that is the only material and cannot load falls back to the
GEAK harness (2a) instead of measuring plain baseline under GEAK's name.
Two traps in that gate, both found by replaying the 26 recorded dispatches:
- Loadable is not the same as installing a kernel. GEAK also emits a
config-only overlay - {"modules": [], "rebinds": [], "note": "config-only
result: no kernel overlay accepted ..."} - which imports cleanly and
installs nothing, so a sitecustomize.py check alone would label a pure
config win as a kernel win. When a manifest is present it must now name at
least one module, rebind or capture. An overlay with no manifest keeps the
old behaviour: absence of evidence is not evidence of an empty overlay.
Latent today - none of the 7 dispatched kernel overlays is empty.
- The manifest names the bind target, not the kernel body, so it does not
identify what would run: three unrelated sessions share one manifest digest
because all three patch sglang.kernels.ops.attention.decode_attention. The
digest now folds in the bodies each entry points at. Over the 7 dispatches
with a loadable overlay this turns 5 distinct digests into 7.
The gain written into the kernel ledger is always the orchestrator-measured
rebench number, never GEAK's self-reported delta. When several kernels rode in
on one rebench, or the overlay was not proven loaded, the row is written with a
null gain and validated: False.
Verified against the campaign, no privilege escalation (all state.json and
reports/decision_trace.json are root mode 600):
- attribution, on the 12 real wins reconstructed from coordinator.db: every
one moves explore -> geak, 217.1 gain-points total (186.3 config, 30.9
joint), and the joint rows name their kernel
- ledger: 9 sessions, 11 rows, no duplicates; null gain whenever the overlay
is not proven loaded
- overlay gate, replayed on the 26 recorded dispatch payloads rather than on
geak/result.json: all 26 declared an overlay, only 7 can load one; 17 drop
to flags-only, 2 fall back to 2a, 0 cfg_hash changes, 7 distinct digests
- dead runs: exactly 1 recovery of 30, 29 correctly stay empty
- geak+breakdown suites: 253 passed / 33 failed, failure set identical to
origin/main; the 10 new overlay-identity tests pass
geak/result.json is a last-write-wins snapshot - GEAK runs several e2e cycles
and overwrites it - so any offline replay keyed on that file reads the last
cycle, not the one that was dispatched. The dispatch payloads in
storage/coordinator.db are the record; the counts above use them.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
PR #1209 — Blocking itemsThe five items below all sit in B1.
|
Five defects found while replaying the claw campaign against this branch. Each is keyed on evidence that holds campaign-wide, not on one session. B1 attribution: a GEAK revalidation dispatches as a plain `explore` task, so its action label says `explore` and only its provenance says GEAK. `_entry_family` now consults provenance, reusing the existing `_CONFIG_REPLAY_PROVENANCE` set rather than adding a second list. Without this every GEAK-driven gain is bucketed as generic exploration. B2 alias twins: GEAK records one acceptance twice -- the candidate-slot row carries `gpu_pct`, the resolved profiler symbol carries `gpu_pct: null`. `_collapse_journey_aliases` keeps the measured row (it is the only one holding the number) and names it by the symbol (it is the id the acceptance ledger keeps). Naming it by the slot tag put one kernel under two names in two tables of the same report. B3 kernel_integrate_attempts: the GEAK path never wrote the ledger, so dead-run recovery read a file that is never produced. `_record_geak_ candidate` now writes it. Recovers `MiniMax-M3-MXFP8/20260731T182731Z` (`_mxfp8_linear_kernel`, +40.626%). B4 status gate: `no_gain` is GEAK's verdict on its own headline number, not on the kernels it accepted. A run can report `no_gain` on the promoted basis while carrying an accepted kernel with a positive parity-checked same-config A/B. The gate now admits on `_geak_has_accepted_kernel`, and the rebench downstream still decides. Of 38 `no_gain` results in the campaign exactly one flips (`GLM-5.2-MXFP4/20260814T163244Z`, `dsa_sparse_attn_prefill_main_kernel`, +29.994%); the other 37 stay declined. B5 overlay gate: a config-only `_overlay_manifest.json` imports cleanly and installs nothing, so "the overlay imported" was never evidence a kernel ran. `_geak_overlay_is_loadable` requires `sitecustomize.py` and at least one module or rebind; `overlay_loaded_2a` also checks the digest matches what the task was dispatched with. 64 runs declare `final_overlay`, 9 are loadable -- 55 false "kernel ran" credits removed. `_geak_has_accepted_kernel`, `_geak_overlay_is_loadable` and `_geak_overlay_digest` live in `loop/coordinator_helpers.py` as the single owner; `phases/kernel.py` and `loop/writeback.py` import them. Tests: three new files (37 tests) plus two assertion updates in `test_geak_breakdown_unit.py` that still expected the pre-B2 spelling. 80 tests pass across the five GEAK test files. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Re-running the collector A/B over all 98 claw sessions showed ``kernels_optimized`` falling 31 -> 26 on the recovery path, and one session (Qwen3-14B-FP8/20260814T163051Z) going 2 -> 0 while discarding a +14.924% KEEP. That read as a regression. It is not. The journey holds an alias twin whose resolved symbol is a CK library GEMM. ``result.json`` ships an empty ``accepted_kernels`` and puts the win in ``accepted_heads`` with ``kind: env``. The collapse names the surviving row by the symbol, the kind join then finds GEAK's own declaration, and the row is excluded from the kernel bucket -- correctly. The e2e gain is untouched; it is attributed to config, which is what it is. ``main`` reported 2 only because it neither collapsed the twin nor read the kind. Checked campaign-wide rather than on this one run: of every symbol ``result.json`` names across the 68 sessions holding one, exactly 2 are not covered by a kept row, and both are declared ``env``. Non-env exclusions: 0. The two tests pin both directions, because the exclusion is only safe while it stays narrow: a declared ``env`` twin collapses to 0, and the same twin declared ``authored`` survives at 1 with ``kind_source: result_json``. A row no lane names stays admitted as ``kind_source: absent`` -- guessing "env" there would delete real kernels from dead runs, which is the loss this collector recovers. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The kind join read only `result.json`. That file is rewritten once per cycle and the last write wins, so a later cycle that accepts nothing blanks the lanes an earlier cycle declared. On Qwen3-14B-FP8/20260816T050457Z it names 0 lanes while the run accepted three rows: every recovered row came out `kind_source: absent` and the `kind == "env"` exclusion could not run on them at all. The `action == "geak_e2e"` entries of `state.optimization_stack` carry the same two lanes in the same spelling -- KernelPhase copies them from the result of *that* cycle -- and the stack is append-only, so it keeps what the file lost. `collect_geak` already receives the parsed state, so the second source costs one argument threaded through, not a new reader. Precedence keeps the run's own words authoritative: a declared kind beats an undeclared one whichever artifact holds it, and between two declarations `result.json` wins. So the stack can only fill gaps, never overwrite a published kind. `kind_source` gains `stack` / `stack_undeclared` beside the existing `result_json` forms, so a stack-sourced kind is never reported as something result.json said, and the untyped residual stays countable. Replay over all 98 claw sessions is byte-identical (26 recover, 23 live): every campaign `state.json` is root mode 600, so the replay synthesises state from `result.json` and cannot reach a stack. The change is strictly additive there; the three new unit tests cover the behaviour. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Re: the five blocking itemsAll five are fixed. Branch HEAD is now B1 — both lanesFixed.
Your diagnosis was right: the label was wrong, not the count. Nine sessions B2 — one name per kernelFixed, in your direction. B3 — the journey rows now have a kind sourceThis is the item that needed new code, and your last sentence named the fix. The kind join read only The Your case, on the branch: Two details, so the record is straight:
Precedence keeps GEAK's own words authoritative. A declared kind beats an
Replay over all 98 sessions is byte-identical to the previous commit. That B4 — 2a holds the same bar as 2bFixed. The 2a path now calls B5 — distinct status valueFixed. The new case writes Tests: 59 pass across |
A/B on the claw CI corpus: what this PR changes about GEAK creditControl = the already-run Hyperloom CI sessions replayed through Corpus: 99 GEAK session dirs under 1. Census
2. Where GEAK actually moves Hyperloom e2eRuns the PR attributes to a kernel, where the overlay can load and the measured pair is a real gain.
6 runs. Measured e2e spans +8.07% to +47.58%; the arithmetic mean is +18.18%. These are six different models, so the figures do not compose into a single number — read them per row. On 3. Named by the PR, but deliberately NOT counted as e2eListed because section 2 alone would overstate the PR.
4. False credit this PR removes
13 rows / 159.68% were duplicate counts of a gain already credited once. 2 rows / 29.85% are the CK GEMM library pick, which stays credited to GEAK but as 5. What did not change
Same corpus, same statuses. This PR reads the record; it does not re-run or re-judge anything. Reproduce with |
The coverage gate failed at 89.98% against fail-under=90.00. The two units
this PR adds to carry a GEAK kernel from an acceptance to a named row were
the largest uncovered blocks it introduced:
* _geak_accepted_kernel_specs -- the whole selection loop (lanes, env
exclusion, non-positive deltas, alias-twin collapse) was reached only
through its early returns.
* KernelPhase._record_geak_adopted_kernels -- the per-kernel ledger writer,
entirely uncovered. by_kernel, kernel_lifecycle.adopted and the
attribution split all read what it writes.
Writing the ledger tests surfaced a real defect: best_gain_pct was set to the
current rebench gain, while the canonical writer in _kernel_decisions.py
computes it as a max over attempts. Both feed the same downstream field, so a
second, worse rebench lowered the kernel's recorded best. Now a max, keeping
None (rather than that writer's 0.0 default) for the unattributable case --
"not attributable" is not the same claim as "no gain".
Measured by combining this run's four CI coverage shards with the new tests:
89.98% -> 90.04% (7485 -> 7438 missing lines).
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Coverage gate cleared, and one defect it surfaced
The cause was this PR's own code, not the baseline.
A real defect the tests surfaced
It is now a max. NumbersMeasured by combining this run's four CI coverage shards with the new tests,
The 18 that remain are the 2b-declined → GEAK-harness fallback block, which No behaviour in the A/B changes: the posted per-kernel and e2e numbers are |
Blocking: the conflict is semantic —
|
Resolves the two conflicts in the GEAK promote path, both created by #1188's correctness sweep landing on the same function this branch extends. phases/kernel.py, _promote_geak_from_candidate: main added a KEEP guard (a rebench that does not beat current_best must not overwrite the headline) and dropped the then-unused `provenance` parameter. This branch needs `provenance` and `overlay_loaded` to record the adoption ledger. Kept both sides: main's KEEP guard is untouched, and the two parameters return -- with `provenance` now DEFAULTED rather than required, so main's callers in writeback.py and test_geak_gain_alignment.py continue to work unchanged. loop/writeback.py: restores the two keyword arguments at the promote call sites, matching the signature above. 143 GEAK tests pass, including main's test_geak_gain_alignment.py, which omits `provenance` and so exercises the new default.
|
Rebased onto
Failure sets: 20 identical, 2 swapped. Both swapped tests are pre-existing flakes in files this branch does not touch (
The 20 common failures are all in The two stale alias-twin assertions now expect the resolved symbol ( |
|
@lishuoshuo-amd both resolved. On #1188 /
On B1 — you were right, and the fix is the label, not the count. Suite, both arms, after the rebase:
Failure sets: 20 identical, 2 swapped, and neither swapped file is touched by this branch. I did not assume they were flakes — |
2 participants
Branch:
fix/geak-backfill-no-gain→mainCloses: #1208
Scope
This PR grew past the issue it closes. #1208 describes three collector defects;
the PR also carries four orchestrator fixes that the same investigation turned
up, because the first three only make the kernels visible and the rest are
what make them credited. Reading the sections in order follows that chain.
no_gain, not onlyokcollectors/geak.pycollectors/geak.pycollectors/geak.pycollectors/attribution.pycollectors/geak.pyphases/kernel.pyloop/writeback.py,loop/coordinator_helpers.pyThe A/B against the live campaign is in
this comment;
the coverage-gate fix and one defect it surfaced are in
this one.
What this fixes
Three defects in
collectors/geak.py, one commit each.1.
cd88067— back-fill onno_gain, not only onokno_gainis a statement about end-to-end throughput. It says nothing aboutwhether a kernel was integrated. It is also the majority status: 35 of the 65
GEAK sessions in
/shared_nfs/hyperloom-claw/. The collector refused to readany of their journeys.
2.
6166943— read every e2e cycle, not the one the pointer nameskernel_journey_pathnames the last cycle only. A new_journey_pathshelperreturns the pointer first, then its sibling
*/kernel_journey.jsonin sortedcycle order. Kernels accumulate deduped on
kernel_id, and the pointer wins ona tie, so the authoritative cycle keeps precedence.
Safe to union, and checked rather than assumed: across all 64 journeys on disk,
0 kernel_ids are KEEP in one cycle and non-KEEP in another.
3.
8d108b0— collapse GEAK's alias twinGEAK writes one acceptance under two ids: an overlay entry named for the
candidate tag (
c0_triton) and a return entry named for the kernel symbol(
dsa_sparse_attn_prefill_main_kernel), with the samee2e_gain_pctandgpu_pct: nullon the second._collapse_journey_aliasesgroups accepted rows by gain and fires only when agroup holds both a measured and an unmeasured row. The measured row
survives and gains an
aliaseslist; the unmeasured row is dropped. Twomeasured rows of equal gain are two kernels and both survive. Unmeasured rows of
distinct gain are untouched.
accepted_kernelsis typedlist[Any](breakdown/schema.py:1033), so addingaliasesis schema-safe.Root cause is in GEAK and is fixed there —
AMD-AGI/GEAK
fix/journey-overlay-return-alias. This collapse still has toship: the 77 journeys already written to disk contain the duplicates and the
dashboard reads those files. Once the GEAK fix lands, this becomes a no-op on
new runs.
Verification 1 — replay over the live campaign
Ran the real collector over all 65
geak/result.jsonunder/shared_nfs/hyperloom-claw/, onmainand on this branch.mainkernels_optimized > 0Newly visible, all
status: "no_gain":/shared_nfs/hyperloom-claw/GLM-5.2-MXFP4/20260814T163244Z/geak— recovered bythe gate alone
/shared_nfs/hyperloom-claw/Qwen3-14B-FP8/20260816T050457Z/geak— needs thegate and the multi-cycle scan (pointer
e2e_cycle1, KEEPs ine2e_cycle0)/shared_nfs/hyperloom-claw/gemma-4-26B-A4B-it/20260816T112750Z/geak— sameThe six sessions whose duplicate was removed are listed in the issue.
Reproduce:
Verification 2 — tests
Six new unit tests in
src/hyperloom/inference_optimizer/tests/test_geak_breakdown_unit.py:test_collect_geak_backfill_scans_earlier_cyclestest_collect_geak_backfill_dedupes_repeated_kernel_across_cyclestest_collect_geak_backfill_collapses_alias_twintest_collect_geak_backfill_collapses_rounded_alias_twintest_collect_geak_backfill_keeps_two_measured_kernels_of_equal_gaintest_collect_geak_backfill_keeps_unmeasured_kernels_of_distinct_gainThe 13 geak/breakdown test files, same interpreter, same machine:
The 27 failures are pre-existing and the failure sets are byte-identical on both
sides (
pytest-asyncioconfiguration;diffof the twoFAILEDlists isempty). The full repository suite does not complete here — it exceeds a 2-minute
budget and hits collection errors in
test_multinode_server_log_dir.pyandtest_gpu_probe_remote.py, identically before and after — so the run is scopedto the files this change touches.
Not addressed here
patches/on GEAK sessions. By design: GEAK never sends the per-kernelapply request that
session_paths.patches_dirserves(
session_paths.py:219-232, one caller atrequest_handlers.py:1038).reports/kernel_optimization_summary.jsonstill reports "No kernels wereattempted". It reads a different field.
gate: "stack"acceptance should count as an integrated kernel. Onecredited kernel (
c1_tilelang, GLM session) is provisional by its ownintegrate_result.json. Worth deciding, separately.Upstream: AMD-AGI/GEAK#407 / AMD-AGI/GEAK#408.
4.
a44a190— carry the credit the rest of the way to the dashboardThe three commits above make the collector see the kernels. This one makes the
rest of the pipeline name them. Same evidence base: 66 sessions with a result
under
/shared_nfs/hyperloom-claw/.GEAK is already credited 12 times, +1.2% to +68.7%, and every credit is filed as
a server flag. Four places separate a number from what produced it.
Change 1 — attribution resolves the GEAK family from provenance. The
revalidation dispatches as
kind="explore"with provenancegeak_revalidate, so_action_family("explore")returnedexploreand all 12wins landed in the explore family — exactly what the comment above
_ACTION_FAMILY_TABLEsays the geak bucket exists to prevent. The provenancenames are reused from
_CONFIG_REPLAY_PROVENANCE, so there is one list in thetree. A new phase bucket splits the gain by what was running: config / kernel /
joint. A joint measurement is reported whole under
joint— the stack rebenchcannot divide it, and an invented split is worse than none.
Change 2 — dead-run recovery reads a file that exists.
_geak_reconstruct_from_diskrecovered kernels fromkernel_journey.json,which is written last: 0 of 30 killed runs have one, while 10 have per-candidate
overlay/*/integrate_result.json. Those are now a second source, same admissiontest, tagged
accepted_kernels_source=integrate_result_backfillandvalidated: False. Recovers_mxfp8_linear_kernelat +40.626% and nothing else.Change 4 — one adoption ledger.
state.kernel_integrate_attemptsis whatby_kernel,kernel_lifecycle.adopted, attribution and the timeline all read;GEAK wrote only
optimization_stack+geak_pending. The promotion now writesboth from the same measurement. Acceptances are read from both lanes —
accepted_headsandaccepted_kernelsdiffer only by which queue proposedthem, and 8 of the 11 sessions with an acceptance carry it in
accepted_headsalone.
kind == "env"entries are excluded: they select an existing library, sothey are config gain, not an authored kernel.
Change 5 — prove the kernel ran before crediting it.
canonical_fingerprintcovers(args, envs)only, so a revalidation thatsilently dropped the overlay still passed its own identity assertion. Of the 62
results carrying a
final_overlay, 24 name a directory that does not exist and29 hold no
sitecustomize.py. The dispatcher now requires a loadable overlay,carries the overlay manifest digest beside the config hash, and names the
accepted kernels on the grid entry.
canonical_fingerprintitself is unchanged.Two traps found only on live data, both now covered by
test_geak_overlay_identity.py:{"modules": [], "rebinds": []}plus a note — which imports cleanly and installs nothing. Counting it as
loadable would label a pure config win as a kernel win.
_overlay_manifest.jsonnames the bind target, not the kernel body. Threeunrelated sessions shared one manifest digest because all three patch
sglang.kernels.ops.attention.decode_attention. The referenced bodies arefolded in; digests went from 5 distinct to 7, with 0 collisions.
Verification 4 — end to end on campaign data
The credit chain run whole, not stage by stage:
geak/result.json+ the orchestrator's real rebench numbers fromcoordinator.db→_record_geak_adopted_kernels→state.kernel_integrate_attempts→_collect_adopted_kernels→collect_kernel_lifecycle["adopted"]→ attribution family and phase.maingeakfamilyvalidated: Truekernel_lifecycle["adopted"][]Replayed on the recorded dispatch payloads in
coordinator.dbrather thangeak/result.json: 26 dispatches, all declared an overlay, 7 loadable → 17drop to flags-only, 2 fall back to the GEAK harness (2a), 0
cfg_hashchanges,7 distinct digests. Four of the 12 wins had the kernel proven loaded.
geak/result.jsonis a last-write-wins snapshot acrosse2e_cycle0/1/2, soan offline replay keyed on it reads the last cycle, not the dispatched one. The
task params in
coordinator.dbare the dispatch, recorded verbatim, and are thecorrect replay input. This is why the result-driven replay above shows 3 and the
dispatch-driven one shows 4.
Verification 5 — tests
10 new tests in
test_geak_overlay_identity.py. Widest scoped run(
-k "geak or breakdown or kernel or attribution or overlay"), sameinterpreter, same machine:
The 252 failures are pre-existing and the failure sets are identical, not
merely the same size.
Deviation from plan — Change 3 is not adopted
The plan prescribed deriving GEAK's status once, using
cold_speedupas theverdict ratio. Replayed over 63 sessions, that flips 5 sessions the independent
director scored 0.9994–1.0035. It is refuted and was not merged. AMD-AGI/GEAK#406
addresses the same code and fixes the GLM case correctly upstream, publishing
1.7456 rather than leaving
okbeside a 0.9877 ratio.Limits
No GPU is present on the machine this was verified on (
/dev/kfdabsent), soeverything above is an offline replay against recorded campaign data. A live
session remains verification 9 and is only possible after merge. All 66
state.jsonandreports/decision_trace.jsonare root mode 600 and were notread; no privilege escalation was used.