diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 8a786f138..cd732c4aa 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -25056,3 +25056,273 @@ narrower f64 vector explains without anything further. SGLang-Omni is still `gateable = no`. Every reference axis in `docs/BENCHMARKS.md` stays `PENDING`. Everything above is an internal two-arm number on one named box. + +## A2-D1 — NemotronH decodes on the single-step recurrent kernels (#1311) + +**Verdict: this change is TOKEN-NEUTRAL on both gated hosts, and its SPEED +hypothesis is REFUTED on both. The GB10 `95/96` is A2-Q1's device mamba arm and +NOT this row and NOT the architecture** — resolved by a three-leg run, see the +sm_121a section. That same run measured **A2-Q1's arm at 6.64x per output +token** (718.1x -> 108.2x vs vLLM), which is the larger finding on this page.** Measured as a same-binary A/B on `thor:gpu0` +(sm_110) and `dgx:gpu0` (GB10, sm_121a); recipe +`scripts/nemotron-h-a2d1-gpu-gate.sh`. + +| host | ON (single-step) | OFF (chunk scan) | per-token move | verdict | +|---|---|---|---|---| +| sm_110 | `96/96 STRICT PASS` | `96/96 STRICT PASS` | +0.388% (slower) | speed REFUTED | +| sm_121a | `95/96 DIVERGENCE` | `95/96 DIVERGENCE` | -1.991% (faster) | speed REFUTED | + +**Each host returns the SAME token verdict on both arms**, which is what makes +this change token-neutral: the arm that predates it diverges identically to the +arm that replaces it. The sm_121a divergence is therefore NOT this row's and is +filed as #1388 -- it is the "sm_121a re-run pending" `STATUS` has carried, and +it fails on both arms. + +Both moves are under #1311's own 3% refutation bar, so the speed hypothesis is +refuted on both hosts -- and note they point in OPPOSITE directions, which is +itself evidence that neither is signal. + +This entry was first written before any GPU lease was obtained and said no +number was claimed. It is superseded by the sections below rather than deleted, +because the order in which the evidence arrived is part of the record. + +### What IS established, and on what + +The decode-vs-prefill equivalence the swap rests on, at NemotronH's own group +count. The only pre-existing case ran `H=4 G=2`, i.e. `heads_per_group = 2`; +this model runs 8. CPU, `test_ops_mamba2_state_update`, 7 cases / 2527 +assertions / `SUCCESS`, of which the new case contributes 58 assertions: + +| shape | out elements | out scale | worst \|diff\| | state elements | state scale | worst \|diff\| | +|---|---|---|---|---|---|---| +| multi-chunk `T=24 chunk=8` | 98,304 | 26.0174 | 1.90735e-05 | 524,288 | 4.58483 | 9.53674e-07 | +| production `T=1 chunk=128` | 4,096 | 32.2456 | 7.62939e-06 | 524,288 | 5.54134 | 4.76837e-07 | + +The scale is printed beside every comparison and asserted `> 0.1`, because the +inherited 5e-3 atol would accept everything if the tensors compared were ~1e-7 +([[count-based-tolerances-bound-nothing]]). + +Mutation **A2D1-M1** — clamp the state-update group index to `min(h/hpg, 1)`: + +``` +existing case (H=4 G=2 hpg=2) RC=0 1 passed | 0 failed 8 | 8 passed | 0 failed +new driver case (H=64 G=8 hpg=8) RC=1 0 passed | 1 failed 58 | 54 passed | 4 failed +``` + +Invisible to the case that existed, caught by the case this row adds, on BOTH +shapes. Tree restored byte-for-byte (`src/vt/cpu/cpu_ops.cpp` md5 +`753ba5c3d0869396c20f2205eb2617d7` before and after). + +### The launch and allocation counts — ARITHMETIC, NOT A PROFILE + +Counted from `cuda_mamba2_ssd.cuh:596-641` at the driver geometry +(`H=64 P=64 N=128 G=8 cs=128 S=1 nchunks=1 T=1`), per TOKEN over 23 mamba +layers: + +| | chunk scan (before) | state update (after) | +|---|---|---| +| SSD kernel launches | 115 | 23 | +| SSD `Alloc`/`Free` | 230 | 0 | +| SSD `cudaMemsetAsync` | 46, zeroing 57.5 MiB | 0 | +| SSD scratch | 104.9 MiB | 0 | +| gather/scatter launches | 92 | 0 | +| small metadata H2D | 138 | 0 | +| `M2ChunkScanKernel` grid | 524,288 elements for 4,096 | n/a | + +Per-call scratch is `dtv` 32 KiB + `dac` 32 KiB + `states` 2.00 MiB + `cb` +512 KiB + `passed` 2.00 MiB = **4.5625 MiB**, of which 2.50 MiB is memset. + +**Disagreement with #1311, recorded rather than reconciled.** The issue put the +gather/scatter state churn at "roughly +414 MiB/token". Counting the three SSM +movements (gather, scatter, `final_states` copy-back) at 4.00 MiB each gives +12.0 MiB per layer per token = ~276 MiB/token for the SSM plus ~7 MiB for the +conv, about **283 MiB/token**. Neither figure is measured and the direction of +the change does not depend on which is right, so the smaller one is carried and +the difference is stated. + +### The A3 e2e token gate — MEASURED, and it is the acceptance condition + +`thor:gpu0` (sm_110) inside an `rc` lease, `ARCH=110`, tree `68a0ff378`, real +`NVIDIA-Nemotron-3.5-Lightning-30B-A3B-NVFP4` at `/workspace/a3/ckpt-stage`, +device mamba arm ON, `VT_NEMOTRON_H_MAMBA_DECODE_STEP=1` (the default): + +``` +RC[a3 on]=0 +[nemotron-h] engine loaded in 654.7s +[nemotron-h] TOKEN MATCH: 96/96 over 3 prompt(s) (full rows=3, short rows=0, mode=decode) +[nemotron-h] STRICT PASS +on: tokens compared 96 ; matched 96 +reference-tier lines in on: 0 +``` + +The `reference-tier` count matters: the portable reference tier is numerically +CORRECT, so a pass obtained on it is invisible in the tokens and only that line +separates them. + +### Which kernels the decode steps LAUNCHED — the reachability evidence + +Every decode step of that run, read off the `vt::` call sites via +`VT_NEMOTRON_H_ARM_TRACE`: + +``` +[NH-DIAG] ARM step T=1 nd=1 np=0 state_update_rows=23 chunk_scan_calls=0 + conv_update_rows=23 conv_fwd_calls=0 + gathers=0 scatters=0 +ARM lines total: 96 +ARM lines with a DECODE row: 93 +ARM lines with a PREFILL row: 3 +``` + +23 is the mamba layer count and `nd=1` is one decode row, so that is ONE +state-update row per mamba layer, ZERO chunk scans, ZERO gathers and ZERO +scatters on a decode step. 96 forwards = 3 prefills + 93 decodes over 3 prompts +x 32 tokens. `vt::Mamba2StateUpdate` went from zero callers under `src/vllm/` to +23 launches per decoded token through a production entry point. + +### The decode window — SAMPLED ON THE DECODE ONLY + +``` +on: decode window 74.511 s (the engine load is OUTSIDE it) +on: engine load 654.7 s, excluded +on: GPU busy in 244 of 555 DECODE samples = 43.96% +on: per output token 0.776159 s +``` + +**NO vLLM ratio is quoted for arch 110.** The 0.014369 s reference is GB10's, +and a ratio against it would compare two pieces of silicon. The only admissible +comparison is the ON/OFF A/B of this same binary on this same box. + +### The same-binary A/B — the SPEED HYPOTHESIS IS REFUTED ON THIS BOX + +One binary, one box, device mamba arm ON in both legs, only +`VT_NEMOTRON_H_MAMBA_DECODE_STEP` differing. Both legs pass the token gate, so +this is a speed comparison between two CORRECT arms: + +| | ON (single-step, default) | OFF (chunk scan at decode) | +|---|---|---| +| A3 verdict | `96/96 mode=decode STRICT PASS` | `96/96 mode=decode STRICT PASS` | +| `state_update_rows` per decode step | 23 | 0 | +| `chunk_scan_calls` per decode step | 0 | 23 | +| `conv_update_rows` / `conv_fwd_calls` | 23 / 0 | 0 / 23 | +| `gathers` / `scatters` per decode step | 0 / 0 | 46 / 46 | +| decode window | 74.511 s | 74.223 s | +| per output token | 0.776159 s | 0.773156 s | +| GPU busy (decode samples) | 244 of 555 = 43.96% | 230 of 562 = 40.93% | +| engine load, EXCLUDED | 654.7 s | 778.2 s | +| `reference-tier` lines | 0 | 0 | + +**Per output token moved +0.388%, and in the SLOWER direction.** Issue #1311's +own stop condition is "Refuted if per-token time moves less than 3%". At 0.388% +the speed hypothesis is **REFUTED on `thor:gpu0` (sm_110) at concurrency 1**, +and this record says so rather than reporting the two counter columns as though +they were a result. + +**What is NOT refuted, and the distinction matters.** The counters are not a +prediction; they are what the run launched. The single-step arm demonstrably +removed 92 of 115 SSD kernel launches, all 230 driver alloc/frees, all 46 +memsets, 104.9 MiB of per-token scratch and all 92 gather/scatter launches per +token — and per-token time did not move. **So at c1 on this box the decode step +is not bound by any of them.** That is a finding, not a null. + +**Limits of this measurement, stated rather than discovered later:** + +- **n = 1 per leg.** No repetitions, so 0.388% is not separable from run-to-run + noise; it is reported as "did not move", not as a regression. +- **The two legs did not see the same box.** Engine load was 654.7 s and 778.2 s + — an 18.8% spread on a phase that is excluded from the window but is evidence + that the host was not in the same state for both. +- **Thor is not GB10.** The 6.31% busy-fraction and 0.014369 s/token references + are GB10's, so NO ratio against them is quoted here. The sm_121a leg is owed + and is the only thing that can answer the GB10 question + ([[negative-results-are-regime-dependent]]). +- **c1 only.** The gather/scatter tax the `qwen3_5.cpp:4730-4746` comment + describes is stated to scale with CONCURRENCY, and G-SAFE pins `num_reqs <= 1` + here, so the regime where it would show has not been measured at all. + +**The next traceable hypothesis, because no ceiling is declared:** an `nsys` +trace of the DECODE WINDOW ONLY on both legs, attributing the 0.776 s/token. +The counters say what the step stopped launching; the trace would say what the +0.776 s is actually spent on. Peak host during the run was 44402 MiB. + +### sm_121a (GB10) — RESOLVED by a three-leg run: the cause is A2-Q1's arm + +`dgx:gpu0`, `rc` lease, `ARCH=121a`, logs `/workspace/a2d1-discriminate/20260819T200231Z`. +THREE legs of ONE binary. Legs 1 and 2 vary the mamba KERNEL; leg 3 removes +A2-Q1's device mamba arm entirely (`VT_NEMOTRON_H_DEVICE_MAMBA=0`): + +| leg | device mamba | tokens | s/token | vs vLLM | decode busy | +|---|---|---|---|---|---| +| `on` (single-step) | 1 | `95/96 DIVERGENCE` | 1.584694 | 110.3x | 108/1052 = 10.27% | +| `off` (chunk scan) | 1 | `95/96 DIVERGENCE` | 1.554233 | 108.2x | 108/1061 = 10.18% | +| **`hostmamba`** | **0** | **`96/96 STRICT PASS`** | 10.318897 | 718.1x | 559/7115 = 7.86% | + +**★ THE CAUSE IS A2-Q1's DEVICE MAMBA ARM, NOT THE ARCHITECTURE.** Removing it +on the same binary and box restores `96/96`. The earlier "arch-specific" reading +was retracted before this ran, and this is the experiment that actually tested +it. + +**★ WHY THIS IS A DISCRIMINATION AND NOT A CORRELATION: leg 3's own counters.** +It reports `state_update_rows=0 chunk_scan_calls=0 conv_update_rows=0 +conv_fwd_calls=0` with `gathers=46 scatters=46` — all four kernel counters at +zero while the host branch's gather/scatter stay non-zero, which is the +signature of the host path and of nothing else. 96 ARM lines, 93 decode rows, 3 +prefill rows, `reference-tier lines: 0`. A leg that could not prove which path +it took would have made this a correlation between a flag and an outcome; the +counters are what make it an attribution. + +**The divergence is ONE token, and both device legs lose the SAME one.** The +repaired verdict grep — added after the first GB10 run discarded it — captured +it: + +``` +got: ...,1044,12837,2505,1261,9943,1307,11286 +exp: ...,1044,12837,2505,1261,9943,1307,3468 +``` + +Prompt 2 (13 prompt tokens, the longest), position 32 of 32, the LAST generated +token; positions 1-31 byte-identical, and identical between the two device legs. +Prompts 0 and 1 are 32/32. **Whether this is a defect or a bf16 near-tie is NOT +settled here** — a fresh implementer is root-causing it from the oracle's top-2 +margin at that position, and a near-tie would mean there is no defect at all. + +### The second result, and it is the larger one + +**A2-Q1's device mamba arm is worth 6.64x per output token on GB10**: 10.318897 +s with the arm off against 1.554233 s with it on, closing the vLLM ratio +**718.1x -> 108.2x**, with decode busy rising 7.86% -> 10.18%. It agrees in +direction and rough magnitude with #1289's independent Thor decode-window A/B +(7.17x less time per output token), which is two boxes and two harnesses +agreeing. + +**Bound honestly. This is ONE run per leg on a contended box**, engine load +excluded from every window (376.3 s on leg 3). It is not a mean, it has no +spread, and the box was carrying foreign jobs all day. The 6.64x is a large +enough effect to survive that; the ~2% on/off deltas are not, and §A2-D1's own +data now proves it — see below. + +**108.2x is an OPEN GAP, not a win.** It is the distance still to vLLM after the +single biggest movement measured on this model. Nothing here claims parity, and +no ceiling is declared. + +### The +-2% on/off delta is NOISE, and this run proves it + +Two runs of the SAME on/off comparison on the SAME box: + +| run | on | off | on-vs-off | +|---|---|---|---| +| first GB10 | 1.513958 | 1.544706 | **-1.991%** (on faster) | +| this GB10 | 1.584694 | 1.554233 | **+1.960%** (on slower) | + +**The sign flipped.** A quantity that reverses direction between two runs of one +comparison on one box is not a measurement of that comparison. This retires any +reading of the ~2% figures as signal and strengthens #1311's refutation, which +was already argued from the 3% bar: the single-step-vs-chunk-scan speed question +is settled negative on both hosts, and the token question is settled neutral +(95/96 vs 95/96 here, 96/96 vs 96/96 on Thor). + +### Evidence + +`tests/vt/test_ops_mamba2_state_update.cpp` (the two driver-group cases), +`tests/vllm/models/test_nemotron_h_paged_forward.cpp` (the arm recorder, driven +through `ModelRegistry::Forward`), `scripts/nemotron-h-a2d1-gpu-gate.sh` (the +recipe), and the Thor run logs under `/workspace/a2d1-thor/20260819T125936Z`. diff --git a/.agents/issue-index.md b/.agents/issue-index.md index fe4b5e576..456f47406 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -457,3 +457,7 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1376](https://github.com/mudler/vllm.cpp/issues/1376) | `ENG-CUDAGRAPH-BREAK` | `main` was red on `tests/scripts/test_check_gate_commands.py`, measured at `601b576c6` in a detached worktree of `origin/main`: 8 failures of 44 tests, every one a comparison between the computed runnable population and `RUNNABLE_BASELINE`. `ENG-CUDAGRAPH-BREAK` was in the first and absent from the second. Cause: W5 of that row ([#1361](https://github.com/mudler/vllm.cpp/issues/1361)) filled its spec's `## Gates` section with runnable evidence, including a named test binary with its case and assertion counts and an exit status, which is exactly what moves a row into the runnable population. The ratchet's own error text instructs a re-pin in the SAME change, and the re-pin was not made. This is the growth case the ratchet exists to force a decision about, not a defect in that row's work. **It landed with no remote verdict**: the continuous integration lane that would have caught it independently has not executed for this repository since roughly 07:43Z on 19 August 2026, with runs queueing and none starting while GitHub reports Actions operational. FIXED IN FLOW while merging `origin/main` into `row/ENG-HF-MODEL-DOWNLOAD` for [#1280](https://github.com/mudler/vllm.cpp/issues/1280), because the fix is small and clear and a red `main` blocks every other row's gate. The entry is added with a justifying comment in the form the neighbouring entries use, no checker semantics change, and no test is weakened. After the re-pin the suite reports 45 tests OK and the audit reads 39 runnable of 119 gated rows | bug | | [#1375](https://github.com/mudler/vllm.cpp/issues/1375) | `MODEL-DIFFUSION-LTX25` | First end-to-end per-forward cost for the FULL 21.004 B LTX-2.5 DiT on GB10, measured on run `20260819T150230Z` with binary `0a43a750` built from [`7b9e207b1`](https://github.com/mudler/vllm.cpp/commit/7b9e207b1) (#1252). At 1024x576/25f (2304 latent tokens) the governor resolved **7 forward starts from the GPU busy/idle edge counter** and measured `per_forward ~162.0 s` with `first_dit = 481.5 s`, so the recipe's fixed 60 forwards (30 steps x 2 CFG legs, `ltx2_pipeline.cpp:521-529`) project **10 803 s against the rung's 7 153 s budget** and the rung was refused rather than run to the wall. The same lease then COMPLETED 768x448/25f (1344 tokens) in 2990 s, so the ceiling is geometry against lease length, not a defect. TWO instrument facts belong with the number, because both have already caused a wrong reading: `gpu_edges=0` means the GPU never went idle long enough to sample an edge (SATURATED), not that no work ran — this rung sampled 85% of 3191 samples above 50% utilisation; and `eu-stack` resolves no frames in the rc worker container, so phase attribution came from the cpu%/rss signature rather than from symbols. Owned by the LTX-2.5 row; spec [`ltx-2-5.md`](specs/ltx-2-5.md) | measurement | | [#1386](https://github.com/mudler/vllm.cpp/issues/1386) | — | `tools/bench/gpu_clock_state.py`'s `QUERY_FIELDS` collects nine fields and **none of them is thermal and none is electrical**, so the driver's own `SwThermalSlowdown` label can never be checked against a die reading on any window this helper has ever recorded. The measured consequence is that the nine windows of 2026-08-19 cannot distinguish a load transition from a thermal excursion. The concrete evidence is `clock-c1-r1.samples.json` in `/mnt/nas_share/rc/q38bf16/out/bench-20260819T035148Z/`: ours c1 r1 dips five times on the same period at the same `utilization.gpu = 96` — 48.83 s / 2177 MHz, 80.60 s / 2320 MHz, 109.28 s / 2210 MHz, 137.98 s / 2359 MHz, 166.07 s / 2268 MHz — and **two of those five carry `0x0000000000000000`**, no throttle bit at all (2210 and 2359), while three carry `0x20`. The 2210 MHz unlabelled dip is deeper than two of the three labelled ones, so the driver labels comparable excursions inconsistently and the bit alone cannot decide it. What would settle it: add `temperature.gpu` and `power.draw` to `QUERY_FIELDS`. That changes the clock-record schema, so it owes its own row and spec. Split out of [#1354](https://github.com/mudler/vllm.cpp/issues/1354) and owed under `## Owed` in [lease-clock-pinning.md](specs/lease-clock-pinning.md) | gap | +| [#1290](https://github.com/mudler/vllm.cpp/issues/1290) | `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | The A3 gate reads `93/96 mode=decode DIVERGENCE` on `thor:gpu0` (sm_110) with the HOST mamba arm and `96/96 STRICT PASS` with the device FP8 W8A8 arm \-- same binary, box, checkpoint, golden and session, `reference-tier lines: 0` on both. On `dgx:gpu0` (sm_121a) the SAME host arm reads `96/96 STRICT PASS` (the #1157 run, `gate_fixed.out`), so the divergence is established on sm_110 ONLY and is NOT a general statement about `main` on a GPU. n=1 per arm. Leading mechanism: the golden's oracle computes these projections W8A8 while the host arm is W8A16 by construction (`DenseBf16` carries and does not apply `input_scale`, `nemotron_h.cpp:419-422`), so the perturbation is marginal and whether it flips a token depends on the rest of the tower \-- `fa2` provides for `12.1a` and not `11.0` (`CudaArchFeatures.cmake:349`), so GB10 runs FlashAttention-2 where Thor runs the fallback. NOT the fp8 GEMM: the host arm dequantizes and runs a CPU bf16 `MatmulBT`, so no fp8 GEMM executes on either box in that configuration. Next: the oracle's top-2 margin at the three moved tokens, then a `NemotronHTrace` layer bisect. A2-Q1 ([#1289](https://github.com/mudler/vllm.cpp/pull/1289)) makes the Thor gate read 96/96; whether anything was wrong on GB10 is open and this issue owns it | bug | +| [#1311](https://github.com/mudler/vllm.cpp/issues/1311) | `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | **NemotronH ran the PREFILL mamba kernels at decode**, and `vt::Mamba2StateUpdate` — the port of `selective_state_update`, registered on CUDA (`cuda_gdn.cu:6672`) and CPU (`cpu_ops.cpp:3408`) and gated since [`mamba2-ssd.md`](specs/mamba2-ssd.md) §W2 landed 2026-08-13 — had ZERO callers under `src/vllm/`, which is a live `AGENTS.md` §"Nothing lands dead" case. vLLM branches its Mamba2 mixer on `has_decode` (`mamba_mixer2.py:981`) and runs `causal_conv1d_update` `:1012` + `selective_state_update` `:1087` on the decode rows, both taking `conv_state_indices`/`state_batch_indices` and updating the cache IN PLACE at the slot, so upstream's decode half performs no gather and no scatter; A2-Q1 ran `vt::CausalConv1dFwd` + `vt::Mamba2ChunkScan` wrapped in `GatherNemotronHState`/`ScatterNemotronHState` over BOTH halves. Arithmetic from `cuda_mamba2_ssd.cuh:577-645`, NOT profiled: at `T=1 chunk_size=128` the chunk scan launches a grid of 524,288 elements to compute 4,096, in 5 kernels with 10 `cudaMallocAsync`/`FreeAsync`, 2 memsets of 2.50 MiB and 4.56 MiB of scratch, per mamba layer per token. A SECOND defect was found while fixing it and is fixed in the same flow: the decode-vs-prefill equivalence the swap rests on was gated ONLY at `H=4 G=2` (`heads_per_group = 2`), while NemotronH runs `heads_per_group = 8` — mutation A2D1-M1 (clamp the group index to `min(h/hpg, 1)`) passes the pre-existing case 8/8 assertions RC=0 and fails the new driver-group case 4 of 58 assertions RC=1, on both the multi-chunk and the production `T=1 chunk=128` shape. FIXED IN FLOW in [`nemotron-h-a2d1-mamba-decode-step.md`](specs/nemotron-h-a2d1-mamba-decode-step.md) | bug | +| [#1371](https://github.com/mudler/vllm.cpp/issues/1371) | `BACKEND-ATTN-REGISTRY` | **CPU has NO attention backend when FLASH_ATTN declines a head_size.** `369ea7fd4` ([#1332](https://github.com/mudler/vllm.cpp/issues/1332), [#1344](https://github.com/mudler/vllm.cpp/pull/1344)) landed `supports_head_size` on the FA backend (`include/vllm/v1/attention/backend.h:394-397`, `head_size % 8 != 0 => false`), a faithful mirror of `flash_attn.py:170-178`. The predicate is not the defect; what it exposed is that the CPU selector no longer FALLS BACK when FA declines. `tests/vllm/models/test_nemotron_h_paged_forward.cpp` uses `kHeadDim = 6`, and against `origin/main` `18f99481e` the suite reads `RC=1`, `13 cases, 2 passed, 11 failed`, every one throwing `No valid attention backend for device type 0 from {FLASH_ATTN: [head_size not supported]}` at fixture construction. Device type 0 is CPU and `src/vllm/platforms/cpu.cpp:45` returns TWO candidates `{CPU_ATTN, FLASH_ATTN}`, so `CPU_ATTN` either never entered the candidate set or was filtered before `validate_configuration`. ★ The run also prints `assertions: 18, 18 passed, 0 failed` beside `11 failed` cases, because a THROWN case runs no assertions -- a gate grepping only `assertions:` reads this as clean ([[doctest-assertions-line-hides-thrown-cases]]). NOT the head_dim: raising the fixture to a multiple of 8 greens one suite and leaves the selector defect for every other CPU consumer. Attributed by counter-measurement, not by inspection: the SAME suite on the SAME row at `68a0ff378`, which lacks `369ea7fd4`, built for CUDA on `thor:gpu0` sm_110 in an `rc` lease, reads `RC=0`, `13 of 13 passed`, `3272 assertions`; and that row `git diff`s EMPTY over `src/vllm/v1/attention/`, `include/vllm/v1/attention/` and `src/vllm/platforms/`. Found while landing [#1311](https://github.com/mudler/vllm.cpp/issues/1311); FILED not fixed in flow because the repair is in the attention-backend selector, which is this row's surface and affects every CPU consumer rather than one model's suite | bug | +| [#1388](https://github.com/mudler/vllm.cpp/issues/1388) | `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | **NemotronH A3 on GB10 sm_121a reads `95/96 DIVERGENCE` on BOTH recurrent decode arms**, which is the "sm_121a re-run pending" `docs/STATUS.md` has carried since A2-P. Measured on `dgx:gpu0` in an `rc` lease, tree `e35c14d52`, real 20.1 GiB NVFP4 checkpoint, device mamba arm ON in both legs, `cutlass-fp8 ENABLED for [121a]` so the build is not VOID, `reference-tier` lines 0 on both. ON (`vt::Mamba2StateUpdate`, counters `state_update_rows=23 chunk_scan_calls=0 gathers=0 scatters=0`) reads `RC=1`, `95/96`, `full rows=3 short rows=0 mode=decode`; OFF (`vt::Mamba2ChunkScan`, counters `0 / 23 / 46 / 46`) reads `RC=1`, `95/96`, the same shape. The counters prove the legs ran DIFFERENT kernels and both lose exactly one token in 96, so it is arch- or host-specific and NOT arm-specific -- `thor:gpu0` sm_110 reads `96/96 STRICT PASS` on both arms of the same code and checkpoint. ★ NOT ESTABLISHED: whether both legs lose the SAME token. The driver emits `got:`/`exp:` ids for a mismatching row (`nemotron_h_gen/main.cpp:385`) and the gate script's verdict grep DISCARDED them, so a wrong recurrent carry and a benign bf16 near-tie are not yet separated; the script is repaired on `row/A2-D1-mamba-decode-step` so the next run captures them, and that is the next traceable step. Also unexplained and recorded so it is not read as an [#1311](https://github.com/mudler/vllm.cpp/issues/1311) result: GB10 runs this arm ~2x SLOWER than Thor, 1.513958 vs 0.776159 s per output token at 11.06% vs 43.96% decode busy. Probably not independent of [#1157](https://github.com/mudler/vllm.cpp/issues/1157) ("GB10 read 4/24") or [#1290](https://github.com/mudler/vllm.cpp/issues/1290). Found while landing #1311; FILED not fixed in flow because it predates that change, needs its own spec and a scarce `dgx` lease | bug | diff --git a/.agents/model-matrix.md b/.agents/model-matrix.md index 3c8bd3ae4..45df9fda9 100644 --- a/.agents/model-matrix.md +++ b/.agents/model-matrix.md @@ -282,7 +282,7 @@ Transformers compatibility is capability-driven and excluded from finite counts. | `MODEL-TEXT-mimo-v2-mi-mo-v2-flash-for-causal-lm` | `MiMoV2FlashForCausalLM` | `registry.py:166`; `vllm/model_executor/models/mimo_v2.py::MiMoV2FlashForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; FusedMoE/grouped GEMM; sliding-window attention | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-mimo-v2-mi-mo-v2-for-causal-lm` | `MiMoV2ForCausalLM` | `registry.py:167`; `vllm/model_executor/models/mimo_v2.py::MiMoV2ForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; FusedMoE/grouped GEMM; sliding-window attention | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-nemotron-nemotron-for-causal-lm` | `NemotronForCausalLM` | `registry.py:168`; `vllm/model_executor/models/nemotron.py::NemotronForCausalLM` | causal generation / text | model loader/forward; paged attention/KV | ☐ required | `INVENTORIED` | none | unassigned | -| `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | `NemotronHForCausalLM`, `NemotronHPuzzleForCausalLM` | `registry.py:169-170`; `vllm/model_executor/models/nemotron_h.py::NemotronHForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; FusedMoE/grouped GEMM; Mamba/SSM state; sliding-window attention. **2026-08-12: row spec committed ([#517](https://github.com/mudler/vllm.cpp/issues/517), [spec](specs/nemotron-h-model.md)); BLOCKED on `KERNEL-SSM-MAMBA` [#496](https://github.com/mudler/vllm.cpp/issues/496)** ([spec](specs/mamba2-ssd.md)) — the Mamba2 SSD core is unported. Beyond that kernel this arch owes three further things, none of which exist locally: non-gated `relu²` MoE (every grouped-MoE op we have is SwiGLU-shaped), ModelOpt `MIXED_PRECISION` per-module loading (`quantization/modelopt.py:2280` — NVFP4 W4A16 g16 experts and FP8 W8A8 mamba projections in one checkpoint), and the DeepSeek-style MTP head. Driver checkpoint `nvidia/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-NVFP4` (20.1 GiB, fits one GB10). **2026-08-17: A2-P lands the PAGED forward** ([#810](https://github.com/mudler/vllm.cpp/issues/810), [spec](specs/nemotron-h-a2p-paged-forward.md)): `ForwardNemotronHForCausalLM` selects `NemotronHPagedForward` whenever the runner supplies paged KV and recurrent state, so K/V go into the runner's pages at `attn_meta.slot_mapping` and the conv + SSM rows are gathered from and scattered to `GdnStateCache` at the metadata's state indices — the first time this architecture carries state between decode steps. G-SAFE narrows from three clauses to `num_reqs <= 1` (batching is A2-B). The row stays `INVENTORIED`: the A3 end-to-end token gate against the released checkpoint is PENDING on `dgx.casa`, and `lm_head` (A2-Q2b), the FP8 mamba projections (A2-Q1, [#940](https://github.com/mudler/vllm.cpp/issues/940)), MTP (W5) and GGUF (W7) are all still owed | ☐ required | `INVENTORIED` | none | unassigned | +| `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | `NemotronHForCausalLM`, `NemotronHPuzzleForCausalLM` | `registry.py:169-170`; `vllm/model_executor/models/nemotron_h.py::NemotronHForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; FusedMoE/grouped GEMM; Mamba/SSM state; sliding-window attention. **2026-08-12: row spec committed ([#517](https://github.com/mudler/vllm.cpp/issues/517), [spec](specs/nemotron-h-model.md)); BLOCKED on `KERNEL-SSM-MAMBA` [#496](https://github.com/mudler/vllm.cpp/issues/496)** ([spec](specs/mamba2-ssd.md)) — the Mamba2 SSD core is unported. Beyond that kernel this arch owes three further things, none of which exist locally: non-gated `relu²` MoE (every grouped-MoE op we have is SwiGLU-shaped), ModelOpt `MIXED_PRECISION` per-module loading (`quantization/modelopt.py:2280` — NVFP4 W4A16 g16 experts and FP8 W8A8 mamba projections in one checkpoint), and the DeepSeek-style MTP head. Driver checkpoint `nvidia/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-NVFP4` (20.1 GiB, fits one GB10). **2026-08-17: A2-P lands the PAGED forward** ([#810](https://github.com/mudler/vllm.cpp/issues/810), [spec](specs/nemotron-h-a2p-paged-forward.md)): `ForwardNemotronHForCausalLM` selects `NemotronHPagedForward` whenever the runner supplies paged KV and recurrent state, so K/V go into the runner's pages at `attn_meta.slot_mapping` and the conv + SSM rows are gathered from and scattered to `GdnStateCache` at the metadata's state indices — the first time this architecture carries state between decode steps. G-SAFE narrows from three clauses to `num_reqs <= 1` (batching is A2-B). The row stays `INVENTORIED`: the A3 end-to-end token gate against the released checkpoint is PENDING on `dgx.casa`, and `lm_head` (A2-Q2b), MTP (W5) and GGUF (W7) are all still owed. **2026-08-18: A2-Q1 lands the FP8 W8A8 MAMBA DEVICE ARM** ([#810](https://github.com/mudler/vllm.cpp/issues/810), [spec](specs/nemotron-h-a2q1-fp8-mamba.md)), on the shared seam [#940](https://github.com/mudler/vllm.cpp/issues/940) extracted for it: `NemotronHMamba2MixerDevice` mirrors the host arm op for op and both forwards select it at runtime, so the 23 Mamba2 layers no longer download the normed hidden, dequantize 890e6 fp8 elements to bf16 per token and upload the result. The recurrent rows the gather produced are advanced IN PLACE, and G-SAFE is untouched. The GB10 per-block numeric gate and the GPU-occupancy measurement the unit is accepted on are PENDING a lease | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-olmo-olmo-for-causal-lm` | `OlmoForCausalLM` | `registry.py:171`; `vllm/model_executor/models/olmo.py::OlmoForCausalLM` | causal generation / text | model loader/forward; paged attention/KV | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-olmo2-olmo2-for-causal-lm` | `Olmo2ForCausalLM`, `Olmo3ForCausalLM` | `registry.py:172-173`; `vllm/model_executor/models/olmo2.py::Olmo2ForCausalLM` | causal generation / text | OLMo-2-0425-1B dense LANDED (W0-W4). Impl: `include/vllm/model_executor/models/olmo2.h` + `src/vllm/model_executor/models/{olmo2,olmo2_weights,olmo2_registry}.cpp` (TWO `REGISTER_VLLM_MODEL` for `Olmo2ForCausalLM`+`Olmo3ForCausalLM`, one shared factory; reuses the shared dense glue). ZERO new compute kernel — the two "distinctives" both reduce to WIRING over landed ops: (1) PURE POST-NORM (`norm_after`, `olmo2.py:261-277`) = standalone `vt::RmsNorm` on each sublayer OUTPUT + plain `vt::Add` residual re-join, NO pre-norm (the GLM-4 standalone-output-norm op, pre-norms dropped); (2) FULL-WIDTH QK-norm (`olmo2.py:113-117,160-172`) = two standalone `vt::RmsNorm` over `[T,q_size]`/`[T,kv_size]` before NeoX rope (NOT the per-head `kAttnQkNormRope`). Reuse: SiLU SwiGLU (`kSiluAndMul`), NeoX RoPE (`RopeFromCache`/`RopeNeox`), GQA/MHA paged glue (`dense_attn_block.h`), merged qkv/gate_up loader. Checkpoint is F32 on-disk → loader downcasts f32→bf16 (round-to-nearest-even, `vt::F32ToBF16`) to match vLLM-bf16. UNTIED lm_head (`tie_word_embeddings:false`). Loader 179 tensors, zero missing/unmapped. Tokenizer: ONE shared-TU touch — `tokenizer.cpp` accepts the OLMo-2 Split `behavior=Removed,invert=true` encoding (equiv. to `Isolated` for the full-cover cl100k regex → `kLlama3`); guarded additive branch, no other checkpoint uses it. NO-BOS-verified real ByteLevel gate. Runs EAGER (bf16, no decode graph). Gate: 16/16 vs vLLM 0.25.0 (STRICT 13/16 + near-tie 3/16, max gap 0.094 nats, 0 forward-divergent) — speed pending. **OLMo-3 W5 (batch3) IMPLEMENTED** (guarded additive edits to `olmo2.{h,cpp,weights}`, diff-inert for OLMo-2 — re-run OLMo-2 gate 16/16 UNCHANGED): per-layer interleaved routing off `config.layer_types` — sliding_attention layers use plain NeoX rope (theta 500000) + finite window (masked at the FA kernel, inert for short contexts), full_attention layers use a precomputed YaRN cos/sin cache (get_rope yarn: factor 8, original 8192, mscale=yarn_get_mscale(8)=1.2079=config attention_factor, indexed by real positions); dtype-aware loader (OLMo-3 is BF16 on-disk vs OLMo-2 F32). **BUT the pinned vLLM 0.25.0 oracle CANNOT run `allenai/OLMo-3-1025-7B`** — its transformers version predates OLMo-3's nested per-layer-type rope schema: `olmo2.py:143` does `rope_parameters["rope_theta"]` → `KeyError: 'rope_theta'` (top-level rope_theta not folded into standardized rope_parameters), and forcing it surfaces `Unrecognized keys {'sliding_attention','full_attention'}` → `TypeError: unhashable type: 'dict'`. So there is NO pinned-oracle SACRED bar for OLMo-3 (DEP-blocked, spec D5); our engine LOADS + RUNS it (bootstrap loaded clean). W5 SACRED gate pending an oracle that can construct OLMo-3's rope config. | ✅ [sweep-olmo2](specs/sweep-olmo2.md) | `PARTIAL` | `test_olmo2_paged_engine` 16/16 (dgx, 92 assertions, batch3 re-run UNCHANGED); `test_olmo3_paged_engine` present (oracle-blocked, skips); loader 179 tensors zero-unmapped; registry resolves both arch strings; registration `src/vllm/model_executor/models/olmo2_registry.cpp:126`; test `tests/vllm/models/test_model_registry.cpp:118` | `CLAIM-SWEEP-OLMO2` (Claude Code opus-4-8) | | `MODEL-TEXT-olmo-hybrid-olmo-hybrid-for-causal-lm` | `OlmoHybridForCausalLM` | `registry.py:174`; `vllm/model_executor/models/olmo_hybrid.py::OlmoHybridForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; Mamba/SSM state; GDN/linear-attention state | ☐ required | `INVENTORIED` | none | unassigned | diff --git a/.agents/specs/nemotron-h-a2d1-mamba-decode-step.md b/.agents/specs/nemotron-h-a2d1-mamba-decode-step.md new file mode 100644 index 000000000..2124ce364 --- /dev/null +++ b/.agents/specs/nemotron-h-a2d1-mamba-decode-step.md @@ -0,0 +1,358 @@ +# A2-D1 — NemotronH decodes on the single-step recurrent kernels + +**Issue:** [#1311](https://github.com/mudler/vllm.cpp/issues/1311). +**Parent row:** `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` +([#517](https://github.com/mudler/vllm.cpp/issues/517)). +**Governing spec:** [`nemotron-h-abi-e2e.md`](nemotron-h-abi-e2e.md); this unit +is residual work on the arm [`nemotron-h-a2q1-fp8-mamba.md`](nemotron-h-a2q1-fp8-mamba.md) +built. +**Base:** `row/A2-Q1-fp8-mamba` @ `56e2dc36182fa3c5aa4adaa4abe45dc453335e54` +(PR [#1289](https://github.com/mudler/vllm.cpp/pull/1289)), merged forward onto +`origin/main` `edbc47ce0`. `NemotronHMamba2MixerDevice`, the function this unit +changes, does not exist on `main`. +**Pinned oracle:** `${VLLM_SOURCE}` @ `5559679229bc961848b121ccdeaa8fa5d79bec98` +(vLLM 0.26.0.dev0), per [`upstream-sync.md`](../upstream-sync.md). +**Lifecycle at this commit:** unchanged. The row stays `INVENTORIED`. + +--- + +## 0. The divergence + +vLLM's Mamba2 mixer branches on `has_decode` (`mamba_mixer2.py:981`) and runs a +different pair of kernels on the decode rows: + +| | vLLM decode | vLLM prefill | +|---|---|---| +| conv | `causal_conv1d_update` `:1012` | `causal_conv1d_fn` `:869` | +| SSD | `selective_state_update` `:1087` | `mamba_chunk_scan_combined` `:890` | +| state I/O | `conv_state_indices` / `state_batch_indices`, IN PLACE at the slot | gather, scan, scatter | + +A2-Q1 ran the **prefill pair over both halves**: `vt::CausalConv1dFwd` + +`vt::Mamba2ChunkScan` wrapped in `GatherNemotronHState` / +`ScatterNemotronHState`, for every row of every step. + +`vt::Mamba2StateUpdate` — the port of `selective_state_update` — has been +registered on CUDA (`cuda_gdn.cu:6672`) and CPU (`cpu_ops.cpp:3408`) and gated +since [`mamba2-ssd.md`](mamba2-ssd.md) §W2 landed on 2026-08-13, with **zero +callers under `src/vllm/`**. That is a live `AGENTS.md` §"Nothing lands dead" +case, and closing it is half of why this unit exists. + +The tree already carried the warning. `qwen3_5.cpp:4730-4746`, on the GDN decode +path: *"Passing the state indices to the op eliminates the per-request +gather+scatter — the two host<->device copies per sequence per layer that +dominate the decode memcpy tax."* NemotronH did the thing that comment warns +against. + +### 0.1 What it costs, as arithmetic + +Counted from `cuda_mamba2_ssd.cuh:577-645` at `T=1`, `chunk_size=128`, +`nchunks=1`. **Not profiled** — #1311 says so and this spec repeats it: + +| per mamba layer per token | chunk scan | state update | +|---|---|---| +| kernel launches | 5 | 1 | +| `cudaMallocAsync`/`FreeAsync` | 5 + 5 | 0 | +| `cudaMemsetAsync` | 2, of 2.50 MiB | 0 | +| scratch | 4.56 MiB | 0 | +| `M2ChunkScanKernel` grid | 524,288 elements to compute 4,096 | 4,096 | + +x23 mamba layers = ~105 MiB scratch, 230 driver alloc/frees and 115 launches per +token in the SSD alone, plus the gather/scatter state churn upstream's decode +half never performs. + +### 0.2 The same arithmetic, re-derived here, and where it disagrees with #1311 + +Re-counted from `cuda_mamba2_ssd.cuh:596-641` at the DRIVER geometry +(`H=64 P=64 N=128 G=8 cs=128 S=1 nchunks=1 T=1`), because a number that is +quoted often enough starts being treated as measured +([[a-number-quoted-often-becomes-treated-as-measured]]). **This is arithmetic, +not a profile.** + +| scratch term | elements | bytes | +|---|---|---| +| `dtv` = `H*nchunks*cs` | 8,192 f32 | 32 KiB | +| `dac` = same | 8,192 f32 | 32 KiB | +| `states` = `nchunks*H*P*N` | 524,288 f32 | 2.00 MiB | +| `cb` = `nchunks*G*cs*cs` | 131,072 f32 | 512 KiB | +| `passed` = `n_states * state_elem` | 524,288 f32 | 2.00 MiB | +| **total per call** | | **4.5625 MiB** | + +Zeroed per call: `cb` + `passed` = 2.50 MiB. `M2ChunkScanKernel`'s grid is +`nchunks*H*cs*P` = 524,288 elements to produce `T*H*P` = 4,096 — a factor of +`cs` = 128. + +Per TOKEN over the 23 mamba layers, decode: + +| | chunk scan (before) | state update (after) | +|---|---|---| +| SSD kernel launches | 115 | 23 | +| SSD `Alloc`/`Free` | 230 | 0 | +| SSD `cudaMemsetAsync` | 46, zeroing 57.5 MiB | 0 | +| SSD scratch | 104.9 MiB | 0 | +| gather/scatter launches | 92 | 0 | +| small metadata H2D | 138 | 0 | + +**Where this disagrees with #1311.** The issue estimated "roughly +414 MiB/token" +of gather/scatter state churn. Counting the three SSM movements — gather +(read page + write working), scatter (read working + write page) and the +`final_states` copy-back — at 4.00 MiB each gives 12.0 MiB per layer per token, +so **~276 MiB/token for the SSM plus ~7 MiB for the conv, about 283 MiB/token**, +not 414. The disagreement is recorded rather than reconciled to the larger +figure: neither number is measured, and the direction of the change does not +depend on which is right. + +--- + +## 1. Scope + +| In A2-D1 | Out of A2-D1 | +|---|---| +| the decode/prefill split of the device mamba arm in `NemotronHPagedForward` | the MoE block (`NemotronHMoeBlockDevice`) — [#1312](https://github.com/mudler/vllm.cpp/issues/1312) | +| routing decode rows to `vt::CausalConv1dUpdate` + `vt::Mamba2StateUpdate` at their cache slots | `runner.cpp` / `sampler.cpp` — [#1313](https://github.com/mudler/vllm.cpp/issues/1313) | +| narrowing the gather/scatter to the prefill rows | the HOST mamba fallback branch, which stays exactly as it was | +| gating the chunk-scan/state-update equivalence at `n_groups=8` | lifting G-SAFE's `num_reqs <= 1` — A2-B | +| a recording seam for which arm ran | speculative decode rows — #810 W5 | + +**Prefill keeps the chunk scan.** The selection is the metadata's own +`num_decodes` / `num_prefills` split, never the token count: a one-token PREFILL +of a fresh request is also `T == 1`, carries no state in, and can have +`prefill_has_initial_state == 0`, which neither single-step kernel can express. + +--- + +## 2. Design + +`NemotronHMamba2MixerDevice` gains one optional parameter, a +`NemotronHMambaDecodeSlots` descriptor carrying the two FULL cache pages and the +per-row slot indices. Non-null replaces steps 3 and 5 of the block with the +single-step kernels and forbids the gathered carry; null leaves every existing +caller byte-identical. + +The paged forward's device-mamba branch becomes: + +1. `nd = gdn_meta.num_decodes`, `np = gdn_meta.num_prefills` (already computed + by A2-P, mirroring `mamba_attn.py:523-532`). +2. if `nd > 0`: ONE batched mixer call over the leading `nd` tokens with the + decode slots. No gather. No scatter. No `final_states` copy-back. +3. if any rows remain: gather **those rows only**, loop as A2-Q1 did, scatter + **those rows only**. + +`MetaSubView` narrows the `[R]` index vectors to a contiguous sub-range, the +same shape `SubView` already has in `qwen3_5.cpp` for the GDN decode arm. + +### 2.1 Why the decode arm needs no `has_initial` mask + +A decode continues a sequence by definition, which is why upstream leaves +`has_initial_state` `None` on a decode-only step (`gdn_attn.py:405`) and why +A2-P already sets the mask to 1 for every row below `num_decodes`. The gather's +zeroing obligation (A2-P's `NemotronHRecurrentIo` note) applies to rows whose +mask is 0, and a decode row never has one. + +### 2.2 The recording seam + +The two arms compute the same recurrence, so they produce the same tokens and a +token gate cannot tell them apart. `NemotronHMambaArmCounts` counts launches at +the `vt::` call sites — not at the branch condition — so a gate entering through +`ModelRegistry::Forward` observes what the step RAN. It mirrors +`RecordGdnOutActivationDTypes` (`qwen3_5.cpp`). + +--- + +## 3. Risks + +**R1 — a wrong carry is silently wrong tokens, not a crash.** The recurrent +state is exactly what a mis-slotted or mis-grouped update corrupts, and the +result is finite, plausible and wrong. Mitigated by the A3 token gate on the +real checkpoint, run on BOTH arms of the same binary. + +**R2 — the equivalence at `n_groups=8` was ASSUMED.** `mamba2-ssd.md` §8.3 +declares an equivalence contract, but §8.3 is device-vs-host; the only +decode-vs-prefill case ran `H=4 G=2`, i.e. `heads_per_group = 2`. NemotronH runs +`heads_per_group = 8`. A head-to-group map correct at 2 and wrong at 8 passed +every existing case. **Gated by this unit before the swap is relied on** — see +§4. + +**R3 — the instrument could report zero for the wrong reason.** A counter that +is never incremented reads exactly like a kernel that is never launched. Gated +on CPU (§4) so a broken recorder reds in seconds rather than on a GPU window. + +**R4 — the decode rows must lead the batch and be one token each.** Asserted at +the call site against `gdn_meta`, refusing rather than mis-slicing. + +--- + +## 4. Tests + +| # | Case | Where | Runs on | +|---|---|---|---| +| T1 | decode == chunked prefill at the driver group count, two shapes incl. `T=1 chunk=128` | `tests/vt/test_ops_mamba2_state_update.cpp` | CPU | +| T2 | the same on device, both arms CUDA, native provider asserted | same file | CUDA | +| T3 | the arm recorder counts what the step launched, through `ModelRegistry::Forward` | `tests/vllm/models/test_nemotron_h_paged_forward.cpp` | CPU | +| T4 | A3 e2e token gate + same-binary `VT_NEMOTRON_H_MAMBA_DECODE_STEP` A/B | `scripts/nemotron-h-a2q1-dgx-gate.sh` | GPU | + +T1/T2 report the comparison's SCALE and worst absolute difference +unconditionally and assert the scale is O(1): the inherited 5e-3 atol would +accept everything if the tensors compared were ~1e-7 +([[count-based-tolerances-bound-nothing]]). + +### 4.1 Mutations + +| id | mutation | must be caught by | +|---|---|---| +| A2D1-M1 | clamp the state-update group index to `min(h/hpg, 1)` | T1 — and NOT by the pre-existing `H=4 G=2` case, which is the point | +| A2D1-M2 | delete the decode branch of the paged forward | T4's counter readout | +| A2D1-M3 | select the decode arm on `T == 1` instead of on `num_decodes` | a one-token prefill of a fresh request | +| A2D1-M4 | stop incrementing the counters | T3 | + +--- + +## 5. Gates + +**The A3 e2e gate must read `96/96 mode=decode STRICT PASS` on the default +(state-update) arm.** That is the acceptance condition; no throughput number is +accepted before it. + +The A/B is same-binary: `VT_NEMOTRON_H_MAMBA_DECODE_STEP=1` (default) against +`=0`, on one build, on an idle leased box, sampling the DECODE WINDOW ONLY. A +sampler started with the process swallows a multi-minute engine load and has +already produced one void number on this row. + +--- + +## 6. Stop conditions + +- The A3 gate does not read `96/96 STRICT PASS` on the default arm → the swap is + wrong; do not tune, do not widen, revert to the chunk scan and reopen #1311 + with the divergence. +- T1 or T2 fails → the equivalence contract does not hold at `n_groups=8` and + the premise of #1311 is refuted. `NEEDS_DECISION`, not a wider tolerance. +- The counters read zero on both arms → the instrument is broken; T3 is the + triage, not the GPU log. +- **The A3 gate does not read `96/96 STRICT PASS` on ANY gated host → the swap + is not accepted there, and a pass on a DIFFERENT host does not substitute.** + This condition FIRED on sm_121a and was DISCHARGED by the same-binary A/B: the + chunk-scan arm diverges identically, so the host owns the token and this row + does not (#1388). The discharge is the A/B, never the argument that the other + host passed. + +### 6.1 The speed stop condition FIRED, and what it does and does not close + +#1311 pre-registered its own refutation criterion: *"Refuted if per-token time +moves less than 3%."* On `thor:gpu0` at c1 the move is **0.388%**, so the speed +hypothesis is **REFUTED on that box**. Recording it is the point of +pre-registering it; a criterion that only ever confirms is not one. + +**What the refutation does NOT close, and this is not consolation.** The +counters are not a prediction, they are what the run launched: the arm +demonstrably removed 92 of 115 SSD kernel launches, all 230 driver alloc/frees, +all 46 memsets, 104.9 MiB of per-token scratch and all 92 gather/scatter +launches, **and per-token time did not move**. The correct reading is therefore +*the c1 decode step on sm_110 is not bound by launch count, driver allocation or +state-copy traffic*, which is a positive finding about where the time is NOT. + +**Three regimes this measurement never entered**, so none of them is refuted: + +1. **GB10 / sm_121a.** The 6.31% busy-fraction and 0.014369 s/token references + are GB10's and no ratio against them is quoted here + ([[negative-results-are-regime-dependent]]). +2. **Concurrency above 1.** `qwen3_5.cpp:4730-4746` says the gather/scatter tax + scales with concurrency; G-SAFE pins `num_reqs <= 1`, so the regime where it + would show was not measured. A2-B owns lifting that. +3. **n > 1.** One run per leg, and the two legs saw engine loads of 654.7 s and + 778.2 s — an 18.8% spread on an excluded phase, which is evidence the host + was not in one state. 0.388% is not separable from that. + +**No ceiling is declared. The next traceable step** is an `nsys` trace of the +DECODE WINDOW ONLY on both legs. The counters say what the step stopped +launching; only a trace says what the 0.776 s/token is spent on. + +**The row still lands.** Correctness is the acceptance condition and it is met; +the "Nothing lands dead" violation is closed by measurement; and the arm is what +vLLM runs, which is repository policy independent of a local speed delta. + +--- + +## Now + +T1, T2 and T3 committed and green. T4 MEASURED on `thor:gpu0` (sm_110), both +legs of one binary: + +- **Correctness ACCEPTED.** The A3 gate reads `96/96 mode=decode STRICT PASS` + on the single-step arm, with `reference-tier lines: 0`. The chunk-scan leg + passes identically, so the swap changes no token. +- **Reachability ESTABLISHED.** Every decode step launches 23 state-update rows + and 0 chunk scans, 0 gathers and 0 scatters, against 23 chunk scans and 46 + gathers + 46 scatters with the arm off. `vt::Mamba2StateUpdate` went from zero + callers under `src/vllm/` to 23 launches per decoded token through a + production entry point. +- **★ SPEED REFUTED on this box.** Per output token 0.776159 s (on) against + 0.773156 s (off) = **+0.388%**, in the slower direction. §6's stop condition + quotes #1311's own bar — "Refuted if per-token time moves less than 3%" — and + 0.388% is under it. See §6.1. + +**★ The `dgx:gpu0` sm_121a leg read `95/96 DIVERGENCE` on BOTH kernel arms, and +a THIRD leg resolved it: the cause is A2-Q1's device mamba arm, NOT this row and +NOT the architecture.** Three legs of one binary: + +| leg | device mamba | tokens | s/token | vs vLLM | +|---|---|---|---|---| +| `on` (single-step) | 1 | `95/96 DIVERGENCE` | 1.584694 | 110.3x | +| `off` (chunk scan) | 1 | `95/96 DIVERGENCE` | 1.554233 | 108.2x | +| **`hostmamba`** | **0** | **`96/96 STRICT PASS`** | 10.318897 | 718.1x | + +**Leg 3's own counters are what make this a discrimination and not a +correlation**: `state_update_rows=0 chunk_scan_calls=0 conv_update_rows=0 +conv_fwd_calls=0` with `gathers=46 scatters=46`, which is the host path's +signature and nothing else's. Without them the leg would only correlate a flag +with an outcome. Filed as +[#1388](https://github.com/mudler/vllm.cpp/issues/1388). + +**The divergence is ONE token and both device legs lose the SAME one** — +prompt 2, position 32 of 32, `11286` against the oracle's `3468`, positions 1-31 +byte-identical. The repaired verdict grep captured it; the first GB10 run threw +it away. **Defect or bf16 near-tie is NOT settled**, and a separate fresh +implementer owns root-causing it from the oracle's top-2 margin. + +**The acceptance statement for THIS row is unchanged.** Each host returns the +same verdict on both KERNEL arms — `96/96` vs `96/96` on sm_110, `95/96` vs +`95/96` on sm_121a — so the swap is token-neutral, and the sm_121a token belongs +to the arm underneath it. + +### The ±2% on/off delta is NOISE, proven by a sign flip + +| run | on | off | on-vs-off | +|---|---|---|---| +| first GB10 | 1.513958 | 1.544706 | **−1.991%** (on faster) | +| this GB10 | 1.584694 | 1.554233 | **+1.960%** (on slower) | + +A quantity that reverses direction between two runs of one comparison on one box +is not a measurement of that comparison. §6.1's refutation was argued from +#1311's 3% bar; this retires the figures outright. + +### The larger finding, which is not this row's + +A2-Q1's device mamba arm is worth **6.64x per output token on GB10** — 10.318897 +s off against 1.554233 s on — closing **718.1x → 108.2x** versus vLLM, decode +busy 7.86% → 10.18%. It agrees in direction and rough magnitude with #1289's +independent Thor A/B (7.17x). **ONE run per leg on a contended box**, engine +load excluded; the 6.64x is large enough to survive that and the ~2% deltas are +not. **108.2x remains an OPEN GAP, not parity**, and no ceiling is declared. + +## Owed + +- **[#1388](https://github.com/mudler/vllm.cpp/issues/1388)** — RESOLVED to + A2-Q1's device mamba arm and no longer this row's. What remains open there is + whether the single token is a defect or a bf16 near-tie; a fresh implementer + owns it, starting from the oracle's top-2 margin at prompt 2 position 32. +- **A repeated A/B (n >= 3 per leg) on an idle box.** The sm_110 legs saw an + 18.8% spread in engine-load time and the two hosts' deltas have opposite + signs; neither 0.388% nor 1.991% is separable from that at n=1. +- **An `nsys` decode-window trace of both legs**, the next traceable step §6.1 + names. GB10 runs this arm ~2x slower than Thor (1.514 vs 0.776 s/token) and + that is unexplained. +- **Concurrency > 1** is where `qwen3_5.cpp:4730-4746` says the gather/scatter + tax shows; G-SAFE pins `num_reqs <= 1`, so A2-B owns that regime. + +## Outcome + +Pending; this section is authored when the row reaches `DONE`. diff --git a/.agents/specs/nemotron-h-a2q1-fp8-mamba.md b/.agents/specs/nemotron-h-a2q1-fp8-mamba.md index 8c4f593f3..e6dc955e8 100644 --- a/.agents/specs/nemotron-h-a2q1-fp8-mamba.md +++ b/.agents/specs/nemotron-h-a2q1-fp8-mamba.md @@ -418,23 +418,238 @@ stated removal condition is the device/paged runner path, which is A2-P's. ## 10. Now -**State at this commit:** spec only. No product code, no lifecycle change. - -A2-Q1 is **BLOCKED on #960** and is not claimable until it lands. When it does, a -fresh implementer claims this file, captures a RED per §5.3 first, and lands the -mamba device arm with G-SAFE untouched. A fresh reviewer — never the implementer -— runs the §5.3 mutations. - -**Three things to read before the first edit:** §3, so the Thor dependency is -consumed rather than re-derived; §1.1, so the conv-state dtype is left to A2-P; -and §5.2, so the bands are measured rather than chosen. +**State at this commit:** the device arm is IMPLEMENTED and REACHED, and its +binding measurement is PENDING a lease. + +`NemotronHMamba2MixerDevice` (`nemotron_h_device.cpp`) runs the whole block on +the device on the shared FP8 W8A8 seam, and BOTH forwards select it at runtime: +`NemotronHDeviceForward` (non-paged, discarding the recurrence) and +`NemotronHPagedForward` (production, advancing the gathered recurrent rows in +place). The selection is a runtime op-table query plus a weight-form predicate, +never a preprocessor guard, so a dense NemotronH or a device without the fp8 pair +keeps the host bounce. + +**Three design points differ from §4 and the reasons are recorded here rather +than left to be re-derived:** + +1. **The `Fp8Weight` pair is built ON FIRST DEVICE USE, not by the loader** + (§4.1/§4.2 put it in the loader). Building it in a `ResidentSlot` the weights + own is A2-Q2a's newer idiom and it is strictly better here: it does not double + the 890 MB fp8 tower in host memory at load, it does not upload anything on a + host-only run, and it leaves `rep.host_bytes` — the literal `18888922112` that + `test_nemotron_h_loader.cpp:310` pins — untouched, so §4.2's re-derivation + obligation does not arise at all. The e4m3 staging copy is released as soon as + `ResidentFp8` has uploaded it, so the peak cost of the conversion is one + projection. +2. **The upload IS accounted** (§4.4 expected the report to say it was short). + `dense_fp8::ResidentFp8` still does not call `load_stats::AddDeviceUpload` + — that is [#974](https://github.com/mudler/vllm.cpp/issues/974), unchanged, and + the shared header is not touched — so A2-Q1 accounts what IT uploads at the + site that causes it, exactly as `ResidentWeight` and `ResidentNvfp4` do. The + counter is then also the instrument the residency case reads, because an arm + that re-uploaded the tower every step returns identical numbers to one that + uploads it once. +3. **The paged selection carries an `ssm_dtype == f32` term.** + `vt::GdnStateGather` widens the page into an f32 working buffer by op + contract, and the HOST arm then narrows it back to `ssm_dtype` before the + mixer sees it. On a checkpoint whose `mamba_ssm_cache_dtype` is not f32 the + two arms would round differently and the per-block numeric gate would be + comparing two different computations. The released checkpoint resolves f32. + +**The comparison is NOT bit-comparable by construction, and §5.2's "measure the +band" is therefore binding rather than cautionary.** The host reference is +W8A16 — `DenseFor` dequantizes the fp8 weight to bf16 and leaves the activation +alone, as `DenseBf16` says outright — while the device arm is W8A8 as vLLM is. +The difference between them is the e4m3 activation quantization, and every band +in `tests/vllm/models/test_nemotron_h_mamba_device.cpp` is measured in the run +against a defect the fixture separates. + +### 10.1 Thor (sm_110) RAN the arm, and it answers §3's question + +Measured 2026-08-18 under an `rc` lease on `thor:gpu0`, product tree behaviourally +identical to this row's head (the later commits touch tests, scripts and one +comment only). `cmake --build -j 4` returned 0. The feature table read +**`ENABLED for [110]: 1 ; DISABLED cells: 7`** — only `marlin-nvfp4`, with +`cutlass-fp8` and both `scaledmm-c3x` cells DISABLED. + +**That configuration is the whole point.** §3.1 measured Thor as having half an +fp8 arm: the GEMM present through `kMatmulFp8CublasLt`, the activation quant +trapped in a CUTLASS-gated TU. #991 moved the registration out. The arm running +here, on a build with no CUTLASS fp8 at all, is what closes that question. + +`test_nemotron_h_mamba_device` reported **49 assertions** where a GPU-less box +reports 4, so the device path executed rather than skipping: + +| case | result | +|---|---| +| fresh block vs host reference, T=1 / 8 / 12 | agreed 0.164 / 0.282 / 0.309 against a measured band of 0.5; 128 / 1024 / 1536 elements examined; 3 widths covered | +| the fp8 tower uploads ONCE | first call 61760 B == expected 61760 B, second call 0 B | +| refuses a dense projection, refuses a missing `input_scale` | both threw | +| the carry across two legs | **FAILED, and the instrument was the defect** — see below | + +Neighbouring suites, same run, all `Status: SUCCESS!`: `test_nemotron_h_forward` +16/16 (5716 assertions), `test_nemotron_h_paged_forward` 12/12 (3256), +`test_nemotron_h_loader` 2/2, `test_nemotron_h_moe_device` 2/2 (29), +`test_ops_mamba2_ssd` 12/12 (2095), `test_ops_fp8_cpu` 5/5 (62). + +### 10.2 The carry gate banded a defect smaller than the noise it had to accept + +The case banded the SECOND LEG'S OUTPUT against the separation of a dropped +carry. Thor measured the second leg agreeing to 0.705 while a dropped carry +separated by only 0.205, so the derived band (0.102) sat BELOW the deviation a +FRESH leg already shows (0.164 at T=1, from the case above). + +That is §5.2's lesson arriving from the other direction. The two arms are W8A8 +against W8A16, so a fresh leg already disagrees by the e4m3 activation +quantization and a second leg compounds that with the same disagreement +propagated through the carried state. **A defect whose separation is smaller than +the noise the comparison must accept is not resolvable from that comparison**, and +widening the band until it passes is what §8.1 says to stop for. + +The repair gates what the carry IS: the STATE. A dropped carry hands the next leg +zeros, so the separation between the advanced state and a zeroed one is 1.0 by +construction — about six times the noise floor. The conv window and the SSM state +are banded separately, each against its own zeroed twin. The noise floor is +measured in the run at the same width and printed beside the separation, and the +second leg's output carries an assertion only when the separation exceeds twice +that floor, with the condition printed either way. + +The failure also exposed two fixture defects. `mamba_ssm_cache_dtype` was unset +and resolved bf16 — NOT the configuration the paged forward selects the device arm +for (`ssm_dtype == f32`), so the cheap arm was gating a path production does not +take. And the whole file was a skip on a GPU-less box, so a CPU-runnable case now +pins the op contract the split depends on. + +### 10.2b The repaired carry gate, re-run on hardware, and what it measured + +Thor re-run at the branch head: `test_nemotron_h_mamba_device` **5 cases, 63 +assertions, 0 failed**, and every neighbouring suite green again. The repair of +§10.2 therefore holds on the silicon that exposed the defect, and the numbers it +prints now justify the repair rather than merely passing: + +| quantity | measured | separation | band | margin | +|---|---|---|---|---| +| W8A8-vs-W8A16 noise floor, T=1 | **0.2465** | — | — | the reference for everything below | +| carried conv window | **0.1746** over 576 elements | 1.0 (zeroed) | 0.5 | 2.9x | +| carried SSM state | **0.0614** over 2048 elements | 1.0 (zeroed) | 0.5 | 8.1x | +| second leg's output | 0.7055 | 0.2045 | — | NO assertion, see below | + +**The data now proves the diagnosis that drove the repair.** A dropped carry +separates the second leg's output by 0.2045, while the noise the comparison must +accept is 0.2465 — the defect is genuinely SMALLER than the noise, so +`separation > 2 * noise_floor` is false and the case makes no assertion there, by +design. The state comparison carries the case instead, at 2.9x and 8.1x margins. +Had the original band survived, it would have been asserting on a quantity it +cannot resolve. + +**One diagnostic was defective and is fixed.** The line that reports whether the +second leg is resolvable printed `1` rather than its prose, because doctest +stringifies a `const char*` as a BOOL and the message streamed a `char*` ternary. +That line's whole job is to make "no assertion was made here" a STATED result +rather than a silent hole, so a version of it that cannot say what it means is +the same class of defect as the band it reports on. It now builds a +`std::string`; reproduced against doctest 2.5.2 both ways before the fix. + +### 10.2c The decode-window sampler works, and it quantifies the dilution + +Thor re-run, A3 gate with the arm ON, sampler starting only after +`engine loaded in Ns`: + +``` +RC[a3 on]=0 +[nemotron-h] engine loaded in 500.9s +[nemotron-h] TOKEN MATCH: 96/96 over 3 prompt(s) (full rows=3, short rows=0, mode=decode) +[nemotron-h] STRICT PASS +on: GPU busy in 240 of 564 DECODE samples = 42.55% busy +on: decode window 75.418 s (the engine load is OUTSIDE it) +on: per output token 0.785606 s +reference-tier lines in on: 0 +``` + +**42.55% over the decode alone against 15.33% over load+decode.** The load is +500.9 s and the decode is 75.418 s, so the old window was 87% load — the +dilution §10.3 diagnosed is now measured rather than argued. + +**The same cross-silicon defect was still live on the per-token line and is +fixed.** That run printed `ratio 54.7x` against vLLM's 0.014369 s, which is a +GB10 figure, for a decode measured on Thor. It is exactly the defect the +busy-fraction reporter carried, and fixing one surface while leaving its twin is +how a wrong comparison survives a correction. The ratio is now quoted only when +`ARCH` is the arch it was measured on; elsewhere the rate still prints and the +comparison is withheld by name. Both arms are pinned in +`tests/scripts/test_nemotron_h_a2q1_per_token.py`, and quoting the ratio +unconditionally reds the suite. + +### 10.3 The A/B, on the corrected instrument: the busy fraction ROSE + +Second Thor lease, fresh build and clone, sampler measuring the DECODE window +alone. Same binary, same checkpoint, same golden, differing only by +`VT_NEMOTRON_H_DEVICE_MAMBA`: + +| flag | mamba arm | A3 | exit | decode GPU busy | per output token | decode wall | +|---|---|---|---|---|---|---| +| `1` (default) | device FP8 W8A8 | `96/96 STRICT PASS` | 0 | **240/564 = 42.55%** | **0.785606 s** | 75.4 s | +| `0` | host, dequant to bf16 | `93/96 DIVERGENCE` | 1 | **700/3808 = 18.38%** | **5.633442 s** | 540.8 s | + +**THE BUSY FRACTION ROSE, WHICH IS THIS UNIT'S ACCEPTANCE TEST: 18.38% to +42.55%, +24.17 points, a 2.31x rise.** A decode token costs 7.17x less. Peak host +44070 MiB. `reference-tier lines: 0` on both arms, so neither took the portable +tier. + +**Read on the box it was taken on, and nowhere else.** These are sm_110 figures. +The 6.31% baseline and the 0.014369 s per-token reference are BOTH GB10's, so +neither supports a ratio against these numbers, and the instrument now withholds +both comparisons by name on any other arch. What is established is the ON/OFF +difference on one box, and that is exactly the comparison the hypothesis needed. + +**The `ratio 54.7x` and `ratio 392.1x` strings in that run's log are stale and +must not be quoted.** That job cloned before the per-token arch gate landed, so it +still printed the GB10 comparison unconditionally; the per-token VALUES are sound +measurements of that box, the ratios beside them are not. + +### 10.4 The divergence reproduces, so it is not n=1 on Thor + +Two independent Thor leases, separate builds and separate clones, agree exactly +on both arms: + +| run | arm ON | arm OFF | +|---|---|---| +| `20260818T222352Z` | 96/96 | 93/96 | +| `20260818T232910Z` | 96/96 | 93/96 | + +So the `n=1` caveat is lifted FOR THOR: the host arm's 93/96 is reproducible +there, and the device arm's 96/96 is too. GB10 remains n=1 in the other +direction (the #1157 run, host arm, 96/96), and no GB10 run of the DEVICE arm +exists yet. [#1290](https://github.com/mudler/vllm.cpp/issues/1290) carries this. + +**What is still owed.** The occupancy hypothesis is now SUPPORTED on sm_110 and +UNMEASURED on sm_121a, and the two are not interchangeable. Specifically owed: + +- the GB10 run, which is the only one that can be read against the 6.31% + baseline and the 0.014369 s per-token reference, and the only place a DEVICE-arm + A3 result does not yet exist at all; +- the §5.1 per-block numeric gate against `trace.mixer[l]` on the real checkpoint + — the A3 gate is token-level and cannot see a per-layer defect whose argmax is + unchanged; +- the §5.3 mutations, which belong to the fresh reviewer; +- the three moved tokens on the OFF arm, whose oracle top-2 margin decides + whether [#1290](https://github.com/mudler/vllm.cpp/issues/1290) is a near-tie + sensitivity or a wrong answer. + +`scripts/nemotron-h-a2q1-dgx-gate.sh` is the recipe, and it now measures the +decode window, withholds both GB10 references off `121a`, and refuses a fraction +outright when the load boundary never appears. ## 11. Owed -- [#974](https://github.com/mudler/vllm.cpp/issues/974) — the fp8 resident - helpers upload without `AddDeviceUpload` / `AdoptDeviceBytesAsHost`. A2-Q1 - consumes them and reports its accounting as short (§4.4); the fix is not this - unit's. +- [#974](https://github.com/mudler/vllm.cpp/issues/974) — `dense_fp8::ResidentFp8` + uploads without `AddDeviceUpload` / `AdoptDeviceBytesAsHost`. A2-Q1 consumes it + and accounts its own upload at its own call site (§10.2); the fix INSIDE the + shared header is not this unit's, and every other caller of the seam is + byte-unchanged. +- The §5.1 real-checkpoint per-block numeric gate, the §5.3 mutation pass and the + GPU-occupancy measurement are PENDING a GB10 lease. They are the unit's + acceptance test, not paperwork. - The device arm has **no production caller** until A2-P wires it through `ModelRegistry::Forward` (§7). Tracked on [#810](https://github.com/mudler/vllm.cpp/issues/810). diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index bf7257e89..c1ed5cae1 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -12,6 +12,9 @@ | **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 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 | | **Record-anchor ratchet** (`ENG-RECORD-ANCHOR-RATCHET`, #632) | **No number owed:** a record checker. At `8daa67b39`, **832 of 867** in-scope citations (**96.0%**) were already parsed; no symbol test and no report ran. Rot **38** (32 stale, 6 broken), **32 in range** | +| **NemotronH FP8 mamba on device** (`A2-Q1`, [#810](https://github.com/mudler/vllm.cpp/issues/810)) | Thor sm_110, 20.1 GiB checkpoint, c1 greedy, same-binary A/B, n=2 | **Decode GPU busy ROSE 18.38% -> 42.55% (2.31x); per output token 5.6334 -> 0.7856 s (7.17x).** A3 `96/96 STRICT PASS` on; `93/96 DIVERGENCE` off ([#1290](https://github.com/mudler/vllm.cpp/issues/1290)) | Thor figures; the 6.31% and 0.014369 s references are GB10's and NOT comparable. GB10 leg pending. Peak host 44070 MiB | +| **NemotronH FP8 mamba on device, GB10 leg** (`A2-Q1`, [#810](https://github.com/mudler/vllm.cpp/issues/810)) | GB10 sm_121a, 20.1 GiB ckpt, c1 greedy, same-binary A/B, **n=1 per leg**, decode window only | **6.64x per output token: 10.318897 s arm OFF -> 1.554233 s arm ON**, closing **718.1x -> 108.2x** vs vLLM; decode busy 7.86% -> 10.18% | Agrees with the Thor 7.17x. **108.2x is an OPEN GAP, not parity.** Single run, contended box, load excluded. Arm ON also reads `95/96` ([#1388](https://github.com/mudler/vllm.cpp/issues/1388)) | +| **NemotronH single-step mamba decode** (`A2-D1`, [#1311](https://github.com/mudler/vllm.cpp/issues/1311)) | Thor sm_110 and GB10 sm_121a, c1, same-binary A/B, n=1 per leg | **TOKEN-NEUTRAL on both hosts** (same verdict on both arms). **SPEED REFUTED:** the +-2% on/off delta FLIPPED SIGN between two GB10 runs, so it is noise | GB10's `95/96` is A2-Q1's mamba arm, not this row and not the arch ([#1388](https://github.com/mudler/vllm.cpp/issues/1388)); removing it restores `96/96`. See [`benchmark-record.md`](../.agents/benchmark-record.md) | | **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 8880dff65..e3bc6b48a 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -244,6 +244,9 @@ Read-only observability; none change output. | `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_NEMOTRON_H_DEVICE_MAMBA` | on | `=0` routes NemotronH's 23 FP8 W8A8 Mamba2 layers back through the host reference instead of the device arm ([#810](https://github.com/mudler/vllm.cpp/issues/810) A2-Q1). The same-binary A/B every measurement of that arm needs, and a bisect lever for a device-vs-host divergence: with it off the identical binary takes the host bounce, in which each projection is dequantized from fp8 to bf16 on every call, so a throughput or GPU-occupancy difference is attributable to the arm and not to a rebuild | +| `VT_NEMOTRON_H_MAMBA_DECODE_STEP` | on | `=0` makes NemotronH's device Mamba2 layers run the CHUNKED prefill kernels on decode rows too, which is what they did before [#1311](https://github.com/mudler/vllm.cpp/issues/1311). On the default the decode rows take `vt::CausalConv1dUpdate` and `vt::Mamba2StateUpdate` at their cache slots in place, mirroring `mamba_mixer2.py:1012,:1087`, so the per-request gather and scatter go away with them. Prefill keeps the chunk scan on both settings. This exists so the two arms can be A/B'd in ONE binary: a decode-window measurement taken against a differently-built binary measures the build. Not a configuration -- the default is what vLLM runs | +| `VT_NEMOTRON_H_ARM_TRACE` | off | `=1` prints one line per `NemotronHPagedForward` step naming which recurrent kernels it launched: `state_update_rows`, `chunk_scan_calls`, `conv_update_rows`, `conv_fwd_calls`, `gathers`, `scatters`, as a delta over that step. The decode and prefill arms compute the same recurrence, so they emit the same tokens and no token gate can say which one ran ([#1311](https://github.com/mudler/vllm.cpp/issues/1311)). Deliberately SEPARATE from `VT_NEMOTRON_H_DIAG`, which downloads the carry and the residual per layer per step: a timed A/B run under that flag measures the diagnostic. This one is a single `fprintf` of six resident counters, so it can stay on for a timed leg | | `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/FEATURES.md b/docs/FEATURES.md index b3de34b47..a02a9f653 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -144,7 +144,7 @@ speed-pending, which [BENCHMARKS.md](BENCHMARKS.md) tracks. | `LagunaForCausalLM` | poolside/Laguna-S-2.1-NVFP4, GGUF-Q4_K, Laguna-XS | byte-exact near-tie (distributional vs vLLM) | vLLM parity+ 1.03x, default on, via the `laguna-gen` CLI; the registered engine forward VT_CHECKs non-bf16 (`ARCH-ONE-SURFACE` fold) | | `KimiLinearForCausalLM` | Kimi-Linear-48B-A3B (KDA + NoPE-MLA + MoE) | **Folded onto the shared paged runner (ROW 7 §21, #122): engine==CLI 128/128 byte-identical; vs golden 122/128 (the intrinsic near-tie profile); FA2 paged MLA default-ON; SACRED post-fold green** | Served via `vllm_engine_load` + `vllm_complete_tokens` (ABI v13); server 19.0 tok/s wall vs vLLM ~21 (~0.90×), speed residual open | | `KimiK3ForConditionalGeneration` | Kimi-K3 (2.8T MoE) | scaffold: registry+config+enumeration gated, forward refuses | HW-infeasible (~1.56 TB); no run | -| `NemotronHForCausalLM` | Nemotron-3.5-Lightning-30B-A3B-NVFP4 (`nvidia` @`29f2d174`) | config+enumeration+KV-shape gated; hybrid forward COMPUTES; loader materializes 18487/18487 as SHIPPED | **PAGED (#810 A2-P): K/V go into the runner's pages and conv+SSM rows are carried at the metadata's state indices.** G-SAFE narrows to `num_reqs <= 1`. Owed: device `lm_head` (A2-Q2b), FP8 mamba (A2-Q1), MTP, GGUF | +| `NemotronHForCausalLM` | Nemotron-3.5-Lightning-30B-A3B-NVFP4 (`nvidia` @`29f2d174`) | config+enumeration+KV-shape gated; hybrid forward COMPUTES; loader materializes 18487/18487 as SHIPPED | **PAGED (#810 A2-P) and the 23 FP8 W8A8 Mamba2 blocks now run ON THE DEVICE (A2-Q1) instead of dequantizing to bf16 on the host every call.** G-SAFE stays `num_reqs <= 1`. Owed: device `lm_head` (A2-Q2b), MTP, GGUF | | `MuseGlimmerForCausalLM` | real tensors, **bf16 depth 4/52 only**: 5 prefill argmax positions match a torch transcription of vllm#51655 and HF. GGUF full depth generates coherently (#347, #359) but is **NOT token-exact** | text forward + loader vs an fp32 reference, per-mechanism property tests, scaffold 11/11, GGUF gate 17/17. An ABSENT config key now takes the architecture's constant (#412): GGUF post-norms ran at 1e-5, not 1e-8 | no vLLM denominator (pin cannot load it); SECONDARY llama.cpp, same GGUF, GB10 CPU: prefill tie **0.997x**, decode 0.232x, RSS 1.92x (#333) | | `MuseGlimmerForConditionalGeneration` | vision: **no reference run of any kind**; enumeration gated vs the released 30B index (1436/1436). Image/video need bf16 safetensors: `mmproj-kquant.gguf` is refused by name | perception encoder loaded and wired, so an image or video prompt runs; `perception_emb_norm` now armed by default (#405). Reachability plus placeholder scatter only, no image or video correctness | not measurable; anchored to open vllm#51655 | | `LlamaModel` | landed tiny synthetic embedding fixture (engine path == direct pooler path, identical vectors; f64 LAST+normalize reference); real checkpoint (e5-mistral class) is a NAMED residual | pooling/embed only, text paths refuse by task; `vllm_embed` + `/v1/embeddings` | n/a (CPU correctness-grade embeddings) | diff --git a/docs/STATUS.md b/docs/STATUS.md index ef8d10d83..fdc4fea13 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -140,7 +140,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) | **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) | +| Nemotron-3.5-Lightning-30B-A3B (Mamba2 + GQA + relu2 MoE) | **Host and Thor sm_110 A3 PASS 96/96 `STRICT PASS`; sm_121a 95/96 traced to the FP8 mamba arm ([#1388](https://github.com/mudler/vllm.cpp/issues/1388)), worth 6.64x/token there** | ABI-only driver; G-SAFE `num_reqs <= 1`. 23 FP8 Mamba2 blocks on device, DECODING on the single-step recurrent kernels ([#1311](https://github.com/mudler/vllm.cpp/issues/1311)); `lm_head` the last host arm | | 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/docs/USAGE.md b/docs/USAGE.md index a50e50e77..855d10821 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -812,7 +812,7 @@ tokens quietly. | Architecture | Why it refuses | |---|---| | `KimiK3ForConditionalGeneration` | Needs ~1.56 TB (MXFP4); no host here can run it | -| `NemotronHForCausalLM` | **Only BATCHED decode still refuses.** A2-P (#810) narrowed this: `ForwardNemotronHForCausalLM` now selects the paged forward whenever the runner supplies paged KV and recurrent state, so K/V go into the runner's pages and the conv/SSM rows are carried across steps, and `examples/nemotron_h_gen` reaches all of it through `include/vllm.h` alone. What is left is `num_reqs > 1`, refused by name because one request's pages and one request's recurrent state are carried per step and a multi-request step would be decoded as ONE concatenated causal sequence — plausible wrong tokens rather than a failure. Owed to A2-B. **The end-to-end token gate against the pinned oracle has NOT run**, so no claim is made here about what this checkpoint emits; `docs/BENCHMARKS.md` records that as pending rather than as silence. `lm_head` and the FP8 Mamba2 projections still compute on the host, and a GGUF file is refused by name since no GGUF arm exists for it. See *Nemotron-3.5-Lightning-30B: the exact weights, and which arms run* below | +| `NemotronHForCausalLM` | **Only BATCHED decode still refuses.** A2-P (#810) narrowed this: `ForwardNemotronHForCausalLM` now selects the paged forward whenever the runner supplies paged KV and recurrent state, so K/V go into the runner's pages and the conv/SSM rows are carried across steps, and `examples/nemotron_h_gen` reaches all of it through `include/vllm.h` alone. What is left is `num_reqs > 1`, refused by name because one request's pages and one request's recurrent state are carried per step and a multi-request step would be decoded as ONE concatenated causal sequence — plausible wrong tokens rather than a failure. Owed to A2-B. **The end-to-end token gate against the pinned oracle has NOT run**, so no claim is made here about what this checkpoint emits; `docs/BENCHMARKS.md` records that as pending rather than as silence. A2-Q1 puts the 23 FP8 W8A8 Mamba2 blocks on the device, so `lm_head` is the last arm that computes on the host, and a GGUF file is refused by name since no GGUF arm exists for it. See *Nemotron-3.5-Lightning-30B: the exact weights, and which arms run* below | This is a deliberate state, not a bug: registering the architecture is what lets the config parse and weight-name mapping be tested before the forward exists. diff --git a/scripts/agent-preflight.sh b/scripts/agent-preflight.sh index a9e555837..8f6c4a515 100755 --- a/scripts/agent-preflight.sh +++ b/scripts/agent-preflight.sh @@ -135,6 +135,7 @@ SUITES=( test_check_model_checklist test_check_supported_models test_check_env_doc + test_nemotron_h_a2q1_per_token test_checker_text test_check_fusion_consistency test_check_fp4_resident_consistency diff --git a/scripts/nemotron-h-a2d1-gpu-gate.sh b/scripts/nemotron-h-a2d1-gpu-gate.sh new file mode 100755 index 000000000..e7857c100 --- /dev/null +++ b/scripts/nemotron-h-a2d1-gpu-gate.sh @@ -0,0 +1,294 @@ +#!/usr/bin/env bash +# A2-D1 (#1311, .agents/specs/nemotron-h-a2d1-mamba-decode-step.md) — the GPU +# gate for the SINGLE-STEP recurrent decode arm. +# +# It is a SIBLING of scripts/nemotron-h-a2q1-dgx-gate.sh, not an edit of it: +# that script owns A2-Q1's device-vs-host A/B (VT_NEMOTRON_H_DEVICE_MAMBA) and +# is a live surface on another branch. This one owns A2-D1's A/B +# (VT_NEMOTRON_H_MAMBA_DECODE_STEP), with the device mamba arm ON in BOTH legs, +# so the only thing that differs between them is which recurrent kernels the +# DECODE rows take. One file per row, read with a glob (AGENTS.md §Records). +# +# RUN IT INSIDE A LEASE, NEVER OVER ssh: +# rc run -d dgx:gpu0 --max-runtime 6h -- bash -lc \ +# 'git clone --depth 50 -b row/A2-D1-mamba-decode-step \ +# https://github.com/mudler/vllm.cpp /root/src \ +# && bash /root/src/scripts/nemotron-h-a2d1-gpu-gate.sh' +# +# The environment facts (sbsa lane, anchored cuda package names, the toolkit +# proved by a LINK and not by --version, /workspace refusing symlinks, -j 4) +# are the a2q1 script's and are re-encoded here rather than re-discovered. +set -u -o pipefail + +LOG_ROOT=${LOG_ROOT:-/workspace/a2d1} +CKPT=${CKPT:-/workspace/a3/ckpt-stage} +SRC=${SRC:-/root/src} +ARCH=${ARCH:-121a} # 121a = GB10, 110 = Thor +BUILD=${BUILD:-/root/build-cuda-a2d1} +CUTLASS=${CUTLASS:-/root/cutlass} +STAMP=$(date -u +%Y%m%dT%H%M%SZ) +RUN=$LOG_ROOT/$STAMP +mkdir -p "$RUN" +echo "A2D1 run dir: $RUN" + +step() { echo; echo "=== $* ==="; } +# Every gate command runs BARE and echoes its own status. Never pipe a command +# whose exit code matters -- a pipeline reports the LAST stage. +rcx() { "$@"; local r=$?; echo "RC[$*]=$r"; return $r; } + +step "0. the box" +rcx uname -m +rcx nvidia-smi --query-gpu=name,compute_cap,memory.total --format=csv +rcx df -h /root /workspace +rcx free -m + +step "1. contention -- a timing number measured beside another job is VOID" +nvidia-smi --query-compute-apps=pid,used_memory --format=csv | tee "$RUN/contention.txt" +echo "RC[nvidia-smi compute-apps]=${PIPESTATUS[0]}" + +step "2. toolchain" +export DEBIAN_FRONTEND=noninteractive +rcx apt-get update -qq +rcx apt-get install -y -qq git cmake ninja-build g++ curl ca-certificates python3 python3-dev +if ! command -v nvcc >/dev/null 2>&1; then + curl -fsSL -o /tmp/cuda-keyring.deb \ + https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2404/sbsa/cuda-keyring_1.1-1_all.deb + echo "RC[curl keyring]=$?" + rcx dpkg -i /tmp/cuda-keyring.deb + rcx apt-get update -qq + rcx apt-get install -y -qq cuda-nvcc-13-0 cuda-cudart-dev-13-0 libcublas-dev-13-0 \ + cuda-nvrtc-dev-13-0 cuda-nvtx-13-0 cuda-profiler-api-13-0 libcurand-dev-13-0 +fi +export PATH=/usr/local/cuda/bin:$PATH +rcx nvcc --version + +step "2b. the toolkit is proved by a LINK, not by --version" +cat > /tmp/probe_a2d1.cu <<'EOF' +#include +#include +int main() { + cublasLtHandle_t h = nullptr; + const auto s = cublasLtCreate(&h); + std::printf("cublasLtCreate=%d\n", static_cast(s)); + return s == CUBLAS_STATUS_SUCCESS ? 0 : 1; +} +EOF +rcx nvcc -arch=sm_$ARCH /tmp/probe_a2d1.cu -o /tmp/probe_a2d1 -lcublasLt +if [ $? -ne 0 ]; then + echo "VOID: the CUDA toolkit does not link cublasLt; every number below would be a lie" + exit 2 +fi +rcx /tmp/probe_a2d1 + +step "3. cutlass" +if [ ! -f "$CUTLASS/include/cutlass/cutlass.h" ]; then + rcx git clone --depth 1 --branch v4.5.0 https://github.com/NVIDIA/cutlass "$CUTLASS" +fi +rcx test -f "$CUTLASS/include/cutlass/cutlass.h" + +step "4. configure" +cmake -S "$SRC" -B "$BUILD" -G Ninja \ + -DCMAKE_BUILD_TYPE=Release \ + -DVLLM_CPP_CUDA=ON \ + -DVLLM_CPP_CUDA_ARCHITECTURES="$ARCH" \ + -DVLLM_CPP_CUTLASS_DIR="$CUTLASS" > "$RUN/configure.log" 2>&1 +echo "RC[cmake configure]=$?" + +step "4b. the feature table -- a DISABLED or wrong-arch cell VOIDS the run" +grep -E "CUDA feature .*(ENABLED|DISABLED)" "$RUN/configure.log" | tee "$RUN/features.txt" +echo "feature cells ENABLED for [$ARCH]: $(grep -cE "ENABLED for \[$ARCH\]" "$RUN/features.txt") ; DISABLED cells: $(grep -cE 'DISABLED' "$RUN/features.txt")" + +step "5. build (-j 4: unconstrained parallelism has OOM-REBOOTED this box)" +# OPS_ONLY=1 builds and runs ONLY the op-level equivalence suite and stops. +# It exists for a small board that can host the CUDA kernels but not the 20.1 +# GiB checkpoint: the driver-group equivalence is the one piece of evidence that +# needs real CUDA and does NOT need the model, so it should not queue behind a +# box that can run the whole gate. +OPS_ONLY=${OPS_ONLY:-0} +if [ "$OPS_ONLY" = "1" ]; then + cmake --build "$BUILD" -j 4 --target test_ops_mamba2_state_update > "$RUN/build.log" 2>&1 +else + cmake --build "$BUILD" -j 4 > "$RUN/build.log" 2>&1 +fi +echo "RC[cmake build]=$?" +tail -5 "$RUN/build.log" + +step "6. the OP-LEVEL equivalence this swap rests on, at n_groups=8" +# ★ THE PREMISE, GATED BEFORE THE SWAP IS TRUSTED. The only pre-existing +# decode-vs-prefill case ran heads_per_group = 2. NemotronH runs 8. +"$BUILD/tests/test_ops_mamba2_state_update" > "$RUN/state_update.log" 2>&1 +echo "RC[test_ops_mamba2_state_update]=$?" +grep -E "test cases:|assertions:|Status:" "$RUN/state_update.log" +grep -E "driver-group equivalence" "$RUN/state_update.log" + +if [ "$OPS_ONLY" = "1" ]; then + echo + echo "OPS_ONLY=1: the op-level equivalence ran and NOTHING ELSE did. The A3 token" + echo "gate and the decode-window numbers are NOT MEASURED on this host. That is a" + echo "stated absence, not a pass." + echo "ALL LOGS: $RUN" + exit 0 +fi + +step "6b. the neighbouring suites this arm can break" +for t in test_nemotron_h_mamba_device test_nemotron_h_paged_forward test_nemotron_h_forward \ + test_nemotron_h_loader test_nemotron_h_moe_device test_ops_mamba2_ssd; do + if [ -x "$BUILD/tests/$t" ]; then + "$BUILD/tests/$t" > "$RUN/$t.log" 2>&1 + echo "RC[$t]=$?" + grep -E "test cases:|assertions:|Status:" "$RUN/$t.log" + else + echo "MISSING BINARY: $t" + fi +done + +if [ ! -d "$CKPT" ]; then + echo + echo "NO CHECKPOINT AT $CKPT -- the A3 gate and the decode-window numbers are NOT MEASURED on this host." + echo "That is a stated absence, not a pass." + echo "ALL LOGS: $RUN" + exit 0 +fi + +# ───────────────────────────────────────────────────────────────────────────── +# The A/B. ONE binary, device mamba ON in both legs, only the decode arm differs. +# ───────────────────────────────────────────────────────────────────────────── +# $1 = label, $2 = VT_NEMOTRON_H_MAMBA_DECODE_STEP, $3 = VT_NEMOTRON_H_DEVICE_MAMBA +# (defaults to 1; leg 3 sets it to 0 -- see the CONFOUND note above step 7). +run_leg() { + local label=$1 flag=$2 devmamba=${3:-1} + local log="$RUN/a3_$label.log" + : > "$log" + + # ★ THE SAMPLING WINDOW IS THE DECODE, NOT THE WHOLE PROCESS. A sampler + # started with the process contains the multi-minute 20.1 GiB engine load, + # which is GPU-IDLE, and dilutes both arms toward each other. That already + # produced one void number on this row. The driver starts FIRST and the + # sampler starts only once it has printed `engine loaded in Ns`. + # ★ VT_NEMOTRON_H_ARM_TRACE, *NOT* VT_NEMOTRON_H_DIAG. The latter downloads + # the carry and the residual per layer per step, so a timed leg under it + # measures the diagnostic. The arm trace is one fprintf of six resident + # counters per step, so the reachability evidence and the timing come from + # the SAME run rather than from two runs that might differ. + VT_NEMOTRON_H_DEVICE_MAMBA=$devmamba VT_NEMOTRON_H_MAMBA_DECODE_STEP=$flag \ + VT_NEMOTRON_H_ARM_TRACE=1 "$BUILD/examples/nemotron-h-gen" \ + --model "$CKPT" \ + --golden "$SRC/tests/parity/goldens/nemotron_35_lightning_greedy/oracle.json" \ + > "$log" 2>&1 & + local pid=$! + local waited=0 loaded=0 + while kill -0 "$pid" 2>/dev/null; do + if grep -q "engine loaded in" "$log" 2>/dev/null; then loaded=1; break; fi + sleep 1 + waited=$((waited + 1)) + if [ "$waited" -ge 5400 ]; then break; fi + done + if [ "$loaded" -ne 1 ]; then + echo "$label: the engine never reported a load in ${waited}s -- the decode window is UNKNOWN, so NO busy fraction is sampled" + fi + local t0=$(date +%s.%N) + ( while true; do nvidia-smi --query-gpu=utilization.gpu --format=csv,noheader,nounits; sleep 0.1; done ) \ + > "$RUN/util_$label.txt" 2>/dev/null & + local sampler=$! + wait "$pid"; local r=$? + local t1=$(date +%s.%N) + kill "$sampler" 2>/dev/null; wait "$sampler" 2>/dev/null + + echo "RC[a3 $label]=$r" + grep -E "STRICT|PASS|FAIL|DIVERGENCE|REFUSING|SHORT|mode=|engine loaded in" "$log" | tail -20 + + # ★ ON A DIVERGENCE, PRINT THE TOKENS. The driver already emits `got:` and + # `exp:` for any row that mismatches (nemotron_h_gen/main.cpp:385), and the + # verdict grep above throws them away -- so a `95/96 DIVERGENCE` arrived with + # no way to tell a wrong recurrent carry from a benign bf16 near-tie without + # spending a second lease. A gate that reports a failure it cannot triage has + # not finished reporting. Also print the per-prompt row lines, which carry + # which prompt was short and by how much. + if grep -q "DIVERGENCE" "$log"; then + echo "--- DIVERGENCE detail ($label): per-prompt rows, then got/exp ---" + grep -E "^\s*prompt |matched=|wall=" "$log" | tail -10 + grep -E "^\s+(got|exp):" "$log" | tail -20 + fi + + # ★ WHICH KERNELS THE DECODE STEPS ACTUALLY LAUNCHED. The two arms compute the + # same recurrence, so the tokens agree and cannot separate them; this is the + # only line that can. A counter that reads 0 on BOTH legs means the recorder + # is broken, not that no kernel ran -- test_nemotron_h_paged_forward's + # "recurrent arm recorder" case is the triage for that. + echo "--- decode-step arm counters ($label) ---" + # Match on the FIELD TEXT, not on a field INDEX. The line is + # `[NH-DIAG] ARM step T=1 nd=1 np=0 state_update_rows=...`, so nd is field 5 + # and an `$4` test silently selects NOTHING -- which prints as a clean empty + # result and reads exactly like "no decode steps ran". + grep -E 'ARM step .* nd=[1-9]' "$log" | head -3 + grep -E 'ARM step .* nd=[1-9]' "$log" | tail -1 + echo "ARM lines total: $(grep -c 'ARM step' "$log")" + echo "ARM lines with a DECODE row: $(grep -cE 'ARM step .* nd=[1-9]' "$log")" + echo "ARM lines with a PREFILL row: $(grep -cE 'ARM step .* np=[1-9]' "$log")" + # A count of ZERO decode lines is an instrument failure, not a result. + if [ "$(grep -cE 'ARM step .* nd=[1-9]' "$log")" -eq 0 ]; then + echo "$label: NO decode-step ARM lines were recorded -- the counters say NOTHING about this run, and that is not a pass" + fi + + python3 - "$RUN/util_$label.txt" "$label" "$ARCH" "$loaded" <<'PY' +import sys +vals = [int(x) for x in open(sys.argv[1]).read().split() if x.strip().isdigit()] +label, arch, loaded = sys.argv[2], sys.argv[3], sys.argv[4] +if loaded != "1": + print(f"{label}: busy fraction NOT REPORTED -- the decode window was never identified") +elif not vals: + print(f"{label}: NO SAMPLES -- the busy fraction is unmeasured, not 0") +else: + busy = sum(1 for v in vals if v > 0) + print(f"{label}: GPU busy in {busy} of {len(vals)} DECODE samples = {100.0*busy/len(vals):.2f}% " + f"(arch {arch}; compare ONLY against the other leg of this same binary)") +PY + python3 "$SRC/scripts/nemotron-h-a2q1-per-token.py" "$log" "$label" "$t0" "$t1" "$ARCH" + echo "reference-tier lines in $label: $(grep -c 'reference-tier' "$log")" +} + +step "7. A3 gate + decode window, SINGLE-STEP decode arm (the default, = vLLM)" +( while true; do free -m | awk '/^Mem:/{print $3}'; sleep 1; done ) > "$RUN/rss.txt" 2>/dev/null & +MEMPID=$! +# ★ THE CONFOUND THIS THIRD LEG EXISTS TO REMOVE. +# +# Legs 1 and 2 vary the mamba KERNEL (single-step vs chunk scan) and BOTH sit on +# top of A2-Q1's FP8 W8A8 device projections, because both run inside the +# `mamba_on_device` branch. `main` carries NO device mamba arm at all -- +# `NemotronHMamba2MixerDevice`, `MambaIsFp8` and even `VT_NEMOTRON_H_DEVICE_MAMBA` +# are absent there -- so a `95/96` seen on BOTH of legs 1 and 2 is equally +# consistent with: +# +# (a) the divergence being the HOST's (arch-specific), or +# (b) the divergence being A2-Q1's FP8 PROJECTIONS, which neither leg turns off. +# +# Leg 3 discriminates by routing the whole mamba block back to the host +# reference on the SAME binary and box. Its counters are the check that it +# really took that path: the four kernel counters must read 0 while +# gathers/scatters stay non-zero, because the host branch still gathers and +# scatters but never enters the instrumented device mixer. +# +# leg 3 `96/96` => the divergence is A2-Q1's FP8 arm, NOT the architecture. +# leg 3 `95/96` => the host owns it and the arch-specific framing stands. +run_leg on 1 + +step "8. the same binary with the decode arm OFF -- the chunk scan on decode rows" +run_leg off 0 + +step "8b. THE DISCRIMINATOR -- device mamba arm OFF, so A2-Q1's FP8 projections are OUT" +run_leg hostmamba 1 0 + +kill "$MEMPID" 2>/dev/null +echo "peak host MiB used during the run: $(sort -n "$RUN/rss.txt" | tail -1)" + +step "8c. the three-way verdict" +for L in on off hostmamba; do + printf '%-10s %s\n' "$L" "$(grep -E 'TOKEN MATCH' "$RUN/a3_$L.log" 2>/dev/null | tail -1)" +done +echo "leg 3 (hostmamba) MUST show the four kernel counters at 0 with gathers/scatters non-zero;" +echo "anything else means it did NOT take the host path and the discrimination is VOID." +grep -E 'ARM step .* nd=[1-9]' "$RUN/a3_hostmamba.log" 2>/dev/null | tail -1 + +step "9. contention, after" +nvidia-smi --query-compute-apps=pid,used_memory --format=csv +echo "ALL LOGS: $RUN" diff --git a/scripts/nemotron-h-a2q1-dgx-gate.sh b/scripts/nemotron-h-a2q1-dgx-gate.sh new file mode 100755 index 000000000..9a74c24a8 --- /dev/null +++ b/scripts/nemotron-h-a2q1-dgx-gate.sh @@ -0,0 +1,231 @@ +#!/usr/bin/env bash +# A2-Q1 (#810, .agents/specs/nemotron-h-a2q1-fp8-mamba.md) — the GB10 gate for +# the FP8 W8A8 Mamba2 device arm, recorded as a script so the recipe is +# reproducible rather than retyped (AGENTS.md §Gates: "Record the exact build and +# run recipe"). +# +# RUN IT INSIDE A LEASE, NEVER OVER ssh: +# rc run -d dgx:gpu0 --max-runtime 6h -- bash -lc \ +# 'git clone --depth 50 -b row/A2-Q1 https://github.com/mudler/vllm.cpp /root/src \ +# && bash /root/src/scripts/nemotron-h-a2q1-dgx-gate.sh' +# +# THREE ENVIRONMENT FACTS THIS SCRIPT ENCODES RATHER THAN REDISCOVERS: +# * the CUDA lane on this box is `sbsa`, not `arm64`, and an UNANCHORED +# `cuda-toolkit-13*` match selects `cuda-toolkit-13-config-common`, which +# installs cleanly, ships no compiler and returns 0; +# * `nvcc --version` is NOT a sufficient postcondition — a partial CUDA install +# printed every feature line and then failed to link `CUDA::cublasLt`, so the +# toolkit is proved by an actual link; +# * `/workspace` is CIFS and refuses symlinks, so the build lives on local disk +# and only the log is copied out. +set -u -o pipefail + +LOG_ROOT=${LOG_ROOT:-/workspace/a2q1} +CKPT=${CKPT:-/workspace/a3/ckpt-stage} +SRC=${SRC:-/root/src} +# 121a is the GB10; 110 is Thor. The spec gates BOTH hosts, and the FP8 W8A8 arm +# on Thor is the whole point of #960/#991, so the arch is a parameter. +ARCH=${ARCH:-121a} +BUILD=${BUILD:-/root/build-cuda} +CUTLASS=${CUTLASS:-/root/cutlass} +STAMP=$(date -u +%Y%m%dT%H%M%SZ) +RUN=$LOG_ROOT/$STAMP +mkdir -p "$RUN" +echo "A2Q1 run dir: $RUN" + +step() { echo; echo "=== $* ==="; } +# Every gate command runs BARE and echoes its own status. Never pipe a command +# whose exit code matters -- a pipeline reports the LAST stage, which is how a +# failing `mount` once reported rc=0. +rc() { "$@"; local r=$?; echo "RC[$*]=$r"; return $r; } + +step "0. the box, before anything is installed" +rc uname -m +rc id -u +rc nvidia-smi --query-gpu=name,compute_cap,memory.total --format=csv +rc df -h /root /workspace +rc free -m + +step "1. contention -- a timing number measured beside another job is VOID" +nvidia-smi --query-compute-apps=pid,used_memory --format=csv | tee "$RUN/contention.txt" +echo "RC[nvidia-smi compute-apps]=${PIPESTATUS[0]}" + +step "2. toolchain" +export DEBIAN_FRONTEND=noninteractive +rc apt-get update -qq +rc apt-get install -y -qq git cmake ninja-build g++ curl ca-certificates python3 python3-dev +if ! command -v nvcc >/dev/null 2>&1; then + # ANCHORED package names on the sbsa lane. `cuda-nvcc-13-0` is the compiler; + # the libraries this build links are named individually so a metapackage that + # ships nothing cannot satisfy the check. + curl -fsSL -o /tmp/cuda-keyring.deb \ + https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2404/sbsa/cuda-keyring_1.1-1_all.deb + echo "RC[curl keyring]=$?" + rc dpkg -i /tmp/cuda-keyring.deb + rc apt-get update -qq + rc apt-get install -y -qq cuda-nvcc-13-0 cuda-cudart-dev-13-0 libcublas-dev-13-0 \ + cuda-nvrtc-dev-13-0 cuda-nvtx-13-0 cuda-profiler-api-13-0 libcurand-dev-13-0 +fi +export PATH=/usr/local/cuda/bin:$PATH +rc nvcc --version + +step "2b. the toolkit is proved by a LINK, not by --version" +cat > /tmp/probe.cu <<'EOF' +#include +#include +int main() { + cublasLtHandle_t h = nullptr; + const auto s = cublasLtCreate(&h); + std::printf("cublasLtCreate=%d\n", static_cast(s)); + return s == CUBLAS_STATUS_SUCCESS ? 0 : 1; +} +EOF +rc nvcc -arch=sm_$ARCH /tmp/probe.cu -o /tmp/probe -lcublasLt +TOOLCHAIN_OK=$? +if [ "$TOOLCHAIN_OK" -ne 0 ]; then + echo "VOID: the CUDA toolkit does not link cublasLt; every number below would be a lie" + exit 2 +fi +rc /tmp/probe + +step "3. cutlass (the fp8/fp4 fast-path cells resolve from it)" +if [ ! -f "$CUTLASS/include/cutlass/cutlass.h" ]; then + rc git clone --depth 1 --branch v4.5.0 https://github.com/NVIDIA/cutlass "$CUTLASS" +fi +rc test -f "$CUTLASS/include/cutlass/cutlass.h" + +step "4. configure" +rc cmake -S "$SRC" -B "$BUILD" -G Ninja \ + -DCMAKE_BUILD_TYPE=Release \ + -DVLLM_CPP_CUDA=ON \ + -DVLLM_CPP_CUDA_ARCHITECTURES="$ARCH" \ + -DVLLM_CPP_CUTLASS_DIR="$CUTLASS" 2>&1 | tee "$RUN/configure.log" +echo "RC[cmake configure]=${PIPESTATUS[0]}" + +step "4b. the feature table -- a DISABLED or wrong-arch cell VOIDS the run" +grep -E "CUDA feature .*(ENABLED|DISABLED)" "$RUN/configure.log" | tee "$RUN/features.txt" +BAD=$(grep -cE "CUDA feature .*DISABLED" "$RUN/features.txt") +ENABLED=$(grep -cE "ENABLED for \[$ARCH\]" "$RUN/features.txt") +echo "feature cells ENABLED for [$ARCH]: $ENABLED ; DISABLED cells: $BAD" + +step "5. build (-j 4: unconstrained parallelism has OOM-REBOOTED this box)" +rc cmake --build "$BUILD" -j 4 2>&1 | tail -40 | tee "$RUN/build.tail" +echo "RC[cmake build]=${PIPESTATUS[0]}" + +step "6. the focused device gate -- A2-Q1's own cases" +"$BUILD/tests/test_nemotron_h_mamba_device" -s 2>&1 | tee "$RUN/mamba_device.log" +echo "RC[test_nemotron_h_mamba_device]=${PIPESTATUS[0]}" +grep -E "test cases:|assertions:|Status:" "$RUN/mamba_device.log" + +step "6b. the neighbouring suites the arm can break" +for t in test_nemotron_h_forward test_nemotron_h_paged_forward test_nemotron_h_loader \ + test_nemotron_h_moe_device test_ops_mamba2_ssd test_ops_fp8_cpu; do + if [ -x "$BUILD/tests/$t" ]; then + "$BUILD/tests/$t" > "$RUN/$t.log" 2>&1 + echo "RC[$t]=$?" + grep -E "test cases:|assertions:|Status:" "$RUN/$t.log" + else + echo "MISSING BINARY: $t" + fi +done + +step "7. the A3 gate + the GPU busy fraction, device mamba ON" +# The acceptance test of this unit is NOT a ratio: the GPU busy fraction must +# RISE from the 6.31% baseline, so it is sampled on a loop for the whole decode +# and the SAMPLE COUNT is reported beside it. A fraction with no denominator is +# not a measurement. +run_gate() { # $1 = label, $2 = VT_NEMOTRON_H_DEVICE_MAMBA value + local label=$1 flag=$2 + local log="$RUN/a3_$label.log" + : > "$log" + + # ★ THE SAMPLING WINDOW IS THE DECODE, NOT THE WHOLE PROCESS. + # + # The first Thor run sampled from process start, so the window contained the + # multi-minute 20.1 GiB engine load, which is GPU-IDLE. That dilutes both arms + # toward each other and toward zero: it read 15.33% with the arm on and 14.73% + # with it off, a 0.59-point difference that says almost nothing about what the + # DECODE does, and it is not comparable to the 6.31% figure either. Mixing a + # load phase into a decode measurement is the same defect as summing prefill + # and decode in one profile. + # + # So the driver starts FIRST, and the sampler starts only once the driver has + # printed `engine loaded in Ns`. Everything after that line is decode. + VT_NEMOTRON_H_DEVICE_MAMBA=$flag "$BUILD/examples/nemotron-h-gen" \ + --model "$CKPT" \ + --golden "$SRC/tests/parity/goldens/nemotron_35_lightning_greedy/oracle.json" \ + > "$log" 2>&1 & + local pid=$! + local waited=0 + local loaded=0 + while kill -0 "$pid" 2>/dev/null; do + if grep -q "engine loaded in" "$log" 2>/dev/null; then loaded=1; break; fi + sleep 1 + waited=$((waited + 1)) + if [ "$waited" -ge 5400 ]; then break; fi + done + if [ "$loaded" -ne 1 ]; then + # Say so rather than sampling a window whose meaning is unknown. A fraction + # over the wrong window is worse than no fraction, because it still formats + # like a measurement. + echo "$label: the engine never reported a load in ${waited}s -- the decode window is UNKNOWN, so NO busy fraction is sampled" + fi + local t0=$(date +%s.%N) + ( while true; do nvidia-smi --query-gpu=utilization.gpu --format=csv,noheader,nounits; sleep 0.1; done ) \ + > "$RUN/util_$label.txt" 2>/dev/null & + local sampler=$! + wait "$pid" + local r=$? + local t1=$(date +%s.%N) + kill "$sampler" 2>/dev/null + wait "$sampler" 2>/dev/null + + echo "RC[a3 $label]=$r" + grep -E "STRICT|PASS|FAIL|DIVERGENCE|REFUSING|SHORT|mode=|engine loaded in" "$log" | tail -20 + # The busy fraction WITH its denominator, over the DECODE window alone. The + # 6.31% baseline is a GB10 number, so it is quoted ONLY on GB10: printing it + # beside a Thor figure invites a cross-silicon comparison that answers a + # different question, and a number quoted often enough becomes treated as + # measured. + python3 - "$RUN/util_$label.txt" "$label" "$ARCH" "$loaded" <<'PY' +import sys +vals = [int(x) for x in open(sys.argv[1]).read().split() if x.strip().isdigit()] +label, arch, loaded = sys.argv[2], sys.argv[3], sys.argv[4] +if loaded != "1": + print(f"{label}: busy fraction NOT REPORTED -- the decode window was never identified") +elif not vals: + print(f"{label}: NO SAMPLES -- the busy fraction is unmeasured, not 0") +else: + busy = sum(1 for v in vals if v > 0) + pct = 100.0 * busy / len(vals) + note = ("baseline 6.31% on this same GB10 workload" if arch == "121a" + else f"NO baseline for arch {arch}: the 6.31% figure is GB10's, so it " + f"is not comparable here; use the ON/OFF A/B on this box") + print(f"{label}: GPU busy in {busy} of {len(vals)} DECODE samples = {pct:.2f}% busy ({note})") +PY + # The per-output-token time over the SAME decode window, with the terms shown. + python3 "$SRC/scripts/nemotron-h-a2q1-per-token.py" "$log" "$label" "$t0" "$t1" "$ARCH" + # The reference tier is numerically CORRECT, so a pass obtained on it is + # invisible in the numbers and only this line separates them (spec R2). + echo "reference-tier lines in $label: $(grep -c 'reference-tier' "$log")" +} + +if [ ! -d "$CKPT" ]; then + echo "NO CHECKPOINT AT $CKPT -- the A3 gate and the GPU busy fraction are NOT MEASURED on this host." + echo "That is a stated absence, not a pass. Steps 7-9 are skipped." + echo "ALL LOGS: $RUN" + exit 0 +fi + +( while true; do free -m | awk '/^Mem:/{print $3}'; sleep 1; done ) > "$RUN/rss.txt" 2>/dev/null & +MEMPID=$! +run_gate on 1 + +step "8. the same binary with the arm OFF -- the A/B this unit is measured by" +run_gate off 0 +kill "$MEMPID" 2>/dev/null +echo "peak host MiB used during the run: $(sort -n "$RUN/rss.txt" | tail -1)" + +step "9. contention, after" +nvidia-smi --query-compute-apps=pid,used_memory --format=csv +echo "ALL LOGS: $RUN" diff --git a/scripts/nemotron-h-a2q1-per-token.py b/scripts/nemotron-h-a2q1-per-token.py new file mode 100755 index 000000000..56baba5b0 --- /dev/null +++ b/scripts/nemotron-h-a2q1-per-token.py @@ -0,0 +1,91 @@ +#!/usr/bin/env python3 +"""Derive the per-output-token time of one A2-Q1 A3 gate run, WITH its terms. + +`examples/nemotron_h_gen` reports neither a rate nor a duration: it prints the +engine load time and a `TOKEN MATCH: m/n` line and nothing else. So the per-token +number this row is measured on has to be derived, and a derived number that hides +its terms is how a rate over an unknown denominator comes to be quoted as if it +were measured. + +THE WINDOW IS THE DECODE, AND THE CALLER BRACKETS IT. `t0`/`t1` are taken AFTER +the driver prints `engine loaded in Ns`, so the 20.1 GiB load is already outside +them and nothing is subtracted here. An earlier version bracketed the whole +process and subtracted the load, which put a multi-minute GPU-idle phase inside +the same window as the decode -- the same defect as summing prefill and decode +into one profile. The load is still read out of the log and printed, because it +is context for the number, not a term in it. + +This prints the window, the load it excludes and the token count it divides by on +separate lines, and it REFUSES rather than printing 0 when a term is missing or +the window is not positive. + +THE vLLM RATIO IS QUOTED ONLY ON THE SILICON IT WAS MEASURED ON. 0.014369 s is a +GB10 figure. Printing it beside a Thor per-token number invites a comparison +across two different pieces of silicon, and the first Thor run did exactly that +-- it printed `ratio 54.7x` for a number that was never measured against vLLM on +that box. This is the same defect the busy-fraction reporter carried, and fixing +one surface while leaving its twin is how a wrong comparison survives. + + nemotron-h-a2q1-per-token.py