fix(#1157): NemotronH decoded with STALE input ids, and the recurrent carry was never the cause

.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

.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 |

.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

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) |