mudler · localai-bot · Aug 17, 2026 · Aug 18, 2026
diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md
@@ -19,6 +19,81 @@ from relative link targets repointed for this file's location.
 
 # Benchmarks
 
+## MODEL-NEMOTRON-H-ABI-A3-E2E — the A3 token gate RAN on real weights, and FAILED 6/96 (2026-08-17, `row/MODEL-NEMOTRON-H-ABI-A3-E2E`, dgx.casa GB10 sm_121a, #810, #1157)
+
+**This supersedes the entry below it, which was written while the gate could not
+run. The gate has now run. It is a FAIL, recorded as a result.**
+
+`GATE_RC=1`. **6 of 96 tokens match.** Driven entirely through `include/vllm.h`
+(`vllm_engine_load` + `vllm_complete_tokens`) by `examples/nemotron_h_gen`
+against the committed oracle golden.
+
+```
+prompt 0  compared=32 matched=2   got 6993,1046,1317,1048,1032,1050,1050,1050,...
+                                  exp 6993,1046,1256,1010,1784,8961,1307,10787,...
+prompt 1  compared=32 matched=3   got 1032,1058,1048,1050,1010,1050,1050,1050,...
+                                  exp 1032,1048,1044,1032,1049,1044,1032,1049,...
+prompt 2  compared=32 matched=1   got 1349,1010,1032,1032,1032,1032,1050,1050,...
+                                  exp 1349,3468,5702,3244,1395,1261,30622,13406,...
+```
+
+**Token 0 is correct on ALL THREE prompts; everything after decays and collapses
+onto the repeated id `1050` by about step 5.** Token 0 comes from prefill. Every
+later token comes from a decode step that must read what the previous step
+wrote. A correct first token followed by convergence onto a fixed point is the
+signature of a broken recurrent carry, not of numeric drift: drift wanders, it
+does not settle on one id and stay.
+
+**The prediction this confirms, written in this row's own documents before the
+run:** the parent spec §6d matched 3/3 FIRST tokens against a forward carrying
+no state at all and warned that is "exactly how little a first token proves".
+This run reproduces 3/3 first tokens and then fails. And
+`nemotron-h-a2p-paged-forward.md` §10 records that on the synthetic fixture the
+token arms could NOT see the recurrent carry — conv and SSM zeroed on every step
+still produced byte-identical output, and P-M1/P-M2/P-M4/P-M9 all survived.
+**The two numeric cases added to compensate were not sufficient: the CPU gate is
+12/12 green at the very commit that produces this.** Green unit gate, wrong
+model — the class works, the capability does not.
+
+**Conditions, none of them degraded.** CUDA 13.3.73 from the `ubuntu2404/sbsa`
+lane; `CFG_RC=0`; `VOID_FLAG=0` with all four cells `ENABLED for [121a]`
+(`cutlass-nvfp4`, `cutlass-fp8`, `marlin-nvfp4`, `fa2` — not `[121]`, not
+DISABLED); `BUILD_RC=0`, `compile_errors=0`, `enospc=0`; binary sha256
+`b2a8f6a986938037d9f17edb5b93550512127e1913d34cfb1c19bcd7e47253ec`.
+
+**The instrument was not mute:** `compared=96` over 3 prompts of width 32,
+`full rows=3, short rows=0` — the expected geometry exactly. A run that compared
+fewer would have exited 4 instead of 1.
+
+**Ungated numbers from the same run, recorded because they were measured, not
+because anything claims them:** engine load **397.0 s**; peak host RAM **43 405
+MB** over 676 samples (well inside 119 GB, so `gpu_memory_utilization`'s
+host-RAM blindness did not bite here); decode ≈**10.3 s/token**, with `lm_head`
+and the FP8 Mamba2 projections still computing on the host.
+
+**Two environment facts this run established, both costly to rediscover:**
+
+- **The `rc` worker container is Ubuntu 24.04 as uid 0 with a full toolchain and
+  network, but NO `nvcc`.** apt's own `nvidia-cuda-toolkit` is 12.0.140, too old
+  for sm_121a. The NVIDIA repo lane is **`ubuntu2404/sbsa`**, NOT
+  `ubuntu2404/arm64`: the arm64 path answers `HTTP/2 200` but is the
+  Jetson/Tegra lane and carries no `cuda-toolkit-13-*` at all. Select
+  `cuda-toolkit-13-<N>` by an anchored pattern — an unanchored "newest
+  `cuda-toolkit-13*`" selects `cuda-toolkit-13-config-common`, which installs
+  cleanly, ships no compiler, and returns 0.
+- **`/workspace` is CIFS and refuses symlinks, so a CMake build directory cannot
+  live there.** A build that compiled the entire tree cleanly
+  (`compile_errors=0`) died on its last step at
+  `cmake_symlink_library libvllm.so.0.0.3 -> Operation not supported`, costing
+  1h33m. Build on the container-local overlay; keep only source and staged CUDA
+  on `/workspace`, which does persist between `rc run` invocations
+  (`WORKSPACE PERSISTED` marker confirmed).
+
+**Repair is owed and is NOT this row's author's**, per the fresh-implementer
+rule: the defect is in `NemotronHPagedForward`'s state carry, and this row wrote
+the driver that found it. Tracked on
+[#1157](https://github.com/mudler/vllm.cpp/issues/1157).
+
 ## 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
@@ -380,3 +380,4 @@ rather than merged. `scripts/check-agent-record.py` gates both.
 | [#1225](https://github.com/mudler/vllm.cpp/issues/1225) | `SPEC-DSPARK-BLOCK-SIZE-GUARD` | A DSpark speculative length below the draft's block is accepted silently. `ResolveDspark` carries upstream's `k >= dspark_block_size` hard error (`include/vllm/config/speculative.h:179-185`, from `vllm/config/speculative.py:1003-1027` @ `555967922`) and both production call sites pass `std::nullopt` for `n_predict` and for `dspark_block_size` (`src/vllm/entrypoints/model_loader.cpp:881-883` and `:1675-1677`), so the floor reaches no user path and only `tests/vllm/config/test_speculative_dspark.cpp:99-107` drives it. Nothing in our draft path reads the checkpoint's block key — the block layout is sized by `k` alone (`include/vllm/v1/worker/gpu/spec_decode/dspark/speculator.h:56`) and no weight is block-shaped — so a short `k` raises no shape error and drafts a structurally wrong block while the tokens keep flowing. A literal port does NOT close it: `dspark_block_size` appears in no pinned file but `speculative.py`, and neither published Qwen3 draft sets it (`deepseek-ai/dspark_qwen3_4b_block7` and `RadixArk/Qwen3.8-27B-DSpark` @ `85ef153b` both carry `block_size: 7`, no `n_predict`), while upstream's `block_size` normalization at `:945-961` is Gemma4-only — so upstream accepts `k=6` on both sides of vllm#52197. Closing it for the lane we ship needs `block_size` as the floor fallback, one tracked divergence argued in the spec and the commit | bug |
 | [#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 |
+| [#1157](https://github.com/mudler/vllm.cpp/issues/1157) | `MODEL-NEMOTRON-H-ABI-A2P` | NemotronH A3 token gate DIVERGES on real weights: 6/96 tokens, `GATE_RC=1`. Token 0 is correct on all 3 prompts and decode then collapses onto a repeated id, which is the paged forward's recurrent carry rather than drift. The A2-P CPU gate is 12/12 green at the same commit, so the synthetic fixture cannot see it \| owned by `MODEL-NEMOTRON-H-ABI-A2P`, repair to a fresh implementer | bug |
diff --git a/.agents/specs/nemotron-h-a2p-paged-forward.md b/.agents/specs/nemotron-h-a2p-paged-forward.md
@@ -863,9 +863,15 @@ nobody routes this architecture through a block that ropes.
 
 ## 11. Owed
 
-- **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).
-  Nothing about the released checkpoint is claimed until it runs. **The recorded
+- **The §5.4 A3 end-to-end token gate HAS RUN and FAILED: 6/96 tokens,
+  `GATE_RC=1`** (2026-08-17, sm_121a leg included,
+  [#1157](https://github.com/mudler/vllm.cpp/issues/1157)). Token 0 is correct
+  on all three prompts and decode then collapses onto a repeated id, which is
+  this unit's state carry rather than drift. **§10's "one measured surprise" is
+  the direct cause of it going undetected**: the synthetic fixture could not see
+  the recurrent carry, and the two numeric cases added to compensate were not
+  sufficient, because the CPU gate is 12/12 green at the commit that produces
+  this. A2-P is NOT done. Repair is owed to a fresh implementer. **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
   119 GB — and the box now answers at **loadavg 0.36 with 115 of 119 GB

diff --git a/.agents/specs/nemotron-h-abi-e2e.md b/.agents/specs/nemotron-h-abi-e2e.md
@@ -1121,34 +1121,35 @@ five malformed-golden shapes each exit 2. That last one is the guard this
 section's own §5.2 is really about — a comparison over too few elements reports
 a perfect score, and here it cannot.
 
-**§5.2's A3 token gate has NOT RUN. Its recorded cause was wrong twice, and both
-corrections are kept here because the second one is a trap this section can save
-the next reader from.**
-
-*First cause, dead:* contention. Re-measured, `dgx.casa` is idle (loadavg 0.36,
-115 of 119 GB available, GPU 0%), the checkpoint is present, and its first shard
-hashes to revision `29f2d174`'s own LFS record.
-
-*Second cause, also dead:* "nothing can build a gate binary". **That was a HOST
-measurement reported as a CONTAINER measurement, and the two are different
-machines for this purpose.** The host genuinely has no `nvcc`/`cmake` since the
-14 Aug reimage ([#1019](https://github.com/mudler/vllm.cpp/issues/1019)) — but
-the host is not where work runs. Inside `rc run` the worker container is Ubuntu
-24.04 as **uid 0**, carrying `gcc`, `g++`, `cmake`, `ninja`, `make`, `python3`,
-`git`, `apt`, with the GB10 visible and DNS working. **Only `nvcc` is missing**,
-and apt's `nvidia-cuda-toolkit` 12.0.140 is too old for sm_121a, so CUDA 13.x
-comes from `developer.download.nvidia.com/…/ubuntu2404/arm64`. Neither `docker`
-nor `sudo` is involved.
-
-> **Rule this cost two cycles to learn: re-derive every environment fact INSIDE
-> `rc run`.** A probe that runs somewhere other than where the work will run
-> answers a question nobody asked, and it fails toward a confident verdict about
-> the code's environment rather than toward an obvious error.
-
-*What is genuinely outstanding:* install `nvcc` in the build container, and
-whether that container can see `$CHECKPOINT_ROOT`, which is **OPEN** — no probe
-has answered it, and none is claimed. `docs/BENCHMARKS.md` records the gate as
-**pending a named resource**, never as a pass.
+**★ §5.2's A3 token gate HAS NOW RUN, ON REAL WEIGHTS, AND IT FAILED: 6/96
+tokens, `GATE_RC=1`** (2026-08-17, dgx.casa GB10 sm_121a,
+[#1157](https://github.com/mudler/vllm.cpp/issues/1157)). Driven entirely
+through `include/vllm.h` by `examples/nemotron_h_gen`.
+
+**Token 0 is correct on all three prompts; decode then decays and collapses onto
+one repeated id by about step 5.** That is the recurrent carry, not numeric
+drift — drift wanders, it does not settle on a fixed point. Full token streams,
+conditions and ungated numbers are in
+[`benchmark-record.md`](../benchmark-record.md).
+
+**This section predicted its own failure and should be read as evidence that the
+prediction was right.** §5.2 says a single-step gate "is structurally blind to
+the entire class of defect A2 introduces", and the parent spec §6d had already
+matched 3/3 FIRST tokens against a forward carrying no state at all. This run
+matched 3/3 first tokens and then failed. The A2-P unit gate is **12/12 green at
+the same commit**, so the CPU arms cannot see this: a green class and a wrong
+capability.
+
+The build was not degraded and cannot explain it: CUDA 13.3.73, `CFG_RC=0`,
+`VOID_FLAG=0` with all four cells `ENABLED for [121a]`, `BUILD_RC=0`,
+`compile_errors=0`. The instrument was not mute: `compared=96`, `full rows=3`,
+`short rows=0`.
+
+*Two environment facts, both established the hard way and both worth reusing:*
+the `rc` container needs CUDA from the **`ubuntu2404/sbsa`** lane (the `arm64`
+lane answers 200 but carries no `cuda-toolkit-13-*`), and **a CMake build
+directory cannot live on `/workspace`**, which is CIFS and refuses the
+`libvllm.so.0` symlink after a fully successful compile.
 
 **§5.2 arm 2 (multi-request) is additionally blocked by design, not by a host.**
 G-SAFE refuses `input.num_reqs > 1` and A2-B owns that clause, so the three