gate(#378): the first DEVICE run of the warp MoE router, and two record corrections - #481
Open
localai-bot wants to merge 3 commits into
Open
gate(#378): the first DEVICE run of the warp MoE router, and two record corrections#481localai-bot wants to merge 3 commits into
localai-bot wants to merge 3 commits into
Conversation
Replayed onto current main from the local-only branch row/KERNEL-MOE-ROUTER-WARP,
which was never pushed and never opened as a PR. This squashes its four authored
commits, whose messages carry the full derivation and review history:
9f7afa7a spec(kernel): MoE router top-k as one warp per token, byte-exact
cce81c7e perf(kernel): MoE router top-k as one warp per token, byte-exact
be6a1f57 fix(test): moe_router_topk device case fed the op an input it rejects
4f4ee725 fix(test): moe_router_topk gates could not see the kernel they gated
They are squashed rather than cherry-picked because the branch carried three
internal merge commits with no protocol trailers, and its base is far enough back
that replaying them individually conflicts with rows landed since. Nothing in the
kernel, the lane map or the tests is modified here; the tree is byte-for-byte the
branch's, rebased.
Adds MoeRouterTopKWarpKernel<Tin,VPT>: one WARP per token with the whole logit row
in registers, no shared memory, no __syncthreads(), and ONE global read of the row.
Dispatched by default (VT_MOE_ROUTER_WARP, =0 rolls back on the same binary) for
the ungrouped softmax router at E in {32,64,128,256}. Every other E, the grouped
noaux_tc router and the Serial byte-exact oracle fall through to the UNCHANGED
block kernel, which at E=256 spends 37 barriers, 3 KiB of dynamic shared memory
and two full reads of the row.
BYTE-IDENTICAL by structural congruence, not by an appeal to associativity: the
block reduction's s=128,64,32 levels are all multiples of the warp width, so they
never cross a lane and are exactly a per-lane halving tree over
`expert = lane + 32*slot`; levels s=16..1 live inside warp 0 and are what
__shfl_down_sync reproduces. SHAPE port only -- the divide, the sum>0 guard, the
isfinite clamp, the -INFINITY mask and seed, the denom<=0 guard and the best<0
sentinel all stay the incumbent's, because vLLM does five of them differently and
porting any would change tokens.
The device evidence this kernel never had is in the commit that follows.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]
…rd corrections Dispatched as punch-list item 5 of roadmap-v1-completion.md §3 (ROAD-V1-A, "close the 35B c1/c2 residual"). Re-verifying the gap first, per AGENTS.md, changed what the task was. THE ROUTER HAD NOT LANDED. The premise -- that the warp-shuffle router landed and took 35B c1/c4 to ~0.98x -- is false in both halves. `git merge-base --is-ancestor` says none of 9f7afa7a, cce81c7e, be6a1f57, 4f4ee725 is an ancestor of bbc482a, and `git log -S` on MoeRouterTopKWarpKernel, MoeRouterWarpEnabled, VT_MOE_ROUTER_WARP, MoeRouterWarpTreeSum and moe_router_warp returns 0 commits on main for every one. It sat on a local-only branch, never pushed, no PR, 9 commits behind, and cce81c7e says plainly the .cu was never compiled or run. The tree never claimed otherwise -- KERNEL-MOE-ROUTER-WARP is in no matrix and no doc on main -- so the drift was in the dispatch, not the repo. The ~0.98x/0.87x/0.92x figures belong to SPEC-DSPARK (#442), a different row. So this commit does the one thing the branch always needed and its authors could not do: run it on a GPU. FIRST DEVICE VERIFICATION (GB10, source 6c3be5c3, one flock per group). It COMPILES -- not previously known. Gates 1/2/3/6 all green; 35B test_qwen36_paged_engine 315/315 on BOTH arms of the same binary, plus six more green legs during the A/B. Counts land exactly on §11's predictions (9/4597, 315/315); gate 2's 1907 is the post-§8.3 doubled count. Mutation RED re-proven: a left-linear fold fails 76 assertions / 2 cases. The 35B gate is confirmed to have RUN, not skipped. WHICH KERNEL RAN, proven by name in nsys --cuda-graph-trace=node rather than inferred: MoeRouterTopKWarpKernel<__nv_bfloat16,(int)8> (VPT=8 = E=256) in the ON arm with no block kernel present, and the converse under =0. That closes the F1 hazard at the model level -- the 315/315 could not have come from either kernel. KERNEL-LEVEL WIN, ESTABLISHED. 3 reps/arm, order-alternated, 1280 router calls every leg: warp 10.422 ms (10.307/10.529) vs block 14.201 ms (14.141/14.275) -- -3.779 ms, -26.6%, 1.363x, bands non-overlapping. STEP-LEVEL: NOT SEPARABLE, and NOT CLAIMED. Total GPU-busy read 1601.890 vs 1619.395 ms, but the kernels this change does not touch -- identical code, identical 155,310 launches -- differ 13.726 ms between arms with a 24.140 ms spread across six legs, 6.4x the effect. The window is prefill+decode aggregated, diluting 40 router calls/step across ~1600 ms. Sized against #378's 16.1863 ms/step decode GPU-busy the saving is ~0.5-0.75% of decode GPU-busy -- BELOW #378's ~1.5% estimate, because the warp kernel closes ~79% of the gap to topkGating (12.98 -> 8.14 us/call vs 6.85), not all of it. Per gate 7 that is diagnostic evidence, NOT default-flip credit; the default stays ON as cce81c7e authored it because that is the reviewed author's decision, not because this earned it. Whether default-ON is right here is a NEEDS_DECISION for the operator. THE CANONICAL 35B GRID IS STALE. docs/BENCHMARKS.md carried "CANONICAL @348c265d: 0.918x-0.972x" as current. 348c265 is an ancestor of main, but 136 src/include commits have landed since -- including 972a152 ("+2.05% c8", one CastF32 per layer-step gone), which merged 1h36m AFTER the grid was captured. Holding vLLM fixed that implies c8 ~0.937 and c4 ~0.980 today, an ESTIMATE not a measurement. BENCHMARKS and STATUS now say stale and regrid-owed rather than asserting a number nobody has measured on current main. Also repaired: the roadmap issue table named KERNEL-MOE-ROUTER-WARP as a stable row, which check-agent-record rejects -- the spec says this is a lever inside KERNEL-MOE-ROUTING and registers no new row. That defect would have been caught the first time the branch met the gate, which is itself evidence it never did. The measurement lands in .agents/benchmark-record.md, not the scoreboard: it has no vLLM arm, and the scoreboard's canonical sections are reference engines and resource axes, not kernel levers. NOT VERIFIED HERE: gate 5 compute-sanitizer memcheck; a decode-only window; and any vLLM denominator -- none was run, so no ratio-vs-vLLM is asserted. The pinned oracle SOURCE is at 5559679229bc (asserted by commit), but the shared ~/venvs/vllm-oracle symlink resolves to v0.25.0-stage, NOT the pin; left untouched for other sessions and recorded as a hazard. Next traceable hypotheses, no ceiling declared: a decode-only window to turn the 0.5-0.75% into a measurement, then the gate-GEMM epilogue fold and batching the 40 per-step launches. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]
FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode] # Conflicts: # .agents/benchmark-record.md
localai-bot
pushed a commit
that referenced
this pull request
Aug 12, 2026
…again Found while doing #460. MAX_README_CHARS = 30000 in check-readme-structure.py is a budget on a whole shared file, and README.md measures 29,965: THIRTY-FIVE characters free, tighter than any of the three budgets already retired for this exact reason. Measured over the last 20 commits touching README.md: 13 sat under 60 characters free, and 031410e landed the file 52 characters OVER the cap. There is a whole merged row whose only purpose was paying rent, row/DOCS-README-BUDGET (#161), plus 44206e4 "restore the landing-page budget after the MANIFESTO link". FILED, NOT FIXED. This is a third checker with its own mutation suite, so per AGENTS.md it needs its own spec and its own red-before in tests/scripts/test_check_readme_structure.py rather than riding along in a PR scoped to the two benchmark-surface checkers. Recorded in the roadmap intake table and as spec W7. Also reconciles two claims the spec made before the code existed and that the diff does not contain: docs/BENCHMARKS.md is UNCHANGED, because #481 is open and rewrites the 35B row in place, so writing a second copy of "regrid owed" would duplicate a keyed row the moment both merge. The surface's ability to accept the row is what this row owes, and it is proven by test_the_shipped_page_can_accept_the_next_measurement_row, which adds the row to the real page and drops it again: 44,832 chars and 162 rows to 45,173 and 163, errors [], every pre-existing row asserted still present. Issue: #498 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]
localai-bot
pushed a commit
that referenced
this pull request
Aug 12, 2026
… does not make Left over from the pre-code spec: it said "the row added here is additive and placed in ## Open gaps" and described how to resolve a conflict with #481 over it. There is no such edit. docs/BENCHMARKS.md is untouched, #481 keeps ownership of the "regrid owed" fact, and the two branches cannot conflict at all. Also records the mid-row base move, 918c568 -> e1087a8 (12 commits), and that every gate was rerun on the merged tree rather than the pre-merge one. The figures in Our baseline are as taken at 918c568 and are deliberately not restated: they are the measurement that justified the change. Spec text only. No checker, test or public document changes. Issue: #460 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]
localai-bot
pushed a commit
that referenced
this pull request
Aug 12, 2026
…d pair fences the way CommonMark does (#460) Review returned FAIL on PR #494 with two blocking findings and seven smaller ones. The design is unchanged: AGENTS.md, Records, says cap the entry, never the file, and the byte cap stays retired. F1. "Strictly tighter than the byte cap" was FALSE, and the gap was the append-log class this row's own stop condition names. `_prose_paragraphs` excluded every line starting with `-`, `*`, `>`, `|` or `#`, so a bulleted or quoted wall was counted by NOTHING: not the paragraph count, not MAX_PARAGRAPH_CHARS, not MAX_CELL_CHARS, not MAX_ROW_CHARS, and not the heading guard, which fires on a DATE. Measured: 3,000 appended bullet lines took docs/BENCHMARKS.md to 113,833 chars with the checker reporting `[]` (BASE rejects the same mutant on size), and 500 UNDATED `### Attempt N` sections with bulleted forensics reached 117,222 chars at exit 0 with 68/68 green. Fixed both ways the finding asked. The claim is corrected wherever it appeared, to the narrower true statement -- tighter than the canonical-section allowlist, which sees `## ` only -- in the checker comment, spec Design 3, spec Risks and the PR body. And the channel is closed: list items and blockquote lines fold into `_prose_paragraphs`, so a contiguous run is ONE paragraph and trips MAX_PARAGRAPH_CHARS while a run per section trips the count. BENCHMARKS carries no list item and is unmoved at 35; FEATURES carries one, so `max_prose_paragraphs` re-baselines 20 to 21; STATUS moves the other way, 82 long paragraphs to 75, so STATUS_RATCHET follows the measurement DOWN rather than banking 7 units of slack. An emphasis-lead paragraph still starts with `*` and is still excluded, so a `**bold**`-led wall remains unbounded. That is stated plainly, not glossed: closing it turns four paragraphs already shipped on docs/BENCHMARKS.md red at 717, 719, 748 and 1,084 chars against MAX_PARAGRAPH_CHARS, so it owes an edit to a page #481 holds open. Filed as #507, deferred as spec W8, and pinned by test_an_EMPHASIS_lead_wall_is_a_KNOWN_residue, which goes red the day it lands. F2. `strip_code_spans` closed a block on any line matching a fence, including one with an info string, which is not CommonMark and does not fail safe: it INVERTS fence phase for the rest of the file. With the one unbalanced fence this tree has, the bare fence at STATE-LEGACY-000001.md:17697 was "closed" by the ```sh at :17948, and ordinary prose at :18297 was blanked, so a live reader-followable link stopped being validated. A closing fence now must match the opener's character, be at least as long, and carry nothing but whitespace after the marker. Re-measured: 4,114 raw targets, 4,110 after the strip, and all 4 losses are genuine samples. The Risk said 4 and was measuring 5. F3. Deleted the `assertGreater(len(grown), RETIRED_PAGE_BUDGET)` on the LIVE page. It stored a measurement of docs/BENCHMARKS.md inside a test file with 173 chars of margin, which is the corollary this row removed; 887e04f shrank the page by 280. F4 to F9, same pass. The second-base docstring now says what it enforces (the evidence still exists under a declared base) instead of "no reader can follow it", which is false of the links it accepts, and names W5 as what would make the archived copy followable. `[`name`](path)` wrapped in stray backticks hid the non-existent target `path` in this row's own spec: fixed and pinned by a test. "Line numbers stay honest" is no longer cited as evidence, because check_links reports none. The PageRules docstring now says the no-eviction property is true of ROWS only, with both pages sitting on the paragraph count. The heading survey is re-measured: 282 of 305 archived titles dated, 0 of 36 live headings, not 278/301 and 0/32. The acceptance test inserts the owed row directly under the last Open gaps row so markdown renders it in that table. Mutation set, three trees, in spec Risks. M7, M8, M8b and M9 flip SURVIVED to CAUGHT. M6 survives by design and M11 is the filed residue. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]
localai-bot
pushed a commit
that referenced
this pull request
Aug 12, 2026
…RELIEF (#494) Merges `row/BENCHMARKS-CAP-RELIEF` at 5f1a7fe into main at ee9bc7a. Closes #460. Row `ENG-RECORD-CONFLICT-SURFACES`. Spec: `.agents/specs/benchmarks-entry-cap.md`. ## What it does `docs/BENCHMARKS.md` was gated by a 45,000-character budget on the WHOLE FILE and measured 44,795, so 205 characters were free. Adding a measurement row meant deleting a row somebody else owns, and the documented way to delete one, moving it byte-for-byte into `.agents/benchmark-record.md`, did not work for exactly the rows carrying evidence links (#460). Worse, the success mode was unsafe: 04b2b9f is a CLEAN automatic merge that landed the page at 45,007 chars, because two PRs each paid by evicting a DIFFERENT row and the three-way merge applied both additions and neither eviction. That is AGENTS.md, Records, verbatim: "merging two such edits cleanly is worse than conflicting". This applies the corollary instead of raising anything. - `max_chars` is removed from `PageRules`, with the measurement recorded in place. Third instance of a defect this row already retired twice: 87308de removed `MAX_CHARS` from check-now-current.py and the `chars` key from `STATUS_RATCHET` under #364, on this exact argument. - `MAX_ROW_CHARS = 600` caps one ENTRY, so an author bounds their own row and never anyone else's. Set from the shipped pages (longest live row 520 on BENCHMARKS, 580 on FEATURES) and tighter than the cell cap it joins. - `DATED_HEADING_RE` rejects a per-attempt heading at EVERY depth, not just the `## ` the canonical-section allowlist sees. - `_prose_paragraphs` folds list items and blockquote lines into the paragraph running through them, which is the regrowth guard that actually holds. - `check_links` stops validating text no reader can follow, and pairs fences the way CommonMark does; `.agents/benchmark-record.md` gains `docs/` as a second link base because it is the declared archive of `docs/BENCHMARKS.md`. ## The independent review returned FAIL, on two blocking findings Both were real, both are fixed on the branch, and neither is a design change. F1: the claim that the entry cap was "strictly tighter than the byte cap" was FALSE. `_prose_paragraphs` excluded every line starting with `-`, `*`, `>`, `|` or `#`, so a bulleted, blockquoted or UNDATED-subsection wall was counted by nothing at all: not the paragraph count, not `MAX_PARAGRAPH_CHARS`, not `MAX_CELL_CHARS`, not `MAX_ROW_CHARS`, and not the heading guard, which fires on a date. The proof was the deleted `test_oversized_page_fails`'s OWN mutant surviving: 3,000 appended bullet lines took the page to 113,833 characters and the checker reported no errors. F2: `strip_code_spans` mis-paired fences and silently dropped a LIVE link from validation. It closed a block on any line matching a fence, including one with an info string, which CommonMark forbids; getting that wrong INVERTS fence phase for the rest of the file. With the one unbalanced fence this tree already has, the bare fence at `.agents/completed/state-events/0000-00/STATE-LEGACY-000001.md:17697` was "closed" by the ```sh at :17948, and ordinary prose at :18297 was blanked, so a reader-followable spec link stopped being checked. ## How the repair answered them F1(a): the false claim is corrected wherever it appeared, in the checker comment, in spec Design and Risks, and in the PR body. The true statement is the narrow one: the guard is tighter than the canonical-section allowlist, which sees `## ` only. It is NOT tighter than the retired byte cap in general. F1(b): the channel is closed entry-shaped. List items and blockquote lines fold into the surrounding paragraph, so a contiguous run is ONE paragraph and trips `MAX_PARAGRAPH_CHARS`, and a run per appended section trips the paragraph COUNT. The re-baselining is a re-measurement under a larger counted population, not slack: BENCHMARKS unmoved at 35 of 35, FEATURES 20 to 21 for its one list item, and `STATUS_RATCHET["long_paragraphs"]` moved the OTHER way, 82 down to 75, because STATUS's 29 list items now join neighbouring paragraphs. Leaving 82 would have banked 7 units of slack this row did not earn. F2: a closing fence must now use the opener's character, be at least as long, and carry nothing but whitespace after the marker. The corrected rule validates strictly MORE than the loose one, and its remaining losses are code samples. F3 to F9, all non-blocking, were taken in the same pass: the stored measurement of BENCHMARKS.md inside a test file removed, the second-base docstring corrected to claim only that the evidence still EXISTS rather than that it is clickable, two stray backticks that hid a non-existent target fixed and pinned, the "line numbers stay honest" justification withdrawn, the row-only scope of the no-eviction property stated in the `PageRules` docstring, the heading census re-measured (18 per page, 36 total; archive 284 of 307 dated), and the acceptance test's owed row inserted where markdown actually renders it. ## Two mutants the byte cap caught are not caught after this, both disclosed M6 SURVIVES BY DESIGN. 200 legal rows of 599 characters each grow the page without limit and the checker stays green. That is the point of the row, not an oversight: total page size is now unbounded for rows, because rows are the growth mode of a keyed table, each one is capped at 600, and nobody pays for a new measurement by deleting a row someone else owns. AGENTS.md, Records: "cap the entry, never the file." #507 IS A DISCLOSED RESIDUE. A line beginning with emphasis still starts with `*` and is still excluded, so a `**bold**`-led prose wall is unbounded. Closing it turns four paragraphs already shipped on `docs/BENCHMARKS.md` red at 717, 719, 748 and 1,084 characters against `MAX_PARAGRAPH_CHARS = 700`, so it owes an edit to a page #481 holds open and it changes what an existing gate counts. It is filed as #507, deferred as spec W8, and pinned by `test_an_EMPHASIS_lead_wall_is_a_KNOWN_residue`, which goes RED the day the residue is closed. Visible debt, not success. #495 and #498 are filed on the same terms. ## Operator gate, rerun by the operator on a scratch copy of 5f1a7fe Not quoted from the implementer's table. Measured here, each mutant applied to `docs/BENCHMARKS.md`, the real `scripts/check-public-doc-tables.py` run, then the tree restored and verified byte-for-byte by sha256 and `git status`. - M8, 3,000 appended bullet lines, the exact mutant the deleted `test_oversized_page_fails` used and the one that SURVIVED before the repair: page 45,024 -> 256,915 bytes, exit 1, CAUGHT. Two errors, 36 paragraphs over the 35 budget AND a 211,889-char paragraph over the 700 cap. - M8b, 3,000 appended blockquote lines: exit 1, CAUGHT, same two errors. - M7, 500 UNDATED `### Attempt N` sections with two bulleted lines each: page 113,305 bytes, exit 1, CAUGHT, 535 prose paragraphs over the 35 budget. - M6, 200 rows of exactly 599 chars inserted under the last table row: page 45,024 -> 165,024 bytes, exit 0, SURVIVED. Intended, per above; the repair did NOT over-tighten. - M11, 500 `**bold**`-lead paragraphs: exit 0, SURVIVED, the #507 residue reproducing exactly as disclosed. - M9, the fence/link case: `scripts/check-agent-record.py` is OK on the clean tree. The corrected rule extracts 87 targets from STATE-LEGACY-000001.md where the loose rule extracts 86, and the one it adds is precisely the live `specs/cpu-llamacpp-floor-remeasure-2026-07-22.md` link at :18297. Removing that target turns the checker RED with a dangling-link error naming that file; restoring it returns exit 0. ## Gate on the merged tree - `python3 -m pytest tests/scripts/ -q --continue-on-collection-errors`: 8 failed, 1274 passed, 3 skipped, 1 error, 1506 subtests passed. A clean `origin/main` worktree was RUN rather than assumed and reports 8 failed, 1238 passed, 3 skipped, 1 error, 1502 subtests passed: the same failure SET name for name (mlx system headers, now-render, 6 vulkan spirv subtests, and the `test_cpu_kernel_bench.py` collection error), with the merged tree's 36 extra passes being this row's new tests. One intermediate run reported 137 failures and was discarded as an ARTEFACT, not a result: the box was at 100% disk with 127 MB free, and the run reproduced the 8-failure set exactly once space was reclaimed. - `check-public-doc-tables.py`, `check-agent-record.py` (ENGINE=152 MODEL=362 QUANT=82 KERNEL=51 BACKEND=81), `check-doc-checkpoint.py`, `check-now-current.py`, `check-commit-trailers.py`, `check-pr-size.py`: OK. - `scripts/agent-preflight.sh --staged`: 68 gates ok, 0 failed. - Main landed a benchmark row while the branch sat, so the merged pages were re-verified against the NEW rules rather than the retired one: `docs/BENCHMARKS.md` 44,979 chars, 163 rows, longest row 520 of 600, paragraphs 35 of 35, longest paragraph 685 of 700; `docs/FEATURES.md` 29,967 chars, 178 rows, longest row 580 of 600, paragraphs 21 of 21; `docs/STATUS.md` on its ratchet at 11 H2, 75 long paragraphs, 44 oversized cells. Under the RETIRED 45,000-char cap the scoreboard would now have 21 characters free, which is the argument this row makes, landing on its own gate. - Main moved from 10bde85 to ee9bc7a mid-gate, gaining NemotronH and two SPEC-DSPARK measurements. Re-fetched, re-merged and re-gated in full rather than forced. - `.agents/roadmap_v1.md` is a keyed record and was not left at git's three-way result: main's version at ee9bc7a was taken WHOLESALE and the branch's four issue rows (#495, #498, #503, #507) reapplied as two pure insertion hunks, each anchor asserted to match exactly once, each issue ID asserted to appear exactly once, and zero pre-existing main lines lost. The reconstruction was then compared against git's three-way output and matched it, which is a verification of the automatic merge rather than a substitute for one. - CPU only. No GPU was used. - A red `windows-msvc-*` is not this row's: main itself does not compile under MSVC, filed as #503, and this branch touches no `src/`, `include/`, `tests/vt/`, `cmake/` or `.ps1` path. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]
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
2 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.
Closes part of #378. Dispatched as punch-list item 5 of
roadmap-v1-completion.md§3 (
ROAD-V1-A, "close the 35B c1/c2 residual"). Re-verifying the gap first, asAGENTS.md requires, changed what the task was — so this PR is a device gate plus
two record corrections, not an optimisation.
The router had not landed
The premise this work was dispatched on — that the warp-shuffle router landed and
took 35B c1/c4 to ~0.98x — is false in both halves.
git merge-base --is-ancestor: none of9f7afa7a,cce81c7e,be6a1f57,4f4ee725is an ancestor ofbbc482a2.git log -SonMoeRouterTopKWarpKernel,MoeRouterWarpEnabled,VT_MOE_ROUTER_WARP,MoeRouterWarpTreeSum,moe_router_warp: 0 commits onmain for every one.
It sat on a local-only branch, never pushed, no PR, 9 commits behind main, and
cce81c7esays plainly that the.cuwas never compiled or run. The tree neverclaimed otherwise —
KERNEL-MOE-ROUTER-WARPis in no matrix and no doc on main —so the drift was in the dispatch, not the repo. The
~0.98x / 0.87x / 0.92xfigures belong to
SPEC-DSPARK(#442), a different row, and main has sincesuperseded those too.
So the 35B residual was, in this part, a landing problem. This PR does the one
thing the branch always needed and its authors could not do: run it on a GPU.
First device verification
It compiles — not previously known. All under one
flock $HOME/gpu.lockpergroup, GB10, source
1d50cafe,RelWithDebInfo, pinned nvcc, arch121a,CUTLASS + Triton on, FA-2
ENABLEDverified in the configure log.test_moe_router_warp_maptest_ops_moe_grouped(self-pins both arms)VT_MOE_ROUTER_WARP=0test_ops_moetest_ops_moe_router_groupedtest_ops_moe_grouped_bf16test_qwen36_paged_enginewarp ONtest_qwen36_paged_engine=0rollbackCounts land exactly on the spec §11 predictions. Mutation RED re-proven: a
left-linear fold fails 76 assertions / 2 cases. The 35B gate is confirmed to
have run, not skipped — real snapshot
@491c2f1e, real continuation.Which kernel ran is proven by name, not inferred, in
nsys --cuda-graph-trace=node:MoeRouterTopKWarpKernel<__nv_bfloat16,(int)8>(VPT=8 =E=256) in the ON arm with no block kernel present, and the converse under
=0.That closes the F1 hazard at the model level — the 315/315 could not have come
from either kernel.
What it is worth, honestly
Kernel-level, ESTABLISHED (3 reps/arm, order-alternated, 1280 router calls
every leg): warp 10.422 ms (10.307/10.529) vs block 14.201 ms
(14.141/14.275) — -3.779 ms, -26.6%, 1.363x, bands non-overlapping.
Step-level: NOT SEPARABLE, and not claimed. Total GPU-busy read 1601.890 vs
1619.395 ms, but the kernels this change does not touch — identical code,
identical 155,310 launches — differ 13.726 ms between arms with a 24.140 ms spread
across the six legs, 6.4x the effect. The window is prefill+decode aggregated,
diluting 40 router calls/step across ~1600 ms. Sized against #378's 16.1863 ms/step
decode GPU-busy, the saving is ~0.5-0.75% of decode GPU-busy — below #378's
~1.5% estimate, because the warp kernel closes ~79% of the gap to
topkGating(12.98 → 8.14 us/call vs 6.85), not all of it.
Per spec gate 7 that is diagnostic evidence, not default-flip credit. The
default is left ON as
cce81c7eauthored it because that is the reviewed author'sdecision and the standing parity-enabler policy — not because this measurement
earned it. Whether default-ON is right here is a
NEEDS_DECISIONfor theoperator.
The canonical 35B grid is stale
docs/BENCHMARKS.mdcarried "CANONICAL @348c265d: 0.918x-0.972x" as current.348c265dis an ancestor of main, but 136src/+include/commits have landedsince, including
972a152c"+2.05% c8, one CastF32 per layer-step gone", whichmerged 1 h 36 min after the grid was captured. Holding vLLM fixed that implies
c8 ~0.937 and c4 ~0.980 today — an estimate, not a measurement. BENCHMARKS and
STATUS now say stale and regrid-owed rather than asserting a number nobody has
measured on current main.
So: where is the 35B gap today? Nobody knows, and a regrid is owed. The last
measured value is stale in our favour by at least one landed lever.
Also repaired
The roadmap issue table named
KERNEL-MOE-ROUTER-WARPas a stable row, whichcheck-agent-recordrejects — the spec says this is a lever insideKERNEL-MOE-ROUTINGand registers no new row. That defect would have been caughtthe first time the branch met the gate, which is itself evidence it never did.
The measurement lands in
.agents/benchmark-record.md, not the scoreboard: no vLLMarm was run, and the scoreboard's canonical sections are reference engines and
resource axes, not kernel levers.
Not verified here
compute-sanitizer memcheck.pinned oracle source is at
5559679229bc(asserted by commit), but the shared~/venvs/vllm-oraclesymlink resolves tov0.25.0-stage, not the pin. Leftuntouched for other sessions; recorded as a hazard.
Correctness gates were re-run on this exact head under load 8.94; the timing
evidence was taken earlier in a quiet window (load 0.50-1.52). No timing claim is
made from the contended run.
Next traceable hypotheses, no ceiling declared: the decode-only window, then
folding the router into the preceding gate-GEMM epilogue and batching the 40
per-step launches.