diff --git a/.agents/issue-index.md b/.agents/issue-index.md index fe4b5e576..bfc3aae6e 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -457,3 +457,4 @@ 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 | 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-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..cdc43d771 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -12,6 +12,8 @@ | **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** | +| **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) over **844** OK, **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 | | **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..ee83c81ea 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -244,6 +244,7 @@ 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_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..141d4e8ed 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 gates PASS 96/96 `STRICT PASS`, the latter with the device FP8 mamba arm; sm_121a re-run pending** | ABI-only driver; G-SAFE `num_reqs <= 1`. The 23 FP8 Mamba2 blocks reach the device (A2-Q1, [#810](https://github.com/mudler/vllm.cpp/issues/810)); `lm_head` is 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-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