From 0f336401f39c928f65207d15c67668a4f6ad0999 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 22:29:42 +0000 Subject: [PATCH 01/12] feat(#810 A2-Q1): NemotronH's 23 Mamba2 blocks reach the device on the FP8 W8A8 seam (#517) The 23 Mamba2 layers were the last host bounce of a decode step, and they were the expensive one. The host arm reaches its two projections through `Linear(..., const NemotronHOwned&)`, whose `DenseFor` calls `NemotronHOwned::DenseBf16()` -- a FULL dequant of the fp8 tower into a fresh bf16 buffer ON EVERY CALL. That is 23 x (10304x2688 + 2688x4096) = 890e6 elements re-expanded per token, plus a download of the normed hidden and an upload of the mixer output per layer. The block moves whole or not at all: `mixer.in_proj` produces the fused `zxbcdt` that the causal conv and the SSD scan both consume (mamba_mixer2.py:550, split :692-696), so there is no intermediate landing in which the conv is on the device and `in_proj` is not. That is why the shared FP8 W8A8 linear seam had to be extracted first (#940, `dense_fp8_gemm.h`), and this change routes through it rather than re-typing its entry points. `NemotronHMamba2MixerDevice` mirrors the host arm statement for statement -- same vt:: ops, same order, same dtypes, different backend, which is the property A2-R established and the numeric gate reads: in_proj (FP8 W8A8) -> QkvSplit(z | xBC | dt) -> CausalConv1dFwd(silu) -> QkvSplit(x | B | C) -> Mamba2ChunkScan -> RmsNormGatedGroup(n_groups) -> out_proj (FP8 W8A8) The one substitution is the split: the host arm copies columns with `SliceCols` because `vt::Mamba2ChunkScan` validates every operand contiguous, and `vt::QkvSplit` is exactly that copy on the device. Both forwards select it at runtime, never by preprocessor: `MambaIsFp8` names only `NemotronHWeightForm` and `vt::OpRegistered` is the op table's own answer, so a dense NemotronH or a device without the fp8 pair keeps the host bounce. The paged arm hands the mixer the recurrent rows `vt::GdnStateGather` already gathered and zeroed and it advances them IN PLACE, so the two downloads, the host mixer and the two uploads per layer are gone; A2-P's gather/scatter pair and its `has_initial = true` reasoning are unchanged. The `ssm_dtype == f32` term in the paged selection is not decoration. The gather widens the page into an f32 working buffer by op contract and the host arm then narrows it back to `ssm_dtype`, so 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 device residency is built on first use and keyed on a `ResidentSlot` the weights own, never on an address (#237): across two engine builds in one process an address-keyed cache returns the previous engine's device pointers, which is plausible wrong values rather than a crash. The e4m3 staging copy is released as soon as `ResidentFp8` has uploaded it, so the conversion peaks at one projection rather than at the whole 890 MB tower. `dense_fp8::ResidentFp8` still does not account its own upload -- that is #974, and this change does not touch the shared header -- so A2-Q1 accounts what IT uploads at the site that causes it, through the same `load_stats::AddDeviceUpload` `ResidentWeight` and `ResidentNvfp4` call. That counter is also the instrument the new residency case reads: an arm that re-uploaded the tower every step returns identical numbers to one that uploads it once. G-SAFE is untouched. All three clauses of the interlock stand, this change creates no paging, no carried state and no batching that A2-P did not already create, and `num_reqs <= 1` remains A2-B's to remove. `tests/vllm/models/test_nemotron_h_mamba_device.cpp` is the cheap arm in front of the real-checkpoint gate, for the same reason A2-Q2a's device MoE file exists: `BuildTiny` is all-`kDense`, so nothing in the existing device suite could reach the FP8 W8A8 arm at all. The two arms are NOT bit-comparable by construction -- the host reference is W8A16 and the device arm is W8A8 as vLLM is -- so every band is measured in the run against a defect the fixture separates, and the element count is asserted against the geometry so a maximum over zero elements cannot read as agreement. `scripts/nemotron-h-a2q1-dgx-gate.sh` records the GB10 recipe rather than leaving it to be retyped, including the three environment facts that have each voided a run before: the CUDA lane is `sbsa`, `nvcc --version` is not a sufficient postcondition, and `/workspace` is CIFS and refuses symlinks. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/model-matrix.md | 2 +- .agents/specs/nemotron-h-a2q1-fp8-mamba.md | 73 ++- docs/BENCHMARKS.md | 1 + docs/ENVIRONMENT.md | 1 + docs/FEATURES.md | 2 +- docs/STATUS.md | 2 +- docs/USAGE.md | 2 +- scripts/nemotron-h-a2q1-dgx-gate.sh | 186 ++++++ .../models/nemotron_h_device.cpp | 584 +++++++++++++++++- .../models/nemotron_h_forward.h | 37 ++ tests/CMakeLists.txt | 15 + .../models/test_nemotron_h_mamba_device.cpp | 439 +++++++++++++ 12 files changed, 1321 insertions(+), 23 deletions(-) create mode 100755 scripts/nemotron-h-a2q1-dgx-gate.sh create mode 100644 tests/vllm/models/test_nemotron_h_mamba_device.cpp diff --git a/.agents/model-matrix.md b/.agents/model-matrix.md index dc01be5be..aa5642853 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..d9f1dd932 100644 --- a/.agents/specs/nemotron-h-a2q1-fp8-mamba.md +++ b/.agents/specs/nemotron-h-a2q1-fp8-mamba.md @@ -418,23 +418,68 @@ 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. + +**What is still owed, and it is the acceptance test rather than a formality:** +the §5.1 per-block numeric gate on the real checkpoint on both hosts, the A3 +`96/96 mode=decode STRICT PASS` re-run, the §5.3 mutations, and the GPU busy +fraction — which must RISE from the measured 6.31% baseline, sampled with its +denominator reported. `scripts/nemotron-h-a2q1-dgx-gate.sh` is the recipe. ## 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 7f965fc7e..fa59e9d2d 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -11,6 +11,7 @@ | **Container images (ACTIVE; arm64 cuda verified on GB10 + Orin 2026-08-11)** | `ENG-RELEASE-CONTAINERS` ([spec](../.agents/specs/container-images.md)) | cpu amd64 783 MB; cuda arm64 **1.71 GB**. GB10 `sm_121a`: `/health`+`/version`+SIGTERM on `--gpus all`. Orin `sm_87` (Tegra): Qwen3-0.6B **generates**, GPU **GR3D 95-97%** | n/a | | **Developer/row protocol** | Contribution entry point; `ENG-NOW-DERIVED` #374 @`dbd0d51c` | Entry-point gates retained. #374 W1-W5 DONE; benchmark/runtime/parity `VOID`; row specs now carry `## Now` | n/a | | **NemotronH paged forward** (`MODEL-NEMOTRON-H-ABI-A2P`, [#810](https://github.com/mudler/vllm.cpp/issues/810)) | **No 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 | +| **NemotronH FP8 mamba on device** (`A2-Q1`, [#810](https://github.com/mudler/vllm.cpp/issues/810)) | GB10 sm_121a, released 20.1 GiB checkpoint, c1 greedy | **PENDING a lease.** The unit's acceptance test is the GPU busy fraction rising from the 6.31% baseline, not a ratio, and it is unmeasured here | Host arm dequantized 890e6 fp8 elements per token; `VT_NEMOTRON_H_DEVICE_MAMBA=0` is the same-binary A/B | | **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 9a046a657..7542f4f1d 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -242,6 +242,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 63f73c228..beda5a523 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 c5c461c39..78b15d9e3 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -131,7 +131,7 @@ token-for-token correctness against the pinned oracle. | GLM-4 dense (sandwich norms, partial rope) | Correctness-complete, speed-pending | Token-exact 16/16 (GLM-4-9B-0414); first GLM-family model; partial interleaved RoPE + Gemma2 sandwich norms + biased qkv | | GLM-4.7-Flash (MLA + GLM MoE) | Correctness-complete, speed-pending | Token-exact 8/8 (GLM-4.7-Flash, 31.2B); reuses the DeepSeek-V2 MLA stack; first e2e coverage of the q_lora query branch + noaux_tc sigmoid router with routed-scaling | | Kimi-Linear-48B-A3B (KDA + NoPE-MLA + MoE hybrid) | **RUNNER FOLD LANDS (ROW 7 §21, #122): engine==CLI 128/128 byte-identical; golden 122/128 (near-tie profile); FA2 MLA default-ON; `vllm_complete_tokens` (ABI v13).** Grouped-router top-k block-parallel (byte-identical); no binding speed number: ckpt is tiktoken-only, so no warm-server harness. STRICT stays CLOSED. Server 19.0 tok/s wall (~0.90× vLLM floor) = speed open | paged suite 8/8·206; SACRED post-fold 35B 315/315 + 27B 235/235; thin ABI client (ratchet 8) | -| Nemotron-3.5-Lightning-30B-A3B (Mamba2 + GQA + relu2 MoE) | **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 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 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 b1fa1f141..a9f037edb 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -725,7 +725,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/nemotron-h-a2q1-dgx-gate.sh b/scripts/nemotron-h-a2q1-dgx-gate.sh new file mode 100755 index 000000000..19ca56dda --- /dev/null +++ b/scripts/nemotron-h-a2q1-dgx-gate.sh @@ -0,0 +1,186 @@ +#!/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 + ( 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=$! + local t0=$(date +%s.%N) + VT_NEMOTRON_H_DEVICE_MAMBA=$flag "$BUILD/examples/nemotron-h-gen" \ + --model "$CKPT" \ + --golden "$SRC/tests/parity/goldens/nemotron_35_lightning_greedy/oracle.json" \ + > "$RUN/a3_$label.log" 2>&1 + local r=$? + local t1=$(date +%s.%N) + kill "$sampler" 2>/dev/null + wait "$sampler" 2>/dev/null + echo "RC[a3 $label]=$r" + echo "wall seconds ($label): $(echo "$t1 - $t0" | bc)" + grep -E "STRICT|PASS|FAIL|mode=|tok/s|per output token" "$RUN/a3_$label.log" | tail -20 + python3 - "$RUN/util_$label.txt" "$label" <<'PY' +import sys +vals = [int(x) for x in open(sys.argv[1]).read().split() if x.strip().isdigit()] +if not vals: + print(f"{sys.argv[2]}: NO SAMPLES -- the busy fraction is unmeasured, not 0") +else: + busy = sum(1 for v in vals if v > 0) + print(f"{sys.argv[2]}: GPU busy in {busy} of {len(vals)} samples = " + f"{100.0*busy/len(vals):.2f}% busy (baseline 6.31%)") +PY + grep -c "reference-tier" "$RUN/a3_$label.log" > /dev/null 2>&1 + echo "reference-tier lines in $label: $(grep -c 'reference-tier' "$RUN/a3_$label.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/src/vllm/model_executor/models/nemotron_h_device.cpp b/src/vllm/model_executor/models/nemotron_h_device.cpp index 4455998d9..d1cc491d1 100644 --- a/src/vllm/model_executor/models/nemotron_h_device.cpp +++ b/src/vllm/model_executor/models/nemotron_h_device.cpp @@ -104,6 +104,13 @@ // header rather than dense_device_glue.h — `ResidentWeight`, the lazy // upload-once seam this row converts NemotronH's dense weights onto. #include "vllm/model_executor/models/dense_attn_block.h" +// A2-Q1 (#810): the SHARED per-tensor FP8 W8A8 dense GEMM glue extracted by +// #940 — `ResidentFp8` (the upload-once device view) and `MatmulFp8CutlassD` +// (static per-tensor activation quant + fp8 GEMM with the folded alpha). This +// header exists BECAUSE of this model (its own note says so), so the mamba +// projections route through it rather than re-typing the entry points here — +// the hand-rolled parallel path AGENTS.md forbids. +#include "vllm/model_executor/models/dense_fp8_gemm.h" #include "vt/backend.h" #include "vt/ops.h" #include "vt/recipes.h" // kFusedAddRmsNormStd @@ -355,9 +362,6 @@ bool MoeIsNvfp4(const NemotronHMoeWeights& w) { return true; } -// DSR-ALLOW(A2-Q2a): TYPES, not behaviour -- the whole arena region names vt::cuda::Marlin* functions declared only in the guarded header above, so it cannot compile on a build without them. vt::OpRegistered answers availability, never declaration. Mirrors laguna.cpp:456, the same CUDA-leg arena. -#ifdef VT_MARLIN_NVFP4 - // Fetch (building on first use) the resident state a weight owns. A VERBATIM // copy of qwen3_5.cpp:680, which is file-local `static` there and so cannot be // called from here. The duplication is deliberate and is the same call this file @@ -374,6 +378,9 @@ R& ResidentIn(const ResidentSlot& slot) { return *static_cast(slot.state.get()); } +// DSR-ALLOW(A2-Q2a): TYPES, not behaviour -- the whole arena region names vt::cuda::Marlin* functions declared only in the guarded header above, so it cannot compile on a build without them. vt::OpRegistered answers availability, never declaration. Mirrors laguna.cpp:456, the same CUDA-leg arena. +#ifdef VT_MARLIN_NVFP4 + // ─── A2-Q2a: one MoE layer's device-resident Marlin arena ─────────────────── // // THE SHAPE IS AN ARENA, NOT A POINTER ARRAY, and that distinction is the whole @@ -763,6 +770,409 @@ DBuf NemotronHMoeBlockDevice(Dev d, const NemotronHMoeWeights& w, #endif // VT_MARLIN_NVFP4 +// ─── A2-Q1: the FP8 W8A8 Mamba2 block on the DEVICE ───────────────────────── +// +// The 23 Mamba2 layers were the last host bounce of the decode step, and they +// were the EXPENSIVE one. The host arm reaches its two projections through +// `Linear(..., const NemotronHOwned&)` (nemotron_h.cpp:291), whose `DenseFor` +// (:245) calls `NemotronHOwned::DenseBf16()` — a FULL dequant of the fp8 tower +// into a fresh bf16 buffer ON EVERY CALL. That is 23 x (10304x2688 + 2688x4096) +// = 890e6 elements re-expanded per TOKEN, plus a download of the normed hidden +// and an upload of the mixer output per layer. It is why a decode step spent +// its time on the host. +// +// ★ THE BLOCK MOVES WHOLE, ON THE FP8 SEAM, OR NOT AT ALL. `mixer.in_proj` +// produces the fused `zxbcdt` that the causal conv and the SSD scan both consume +// (mamba_mixer2.py:550, split :692-696), so there is no intermediate landing in +// which the conv is on the device and `in_proj` is not. That is the whole reason +// the shared FP8 W8A8 linear seam (#940, `dense_fp8_gemm.h`) had to be extracted +// first: without it this block has no device path to build. +// +// WHAT IS MIRRORED, statement for statement, from the host arm +// (nemotron_h.cpp:451-625) — same vt:: ops, same order, same dtypes, different +// backend, which is the property A2-R established and the numeric gate reads: +// +// in_proj (FP8 W8A8) -> QkvSplit(z | xBC | dt) -> CausalConv1dFwd(silu) +// -> QkvSplit(x | B | C) -> Mamba2ChunkScan -> RmsNormGatedGroup(n_groups) +// -> out_proj (FP8 W8A8) +// +// The ONE substitution is the split: the host arm copies columns with +// `SliceCols` (nemotron_h.cpp:333) because `vt::Mamba2ChunkScan` validates every +// operand contiguous (ops.cpp CheckMamba2Operand). `vt::QkvSplit` is exactly +// that copy on the device — a three-way column split with INDEPENDENT widths +// into three contiguous outputs (ops.h:3512) — so the two arms produce the same +// bytes and neither hands the scan a strided view. + +// Every projection of one Mamba2 block on the FP8 W8A8 seam, plus the six small +// recurrence parameters, uploaded ONCE and held for the model lifetime. Built on +// first device use and keyed on the weights' own `ResidentSlot`, never on an +// address (nemotron_h_forward.h, `device_fp8`). +struct NemotronHMambaDeviceResident { + Fp8Weight in_proj; // [in_proj_out_features, hidden_size] e4m3 + Fp8Weight out_proj; // [hidden_size, mamba_intermediate_size] e4m3 + // The device buffers this arm owns for the model lifetime. Raw `Backend::Alloc` + // with a freeing deleter rather than a pooled `DBuf`: a `DBuf` held forever + // would take its block out of the shared scratch pool for good, which is the + // opposite of what the pool is for. + std::shared_ptr conv_w; // act dtype [conv_dim, conv_kernel] + std::shared_ptr conv_b; // act dtype [conv_dim] (use_conv_bias only) + std::shared_ptr a_neg; // f32 [num_heads] = -exp(A_log) + std::shared_ptr d_term; // f32 [num_heads] + std::shared_ptr dt_bias; // f32 [num_heads] + std::shared_ptr norm_w; // act dtype [mamba_intermediate_size] + bool has_conv_bias = false; + bool ready = false; +}; + +// Refuse by name unless BOTH projections are the FP8 W8A8 static form this arm +// is built from. `BuildTiny` and any future unquantized NemotronH ship them +// `kDense`, and those layers keep the host bounce — stated here rather than +// discovered later as a silent slow path. +bool MambaIsFp8(const NemotronHMambaWeights& w) { + auto q = [](const NemotronHOwned& t) { + return t.form == NemotronHWeightForm::kFp8W8A8Static && !t.bytes.empty() && + t.shape.size() == 2; + }; + return q(w.in_proj) && q(w.out_proj); +} + +// VT_NEMOTRON_H_DEVICE_MAMBA, default ON. The same-binary A/B switch every +// measurement of this row needs: with it OFF the identical build takes the host +// bounce, so a throughput or GPU-occupancy difference is attributable to THIS +// arm and not to a rebuild. Mirrors `NemotronHDeviceMoeEnabled`, which A2-Q2a +// introduced for its own A/B. +bool NemotronHDeviceMambaEnabled() { + static const bool on = [] { + const char* e = std::getenv("VT_NEMOTRON_H_DEVICE_MAMBA"); + return !(e != nullptr && e[0] == '0'); + }(); + return on; +} + +// A lifetime-resident device copy of `nbytes` host bytes. Accounted through +// `load_stats::AddDeviceUpload` AT THE SITE THAT CAUSES THE UPLOAD, which is +// what `ResidentWeight` (dense_attn_block.h:197) and `ResidentNvfp4` +// (dense_nvfp4_gemm.h:306) already do. `dense_fp8::ResidentFp8` does NOT account +// its own upload — that is issue #974, and A2-Q1 does not fix it inside the +// shared header; it accounts what IT uploads, here, so this row's device-upload +// total is honest and every other caller of the seam is byte-unchanged. +std::shared_ptr ResidentBytes(Dev d, const void* src, size_t nbytes) { + VT_CHECK(nbytes > 0, "NemotronH device mamba: refusing a zero-byte residency"); + void* p = d.b.Alloc(nbytes); + d.b.Copy(d.q, p, src, nbytes); + vllm::load_stats::AddDeviceUpload(nbytes); + Backend* bk = &d.b; + return std::shared_ptr(p, [bk](void* q) { bk->Free(q); }); +} + +// The same three properties `RequireWeight` checks on the host arm +// (nemotron_h.cpp:180), so a defect that refuses there refuses here rather than +// reaching a kernel with a null pointer or a transposed extent. +std::shared_ptr ResidentDense(Dev d, const NemotronHOwned& w, const char* what, + DType want, const std::vector& shape) { + VT_CHECK(!w.Empty(), std::string("NemotronH device mamba: weight '") + what + + "' is not materialized"); + VT_CHECK(w.IsDense(), std::string("NemotronH device mamba: weight '") + what + + "' is not dense; only the two projections are quantized"); + VT_CHECK(w.dtype == want, std::string("NemotronH device mamba: weight '") + what + + "' has the wrong dtype for this arm"); + VT_CHECK(w.shape == shape, std::string("NemotronH device mamba: weight '") + what + + "' has the wrong shape"); + return ResidentBytes(d, w.bytes.data(), w.bytes.size()); +} + +// One `NemotronHOwned` FP8 W8A8 projection as the SHARED `Fp8Weight` the seam +// consumes. `packed` is a staging copy of the e4m3 bytes and is RELEASED as soon +// as `ResidentFp8` has uploaded them (see the build below), so the peak cost of +// the conversion is one projection rather than the whole 890 MB tower. +Fp8Weight MambaFp8Weight(const NemotronHOwned& t, const char* what) { + VT_CHECK(t.form == NemotronHWeightForm::kFp8W8A8Static, + std::string("NemotronH device mamba: '") + what + + "' is not the FP8 W8A8 static form this arm is built from"); + VT_CHECK(t.shape.size() == 2, std::string("NemotronH device mamba: '") + what + + "' is not a rank-2 [out, in] projection"); + Fp8Weight f; + f.n = t.shape[0]; + f.k = t.shape[1]; + VT_CHECK(t.bytes.size() == static_cast(f.n) * static_cast(f.k), + std::string("NemotronH device mamba: '") + what + + "' does not carry exactly one e4m3 byte per [out, in] element"); + // ★ `Fp8Weight` HAS NO `has_input_scale` (qwen3_5_weights.h:318-325): a 0.0 + // `input_scale` IS "no scale shipped" there, while `NemotronHOwned` carries the + // distinction explicitly (nemotron_h_forward.h:174). The A2-Q1 spec §4.1 says + // the distinction survives the conversion because the default is 0.0 — ASSERT + // that mapping rather than assume it. A shipped 1.0 is a real value and passes; + // an unshipped scale reaching the GEMM as 1.0 would quantize the activation + // against the wrong divisor and produce a finite, plausible, wrong answer. + VT_CHECK(t.has_input_scale, + std::string("NemotronH device mamba: '") + what + + "' ships no input_scale, so the static per-tensor activation quant " + "this arm performs has no divisor; refusing rather than defaulting to 1.0"); + VT_CHECK(t.input_scale > 0.0F && t.global_scale > 0.0F, + std::string("NemotronH device mamba: '") + what + + "' has a non-positive input_scale or weight_scale"); + f.weight_scale = t.global_scale; + f.input_scale = t.input_scale; + // The folded per-tensor GEMM scalar, mirroring vLLM's per-tensor + // `ScaledEpilogue` (dense_fp8_gemm.h:20-39 cites the apply chain). + f.alpha = t.input_scale * t.global_scale; + f.packed.dtype = DType::kI8; + f.packed.rank = 2; + f.packed.shape[0] = f.n; + f.packed.shape[1] = f.k; + f.packed.nk = true; + f.packed.bytes.assign(t.bytes.begin(), t.bytes.end()); + return f; +} + +void BuildNemotronHMambaDeviceResident(Dev d, const NemotronHMambaWeights& w, + const NemotronHParams& params, DType adt, + NemotronHMambaDeviceResident& mr) { + if (mr.ready) return; + const int64_t H = params.hidden_size; + const int64_t I = params.mamba_intermediate_size(); + const int64_t Cd = params.conv_dim(); + const int64_t Hh = params.mamba_num_heads; + const int64_t K = params.conv_kernel; + const int64_t proj = params.in_proj_out_features(); + + mr.in_proj = MambaFp8Weight(w.in_proj, "mixer.in_proj"); + VT_CHECK(mr.in_proj.n == proj && mr.in_proj.k == H, + "NemotronH device mamba: 'mixer.in_proj' is not " + "[in_proj_out_features, hidden_size]"); + mr.out_proj = MambaFp8Weight(w.out_proj, "mixer.out_proj"); + VT_CHECK(mr.out_proj.n == H && mr.out_proj.k == I, + "NemotronH device mamba: 'mixer.out_proj' is not " + "[hidden_size, mamba_intermediate_size]"); + // Upload through the SHARED seam, which caches the device copy on the weight + // (`Fp8Weight::d_packed`). Synchronize before releasing the staging bytes: the + // copy `ResidentFp8` issues is asynchronous, so freeing the source first is a + // use-after-free that pageable-memory semantics merely hide most of the time. + (void)dense_fp8::ResidentFp8(d, mr.in_proj); + (void)dense_fp8::ResidentFp8(d, mr.out_proj); + d.b.Synchronize(d.q); + vllm::load_stats::AddDeviceUpload(mr.in_proj.packed.bytes.size() + + mr.out_proj.packed.bytes.size()); + // The device copy is now the only one. `ResidentFp8` never re-reads `packed` + // once `d_packed` is set, so dropping the host staging here is what keeps the + // conversion from doubling the 890 MB fp8 tower on a unified-memory box. + mr.in_proj.packed.bytes.Reset(); + mr.out_proj.packed.bytes.Reset(); + + VT_CHECK(params.mamba_hidden_act == "silu", + "NemotronH device mamba: only mamba_hidden_act=silu is ported (the " + "checkpoint ships silu); an other activation is refused rather than " + "silently substituted"); + mr.conv_w = ResidentDense(d, w.conv1d_weight, "mixer.conv1d.weight", adt, {Cd, K}); + if (params.use_conv_bias) { + // The conv bias is a MODEL-DTYPE tensor (ColumnParallelLinear's bias), not + // one of the three f32 SSM parameters below — the host arm says the same + // (nemotron_h.cpp:516-519). + mr.conv_b = ResidentDense(d, w.conv1d_bias, "mixer.conv1d.bias", adt, {Cd}); + mr.has_conv_bias = true; + } + // `A = -exp(A_log)` in f32. THE f32 IS UPSTREAM'S OWN POLARITY, not a local + // widening: `self.A = -torch.exp(self.A_log.float())` keeps A in f32 whatever + // the model dtype, and `vt::Mamba2ChunkScan` validates A/D/dt_bias as f32 + // (ops.cpp:2162-2180). Evaluated ONCE here rather than per step, and in the + // same `std::exp` over the same f32 input the host arm uses + // (nemotron_h.cpp:539-543), so the two arms feed the scan bit-identical A. + VT_CHECK(!w.A_log.Empty() && w.A_log.dtype == DType::kF32 && + w.A_log.shape == std::vector{Hh}, + "NemotronH device mamba: 'mixer.A_log' is absent or is not f32 [num_heads]"); + std::vector a_neg(static_cast(Hh)); + { + const auto* src = reinterpret_cast(w.A_log.bytes.data()); + for (int64_t h = 0; h < Hh; ++h) a_neg[static_cast(h)] = -std::exp(src[h]); + } + mr.a_neg = ResidentBytes(d, a_neg.data(), a_neg.size() * sizeof(float)); + mr.d_term = ResidentDense(d, w.D, "mixer.D", DType::kF32, {Hh}); + mr.dt_bias = ResidentDense(d, w.dt_bias, "mixer.dt_bias", DType::kF32, {Hh}); + mr.norm_w = ResidentDense(d, w.norm_weight, "mixer.norm.weight", adt, {I}); + d.b.Synchronize(d.q); // every upload has landed -> the host stagings are dead + mr.ready = true; +} + +// ONE NemotronH Mamba2 block on the device. +// +// `normed` is the already-normed hidden [T,H] in `adt` on the device; the return +// is the `out_proj` output [T,H] in `adt` on the device. +// +// `conv_state` / `ssm_state` are the recurrence buffers for this one request, +// device-resident and UPDATED IN PLACE — the paged forward hands them the rows +// `vt::GdnStateGather` already gathered and zeroed, so nothing round-trips +// through the host. Both null is the discard arm the non-paged forward takes, +// the exact analogue of `state == nullptr` on the host arm. +// +// `has_initial` is CARRIED SEPARATELY rather than derived from the pointers, +// because the host arm distinguishes the two (nemotron_h.cpp:500-506): a caller +// that wants the ADVANCED state back but starts from zeros passes buffers with +// `has_initial = false`, and the conv window then reads zeros and the scan gets +// NO `initial_states` — precisely what the host arm does. Deriving the flag from +// pointer-ness would make those two cases indistinguishable. +DBuf NemotronHMamba2MixerDevice(Dev d, const NemotronHMambaWeights& w, + const NemotronHParams& params, const Tensor& normed, + int64_t T, DType adt, Tensor* conv_state, + Tensor* ssm_state, bool has_initial) { + const int64_t H = params.hidden_size; + const int64_t I = params.mamba_intermediate_size(); + const int64_t Cd = params.conv_dim(); + const int64_t P = params.mamba_head_dim; + const int64_t Hh = params.mamba_num_heads; + const int64_t G = params.n_groups; + const int64_t N = params.ssm_state_size; + const int64_t Kw = params.conv_kernel; + const int64_t proj = params.in_proj_out_features(); + VT_CHECK(T > 0, "NemotronH device mamba: empty token sequence"); + VT_CHECK(normed.rank == 2 && normed.shape[0] == T && normed.shape[1] == H, + "NemotronH device mamba: the normed hidden is not [T, hidden_size]"); + VT_CHECK(Hh * P == I, "NemotronH device mamba: num_heads*head_dim != intermediate"); + VT_CHECK(I + 2 * G * N == Cd, "NemotronH device mamba: conv_dim mismatch"); + VT_CHECK((conv_state == nullptr) == (ssm_state == nullptr), + "NemotronH device mamba: the conv and SSM carries are one unit — pass " + "both or neither"); + VT_CHECK(vt::OpRegistered(vt::OpId::kQuantFp8Static, d.q.device.type), + "NemotronH device mamba: this device has no static per-tensor fp8 " + "activation quant, so the FP8 W8A8 arm cannot run (issue #960)"); + + // `ResidentIn` locks only the slot's creation, and the build below runs + // outside that lock -- the same shape `NemotronHMoeBlockDevice` has for its + // arena. Two threads entering one layer for the first time would both build. + // Nothing in this tree drives one model's forward from two threads, and + // diverging from the arena's idiom here would be an unrelated change; stated + // rather than silently inherited. + NemotronHMambaDeviceResident& mr = + ResidentIn(w.device_fp8); + BuildNemotronHMambaDeviceResident(d, w, params, adt, mr); + + // 1. the fused zxbcdt projection (mamba_mixer2.py:550), on the shared FP8 + // W8A8 seam: static per-tensor activation quant against `input_scale`, then + // the fp8 GEMM with the folded `alpha = input_scale * weight_scale`. + DBuf zxbcdt = dense_fp8::MatmulFp8CutlassD(d, normed, mr.in_proj, adt); + + // 2. split: z | xBC | dt (mamba_mixer2.py:692-696 reads xBC/dt off the tail, + // :583 reads the gate off the head). The device twin of the host arm's + // three `SliceCols` copies. + DBuf z(d, adt, {T, I}); + DBuf xbc(d, adt, {T, Cd}); + DBuf dt(d, adt, {T, Hh}); + VT_CHECK(I + Cd + Hh == proj, + "NemotronH device mamba: the zxbcdt widths do not sum to " + "in_proj_out_features"); + vt::QkvSplit(d.q, z.t(), xbc.t(), dt.t(), zxbcdt.t()); + + // 3. the causal depthwise conv with the silu activation + // (`activation=config.mamba_hidden_act` = "silu", mamba_mixer2.py:832-846). + // The conv state is f32 BY OP CONTRACT and, when the caller carries none, + // is a TRANSIENT per-call buffer exactly as the host reference's is + // (nemotron_h.cpp:493-498) — A2-Q1 is non-paged in its own right and reads + // no persistent page it did not receive. + const bool carry_in = conv_state != nullptr && has_initial; + DBuf conv_fresh(d, DType::kF32, {1, Cd, Kw - 1}); + if (conv_state == nullptr) conv_fresh.Zero(d); + Tensor cst = conv_state != nullptr ? *conv_state : conv_fresh.t(); + const int32_t qsl[2] = {0, static_cast(T)}; + const int32_t hinit[1] = {carry_in ? 1 : 0}; + DBuf dqsl(d, DType::kI32, {2}, qsl); + DBuf dhinit(d, DType::kI32, {1}, hinit); + DBuf xbc_out(d, adt, {T, Cd}); + { + Tensor cw = MakeTensor(mr.conv_w.get(), adt, d.q.device, {Cd, Kw}); + Tensor cb = MakeTensor(mr.conv_b.get(), adt, d.q.device, {Cd}); + vt::CausalConv1dArgs cargs; + cargs.silu_activation = true; + vt::CausalConv1dFwd(d.q, xbc_out.t(), xbc.t(), cw, mr.has_conv_bias ? &cb : nullptr, + cst, dqsl.t(), dhinit.t(), cargs); + } + + // 4. split the conv output into x | B | C (mamba_mixer2.py:535-543). + DBuf ssm_x(d, adt, {T, Hh, P}); + DBuf ssm_b(d, adt, {T, G, N}); + DBuf ssm_c(d, adt, {T, G, N}); + { + Tensor xf = Reshape(ssm_x.t(), {T, I}); + Tensor bf = Reshape(ssm_b.t(), {T, G * N}); + Tensor cf = Reshape(ssm_c.t(), {T, G * N}); + vt::QkvSplit(d.q, xf, bf, cf, xbc_out.t()); + } + + // 5. the SSD scan. The SSM state dtype is resolved INDEPENDENTLY of every + // activation dtype above — `mamba_ssm_cache_dtype` is "float32" on this + // checkpoint while the tower is bf16, and deriving it from `adt` would + // halve the recurrent state invisibly to a token gate (nemotron_h.h, + // NemotronHSsmCacheDType records why). + const DType ssm_dtype = NemotronHSsmCacheDType(params, adt); + if (ssm_state != nullptr) { + VT_CHECK(ssm_state->dtype == ssm_dtype, + "NemotronH device mamba: the carried SSM state is not the cache " + "dtype this model resolves"); + } + // One sequence, chunked on the GLOBAL token position — the single-sequence + // case of `compute_varlen_chunk_metadata` (v1/attention/backends/mamba2_attn.py + // :22-88), built exactly as the host arm builds it (nemotron_h.cpp:559-571). + const int64_t chunk = params.chunk_size; + const int32_t cu_seqlens[2] = {0, static_cast(T)}; + std::vector cu_chunk = {0}; + std::vector seq_idx; + for (int64_t pos = 0; pos < T; pos += chunk) { + cu_chunk.push_back(static_cast(std::min(pos + chunk, T))); + seq_idx.push_back(0); + } + const int32_t last_chunk[1] = {static_cast(seq_idx.size()) - 1}; + // These five metadata uploads are NOT followed by a `Synchronize`, unlike + // `UploadAs` above, and the difference is deliberate. Every one is a few + // hundred bytes, and a pageable H2D copy that small is staged by the driver + // before `cudaMemcpyAsync` returns, so the host buffers below may die at the + // end of this scope. That is the same reliance qwen3_5.cpp:3936 already makes + // for the identical GDN conv metadata. A per-layer stream synchronize here + // would reintroduce exactly the host/GPU lockstep this unit exists to remove. + DBuf dcu(d, DType::kI32, {2}, cu_seqlens); + DBuf dcc(d, DType::kI32, {static_cast(cu_chunk.size())}, cu_chunk.data()); + DBuf dlc(d, DType::kI32, {1}, last_chunk); + DBuf dsi(d, DType::kI32, {static_cast(seq_idx.size())}, seq_idx.data()); + DBuf y(d, adt, {T, Hh, P}); + DBuf final_states(d, ssm_dtype, {1, Hh, P, N}); + { + Tensor At = MakeTensor(mr.a_neg.get(), DType::kF32, d.q.device, {Hh}); + Tensor Dt = MakeTensor(mr.d_term.get(), DType::kF32, d.q.device, {Hh}); + Tensor dbt = MakeTensor(mr.dt_bias.get(), DType::kF32, d.q.device, {Hh}); + vt::Mamba2Args args; + args.chunk_size = chunk; + // mamba_mixer2.py:888-889: dt_softplus=True, dt_limit=(0.0, +inf). `z` is + // NOT passed to the scan — upstream gates in the norm below (:583-585), and + // passing it here would apply silu(z) twice. + args.dt_softplus = true; + args.dt_min = 0.0F; + args.dt_max = std::numeric_limits::infinity(); + vt::Mamba2ChunkScan(d.q, y.t(), final_states.t(), ssm_x.t(), dt.t(), At, ssm_b.t(), + ssm_c.t(), &Dt, /*z=*/nullptr, &dbt, carry_in ? ssm_state : nullptr, + dcu.t(), dcc.t(), dlc.t(), dsi.t(), args); + } + if (ssm_state != nullptr) { + // The scan READS `ssm_state` as its initial state, so the final state lands + // in its own buffer and is copied back afterwards rather than aliasing the + // operand the kernel is still reading. + const size_t nb = static_cast(Hh * P * N) * vt::SizeOf(ssm_dtype); + d.b.Copy(d.q, ssm_state->data, final_states.t().data, nb); + } + + // 6. the silu-gated GROUP RMS norm (Mixer2RMSNormGated, mamba_mixer2.py:478-480, + // :583-585). n_groups is the mixer's, NOT 1. + DBuf normed_gated(d, adt, {T, I}); + { + Tensor yt = Reshape(y.t(), {T, I}); + Tensor gt = z.t(); + Tensor nw = MakeTensor(mr.norm_w.get(), adt, d.q.device, {I}); + vt::RmsNormGatedGroupArgs args; + args.eps = static_cast(params.layer_norm_epsilon); + args.n_groups = G; + vt::RmsNormGatedGroup(d.q, normed_gated.t(), yt, gt, &nw, args); + } + + // 7. out_proj (mamba_mixer2.py:586), the second FP8 W8A8 projection. + return dense_fp8::MatmulFp8CutlassD(d, normed_gated.t(), mr.out_proj, adt); +} + } // namespace // ─── the per-block equivalence seam ───────────────────────────────────────── @@ -833,6 +1243,78 @@ std::vector NemotronHMoeBlockDeviceHostIO(const NemotronHMoeWeights& w, #endif } +std::vector NemotronHMamba2MixerDeviceHostIO(const NemotronHMambaWeights& w, + const NemotronHParams& params, + const std::vector& hidden_normed, + int64_t num_tokens, DType act_dtype, + Queue& dev_queue, + NemotronHMambaState* state) { + const int64_t T = num_tokens; + const int64_t H = params.hidden_size; + const int64_t Cd = params.conv_dim(); + const int64_t Kw = params.conv_kernel; + const int64_t Hh = params.mamba_num_heads; + const int64_t P = params.mamba_head_dim; + const int64_t N = params.ssm_state_size; + VT_CHECK(T > 0, "NemotronH device mamba: empty token sequence"); + VT_CHECK(static_cast(hidden_normed.size()) == T * H, + "NemotronH device mamba: hidden size mismatch"); + VT_CHECK(act_dtype == DType::kBF16 || act_dtype == DType::kF32, + "NemotronH device mamba: the model dtype must be bf16 or f32"); + VT_CHECK(dev_queue.device.type != vt::DeviceType::kCPU, + "NemotronH device mamba: this is the DEVICE arm and requires a " + "non-CPU queue; the host reference is NemotronHMamba2Mixer"); + VT_CHECK(MambaIsFp8(w), + "NemotronH device mamba: this layer's projections are not FP8 W8A8 " + "static, which is the only form A2-Q1's device arm is built from"); + + Dev d{vt::GetBackend(dev_queue.device.type), dev_queue}; + // Round the input through `act_dtype` on the way in, exactly as the host arm + // does with `PackF32` (nemotron_h.cpp:466). Feeding the device f32 values the + // host arm would have rounded first is the deviation that makes an + // equivalence gate quietly meaningless. + DBuf x = UploadAs(d, hidden_normed, act_dtype, {T, H}); + + const DType ssm_dtype = NemotronHSsmCacheDType(params, act_dtype); + const bool carry = state != nullptr && state->has_initial; + // Allocated whenever the caller wants the advanced state back, carrying or + // not — the three cases the host arm distinguishes (nemotron_h.cpp:500-506, + // :600-606) are `state == nullptr` (discard), `state` fresh (start from zeros, + // report the advance) and `state` carrying. + DBuf conv(d, DType::kF32, {1, Cd, Kw - 1}); + DBuf ssm(d, ssm_dtype, {1, Hh, P, N}); + if (carry) { + VT_CHECK(static_cast(state->conv.size()) == Cd * (Kw - 1), + "NemotronH device mamba: carried conv state has the wrong extent"); + VT_CHECK(state->ssm.dtype == ssm_dtype && state->ssm.Numel() == Hh * P * N, + "NemotronH device mamba: carried SSM state has the wrong dtype or extent"); + conv = UploadAs(d, state->conv, DType::kF32, {1, Cd, Kw - 1}); + DBuf up(d, ssm_dtype, {1, Hh, P, N}, state->ssm.bytes.data()); + d.b.Synchronize(d.q); + ssm = std::move(up); + } else if (state != nullptr) { + conv.Zero(d); + ssm.Zero(d); + } + Tensor ct = conv.t(); + Tensor st = ssm.t(); + const bool keep = state != nullptr; + DBuf out = NemotronHMamba2MixerDevice(d, w, params, x.t(), T, act_dtype, + keep ? &ct : nullptr, keep ? &st : nullptr, carry); + if (state != nullptr) { + // The advanced state, reported exactly as the host arm reports it + // (nemotron_h.cpp:600-606), so a multi-leg gate can compare the carry too. + state->conv = DownloadF32(d, conv, DType::kF32, Cd * (Kw - 1)); + std::vector sb(static_cast(Hh * P * N) * vt::SizeOf(ssm_dtype)); + ssm.Download(d, sb.data()); + state->ssm.dtype = ssm_dtype; + state->ssm.shape = {Hh, P, N}; + state->ssm.bytes = std::move(sb); + state->has_initial = true; + } + return DownloadF32(d, out, act_dtype, T * H); +} + // ─── the hybrid forward ───────────────────────────────────────────────────── std::vector NemotronHDeviceForward(const NemotronHHostWeights& host, @@ -935,9 +1417,20 @@ std::vector NemotronHDeviceForward(const NemotronHHostWeights& host, const bool moe_on_device = lw.block == NemotronHBlock::kMoe && adt == DType::kBF16 && MoeIsNvfp4(lw.moe) && vt::OpRegistered(vt::OpId::kMoeGroupedGemmNvfp4Marlin, d.q.device.type); + // A2-Q1: a Mamba2 layer whose two projections are FP8 W8A8 static runs on + // the DEVICE now, so it needs no host bounce. Same shape of selection as the + // MoE one above and for the same reason: every term is a runtime query + // (`MambaIsFp8` names only NemotronHWeightForm, `vt::OpRegistered` IS the + // op table's own answer), never a preprocessor guard, so a build or a device + // without the fp8 pair resolves false and keeps the host arm. + const bool mamba_on_device = + lw.block == NemotronHBlock::kMamba && MambaIsFp8(lw.mamba) && + NemotronHDeviceMambaEnabled() && + vt::OpRegistered(vt::OpId::kMatmulFp8CublasLt, d.q.device.type) && + vt::OpRegistered(vt::OpId::kQuantFp8Static, d.q.device.type); std::vector nvec; const bool needs_host = - lw.block != NemotronHBlock::kAttention && !moe_on_device; + lw.block != NemotronHBlock::kAttention && !moe_on_device && !mamba_on_device; if (needs_host || (trace != nullptr && trace->capture)) { nvec = DownloadF32(d, normed, adt, T * H); } @@ -954,6 +1447,16 @@ std::vector NemotronHDeviceForward(const NemotronHHostWeights& host, if (trace != nullptr && trace->capture) { mvec = DownloadF32(d, carry, adt, T * H); } + } else if (mamba_on_device) { + // NON-PAGED, exactly as this whole forward is: the recurrence is discarded + // on return, which is what the host reference does when it is handed no + // state (nemotron_h.cpp:500-506). A2-P's paged forward below is the arm + // that carries it. + carry = NemotronHMamba2MixerDevice(d, lw.mamba, params, normed.t(), T, adt, + nullptr, nullptr, false); + if (trace != nullptr && trace->capture) { + mvec = DownloadF32(d, carry, adt, T * H); + } // DSR-ALLOW(A2-Q2a): TYPES, not behaviour -- same call, same reason: the symbol is absent without the guarded region. The SELECTION is already a runtime op-table query (moe_on_device above); only the call site needs the build guard. Mirrors laguna.cpp:1187, a dispatch branch with an else fallback. #ifdef VT_MARLIN_NVFP4 } else if (moe_on_device) { @@ -1609,7 +2112,23 @@ ForwardLogits NemotronHPagedForward(const NemotronHHostWeights& host, lw.block == NemotronHBlock::kMoe && adt == DType::kBF16 && MoeIsNvfp4(lw.moe) && NemotronHDeviceMoeEnabled() && vt::OpRegistered(vt::OpId::kMoeGroupedGemmNvfp4Marlin, d.q.device.type); - const bool needs_host = lw.block != NemotronHBlock::kAttention && !moe_on_device; + // A2-Q1 (#810): the 23 Mamba2 layers on the FP8 W8A8 seam. THE `ssm_dtype == + // f32` TERM IS NOT DECORATION. `vt::GdnStateGather` widens the page into an + // f32 working buffer by op contract, and the HOST arm below then NARROWS + // that f32 back to `ssm_dtype` before the mixer sees it + // (`NemotronHOwned::FromF32`, :below). On a checkpoint whose + // `mamba_ssm_cache_dtype` is not f32 the two arms would therefore round + // differently, and the per-block numeric gate would be comparing two + // different computations. The released checkpoint resolves f32, so the + // device arm runs; anything else keeps the host bounce rather than silently + // dropping a rounding step. + const bool mamba_on_device = + lw.block == NemotronHBlock::kMamba && MambaIsFp8(lw.mamba) && + NemotronHDeviceMambaEnabled() && ssm_dtype == DType::kF32 && + vt::OpRegistered(vt::OpId::kMatmulFp8CublasLt, d.q.device.type) && + vt::OpRegistered(vt::OpId::kQuantFp8Static, d.q.device.type); + const bool needs_host = + lw.block != NemotronHBlock::kAttention && !moe_on_device && !mamba_on_device; std::vector nvec; if (needs_host || (trace != nullptr && trace->capture)) { nvec = DownloadF32(d, normed, adt, T * H); @@ -1627,6 +2146,61 @@ ForwardLogits NemotronHPagedForward(const NemotronHHostWeights& host, carry = NemotronHMoeBlockDevice(d, lw.moe, params, normed.t(), T); if (trace != nullptr && trace->capture) mvec = DownloadF32(d, carry, adt, T * H); #endif + } else if (mamba_on_device) { + // ── A2-Q1: the CARRY, without leaving the device. ── + // + // The gather/scatter pair and the `has_initial = true` reasoning are A2-P's + // and are UNCHANGED (see the block comment above `NemotronHRecurrentIo`). + // What changes is everything between them: the two D2H downloads, the host + // mixer with its per-call fp8 dequant, and the two H2D uploads are gone, + // and the gathered rows are handed straight to the device mixer, which + // advances them IN PLACE in the same buffers the scatter then writes back. + const GdnStateCache& cache = input.gdn_state[mamba_i]; + NemotronHRecurrentIo io = GatherNemotronHState(d, cache, params, R, sdi); + const int64_t conv_row = params.conv_dim() * (params.conv_kernel - 1); + const int64_t ssm_row = + params.mamba_num_heads * params.mamba_head_dim * params.ssm_state_size; + DBuf mixed(d, adt, {T, H}); + // Zeroed first, exactly as the host arm's `mvec.assign(T * H, 0.0F)` + // below does, so a token no request's query range covers reads as zero + // rather than as whatever the scratch pool last held there. + mixed.Zero(d); + const size_t esz = vt::SizeOf(adt); + for (int64_t r = 0; r < R; ++r) { + // The RECURRENT half's own query offsets, read exactly as the host arm + // below reads them, so a mixed batch stays correct when A2-B lifts the + // request count. + const std::vector& qsl = *input.gdn_meta.non_spec_query_start_loc; + const int64_t t0 = qsl[static_cast(r)]; + const int64_t t1 = qsl[static_cast(r + 1)]; + VT_CHECK(t1 > t0 && t1 <= T, + "NemotronH paged forward: a request's query range is empty or " + "runs past the step's tokens"); + Tensor rows = MakeTensor(static_cast(normed.t().data) + + static_cast(t0 * H) * esz, + adt, d.q.device, {t1 - t0, H}); + Tensor cr = MakeTensor(static_cast(io.conv.t().data) + + static_cast(r * conv_row) * sizeof(float), + DType::kF32, d.q.device, + {1, params.conv_dim(), params.conv_kernel - 1}); + Tensor sr = MakeTensor(static_cast(io.ssm.t().data) + + static_cast(r * ssm_row) * sizeof(float), + DType::kF32, d.q.device, + {1, params.mamba_num_heads, params.mamba_head_dim, + params.ssm_state_size}); + // `has_initial = true` in EVERY case, over a row the gather has already + // zeroed when the mask said fresh — A2-P's property, restated here + // because this arm is the one that now consumes it. + DBuf got = NemotronHMamba2MixerDevice(d, lw.mamba, params, rows, t1 - t0, adt, + &cr, &sr, /*has_initial=*/true); + d.b.Copy(d.q, static_cast(mixed.t().data) + + static_cast(t0 * H) * esz, + got.t().data, static_cast((t1 - t0) * H) * esz); + } + ScatterNemotronHState(d, cache, io, sdi); + ++mamba_i; + carry = std::move(mixed); + if (trace != nullptr && trace->capture) mvec = DownloadF32(d, carry, adt, T * H); } else if (lw.block == NemotronHBlock::kMamba) { // ── the CARRY. This is the unit. ── const GdnStateCache& cache = input.gdn_state[mamba_i]; diff --git a/src/vllm/model_executor/models/nemotron_h_forward.h b/src/vllm/model_executor/models/nemotron_h_forward.h index 25360fa43..d07a520b8 100644 --- a/src/vllm/model_executor/models/nemotron_h_forward.h +++ b/src/vllm/model_executor/models/nemotron_h_forward.h @@ -240,6 +240,20 @@ struct NemotronHMambaWeights { NemotronHOwned dt_bias; // [mamba_num_heads] f32 // Mixer2RMSNormGated weight over the SSM intermediate width. NemotronHOwned norm_weight; // [mamba_intermediate_size] + + // A2-Q1 (#810): this layer's device-resident mamba state — the two FP8 W8A8 + // towers uploaded through the shared `dense_fp8::ResidentFp8` seam plus the + // six small recurrence parameters, built ONCE on first device use and owned + // BY THE WEIGHTS. Opaque here for the same reason `moe_marlin` below is: the + // resident type is an implementation detail of nemotron_h_device.cpp. + // + // KEYED ON THE SLOT, NEVER ON AN ADDRESS (issue #237, qwen3_5_weights.h:183). + // Two engine builds in one process can hand the second engine's weights the + // address the first engine's had, and an address-keyed cache then returns the + // PREVIOUS engine's device pointers: plausible, wrong values rather than a + // crash. Holding the state in the weights it describes makes that + // unrepresentable. + ResidentSlot device_fp8; }; // GQA attention weights. q/k/v ship SEPARATE on disk (upstream fuses them into @@ -516,6 +530,29 @@ std::vector NemotronHMoeBlockDeviceHostIO(const NemotronHMoeWeights& w, int64_t num_tokens, vt::DType act_dtype, vt::Queue& dev_queue); +// A2-Q1 (#810): ONE NemotronH Mamba2 block on the device, with host-side input +// and output so a gate can drive a single block in isolation — the same +// per-block equivalence seam `NemotronHMoeBlockDeviceHostIO` is, and for the +// same reason. +// +// WHY THE GATE IS NUMERIC AND NOT TOKENS. A token comparison cannot see a +// dropped `input_scale` (the activation quantized against 1.0), an `alpha` +// folded as `weight_scale` alone, a transposed `in_proj` operand, a `zxbcdt` +// split offset by one column, or a projection that quietly stayed on the host +// and dequantized. Every one of those is finite, correctly shaped and +// plausible, and this row has already been bitten by three of them. +// +// `state` is the OPTIONAL carried recurrence, exactly as `NemotronHMamba2Mixer` +// takes it: null runs the fresh-state arm (zero conv window, no initial SSM +// state), non-null carries in and is UPDATED IN PLACE, so the two arms are +// comparable on the carrying path as well as the fresh one. +std::vector NemotronHMamba2MixerDeviceHostIO(const NemotronHMambaWeights& w, + const NemotronHParams& params, + const std::vector& hidden_normed, + int64_t num_tokens, vt::DType act_dtype, + vt::Queue& dev_queue, + NemotronHMambaState* state = nullptr); + // The final output projection, on the HOST, over `num_rows` already-gathered // and already-final-normed rows `[num_rows, hidden_size]` (f32 in, f32 logits // `[num_rows, vocab_size]` out). diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 83faa1da0..36cf3698c 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -847,6 +847,21 @@ vllm_cpp_add_test(test_nemotron_h_moe_device target_include_directories(test_nemotron_h_moe_device PRIVATE ${CMAKE_SOURCE_DIR}/src) +# A2-Q1 (#810, .agents/specs/nemotron-h-a2q1-fp8-mamba.md): the DEVICE Mamba2 arm +# on a SYNTHETIC FP8 W8A8 fixture — the cheap RED in front of the +# real-checkpoint per-block gate, and the same reason A2-Q2a's file above exists: +# `BuildTiny` is all-`kDense`, so nothing in the existing device suite can reach +# the FP8 W8A8 arm at all. The two arms are NOT bit-comparable by construction — +# the host reference dequantizes the weight and leaves the activation in bf16 +# (W8A16), the device arm is W8A8 as vLLM is — so every band is measured in the +# run against a defect the fixture separates. Every GEMM extent is a multiple of +# 16 so a cuBLASLt fp8 plan resolves and a refusal is about VALUES. Skips LOUDLY +# with a reason on a GPU-less box. +vllm_cpp_add_test(test_nemotron_h_mamba_device + vllm/models/test_nemotron_h_mamba_device.cpp) +target_include_directories(test_nemotron_h_mamba_device PRIVATE + ${CMAKE_SOURCE_DIR}/src) + # A2-P (#810, .agents/specs/nemotron-h-a2p-paged-forward.md): the PAGED forward, # gated through the PRODUCTION entry point. Every case builds a real # NemotronHLoadedModel from a synthetic checkpoint, hands it to a real diff --git a/tests/vllm/models/test_nemotron_h_mamba_device.cpp b/tests/vllm/models/test_nemotron_h_mamba_device.cpp new file mode 100644 index 000000000..30d4f5c17 --- /dev/null +++ b/tests/vllm/models/test_nemotron_h_mamba_device.cpp @@ -0,0 +1,439 @@ +// NemotronH A2-Q1 (#810, parent #517) — the DEVICE Mamba2 arm on a SYNTHETIC +// FP8 W8A8 fixture. +// +// ─── WHY THIS FILE EXISTS AT ALL ──────────────────────────────────────────── +// +// The A2-Q1 spec's §5.1 gate is per-block numeric equivalence on the REAL 21 GB +// checkpoint. This file is the CHEAP arm in front of it, and it is not a +// substitute: +// +// * `test_nemotron_h_forward.cpp`'s `BuildTiny` fixture is all `kDense`, so +// nothing in the existing device suite can reach the FP8 W8A8 arm at all. +// Without this file the real checkpoint would be the ONLY instrument, and +// every RED would cost a GB10 window plus a 21 GB load. +// * The same reasoning that put `test_nemotron_h_moe_device.cpp` in front of +// A2-Q2a's real-checkpoint gate, for the same row. +// +// ─── WHAT THE COMPARISON ACTUALLY IS, WHICH IS NOT "TWO WAYS TO DO bf16" ──── +// +// The HOST reference is W8A16: `Linear(..., const NemotronHOwned&)` +// (nemotron_h.cpp:291) reaches `DenseFor` (:245), which calls +// `NemotronHOwned::DenseBf16()` — a full dequant of the fp8 weight into bf16 — +// and then runs a bf16 GEMM on an UNQUANTIZED activation. `DenseBf16` says so +// itself (nemotron_h.cpp:419-422: "Weight-only: `input_scale` is carried, not +// applied"). +// +// The DEVICE arm is W8A8, which is what vLLM does: static per-tensor activation +// quant against `input_scale` (`quant_utils.py:124`) then an fp8 GEMM with the +// folded `alpha = input_scale * weight_scale`. +// +// So the two arms are NOT bit-comparable by construction, and the difference is +// the e4m3 activation quantization — three mantissa bits, averaged down over the +// K-sum. That is exactly why the band below is MEASURED IN THE RUN and the +// acceptance is derived from a defect this fixture can separate, rather than +// chosen. A band chosen up front is how a 3e-2 bf16 tolerance came to sit ABOVE +// a 2.11e-2 defect on this very row and accept a wrong answer. +// +// ─── THE GEOMETRY IS NOT ARBITRARY ────────────────────────────────────────── +// +// Every GEMM extent is a multiple of 16, because a cuBLASLt fp8 plan needs +// 16-byte-aligned leading dimensions and a shape it declines to run would make a +// refusal here a statement about alignment rather than about values: +// in_proj [N=336, K=128] 336 = I(128) + conv_dim(192) + num_heads(16) +// out_proj [N=128, K=128] +// and the recurrence divides: num_heads*head_dim = 16*8 = 128 = intermediate, +// intermediate + 2*n_groups*state = 128 + 2*2*16 = 192 = conv_dim, and +// n_groups(2) divides intermediate(128) for the group RMS norm. +#include + +#include +#include +#include +#include +#include + +#include "vllm/model_executor/model_loader/safetensors_reader.h" // load_stats +#include "vllm/model_executor/models/nemotron_h.h" +#include "vllm/model_executor/models/nemotron_h_forward.h" +#include "vt/backend.h" +#include "vt/dtype.h" +#include "vt/fp8_kv.h" // vt::F32ToF8E4M3 — the fixture's e4m3 encoder + +namespace { + +using vllm::NemotronHBlock; +using vllm::NemotronHMambaState; +using vllm::NemotronHMambaWeights; +using vllm::NemotronHOwned; +using vllm::NemotronHParams; +using vllm::NemotronHWeightForm; +using vt::Device; +using vt::DeviceType; +using vt::DType; +using vt::Queue; + +// NOT 1.0, either of them. A `weight_scale` of 1 is invisible to a mutation that +// ignores it, and an `input_scale` of 1 is invisible to the mutation that drops +// it — which is Q1-M1, the first row of the spec's mutation table. +constexpr float kWeightScale = 0.125F; +constexpr float kInputScale = 0.5F; + +// The same deterministic generator the forward and MoE-device suites use, so a +// value here is reproducible and independent of any RNG seeding order. +std::vector SynthVec(size_t n, uint32_t salt, float k) { + std::vector v(n); + for (size_t i = 0; i < n; ++i) { + const double a = static_cast(i); + const double b = static_cast(salt); + v[i] = static_cast(k * std::sin(0.7 * a + 1.3 * b + 0.21 * a * b)); + } + return v; +} + +// An FP8 W8A8 static projection [rows, cols]. The stored e4m3 bytes are the +// QUANTIZED weight; the logical value is `byte * weight_scale`, which is what +// `DenseBf16` reconstructs on the host arm and what the folded alpha applies on +// the device arm. +NemotronHOwned MakeFp8(int64_t rows, int64_t cols, uint32_t salt, DType logical) { + NemotronHOwned w; + w.form = NemotronHWeightForm::kFp8W8A8Static; + w.dtype = logical; + w.shape = {rows, cols}; + const std::vector v = SynthVec(static_cast(rows * cols), salt, 3.0F); + w.bytes.resize(v.size()); + for (size_t i = 0; i < v.size(); ++i) w.bytes[i] = vt::F32ToF8E4M3(v[i]); + w.global_scale = kWeightScale; + w.input_scale = kInputScale; + w.has_input_scale = true; + return w; +} + +NemotronHOwned OwnF32(const std::vector& v, DType dt, std::vector shape) { + return NemotronHOwned::FromF32(v, dt, std::move(shape)); +} + +// A NemotronH mamba geometry whose every GEMM extent is a multiple of 16. +NemotronHParams MambaParams() { + NemotronHParams p; + p.hidden_size = 128; + p.mamba_num_heads = 16; + p.mamba_head_dim = 8; + p.n_groups = 2; + p.ssm_state_size = 16; + p.conv_kernel = 4; + p.chunk_size = 8; + p.use_conv_bias = true; + p.mamba_hidden_act = "silu"; + p.layer_norm_epsilon = 1e-5; + p.vocab_size = 32; + p.layers_block_type = {NemotronHBlock::kMamba}; + return p; +} + +NemotronHMambaWeights MakeFp8Mamba(const NemotronHParams& p, DType dt) { + NemotronHMambaWeights w; + const int64_t H = p.hidden_size; + const int64_t I = p.mamba_intermediate_size(); + const int64_t Cd = p.conv_dim(); + const int64_t K = p.conv_kernel; + const int64_t Hh = p.mamba_num_heads; + w.in_proj = MakeFp8(p.in_proj_out_features(), H, 11, dt); + w.out_proj = MakeFp8(H, I, 12, dt); + w.conv1d_weight = OwnF32(SynthVec(static_cast(Cd * K), 13, 0.4F), dt, {Cd, K}); + w.conv1d_bias = OwnF32(SynthVec(static_cast(Cd), 14, 0.2F), dt, {Cd}); + // f32 BY CONTRACT on both arms — `vt::Mamba2ChunkScan` validates A/D/dt_bias + // as f32, mirroring upstream's `-torch.exp(self.A_log.float())`. + w.A_log = OwnF32(SynthVec(static_cast(Hh), 15, 0.5F), DType::kF32, {Hh}); + w.D = OwnF32(SynthVec(static_cast(Hh), 16, 0.6F), DType::kF32, {Hh}); + w.dt_bias = OwnF32(SynthVec(static_cast(Hh), 17, 0.3F), DType::kF32, {Hh}); + w.norm_weight = OwnF32(SynthVec(static_cast(I), 18, 0.7F), dt, {I}); + return w; +} + +bool TryCudaQueue(Queue* q) { + try { + *q = vt::GetBackend(DeviceType::kCUDA).CreateQueue(); + return true; + } catch (...) { + return false; + } +} + +// A GPU-less box must SKIP LOUDLY. A device case that silently reports a pass +// over zero device work is indistinguishable from a real one. +void NoteDeviceSkip(const std::string& case_name) { + MESSAGE("SKIPPED '" << case_name + << "': no CUDA device on this box. A2-Q1's device Mamba2 arm " + "gates on dgx.casa (GB10 sm_121a) and 192.168.68.23 (Thor " + "sm_110); a result from a GPU-less box is not an A2-Q1 result."); + CHECK(true); // the skip path ran and said so +} + +double MaxAbs(const std::vector& v) { + double m = 0.0; + for (float x : v) m = std::max(m, std::abs(static_cast(x))); + return m; +} + +// The largest relative deviation between two answers, against a scale-aware +// denominator. ★ IT REPORTS HOW MANY ELEMENTS IT EXAMINED, AND EVERY CALLER +// ASSERTS THAT. A maximum over ZERO elements is 0.0 — and 0.0 is also exactly +// what a bit-exact comparison prints, so the deviation ALONE cannot distinguish +// "the two arms agree exactly over 512 elements" from "the loop ran over +// nothing". +double MaxRel(const std::vector& got, const std::vector& want, + int64_t* examined) { + const double scale = std::max(MaxAbs(want), 1e-30); + double worst = 0.0; + int64_t n = 0; + for (size_t i = 0; i < got.size() && i < want.size(); ++i) { + const double d = std::abs(static_cast(got[i]) - static_cast(want[i])); + worst = std::max(worst, d / scale); + ++n; + } + if (examined != nullptr) *examined = n; + return worst; +} + +} // namespace + +// ─── the gate ─────────────────────────────────────────────────────────────── +// +// NO COMMA IN ANY CASE NAME, EVER. doctest's `-tc` filter splits on commas, so a +// comma here makes a targeted run select ZERO cases, print `SUCCESS!` and exit +// 0 — a whole mutation pass once read GREEN that way on this repository, +// including the row that deleted the guard. +TEST_CASE("NemotronH A2-Q1: the device Mamba2 block matches the host reference on FP8 W8A8") { + Queue dq{Device{DeviceType::kCPU, 0}, nullptr}; + if (!TryCudaQueue(&dq)) { + NoteDeviceSkip("device Mamba2 block vs host reference"); + return; + } + const NemotronHParams p = MambaParams(); + Queue hq{Device{DeviceType::kCPU, 0}, nullptr}; + const DType dt = DType::kBF16; // the released checkpoint's model dtype + const int64_t H = p.hidden_size; + const NemotronHMambaWeights w = MakeFp8Mamba(p, dt); + + // T == 1 IS THE DECODE SHAPE and it is the width this model spends its whole + // decode in; 8 is one full chunk and 12 straddles two, which is where the + // varlen chunk metadata stops being degenerate. A width loop rather than three + // copies, and the count is asserted afterwards so a loop that ran over nothing + // cannot report a clean pass. + int64_t widths_covered = 0; + for (const int64_t T : {int64_t{1}, int64_t{8}, int64_t{12}}) { + INFO("token count T=" << T); + const std::vector x = SynthVec(static_cast(T * H), 77, 0.5F); + + const std::vector host = vllm::NemotronHMamba2Mixer(w, p, x, T, dt, hq); + const std::vector dev = + vllm::NemotronHMamba2MixerDeviceHostIO(w, p, x, T, dt, dq); + + REQUIRE(host.size() == static_cast(T * H)); + REQUIRE(dev.size() == host.size()); + + int64_t examined = 0; + const double agreed = MaxRel(dev, host, &examined); + MESSAGE("device-vs-host worst relative deviation: " << agreed << " over " << examined + << " elements"); + // THE COUNT IS THE REPORT. Asserted against the GEOMETRY rather than against + // `dev.size()`, so a comparison that silently shortened would still red. + REQUIRE(examined == T * H); + REQUIRE(examined > 0); + + // ── THE SEPARATION IS A REAL DEFECT OF THIS BLOCK ──────────────────────── + // Q1-M2 is "alpha folded as weight_scale alone", which scales every output + // of both fp8 GEMMs by 1/input_scale. Applied to the host answer through the + // SAME arithmetic and the SAME element count, that is what this fixture can + // resolve, and the band is derived from it rather than chosen. + std::vector perturbed = host; + for (float& v : perturbed) v = static_cast(static_cast(v) / kInputScale); + int64_t sep_examined = 0; + const double separation = MaxRel(perturbed, host, &sep_examined); + MESSAGE("separation of a dropped input_scale: " << separation << " over " << sep_examined + << " elements"); + REQUIRE(sep_examined == examined); + REQUIRE(separation > 0.0); + + // Half the measured separation: strictly above any agreement this fixture + // can show, strictly below the defect it must reject, and it keeps that + // property at `agreed == 0` (unlike a geometric mean, which degenerates to 0 + // and then fails on the best possible outcome). + const double band = separation / 2.0; + MESSAGE("accepting at band " << band); + CHECK(agreed < band); + + // ── THE GUARD IS A PROPERTY, NOT A TWIN ───────────────────────────────── + // The perturbed answer, run through the SAME comparison that accepted the + // real one, must come out REJECTED. Widen `band` past the separation and + // this line reds, which is the historical hole reproduced by construction. + int64_t guard_examined = 0; + const double guard = MaxRel(perturbed, host, &guard_examined); + REQUIRE(guard_examined == examined); + INFO("does the band " << band << " REJECT a dropped input_scale?"); + CHECK(guard >= band); + ++widths_covered; + } + REQUIRE(widths_covered == 3); +} + +// The CARRY. With one leg and fresh state the recurrence is unobservable — the +// conv window and the SSM state are both zero, so a dropped carry computes the +// identical answer. Two legs is the smallest arrangement in which the device +// arm's in-place state advance is visible at all. +TEST_CASE("NemotronH A2-Q1: the device Mamba2 arm carries conv and SSM state across two legs") { + Queue dq{Device{DeviceType::kCPU, 0}, nullptr}; + if (!TryCudaQueue(&dq)) { + NoteDeviceSkip("device Mamba2 carry across two legs"); + return; + } + const NemotronHParams p = MambaParams(); + Queue hq{Device{DeviceType::kCPU, 0}, nullptr}; + const DType dt = DType::kBF16; + const int64_t H = p.hidden_size; + const NemotronHMambaWeights w = MakeFp8Mamba(p, dt); + + const int64_t T0 = 8; // the prefill leg + const int64_t T1 = 1; // the decode leg that can only be right if the carry is + const std::vector x0 = SynthVec(static_cast(T0 * H), 31, 0.5F); + const std::vector x1 = SynthVec(static_cast(T1 * H), 32, 0.5F); + + NemotronHMambaState hs; + NemotronHMambaState ds; + (void)vllm::NemotronHMamba2Mixer(w, p, x0, T0, dt, hq, &hs); + (void)vllm::NemotronHMamba2MixerDeviceHostIO(w, p, x0, T0, dt, dq, &ds); + REQUIRE(hs.has_initial); + REQUIRE(ds.has_initial); + REQUIRE(ds.conv.size() == hs.conv.size()); + REQUIRE(ds.conv.size() == static_cast(p.conv_dim() * (p.conv_kernel - 1))); + REQUIRE(ds.ssm.dtype == hs.ssm.dtype); + REQUIRE(ds.ssm.Numel() == hs.ssm.Numel()); + REQUIRE(ds.ssm.Numel() == + p.mamba_num_heads * p.mamba_head_dim * p.ssm_state_size); + + // The conv window is the RAW pre-activation input the block just consumed, so + // it is the one carried tensor the two arms compute the same way up to the fp8 + // activation quant of `in_proj`. It is compared here because a device arm that + // advanced the SSM state but left the conv window unwritten would still return + // a plausible first-leg answer. + int64_t conv_examined = 0; + const double conv_agreed = MaxRel(ds.conv, hs.conv, &conv_examined); + MESSAGE("carried conv window device-vs-host: " << conv_agreed << " over " << conv_examined + << " elements"); + REQUIRE(conv_examined == static_cast(hs.conv.size())); + REQUIRE(conv_examined > 0); + + const std::vector host1 = vllm::NemotronHMamba2Mixer(w, p, x1, T1, dt, hq, &hs); + const std::vector dev1 = + vllm::NemotronHMamba2MixerDeviceHostIO(w, p, x1, T1, dt, dq, &ds); + REQUIRE(host1.size() == static_cast(T1 * H)); + REQUIRE(dev1.size() == host1.size()); + + int64_t examined = 0; + const double agreed = MaxRel(dev1, host1, &examined); + MESSAGE("second-leg device-vs-host worst relative deviation: " + << agreed << " over " << examined << " elements"); + REQUIRE(examined == T1 * H); + REQUIRE(examined > 0); + + // THE SEPARATION IS THE DROPPED CARRY ITSELF: what the second leg would return + // if the device arm had started from zeros. Measured in this run, over the same + // population, through the same arithmetic. + const std::vector fresh1 = vllm::NemotronHMamba2Mixer(w, p, x1, T1, dt, hq); + int64_t sep_examined = 0; + const double separation = MaxRel(fresh1, host1, &sep_examined); + MESSAGE("separation of a DROPPED carry: " << separation << " over " << sep_examined + << " elements"); + REQUIRE(sep_examined == examined); + REQUIRE(separation > 0.0); + const double band = separation / 2.0; + MESSAGE("accepting at band " << band); + CHECK(agreed < band); + INFO("does the band " << band << " REJECT a dropped carry?"); + CHECK(separation >= band); +} + +// A separate case so a `-tc` run can select it alone. Same no-comma rule. +TEST_CASE("NemotronH A2-Q1: the device Mamba2 arm refuses a dense projection rather than reading garbage") { + Queue dq{Device{DeviceType::kCPU, 0}, nullptr}; + if (!TryCudaQueue(&dq)) { + NoteDeviceSkip("device Mamba2 refuses a dense projection"); + return; + } + const NemotronHParams p = MambaParams(); + const DType dt = DType::kBF16; + const int64_t T = 4; + const int64_t H = p.hidden_size; + const std::vector x = SynthVec(static_cast(T * H), 78, 0.5F); + + // A layer whose `in_proj` is DENSE, which is what `BuildTiny` produces and what + // an unquantized NemotronH would load. The whole-forward path takes the host + // bounce for exactly this case; the block entry point refuses, because a direct + // caller asked for the device. + NemotronHMambaWeights dense = MakeFp8Mamba(p, dt); + dense.in_proj = OwnF32(SynthVec(static_cast(p.in_proj_out_features() * H), 5, 0.3F), + dt, {p.in_proj_out_features(), H}); + CHECK_THROWS(vllm::NemotronHMamba2MixerDeviceHostIO(dense, p, x, T, dt, dq)); + + // And an `input_scale` the checkpoint never shipped: `Fp8Weight` has no + // `has_input_scale` field, so the conversion must REFUSE rather than let a + // defaulted divisor reach the activation quant (spec §4.1). + NemotronHMambaWeights noscale = MakeFp8Mamba(p, dt); + noscale.in_proj.has_input_scale = false; + CHECK_THROWS(vllm::NemotronHMamba2MixerDeviceHostIO(noscale, p, x, T, dt, dq)); +} + +// THE RESIDENCY IS THE POINT OF THE WHOLE UNIT, and it is not visible in any +// numeric comparison: an arm that re-uploaded the 890 MB fp8 tower on every step +// would return exactly the same numbers as one that uploads it once. So it is +// asserted directly, against the geometry, through the ONE counter the tree +// already keeps for this (`load_stats::AddDeviceUpload`, which +// `dense_attn::ResidentWeight` and `ResidentNvfp4` both call). +// +// This is also the reachability instrument this row needs. `dense_fp8:: +// ResidentFp8` does NOT account its own upload — that is issue #974 and A2-Q1 +// does not change the shared header — so A2-Q1 accounts what IT uploads at the +// site that causes it. A step that took the host bounce instead uploads nothing, +// and the delta is 0. +TEST_CASE("NemotronH A2-Q1: the FP8 mamba tower uploads ONCE and the second call uploads nothing") { + Queue dq{Device{DeviceType::kCPU, 0}, nullptr}; + if (!TryCudaQueue(&dq)) { + NoteDeviceSkip("device Mamba2 residency and upload accounting"); + return; + } + const NemotronHParams p = MambaParams(); + const DType dt = DType::kBF16; + const int64_t T = 4; + const int64_t H = p.hidden_size; + const int64_t I = p.mamba_intermediate_size(); + const int64_t Cd = p.conv_dim(); + const int64_t Kw = p.conv_kernel; + const int64_t Hh = p.mamba_num_heads; + const NemotronHMambaWeights w = MakeFp8Mamba(p, dt); + const std::vector x = SynthVec(static_cast(T * H), 79, 0.5F); + + // What the arm OWES the device, derived from the geometry rather than read + // back from the run: the two e4m3 towers, the conv weight and bias, the three + // f32 [num_heads] recurrence parameters, and the gated-norm weight. + const uint64_t esz = static_cast(vt::SizeOf(dt)); + const uint64_t expect = + static_cast(p.in_proj_out_features() * H) + // in_proj e4m3 + static_cast(H * I) + // out_proj e4m3 + static_cast(Cd * Kw) * esz + // conv1d.weight + static_cast(Cd) * esz + // conv1d.bias + static_cast(Hh) * 4U * 3U + // A, D, dt_bias (f32) + static_cast(I) * esz; // norm.weight + + const uint64_t before = vllm::load_stats::Snapshot().device_upload_bytes; + (void)vllm::NemotronHMamba2MixerDeviceHostIO(w, p, x, T, dt, dq); + const uint64_t after_first = vllm::load_stats::Snapshot().device_upload_bytes; + (void)vllm::NemotronHMamba2MixerDeviceHostIO(w, p, x, T, dt, dq); + const uint64_t after_second = vllm::load_stats::Snapshot().device_upload_bytes; + + const uint64_t first = after_first - before; + const uint64_t second = after_second - after_first; + MESSAGE("device upload: first call " << first << " B, second call " << second + << " B, expected " << expect << " B"); + CHECK(first == expect); + CHECK(second == 0U); +} From 681504c995f3442c6d724a57650bdd67b072ad51 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 22:45:58 +0000 Subject: [PATCH 02/12] test(#810 A2-Q1): the A3 gate reports its per-token time with the terms it divided, or refuses `examples/nemotron_h_gen` prints neither a rate nor a duration. It prints the engine load time and a `TOKEN MATCH: m/n` line, so the per-output-token number this row is measured against has to be derived from the wall clock, and a derived number that hides its terms is exactly how a rate over an unknown denominator comes to be quoted as if it had been measured. `scripts/nemotron-h-a2q1-per-token.py` prints the wall time, the load it subtracts and the token count it divides by on separate lines, and it refuses in the two cases where a plausible zero would otherwise be printed: when either term is absent from the log, and when the wall clock is not greater than the load it contains. The second case is not hypothetical -- the terms coming from different runs yields a NEGATIVE rate, which still formats like a measurement. The gate script calls it once per arm, so the device-on and device-off runs report the same shape. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- scripts/nemotron-h-a2q1-dgx-gate.sh | 13 +++-- scripts/nemotron-h-a2q1-per-token.py | 71 ++++++++++++++++++++++++++++ 2 files changed, 80 insertions(+), 4 deletions(-) create mode 100755 scripts/nemotron-h-a2q1-per-token.py diff --git a/scripts/nemotron-h-a2q1-dgx-gate.sh b/scripts/nemotron-h-a2q1-dgx-gate.sh index 19ca56dda..9d92f4b51 100755 --- a/scripts/nemotron-h-a2q1-dgx-gate.sh +++ b/scripts/nemotron-h-a2q1-dgx-gate.sh @@ -139,17 +139,16 @@ run_gate() { # $1 = label, $2 = VT_NEMOTRON_H_DEVICE_MAMBA value ( 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=$! - local t0=$(date +%s.%N) + t0=$(date +%s.%N) VT_NEMOTRON_H_DEVICE_MAMBA=$flag "$BUILD/examples/nemotron-h-gen" \ --model "$CKPT" \ --golden "$SRC/tests/parity/goldens/nemotron_35_lightning_greedy/oracle.json" \ > "$RUN/a3_$label.log" 2>&1 local r=$? - local t1=$(date +%s.%N) + t1=$(date +%s.%N) kill "$sampler" 2>/dev/null wait "$sampler" 2>/dev/null echo "RC[a3 $label]=$r" - echo "wall seconds ($label): $(echo "$t1 - $t0" | bc)" grep -E "STRICT|PASS|FAIL|mode=|tok/s|per output token" "$RUN/a3_$label.log" | tail -20 python3 - "$RUN/util_$label.txt" "$label" <<'PY' import sys @@ -161,7 +160,13 @@ else: print(f"{sys.argv[2]}: GPU busy in {busy} of {len(vals)} samples = " f"{100.0*busy/len(vals):.2f}% busy (baseline 6.31%)") PY - grep -c "reference-tier" "$RUN/a3_$label.log" > /dev/null 2>&1 + # The per-output-token time, DERIVED and shown with its terms, because the + # driver reports neither a rate nor a duration. Wall minus the engine load, + # over the tokens the run actually COMPARED -- printing a rate without the + # token count is the mute-instrument shape this row has already been bitten by. + python3 "$SRC/scripts/nemotron-h-a2q1-per-token.py" "$RUN/a3_$label.log" "$label" "$t0" "$t1" + # 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' "$RUN/a3_$label.log")" } diff --git a/scripts/nemotron-h-a2q1-per-token.py b/scripts/nemotron-h-a2q1-per-token.py new file mode 100755 index 000000000..6e6a5d30b --- /dev/null +++ b/scripts/nemotron-h-a2q1-per-token.py @@ -0,0 +1,71 @@ +#!/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. + +This prints the wall time, the load it subtracts, and the token count it divides +by, on separate lines, and it REFUSES rather than printing 0 when either term is +missing from the log. + + nemotron-h-a2q1-per-token.py