diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 91b3578a5..29a679f6f 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -19,6 +19,70 @@ from relative link targets repointed for this file's location. # Benchmarks +## MODEL-NEMOTRON-H-ABI-A2P — the A3 gate PASSES on the host, and the device divergence was the STALE input ids (2026-08-18, `row/MODEL-NEMOTRON-H-ABI-A2P-1157`, #1157, #1217, #810) + +**This supersedes the entry that recorded the A3 gate as a 6/96 device failure +with an unknown cause. The cause is known, it is not the recurrent carry, and +the host leg of the gate is a PASS.** + +**Host leg, `STRICT PASS`.** `TOKEN MATCH: 96/96 over 3 prompts (full rows=3, +short rows=0)`, driven through `include/vllm.h` alone by +`examples/nemotron_h_gen` against the committed oracle golden, on the released +`nemotron-3.5-lightning-30b-nvfp4` at revision `29f2d174`. Engine load 209.0 s; +peak RSS 20 142 392 KB; per-prompt wall 928.93 / 839.42 / 1081.16 s. +`--max-model-len 512`, greedy, `ignore_eos`. **No number on any speed axis is +claimed or implied** — this is a correctness result on a host that is not the +performance target. + +**Device leg, GB10 sm_121a, and the three measurements that name the mechanism.** +The build was not degraded: CUDA 13.x from the `ubuntu2404/sbsa` lane, +`CFG_RC=0`, `cutlass-nvfp4` / `cutlass-fp8` / `marlin-nvfp4` / `fa2` all +`ENABLED for [121a]`, `BUILD_RC=0`, `compile_errors=0`, binary sha256 +`b4677cdb7cf521250c5325fa10e5eadc80134763621d187af1f9b380c7d70140`. + +| arm | same binary, same weights, same golden | result | +|---|---|---| +| decode | the shipped path | **4/24** over the first 8 tokens of 3 prompts | +| fresh-prefill | one token per completion, so no decode step is ever taken | **24/24** | +| host | the same driver on a CPU queue, where `device_token_ids` is always null | **96/96** | + +The `got` streams in the decode arm are byte-identical to the earlier recorded +run, so the failure reproduced on a fresh build rather than drifting. + +**The per-layer trace (`VT_NEMOTRON_H_DIAG`) localises it to the first +operation of the decode step.** At the prefill step CPU and GB10 agree to six +digits on every one of the 52 layers, including bit-identical layer-0 numbers. +At the first decode step the gathered conv/SSM state is IDENTICAL on the two +(`|conv|=310.374`, `|ssm|=3985.8` on both), so the recurrent carry is exact — +and layer 0's embedding row differs. It reads `0.228135` on GB10 at BOTH decode +steps, which consumed different tokens. A constant embedding is a constant +input id. + +**Cause:** `NemotronHPagedForward` embedded the host `input.token_ids` while +`ModelForwardInput::device_token_ids` was non-null. That field's contract is +that the host vector is STALE for decode rows (`model_registry.h:314-324`), +because not materialising it on the host is the synchronize ENG-ASYNC-SCHED W4 +exists to remove. Kimi-Linear was cut from the same divergence +(`kimi_linear_device.cpp:2270-2280`); the seam that allows a third is +[#1217](https://github.com/mudler/vllm.cpp/issues/1217). + +**What this REFUTES, recorded because the wrong cause was on the record for a +day.** [#1157](https://github.com/mudler/vllm.cpp/issues/1157) reasoned that +`gm.num_decodes` might classify a decode as a prefill so the gather would hand +the mixer zeros. On real weights the trace reports `nd=1 np=0 init=[1]` on every +decode step, and mutating that mask to 0 turns the A2-P CPU gate RED (1 case, +6 assertions) — so that gate was never blind to that defect. It was blind to the +real one for a structural reason: the runner sets `device_token_ids` only under +`VLLM_CPP_CUDA` with a live device mirror, so no CPU gate reaches the branch. + +**Two things this run established about the environment, both cheap to lose.** +The released checkpoint LOADS AND DECODES ON A CPU-ONLY BOX — 20.1 GB peak RSS, +209-304 s from a CIFS mount — which is what made a same-binary host/device A/B +affordable at all and should be the first instrument reached for the next +device-only divergence on this model. And `/workspace` on the `rc` worker +persists between runs, so a cloned source tree and a CMake build directory under +`/root` survive long enough for an incremental rebuild between arms. + ## MODEL-NEMOTRON-H-ABI-A3-E2E — the A3 token gate did NOT run, and the cause on record was NOT the cause (2026-08-17, `row/MODEL-NEMOTRON-H-ABI-A3-E2E`, base `origin/main` `a6df72777`, #810) **No number is recorded, on any axis. This entry exists so the pending cause is diff --git a/.agents/issue-index.md b/.agents/issue-index.md index ca5256d52..21b09bbcd 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -384,3 +384,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | `ENG-CUDAGRAPH-DEDUP` | The owed device A/B for #1162 RAN on `dgx:gpu0` 2026-08-18 (GB10, driver 580.173.02, nvcc 13.0.88, `rc` job f88d484b, gated commit `72de552c8`, whose four dedup sources are byte-identical to the merged `2a976eb9f`) and it SPLIT. CORRECTNESS PASSED and closes [#1184](https://github.com/mudler/vllm.cpp/issues/1184): 12/12 cells exit 0, zero `invalid device function` and zero `engine-fatal` where the pre-fix head `e4ce5571a` died after exactly one replay, ON replays as often as OFF (60=60, 33=33, 43=43), and `--output-token-ids` is IDENTICAL over 10/10 comparisons with the three OFF/OFF controls passing FIRST and three workloads hashing to three DIFFERENT values. THE BENEFIT IS REFUTED for the case the row was filed for: `N == M` in every `VT_CUDA_GRAPH_DEDUP=1` cell — 3 graphs to 3 execs on sizes [24 16 8], 2 to 2 on [16 8], 2 to 2 on [32 24] — with the registry count CLIMBING 1→1, 2→2, 3→3, so more than one capture reached it and the 1:1 is a measurement rather than the single-capture artefact of the first attempt. The cause is structural and was pre-registered before the run: `AppendKernelPayload` hashes (`func`, `gridDim.{x,y,z}`, `blockDim.{x,y,z}`, `sharedMemBytes`) at `src/vt/graph_dedup_runtime.h:121-128` and the memcpy payload hashes the copy extent, so the padded batch dimension is IN THE KEY, no candidate group forms and `cudaGraphExecUpdate` is never attempted. That contradicts the row's own premise, and SGLang keys the same fields (`cuda_graph_dedup_mixin.py:105-114`), so whatever folds upstream is not decode buckets either. NO throughput or memory number is recorded — clocks unpinned AND the ON arm allocated exactly as many executables as OFF. `VT_CUDA_GRAPH_DEDUP` therefore stays OFF and the default flip is NOT JUSTIFIED on this evidence. THE OPEN HYPOTHESIS this issue owns, deliberately not decided by the record that filed it: a COARSER key keeping the function addresses and topology but dropping the launch dimensions and memcpy extents would let two padded buckets form a candidate group at all, and the probe-before-fold design means a wrong grouping costs one wasted `cudaGraphExecUpdate` probe and a private executable rather than a wrong replay, so it is a cost question needing its own spec, a red-first discrimination test and a device probe-refusal rate. Also owed here: whether ANY other capture site or two models sharing the process-singleton registry can fold (only the Qwen3 dense decode driver was exercised), and a supporting `orin:gpu0` leg, BLOCKED because the Jetson 540.4.0 driver cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`; a CUDA 12.x toolkit is the untried route). Honest gaps carried into the record: per-shape replay counts are unavailable, so B's ~30-per-shape is arithmetic over a printed TOTAL; the driver's "N captured size(s)" counts SLOTS not captures; the container's own cuBLASLt was never re-tested at CUDA 13.0. Full evidence [`benchmark-record.md`](benchmark-record.md) entry `ENG-CUDAGRAPH-DEDUP W4`; spec [`eng-cudagraph-dedup.md`](specs/eng-cudagraph-dedup.md) `## Outcome` | perf | | [#1229](https://github.com/mudler/vllm.cpp/issues/1229) | `ENG-CUDAGRAPH-DEDUP` | `origin/main` went RED at `fe24a3029` because `tests/scripts/test_check_gate_commands.py:670` pinned the literal string `Device byte-identity A/B (owed` inside `.agents/specs/eng-cudagraph-dedup.md`, and that commit legitimately rewrote the sentence from recording the device A/B as owed to recording that it RAN and what it found. The record edit is correct and the assertion described a state the record left, so every branch that merged `origin/main` inherited the failure. The protected property is stated by the assertion's own comment: the row is credited in `RUNNABLE_BASELINE` for `ctest -R test_graph_dedup` and `./scripts/agent-preflight.sh`, both CPU-tier, while gate item 6 of the spec contributes NO runnable command, so the credit covers the CPU tier and the record must not go silent about the leg it does not cover. That property is STATE-INDEPENDENT: `owed` was the leg's value on the day, not the property, and silence is the defect whether the leg was never run or the line was deleted after it ran. Asserting the NEW wording would only move the pin from `(owed` to `RAN 2026-08-18, PASS`, and the rerun this row already owes under [#1226](https://github.com/mudler/vllm.cpp/issues/1226) against a coarser signature key would red it again. FIXED by keying on structure: `gate_items`, `item_lead` and `gate_disposition` land in `scripts/check-gate-commands.py`, and the test now asserts that exactly one gate item names the device byte-identity A/B as its SUBJECT, that `runnable_commands` extracts nothing from it, and that its bold lead declares a disposition from a closed vocabulary. Lead-scoping is measured rather than assumed: a whole-item search credits item 5 with a disposition it does not declare, because its body contains `fail` in the phrase "proving the instrument can fail". Both historical wordings of item 6 are pinned as fixtures, which is what proves the new assertion is state-independent rather than re-pinned to today's sentence. The rule is NOT swept over every gate item and the reason is measured, not assumed: items 1 and 5 yield no command and declare no disposition, so a sweep is red on arrival, and widening the vocabulary until they pass would make it match ordinary English. That general rule is carried under `## Owed`. Also established and NOT a defect: `ERROR: these baseline rows left the gated population: ROW-THAT-IS-NOT-THERE` on the same run is the expected negative-path output of `test_check_mode_is_never_silently_swallowed_by_json`, confirmed by running that fixture alone. Spec [`fix-gate-commands-prose-pin.md`](specs/fix-gate-commands-prose-pin.md) | bug | | [#785](https://github.com/mudler/vllm.cpp/issues/785) | `BACKEND-ROCM` | ROCm: `PagedAttnPrefillSharedKWmma` host launches were behind `#if defined(VT_ROCWMMA_OK)`, which is never true on the host pass. This row repairs d=256 host dispatch only (`<2,8,16,32,false>`). d=512 stays scalar SharedK because shipping-f58b WMMA spills (52/212) and is a separate kernel repair. Spec [`rocm-sharedk-wmma-host-dispatch.md`](specs/rocm-sharedk-wmma-host-dispatch.md) | bug | +| [#1217](https://github.com/mudler/vllm.cpp/issues/1217) | `MODEL-NEMOTRON-H-ABI-A2P` | `ModelForwardInput::device_token_ids` carries the async runner's device-combined ids and its contract is that `token_ids` is STALE for decode rows whenever the pointer is non-null (`model_registry.h:314-324`). A registered forward that embeds the host vector then embeds the same placeholder id on EVERY decode step. The field's own comment says a model that ignores it "is simply never given one", but `runner.cpp:1408` sets the pointer for whatever model the step routes to, with no per-model opt-in and no check \| two models have now been cut from the identical divergence: Kimi-Linear (`kimi_linear_device.cpp:2270-2280`, the GB10 9/128 case) and NemotronH's paged forward under [#1157](https://github.com/mudler/vllm.cpp/issues/1157), whose A3 gate read 4/24 on GB10 against 96/96 for the same binary on CPU where the pointer is always null \| invisible because the runner sets it only under `VLLM_CPP_CUDA` with a live device mirror, so no CPU gate reaches the branch, and the failure is fluent wrong tokens rather than an error \| two closes: give `ModelFactory` an explicit `honors_device_token_ids` and have the runner fall back to the synchronous host path for a forward that has not declared it, or add a checker over the registered `.forward` entry points (a file-level grep flags ~25 false positives because several models delegate through `detail::DeviceTokenIdsScope` or the shared dense block) \| NOT fixed in the #1157 flow because one close changes a shared seam and every model factory and the other changes checker semantics, which is the "needs its own spec" case rather than the in-flow case. Listed under `## Owed` in [`nemotron-h-a2p-paged-forward.md`](specs/nemotron-h-a2p-paged-forward.md) | bug | diff --git a/.agents/specs/nemotron-h-a2p-paged-forward.md b/.agents/specs/nemotron-h-a2p-paged-forward.md index de0198134..e2cb1ab57 100644 --- a/.agents/specs/nemotron-h-a2p-paged-forward.md +++ b/.agents/specs/nemotron-h-a2p-paged-forward.md @@ -794,6 +794,41 @@ lifecycle write. ## 10. Now +**State at this commit: A2-P's PRODUCT CODE HAS LANDED, its A3 end-to-end token +gate HAS RUN on the released checkpoint, and the divergence it found has a +measured cause and a fix (#1157).** + +**The A3 gate PASSES on the host: `TOKEN MATCH: 96/96 over 3 prompts, full +rows=3, short rows=0`, `STRICT PASS`** (2026-08-18, the released +`nemotron-3.5-lightning-30b-nvfp4` at revision `29f2d174`, the committed oracle +golden, through `include/vllm.h` alone). Engine load 209.0 s, peak RSS +20 142 392 KB, per-prompt wall 928.93 / 839.42 / 1081.16 s. That is the whole +paged forward — the recurrent carry, the paged KV, the FP8 Mamba2 projections, +the NVFP4 MoE and `lm_head` — token-exact against the pinned oracle. + +**On GB10 the same binary read 4/24, and the cause was NOT the recurrent +carry.** `NemotronHPagedForward` embedded the HOST `input.token_ids` while +`ModelForwardInput::device_token_ids` was non-null, and that field's contract is +that the host vector is STALE for decode rows (`model_registry.h:314-324`). Every +decode step therefore embedded the same placeholder id. Three measurements name +it rather than infer it: fresh-prefill mode on the SAME GB10 binary, which takes +no decode step at all, read 24/24; the per-layer trace agrees to six digits +between CPU and GB10 at the prefill step across all 52 layers; and at the first +decode step the gathered conv/SSM state is IDENTICAL on the two while layer 0's +embedding row reads 0.228135 on GB10 at BOTH decode steps, which consumed +different tokens. + +**The cause recorded on #1157 was wrong, and this is the part worth keeping.** +It reasoned that `gm.num_decodes` might classify a decode as a prefill so the +gather would hand the mixer zeros. The trace reports `nd=1 np=0 init=[1]` on +every decode step on real weights, and mutating that mask to 0 turns this row's +CPU gate RED (1 case, 6 assertions) — so the gate was never blind to that +defect. It was blind to this one for a structural reason: the runner sets +`device_token_ids` only under `VLLM_CPP_CUDA` with a live device mirror, so no +CPU gate can reach the branch at all. + +**Superseded record below, kept because its corrections are still useful.** + **State at this commit: A2-P's PRODUCT CODE HAS LANDED, and its A3 end-to-end token gate has NOT RUN.** `ForwardNemotronHForCausalLM` selects `NemotronHPagedForward` whenever the runner supplies paged KV and recurrent @@ -863,8 +898,20 @@ nobody routes this architecture through a block that ropes. ## 11. Owed +- **[#1217](https://github.com/mudler/vllm.cpp/issues/1217) — the seam that let + #1157 land.** `ModelForwardInput::device_token_ids` says a forward that ignores + it "is simply never given one", but `runner.cpp:1408` sets the pointer for + whatever model the step routes to, with no per-model opt-in and no check. Two + models have now been cut from the identical divergence: Kimi-Linear + (`kimi_linear_device.cpp:2270-2280`) and this one. Not fixed in the #1157 flow + because both closes — an explicit `ModelFactory::honors_device_token_ids` with + a runner fallback, or a checker over the registered `.forward` entry points — + change a shared seam or checker semantics, which is the case AGENTS.md sends + through its own spec rather than in flow. - **The §5.4 A3 end-to-end token gate**, and the §5.7 sm_121a leg with it. Owned by this row, tracked on [#810](https://github.com/mudler/vllm.cpp/issues/810). + **The HOST leg is now GREEN (96/96, `STRICT PASS`, §10); the sm_121a leg's + green-after re-run under the #1157 fix is what remains.** Nothing about the released checkpoint is claimed until it runs. **The recorded PENDING CAUSE IS NO LONGER TRUE and was re-measured rather than inherited** (2026-08-17): §10 records contention — `dgx.casa` at loadavg 211 with 3 of diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index a560f0af4..a0bf208db 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -10,7 +10,7 @@ | **Binary release (ACTIVE; Windows pre-alpha pending)** | v0.0.2 shipped eight primary archive/checksum/provenance triplets + two indexes (26 assets) from source SHA `7020de93652ca920424a10ac5255b34810dd2f24`, run `31466516224` | Windows W14-W16 implemented. **PENDING:** native hosted gates, merged-SHA ten-tuple dry run, matching-hardware evidence, v0.0.3-pre.1 publication, 32-asset audit | W12 optional/non-primary | | **Container images (ACTIVE; arm64 cuda verified on GB10 + Orin 2026-08-11)** | `ENG-RELEASE-CONTAINERS` ([spec](../.agents/specs/container-images.md)) | cpu amd64 783 MB; cuda arm64 **1.71 GB**. GB10 `sm_121a`: `/health`+`/version`+SIGTERM on `--gpus all`. Orin `sm_87` (Tegra): Qwen3-0.6B **generates**, GPU **GR3D 95-97%** | n/a | | **Developer/row protocol** | Contribution entry point; `ENG-NOW-DERIVED` #374 @`dbd0d51c` | Entry-point gates retained. #374 W1-W5 DONE; benchmark/runtime/parity `VOID`; row specs now carry `## Now` | n/a | -| **NemotronH paged forward** (`MODEL-NEMOTRON-H-ABI-A2P`, [#810](https://github.com/mudler/vllm.cpp/issues/810)) | **No number on any axis, by the unit's own rule** ([spec](../.agents/specs/nemotron-h-a2p-paged-forward.md) §5) | **A3 gate PENDING: `nvcc` + checkpoint visibility in the rc container.** Two earlier causes here (contention, then "cannot build") were measured FALSE, see [benchmark-record](../.agents/benchmark-record.md) | CPU gate 12/12, 9/9 mutations RED. A3 driver `examples/nemotron_h_gen` exists; guards armed on a real engine (pass 0, divergence 1, short 4) | +| **NemotronH paged forward** (`MODEL-NEMOTRON-H-ABI-A2P`, [#810](https://github.com/mudler/vllm.cpp/issues/810)) | **No speed number, by the unit's own rule** ([spec](../.agents/specs/nemotron-h-a2p-paged-forward.md) §5) | **A3 host gate PASSES 96/96 `STRICT PASS`.** GB10 read 4/24; cause and fix [#1157](https://github.com/mudler/vllm.cpp/issues/1157), sm_121a re-run pending a lease | CPU gate 12/12. Load 209.0 s, peak RSS 20 142 392 KB | | **LoRA runtime W2** (`LORA-RUNTIME`, #278) | **No number owed:** correctness-only; a grid PENDS the W7 model gate | | **ARCH audit: ABI is text-only** | 4 capabilities (H3 video, Laguna, Kimi-Linear, DeepSeek-V4) reachable only from `examples/`, none registry-backed. No gate asks whether a CONSUMER can reach a capability. Documentation only | | **CUDA-graph break seam W1** (`ENG-CUDAGRAPH-BREAK`, [#1192](https://github.com/mudler/vllm.cpp/issues/1192)) | **No number owed, and none taken:** coverage and correctness row, no throughput gate declared | Capability, not a rate: mid-forward capture re-begin holds on a leased GPU; scoped forward matches eager, 500 logits, 0 differing. Probe committed, recipe and sha256 in the [record](../.agents/benchmark-record.md) | diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index e123d1335..341e68070 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -238,6 +238,8 @@ Read-only observability; none change output. | `VT_TT_TRACE_DEBUG` | unset | `=1` prints the Tenstorrent capture bisection traces to stderr: op entries (`[TT-OP]`), host readbacks (`to_vector`/`EnsureHostBytes`), device->device copies/zero-fills, and rope cos/sin cache lookups — all gated to fire only while a mesh-trace capture is active. Read-only diagnostics for the host-free decode investigation; byte-identical output when unset | | `VT_OP_PROVIDER_DISABLE` | (none) | Comma-separated provider names to disable, forcing fallback (diagnostic) | | `VT_SERVER_PREFILL_PROGRESS` | off | `=1` prints chunked-prefill progress to stderr, rate-limited to roughly 2 Hz per request. `=0` explicitly disables it even when `VT_SERVER_VERBOSE=1` | +| `VT_NEMOTRON_H_DIAG` | off | `=1` prints one line per step and per layer of `NemotronHPagedForward` to stderr: the decode/prefill split, the recurrent slot indices and the has-initial mask, then per Mamba2 layer the L2 norm of the state gathered IN and the state written OUT, and per layer the last row's mixer and residual norms. It is what separated a broken recurrent carry from a stale input id on this model ([#1157](https://github.com/mudler/vllm.cpp/issues/1157)): the carry read identical on host and GB10 while layer 0's embedding row did not. Every download it does is inside the guard, so an unset run pays nothing | +| `VT_NEMOTRON_H_DEVICE_MOE` | on | `=0` routes NemotronH's NVFP4 MoE layers back through the host reference instead of the device Marlin arm. A bisect lever for a device-vs-host divergence on this model, not a configuration: it trades the whole MoE tower's throughput for the arm the CPU gate proves token-exact | | `VT_GDN_VALIDATE` | off | Run the GDN validation/cross-check path (slower; for kernel debugging) | | `VT_FP4_AUTOTUNE_VERBOSE` | off | Log the NVFP4 GEMM autotuner's tactic selection | | `VT_H3_PROGRESS` | unset | Trace the MiniMax-H3 denoise loop's phases to stderr: which forward path was taken (device vs the CPU reference), how long the ONE-TIME device weight staging took, and per-step forward seconds with the sequence length. A real-checkpoint run spends its minutes in exactly one of those phases, and this says which without guessing — it was added after GPU-utilization counters proved unreliable on Tegra-class boards | diff --git a/docs/STATUS.md b/docs/STATUS.md index 42afdb933..94a9223dd 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -131,7 +131,7 @@ token-for-token correctness against the pinned oracle. | GLM-4 dense (sandwich norms, partial rope) | Correctness-complete, speed-pending | Token-exact 16/16 (GLM-4-9B-0414); first GLM-family model; partial interleaved RoPE + Gemma2 sandwich norms + biased qkv | | GLM-4.7-Flash (MLA + GLM MoE) | Correctness-complete, speed-pending | Token-exact 8/8 (GLM-4.7-Flash, 31.2B); reuses the DeepSeek-V2 MLA stack; first e2e coverage of the q_lora query branch + noaux_tc sigmoid router with routed-scaling | | Kimi-Linear-48B-A3B (KDA + NoPE-MLA + MoE hybrid) | **RUNNER FOLD LANDS (ROW 7 §21, #122): engine==CLI 128/128 byte-identical; golden 122/128 (near-tie profile); FA2 MLA default-ON; `vllm_complete_tokens` (ABI v13).** Grouped-router top-k block-parallel (byte-identical); no binding speed number: ckpt is tiktoken-only, so no warm-server harness. STRICT stays CLOSED. Server 19.0 tok/s wall (~0.90× vLLM floor) = speed open | paged suite 8/8·206; SACRED post-fold 35B 315/315 + 27B 235/235; thin ABI client (ratchet 8) | -| Nemotron-3.5-Lightning-30B-A3B (Mamba2 + GQA + relu2 MoE) | **Paged forward + ABI driver land (#810 A2-P, A3); e2e token gate PENDING on `nvcc` + checkpoint visibility in the rc container** | `examples/nemotron_h_gen` reaches it through `include/vllm.h` alone; G-SAFE narrows to `num_reqs <= 1`. Two earlier pending causes on this row were measured false; the box is idle and the checkpoint is revision-verified | +| Nemotron-3.5-Lightning-30B-A3B (Mamba2 + GQA + relu2 MoE) | **Host gate PASSES 96/96 `STRICT PASS`; GB10 fixed ([#1157](https://github.com/mudler/vllm.cpp/issues/1157)), sm_121a re-run pending** | ABI-only driver; G-SAFE `num_reqs <= 1`. The paged forward now honours `device_token_ids`; seam [#1217](https://github.com/mudler/vllm.cpp/issues/1217) | | Gemma-3 dense (GeGLU, dual rope, sandwich norms) | Correctness-complete, speed-pending | STRICT token-exact 48/48 greedy (gemma-3-1b-it); first Gemma-family model; GeGLU (gelu_pytorch_tanh) + dual per-layer RoPE theta + Gemma-RMSNorm sandwich norms + sqrt(hidden) embed-scale + query_pre_attn_scalar scaling | | Gemma-2 dense (attn + final logit soft-cap) | Correctness-complete, speed-pending | Near-tie-band 48/48 (gemma-2-2b-it): 44/48 strict on vLLM's greedy + 4/48 at 0.0-nat ties in vLLM's own logits; proves the attention + final logit soft-cap primitives (attn_logit_softcapping 50 + final 30); the inverse of Gemma-3 (both soft-caps, no QK-norm) | | Gemma-1 dense (the original Gemma) | Correctness-complete, speed-pending | STRICT token-exact 48/48 greedy (gemma-2b); two fused norms/layer, head_dim scale, GeGLU + sqrt(hidden) embed-scale, tied lm_head; no soft-cap/QK-norm/sliding. **D1 (2026-07-31): the whole Gemma family (1/2/3/4) folded to the default-ON bf16 merged-QKV descriptor (`MergedQkvEnabled`); re-gated Gemma-2 SACRED 48/48 (global+sliding) + Gemma-4 STRICT 32/32 — its existing gate held** | diff --git a/examples/nemotron_h_gen/main.cpp b/examples/nemotron_h_gen/main.cpp index 605f29bca..448c972f9 100644 --- a/examples/nemotron_h_gen/main.cpp +++ b/examples/nemotron_h_gen/main.cpp @@ -190,6 +190,22 @@ int main(int argc, char** argv) { // entry count and per-entry widths — is the number every count assertion // below is measured against. bool golden_info = false; + // ── #1157 DISCRIMINATOR ──────────────────────────────────────────────────── + // Generate the SAME token stream WITHOUT ever taking a decode step: ask for + // one token at a time, each from a FRESH request whose prompt is the original + // prompt plus every token generated so far. Every token then comes out of a + // PREFILL, and the recurrent state and paged KV each start empty. + // + // It is the same engine, the same weights and the same public entry point, so + // the two runs differ in exactly one thing: whether a token was produced by + // continuing a sequence or by recomputing it. A stream that is right this way + // and wrong the normal way localises the defect to the CARRY and rules out + // the tower; a stream that is wrong BOTH ways rules the carry out instead. + bool fresh_prefill = false; + // Run BOTH modes over one engine load, which is the only affordable shape + // when a load is minutes long: the two streams then differ in nothing except + // whether a token came from a decode step or from a re-prefill. + bool both_modes = false; for (int i = 1; i < argc; ++i) { const std::string a = argv[i]; auto next = [&]() -> const char* { return (i + 1 < argc) ? argv[++i] : ""; }; @@ -200,6 +216,8 @@ int main(int argc, char** argv) { else if (a == "--max-model-len") max_model_len = std::atoi(next()); else if (a == "--load-only") load_only = true; else if (a == "--golden-info") golden_info = true; + else if (a == "--fresh-prefill") fresh_prefill = true; + else if (a == "--both-modes") both_modes = true; else { std::fprintf(stderr, "unknown arg %s\n", a.c_str()); return 2; } } if (model.empty() && !golden_info) { @@ -291,22 +309,57 @@ int main(int argc, char** argv) { : static_cast(gold.entries.size()); int total_compared = 0, total_matched = 0, rows_full = 0, rows_short = 0; + const int n_modes = both_modes ? 2 : 1; + for (int mi = 0; mi < n_modes; ++mi) { + if (both_modes) { + fresh_prefill = (mi == 1); + total_compared = 0; total_matched = 0; rows_full = 0; rows_short = 0; + std::fprintf(stderr, "\n[nemotron-h] ===== MODE %s =====\n", + fresh_prefill ? "fresh-prefill" : "decode"); + } for (int pi = 0; pi < n_prompts; ++pi) { const GoldenEntry& e = gold.entries[static_cast(pi)]; std::vector gen(static_cast(steps), 0); int32_t n_gen = 0; const auto ts = std::chrono::steady_clock::now(); + if (fresh_prefill) { + // One token per completion, from a prompt that grows by the token the + // previous completion returned. `n_gen` still counts what THIS driver + // produced, so every count assertion below is unchanged. + std::vector ctx = e.prompt_token_ids; + vllm_sampling_params one = sp; + one.max_tokens = 1; + for (int t = 0; t < steps; ++t) { + int32_t got = 0, n_one = 0; + const vllm_status s1 = vllm_complete_tokens( + eng, ctx.data(), static_cast(ctx.size()), &one, &got, 1, + &n_one, nullptr); + if (s1 != VLLM_OK || n_one != 1) { + std::fprintf(stderr, + "[nemotron-h] prompt %d fresh-prefill step %d FAILED " + "(status=%d n=%d): %s\n", + pi, t, static_cast(s1), static_cast(n_one), + vllm_last_error()); + vllm_engine_free(eng); + return 1; + } + gen[static_cast(t)] = got; + ctx.push_back(got); + ++n_gen; + } + } else { const vllm_status st = vllm_complete_tokens( eng, e.prompt_token_ids.data(), static_cast(e.prompt_token_ids.size()), &sp, gen.data(), static_cast(gen.size()), &n_gen, nullptr); - const auto te = std::chrono::steady_clock::now(); if (st != VLLM_OK) { std::fprintf(stderr, "[nemotron-h] prompt %d FAILED: %s\n", pi, vllm_last_error()); vllm_engine_free(eng); return 1; } + } + const auto te = std::chrono::steady_clock::now(); const int expected = static_cast(e.token_ids.size()); const int n = std::min(expected, static_cast(n_gen)); @@ -336,8 +389,10 @@ int main(int argc, char** argv) { std::fprintf(stderr, "\n[nemotron-h] TOKEN MATCH: %d/%d over %d prompt(s) " - "(full rows=%d, short rows=%d)\n", - total_matched, total_compared, n_prompts, rows_full, rows_short); + "(full rows=%d, short rows=%d, mode=%s)\n", + total_matched, total_compared, n_prompts, rows_full, rows_short, + fresh_prefill ? "fresh-prefill" : "decode"); + } vllm_engine_free(eng); // A pass needs three things to be true at once, and each is checked here diff --git a/src/vllm/model_executor/models/nemotron_h_device.cpp b/src/vllm/model_executor/models/nemotron_h_device.cpp index 4bd6ad1cd..4455998d9 100644 --- a/src/vllm/model_executor/models/nemotron_h_device.cpp +++ b/src/vllm/model_executor/models/nemotron_h_device.cpp @@ -79,6 +79,7 @@ #include #include +#include #include #include #include @@ -1111,6 +1112,49 @@ std::vector OwnedToF32(const NemotronHOwned& w) { return out; } +// ─── #1157 DIAGNOSTIC (VT_NEMOTRON_H_DIAG, documented in ENVIRONMENT.md) ──── +// +// Off unless `VT_NEMOTRON_H_DIAG` is set to something other than "0", and every +// download it does is inside that guard, so a production step pays nothing. +// +// It exists to answer the question no CPU gate on this model can: the runner +// hands a decode step a device-resident input id and a recurrent page, and when +// the tokens come out wrong, only the per-layer numbers say WHICH of the two the +// step actually read. On #1157 they said the carry was exact — the state +// gathered at step k+1 equalled the state written at step k, on host and on +// GB10 alike — and that layer 0's embedding row was constant across two decode +// steps that consumed different tokens. It stays for the next reader of this +// model, because the next divergence here will be diagnosed the same way. +bool NemotronHDiagEnabled() { + static const bool on = [] { + const char* e = std::getenv("VT_NEMOTRON_H_DIAG"); + return e != nullptr && e[0] != '0'; + }(); + return on; +} + +// #1157 BISECT SWITCH. The device MoE arm is 23 of this model's 52 layers and +// has never run at T=1 anywhere: its own gate (test_nemotron_h_moe_device.cpp) +// exercises T=4 and T=2. Setting `VT_NEMOTRON_H_DEVICE_MOE=0` routes those +// layers back through the host reference the CPU arm already proves token-exact +// on this checkpoint, so one run says whether the device MoE is the difference. +bool NemotronHDeviceMoeEnabled() { + static const bool on = [] { + const char* e = std::getenv("VT_NEMOTRON_H_DEVICE_MOE"); + return e == nullptr || e[0] != '0'; + }(); + return on; +} + +double DiagL2(const std::vector& v, int64_t off, int64_t n) { + double acc = 0.0; + for (int64_t i = 0; i < n; ++i) { + const double x = v[static_cast(off + i)]; + acc += x * x; + } + return std::sqrt(acc); +} + // ─── the per-step device inputs ───────────────────────────────────────────── // // Uploaded ONCE per step and shared by all 6 attention layers and all 23 @@ -1213,6 +1257,24 @@ NemotronHPagedStep BuildNemotronHPagedStep(Dev d, const ModelForwardInput& input // for the same hazard and the same remedy). Waiting here costs nothing this // unit measures: A2-P records no throughput number on any axis (spec §5). d.b.Synchronize(d.q); + if (NemotronHDiagEnabled()) { + std::fprintf(stderr, "[NH-DIAG] step T=%lld R=%lld nd=%lld np=%lld idx=[", + static_cast(T), static_cast(R), + static_cast(nd), static_cast(np)); + for (int64_t r = 0; r < R; ++r) + std::fprintf(stderr, "%d%s", idx[static_cast(r)], r + 1 < R ? "," : ""); + std::fprintf(stderr, "] init=["); + for (int64_t r = 0; r < R; ++r) + std::fprintf(stderr, "%d%s", init[static_cast(r)], r + 1 < R ? "," : ""); + std::fprintf(stderr, "] qsl_attn=["); + for (size_t i = 0; i < am.query_start_loc.size(); ++i) + std::fprintf(stderr, "%d%s", am.query_start_loc[i], + i + 1 < am.query_start_loc.size() ? "," : ""); + std::fprintf(stderr, "] seq_lens=["); + for (size_t i = 0; i < am.seq_lens.size(); ++i) + std::fprintf(stderr, "%d%s", am.seq_lens[i], i + 1 < am.seq_lens.size() ? "," : ""); + std::fprintf(stderr, "]\n"); + } return sdi; } @@ -1467,14 +1529,47 @@ ForwardLogits NemotronHPagedForward(const NemotronHHostWeights& host, RequireDeviceWeight(host.norm_f, "backbone.norm_f.weight", adt, {H}); DBuf residual(d, adt, {T, H}); { - std::vector ids = input.token_ids; - for (int32_t id : ids) { - VT_CHECK(id >= 0 && id < V, "NemotronH paged forward: token id out of range"); - } - DBuf it(d, DType::kI32, {T}, ids.data()); - d.b.Synchronize(d.q); // `ids` is a local; see UploadAs. Tensor tab = ResidentWeight(d, host.embeddings); - vt::Embedding(d.q, residual.t(), tab, it.t()); + Tensor rt = residual.t(); + if (input.device_token_ids != nullptr) { + // ★ ENG-ASYNC-SCHED W4 (#1157). `ModelForwardInput::device_token_ids` is + // non-null exactly when the async runner's device combine has already + // spliced each DECODE row's sampled token into ITS device buffer and left + // the host `token_ids` STALE on purpose — materializing it on the host is + // the synchronize W4 exists to remove (model_registry.h:314-324, + // runner.cpp:1175-1194). A forward that embeds the host vector therefore + // embeds the same placeholder id on every decode step. + // + // That is not a hypothesis. With the host vector, this model's A3 token + // gate read 4/24 on GB10 while the SAME binary read 24/24 in + // fresh-prefill mode (no decode step is ever taken) and 96/96 on CPU + // (where this pointer is always null), and the per-layer trace showed the + // layer-0 embedding row identical across two consecutive decode steps + // that consumed different tokens. + // + // Kimi-Linear was cut from this same divergence + // (kimi_linear_device.cpp:2270-2280) and every other registered forward + // already honours the field. This one did not, and nothing could see it: + // the runner sets the pointer only under VLLM_CPP_CUDA with a live device + // mirror, so no CPU gate can reach the branch at all. + // + // The host-side range check below is deliberately NOT repeated here. The + // ids live on the device and validating them would need the D2H + // synchronize this path exists to delete; `LaunchCombineSampledAndDraft + // Tokens` produces them from the sampler's own output, and vt::Embedding + // bounds-checks the gather. + Tensor ids = MakeTensor(const_cast(input.device_token_ids), + DType::kI32, d.q.device, {T}); + vt::Embedding(d.q, rt, tab, ids); + } else { + std::vector ids = input.token_ids; + for (int32_t id : ids) { + VT_CHECK(id >= 0 && id < V, "NemotronH paged forward: token id out of range"); + } + DBuf it(d, DType::kI32, {T}, ids.data()); + d.b.Synchronize(d.q); // `ids` is a local; see UploadAs. + vt::Embedding(d.q, rt, tab, it.t()); + } } vt::RmsNormArgs nargs; @@ -1512,6 +1607,7 @@ ForwardLogits NemotronHPagedForward(const NemotronHHostWeights& host, const bool moe_on_device = lw.block == NemotronHBlock::kMoe && adt == DType::kBF16 && MoeIsNvfp4(lw.moe) && + NemotronHDeviceMoeEnabled() && vt::OpRegistered(vt::OpId::kMoeGroupedGemmNvfp4Marlin, d.q.device.type); const bool needs_host = lw.block != NemotronHBlock::kAttention && !moe_on_device; std::vector nvec; @@ -1540,6 +1636,12 @@ ForwardLogits NemotronHPagedForward(const NemotronHHostWeights& host, params.mamba_num_heads * params.mamba_head_dim * params.ssm_state_size; std::vector conv_all = DownloadF32(d, io.conv, DType::kF32, R * conv_row); std::vector ssm_all = DownloadF32(d, io.ssm, DType::kF32, R * ssm_row); + if (NemotronHDiagEnabled()) { + std::fprintf(stderr, + "[NH-DIAG] L%lld mamba GATHERED |conv|=%.6g |ssm|=%.6g\n", + static_cast(l), DiagL2(conv_all, 0, conv_row), + DiagL2(ssm_all, 0, ssm_row)); + } // At `num_reqs == 1` this loop runs once, and it is written as a loop for // the reason §4.1 gives: the indexing machinery lands here, only the @@ -1587,6 +1689,14 @@ ForwardLogits NemotronHPagedForward(const NemotronHHostWeights& host, ssm_all.begin() + static_cast(r * ssm_row)); } + if (NemotronHDiagEnabled()) { + std::fprintf(stderr, + "[NH-DIAG] L%lld mamba WROTE |conv|=%.6g |ssm|=%.6g " + "|out|=%.6g\n", + static_cast(l), DiagL2(conv_all, 0, conv_row), + DiagL2(ssm_all, 0, ssm_row), + DiagL2(mvec, (T - 1) * H, H)); + } io.conv = UploadAs(d, conv_all, DType::kF32, {R, params.conv_dim(), params.conv_kernel - 1}); io.ssm = UploadAs(d, ssm_all, DType::kF32, @@ -1610,6 +1720,18 @@ ForwardLogits NemotronHPagedForward(const NemotronHHostWeights& host, carry = UploadAs(d, mvec, adt, {T, H}); } + if (NemotronHDiagEnabled()) { + const std::vector cv = DownloadF32(d, carry, adt, T * H); + const std::vector rs = DownloadF32(d, residual, adt, T * H); + const char* kind = lw.block == NemotronHBlock::kAttention ? "attn" + : lw.block == NemotronHBlock::kMamba ? "mamba" + : lw.block == NemotronHBlock::kMoe ? "moe" + : "mlp"; + std::fprintf(stderr, + "[NH-DIAG] L%lld %-5s |mixer_last|=%.6g |resid_last|=%.6g\n", + static_cast(l), kind, DiagL2(cv, (T - 1) * H, H), + DiagL2(rs, (T - 1) * H, H)); + } if (trace != nullptr && trace->capture) { trace->normed[static_cast(l)] = std::move(nvec); std::vector h = DownloadF32(d, residual, adt, T * H); diff --git a/tests/vllm/models/test_nemotron_h_moe_device.cpp b/tests/vllm/models/test_nemotron_h_moe_device.cpp index 95fcce4c5..181f070cb 100644 --- a/tests/vllm/models/test_nemotron_h_moe_device.cpp +++ b/tests/vllm/models/test_nemotron_h_moe_device.cpp @@ -227,10 +227,22 @@ TEST_CASE("NemotronH A2-Q2a: the device MoE block matches the host reference on const NemotronHParams p = MoeParams(); Queue hq{Device{DeviceType::kCPU, 0}, nullptr}; const DType dt = DType::kBF16; // Marlin's a/c contract (ops.cpp:879) - const int64_t T = 4; const int64_t H = p.hidden_size; const NemotronHMoeWeights w = MakeNvfp4Moe(p, dt); + + // ★ T == 1 IS THE DECODE SHAPE, AND UNTIL #1157 THIS CASE NEVER RAN IT. + // The widths here were 4 and 2, both of them PREFILL shapes. Every token + // after the first comes out of a step with exactly one token, so the arm this + // model spends its whole decode in was the one width the gate did not cover — + // and `MarlinMoeAlignBlockSizeSelect` / `MarlinMoeAlignSizes` take different + // branches at a token count below the expert count, which is what T=1 with + // 128 experts is. A width loop rather than a third copy, so the three cannot + // drift apart, and the count is asserted afterwards: a loop that ran over + // nothing would otherwise report a clean pass. + int64_t widths_covered = 0; + for (const int64_t T : {int64_t{1}, int64_t{2}, int64_t{4}}) { + INFO("token count T=" << T); const std::vector x = SynthVec(static_cast(T * H), 77, 0.5F); // The HOST arm dequantizes each touched expert to bf16 and runs the per-pair @@ -301,6 +313,10 @@ TEST_CASE("NemotronH A2-Q2a: the device MoE block matches the host reference on REQUIRE(guard_examined == examined); INFO("does the band " << band << " REJECT a routed-scale defect?"); CHECK(guard >= band); + ++widths_covered; + } + // Three widths, or the loop did not run the one this case was extended for. + REQUIRE(widths_covered == 3); } // A separate case so a `-tc` run can select it alone. Same no-comma rule. diff --git a/tests/vt/test_ops_paged_attn.cpp b/tests/vt/test_ops_paged_attn.cpp index f29deed4f..3fa0ab395 100644 --- a/tests/vt/test_ops_paged_attn.cpp +++ b/tests/vt/test_ops_paged_attn.cpp @@ -1196,8 +1196,18 @@ TEST_CASE("paged_attention CUDA FA-2 prefill (bf16 q/kv/out) matches f32 ref at // kBlockN=128 for d128: covers up to ceil(140/128)*128 = 256 keys = 16 pages; // size the block table to that so no column is read past its row (memcheck). const int64_t max_blocks = 16, num_blocks = 64; + // ★ {32, 2} IS NEMOTRON-3.5-LIGHTNING, AND IT WAS NOT HERE (#1157). + // Every case in this file measured ngroups 2 (16/8) and 4 (32/8), the two + // Qwen3-dense gate configs the varlen d128 decode path was written for. The + // path is DEFAULT ON for ANY bf16 causal pure-decode at head_dim 128, so the + // first model to arrive with a different ratio reaches an untested grid — and + // NemotronH-3.5-Lightning-30B is 32 query heads over 2 KV heads, ngroups 16, + // four times the widest group count ever measured here. Its A3 token gate + // failed 6/96 on the device while the same binary decodes the same checkpoint + // token-exact on CPU, which is what sent the search here. for (const auto& ratio : {std::pair{16, 8}, - std::pair{32, 8}}) { + std::pair{32, 8}, + std::pair{32, 2}}) { const int64_t Hq = ratio.first, Hk = ratio.second, page = Hk * D; CAPTURE(Hq); auto qf = RandF32(static_cast(num_tokens * Hq * D), @@ -1569,8 +1579,18 @@ TEST_CASE("paged_attention CUDA FA-2 varlen d128 decode matches composed referen MESSAGE("no CUDA backend; skipping FA-2 varlen d128 decode parity (dgx-pending)"); return; } + // ★ {32, 2} IS NEMOTRON-3.5-LIGHTNING, AND IT WAS NOT HERE (#1157). + // Every case in this file measured ngroups 2 (16/8) and 4 (32/8), the two + // Qwen3-dense gate configs the varlen d128 decode path was written for. The + // path is DEFAULT ON for ANY bf16 causal pure-decode at head_dim 128, so the + // first model to arrive with a different ratio reaches an untested grid — and + // NemotronH-3.5-Lightning-30B is 32 query heads over 2 KV heads, ngroups 16, + // four times the widest group count ever measured here. Its A3 token gate + // failed 6/96 on the device while the same binary decodes the same checkpoint + // token-exact on CPU, which is what sent the search here. for (const auto& ratio : {std::pair{16, 8}, - std::pair{32, 8}}) { + std::pair{32, 8}, + std::pair{32, 2}}) { for (const int batch : {1, 2, 4, 8}) { for (const int base_len : {5, 21, 1024}) { // short => num_splits==1; long => split CAPTURE(ratio.first); @@ -1615,8 +1635,18 @@ TEST_CASE("paged_attention CUDA FA-2 varlen d128 decode GQA group-swap matches c MESSAGE("no CUDA backend; skipping FA-2 varlen d128 group-swap parity (dgx-pending)"); return; } + // ★ {32, 2} IS NEMOTRON-3.5-LIGHTNING, AND IT WAS NOT HERE (#1157). + // Every case in this file measured ngroups 2 (16/8) and 4 (32/8), the two + // Qwen3-dense gate configs the varlen d128 decode path was written for. The + // path is DEFAULT ON for ANY bf16 causal pure-decode at head_dim 128, so the + // first model to arrive with a different ratio reaches an untested grid — and + // NemotronH-3.5-Lightning-30B is 32 query heads over 2 KV heads, ngroups 16, + // four times the widest group count ever measured here. Its A3 token gate + // failed 6/96 on the device while the same binary decodes the same checkpoint + // token-exact on CPU, which is what sent the search here. for (const auto& ratio : {std::pair{16, 8}, - std::pair{32, 8}}) { + std::pair{32, 8}, + std::pair{32, 2}}) { for (const int batch : {1, 2, 4, 8}) { for (const int base_len : {5, 21, 1024}) { // short => num_splits==1; long => split CAPTURE(ratio.first); @@ -1645,8 +1675,18 @@ TEST_CASE("paged_attention CUDA FA-2 varlen d128 decode swap near-ties the plain MESSAGE("no CUDA backend; skipping FA-2 varlen d128 swap-vs-plain near-tie (dgx-pending)"); return; } + // ★ {32, 2} IS NEMOTRON-3.5-LIGHTNING, AND IT WAS NOT HERE (#1157). + // Every case in this file measured ngroups 2 (16/8) and 4 (32/8), the two + // Qwen3-dense gate configs the varlen d128 decode path was written for. The + // path is DEFAULT ON for ANY bf16 causal pure-decode at head_dim 128, so the + // first model to arrive with a different ratio reaches an untested grid — and + // NemotronH-3.5-Lightning-30B is 32 query heads over 2 KV heads, ngroups 16, + // four times the widest group count ever measured here. Its A3 token gate + // failed 6/96 on the device while the same binary decodes the same checkpoint + // token-exact on CPU, which is what sent the search here. for (const auto& ratio : {std::pair{16, 8}, - std::pair{32, 8}}) { + std::pair{32, 8}, + std::pair{32, 2}}) { for (const int batch : {2, 8}) { for (const int base_len : {21, 1024}) { CAPTURE(ratio.first); @@ -1691,8 +1731,18 @@ TEST_CASE("paged_attention CUDA FA-2 varlen d128 decode num_splits cap engages a MESSAGE("no CUDA backend; skipping FA-2 varlen d128 num_splits-cap check (dgx-pending)"); return; } + // ★ {32, 2} IS NEMOTRON-3.5-LIGHTNING, AND IT WAS NOT HERE (#1157). + // Every case in this file measured ngroups 2 (16/8) and 4 (32/8), the two + // Qwen3-dense gate configs the varlen d128 decode path was written for. The + // path is DEFAULT ON for ANY bf16 causal pure-decode at head_dim 128, so the + // first model to arrive with a different ratio reaches an untested grid — and + // NemotronH-3.5-Lightning-30B is 32 query heads over 2 KV heads, ngroups 16, + // four times the widest group count ever measured here. Its A3 token gate + // failed 6/96 on the device while the same binary decodes the same checkpoint + // token-exact on CPU, which is what sent the search here. for (const auto& ratio : {std::pair{16, 8}, - std::pair{32, 8}}) { + std::pair{32, 8}, + std::pair{32, 2}}) { for (const int batch : {1, 2, 4}) { CAPTURE(ratio.first); CAPTURE(batch);