diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 92df143aa..1d29b270f 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -467,6 +467,10 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1394](https://github.com/mudler/vllm.cpp/issues/1394) | `ENG-CUDAGRAPH-BREAK` | The CPU paged attention reads `btab[r * bt_row + (j / block_size) * bt_col]` for every `j < seq_lens[r]` without checking that the block table has that many columns, so a caller with a short table gets an out-of-bounds read, a plausible block index out of it, and attention over the WRONG page — silently. `tests/vllm/models/test_qwen3_5_decode_graph_seam.cpp`'s `SpecAttnMeta` supplies one: hardcoded `block_table_num_cols = 1` against shape C's `seq_lens = 24` at `block_size = 16`. Found while fixing [#1380](https://github.com/mudler/vllm.cpp/issues/1380), whose `DevicePool` change moved the bytes after the table and turned the same read into a SIGSEGV on `thor:gpu0` (`gdb` at `src/vt/cpu/cpu_paged_attn.cpp:224` under `FullAttnBlockPaged`). PRE-EXISTING: the case passes at `origin/main` only because the read landed on bytes that decoded to an in-range index. FIXED IN FLOW in both halves — the kernel refuses a short table with one compare per request outside the token loop, and the helper sizes its table for the sequence length it declares. Owned by row `ENG-CUDAGRAPH-BREAK` | bug | | [#1387](https://github.com/mudler/vllm.cpp/issues/1387) | `ENG-EXPERT-STREAM-DEVICE` | `docs/FEATURES.md`'s routed-expert-streaming row still read "CPU keep-quant towers only" after W0c made the seam take the slot arm on `is_cpu()` OR `host_memory_is_device_addressable()` and W0d made the load-time fit refusal drop those towers from its bound. `AGENTS.md` routes a feature-surface change to that page, and the change that moved the surface did not write it. `scripts/check-doc-checkpoint.py` said so, on commit `939755f99` of `row/ENG-EXPERT-STREAM-DEVICE-W0`: a measurement was appended to `.agents/benchmark-record.md` with no `docs/FEATURES.md` edit beside it. The PAGE is FIXED IN FLOW while repairing the fresh review of [#1377](https://github.com/mudler/vllm.cpp/pull/1377): the row now names both the device arm ([#1124](https://github.com/mudler/vllm.cpp/issues/1124)) and the residency condition ([#1378](https://github.com/mudler/vllm.cpp/issues/1378)), and the observability detail it displaced to stay inside the 220-character cell budget is stated at `docs/USAGE.md:4598-4620`. The GATE is NOT fixed and needs a decision: the checker walks a range one COMMIT at a time, so once a commit is published on a branch that may not be force-pushed, no later commit can make it green, and `scripts/agent-preflight.sh` keeps reporting `doc-checkpoint range` red on this branch until it merges. The squashed commit that lands on `main` carries both paths and passes. Changing the walk is checker semantics and needs its own row, spec and red-first evidence per `AGENTS.md` "Changing the rules or a checker", so it is not folded in here. Spec [`expert-stream-device-slots.md`](specs/expert-stream-device-slots.md) | bug | | [#1386](https://github.com/mudler/vllm.cpp/issues/1386) | — | `tools/bench/gpu_clock_state.py`'s `QUERY_FIELDS` collects nine fields and **none of them is thermal and none is electrical**, so the driver's own `SwThermalSlowdown` label can never be checked against a die reading on any window this helper has ever recorded. The measured consequence is that the nine windows of 2026-08-19 cannot distinguish a load transition from a thermal excursion. The concrete evidence is `clock-c1-r1.samples.json` in `/mnt/nas_share/rc/q38bf16/out/bench-20260819T035148Z/`: ours c1 r1 dips five times on the same period at the same `utilization.gpu = 96` — 48.83 s / 2177 MHz, 80.60 s / 2320 MHz, 109.28 s / 2210 MHz, 137.98 s / 2359 MHz, 166.07 s / 2268 MHz — and **two of those five carry `0x0000000000000000`**, no throttle bit at all (2210 and 2359), while three carry `0x20`. The 2210 MHz unlabelled dip is deeper than two of the three labelled ones, so the driver labels comparable excursions inconsistently and the bit alone cannot decide it. What would settle it: add `temperature.gpu` and `power.draw` to `QUERY_FIELDS`. That changes the clock-record schema, so it owes its own row and spec. Split out of [#1354](https://github.com/mudler/vllm.cpp/issues/1354) and owed under `## Owed` in [lease-clock-pinning.md](specs/lease-clock-pinning.md) | gap | +| [#1290](https://github.com/mudler/vllm.cpp/issues/1290) | `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | The A3 gate reads `93/96 mode=decode DIVERGENCE` on `thor:gpu0` (sm_110) with the HOST mamba arm and `96/96 STRICT PASS` with the device FP8 W8A8 arm \-- same binary, box, checkpoint, golden and session, `reference-tier lines: 0` on both. On `dgx:gpu0` (sm_121a) the SAME host arm reads `96/96 STRICT PASS` (the #1157 run, `gate_fixed.out`), so the divergence is established on sm_110 ONLY and is NOT a general statement about `main` on a GPU. n=1 per arm. Leading mechanism: the golden's oracle computes these projections W8A8 while the host arm is W8A16 by construction (`DenseBf16` carries and does not apply `input_scale`, `nemotron_h.cpp:419-422`), so the perturbation is marginal and whether it flips a token depends on the rest of the tower \-- `fa2` provides for `12.1a` and not `11.0` (`CudaArchFeatures.cmake:349`), so GB10 runs FlashAttention-2 where Thor runs the fallback. NOT the fp8 GEMM: the host arm dequantizes and runs a CPU bf16 `MatmulBT`, so no fp8 GEMM executes on either box in that configuration. Next: the oracle's top-2 margin at the three moved tokens, then a `NemotronHTrace` layer bisect. A2-Q1 ([#1289](https://github.com/mudler/vllm.cpp/pull/1289)) makes the Thor gate read 96/96; whether anything was wrong on GB10 is open and this issue owns it | bug | +| [#1416](https://github.com/mudler/vllm.cpp/issues/1416) | `ENV-ORACLE-WHEEL-IN-LEASE` | The durable pinned-oracle wheel on `/workspace/oracle-vllm` cannot be `pip install`ed under the name it was staged with. `vllm-0.1.dev1+g555967922-FLASHINFER-ONLY-cp312-cp312-linux_aarch64.whl` has SEVEN `-`-separated parts where PEP 427 allows five (six with a build tag), so `packaging.utils.parse_wheel_filename` raises `InvalidWheelFilename` before pip opens the file — verified against that parser directly, and the conforming `vllm-0.1.dev1+g555967922-cp312-cp312-linux_aarch64.whl` parses. The bytes are fine: `METADATA` says `0.1.dev1+g555967922`, which the conforming name carries exactly, so only the marker is unrepresentable. `README-WHEELS.md` documents the usage as `pip install `, which cannot work as written. It cost a GB10 lease on 2026-08-19 (`/workspace/a2q1-neartie/20260819T215514Z`), and the failure is badly placed in time: the ~13-minute `pip install torch==2.13.0` runs first, so `RC[pip wheel]=1` lands long after the job looks healthy and a driver that does not stop there reports only `ModuleNotFoundError: No module named 'vllm'` at the measurement step — an instrument failure wearing the shape of a verdict about the model. FIXED IN FLOW: the driver copies to a conforming name and exits on a failed identity assert, the shared `README-WHEELS.md` carries a correction, and [`oracle-wheel-in-lease.md`](specs/oracle-wheel-in-lease.md) records the measurement. Found while measuring the oracle top-2 margin for [#1388](https://github.com/mudler/vllm.cpp/issues/1388) | bug | +| [#1431](https://github.com/mudler/vllm.cpp/issues/1431) | `ENV-ORACLE-WHEEL-IN-LEASE` | The pinned vLLM oracle loads the 20.1 GiB NemotronH-3.5-Lightning NVFP4 checkpoint on `dgx:gpu0` and is then killed by the host-memory watchdog during engine START-UP, before generating a token, on all five configurations tried (logs `/workspace/a2q1-neartie/`, box idle at start: `available 117456` MB, no compute apps). Host memory sits flat through the weight load and then falls **~76 GB in ~60 s at ~1.2 GB/s** immediately after the mamba page-size print. **The watchdog floor was NOT constant:** runs 1, 4 and 5 used `< 15000MB` and runs 2 and 3 used `< 20000MB`, so runs 2 and 3 were killed ~5 GB earlier in the drawdown (`watchdog.log` per run: 12597, 19433, 19797, 13941, 14846 MB). That runs 2 and 3 would also have crossed 15000 MB is an INFERENCE, not a measurement — the trace was still falling at ~1.18 GB/s with no arrest — and `job.sh` hard-codes a single `FLOOR_MB=15000` on unversioned CIFS after being overwritten between runs, so the artifact cannot reconstruct three of the five runs it produced. The runs EXCLUDE three of the four candidates the record had bundled together: not `torch.compile` (run 5 hit the AOT cache, `torch.compile took 0.30 s`), not CUDA graph capture (run 2 ran `enforce_eager=True`), and not KV sizing (run 3 `num_gpu_blocks_override=8`, run 5 `kv_cache_memory_bytes=4 GiB`). What is left is the first forward, consistent with the NVFP4 Marlin repack the record already names, and run 5 shows it POSITIVELY rather than by elimination: under an absolute KV budget `v1/worker/gpu_worker.py:465-468` runs `profile_run()` FIRST and only then logs `Initial free memory … skipped memory profiling` (`:470-482`), and run 5's log reaches NEITHER string, so the process died inside that forward. (Those two strings are quoted from the pinned SOURCE; they appear in no log, which is the point.) At this pin no environment knob forces a different NvFp4 MoE backend (`fused_moe/oracle/nvfp4.py:275-292` returns the first supported). **It is a REGRESSION:** run 4 runs `oracle_only.sh` attempt `a`'s ENGINE KWARGS unchanged — byte-for-byte of the kwargs, NOT of the environment — and on 2026-08-18 that configuration bottomed at `minMemAvailable_MB=51528` and generated all three prompts, so the evidence behind "the oracle DOES run a model in a lease" no longer reproduces. The scale of the change is the strongest single number here: 08-18 reported `Available RAM: 66.97 GiB` and peaked ~17 GB of host use, while run 4 reported `94.86 GiB` and burned **~83 GB** — same box, same engine kwargs, **5x the host footprint**. The venvs differ (`/tmp/nhspeed-oracle` against `/tmp/a2q1-oracle`, with a fresh `pip install torch==2.13.0` and a reinstalled wheel) and an identity assert on the vLLM COMMIT cannot see a toolchain delta, so the torch/flashinfer difference is an UNEXCLUDED candidate for the regression. Blocks the top-2 margin [#1388](https://github.com/mudler/vllm.cpp/issues/1388) needs, which holds [#1289](https://github.com/mudler/vllm.cpp/pull/1289). Found while measuring that margin | bug | +| [#926](https://github.com/mudler/vllm.cpp/issues/926) | `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | INDEXED LATE, and the index had never carried it: the issue has been open since 2026-08-15 and no row referenced it. **The Nemotron golden cannot be re-derived.** `tests/parity/goldens/nemotron_35_lightning_greedy/oracle.json` records the model, the revision, the sampling parameters and the vllm/transformers/flashinfer versions, and records NO engine configuration; `af8170154` committed no generator, because that capture ran from `$HOME/venvs/vllm-oracle-next` on GB10. So every comparison against it measures a known configuration against an unrecorded one. #926 already names `enforce_eager`, `gpu_memory_utilization`, `max_model_len` and the batch shape as unrecorded; the 2026-08-18/20 lease logs add one more with a named accuracy mechanism, `kv_cache_dtype=fp8_e4m3` AUTO-SELECTED on this checkpoint (`cache.py:296`, "it may cause accuracy drop without a proper scaling factor") over a checkpoint that carries no q scale, so vLLM imputes it (`kv_cache.py:134`, "Checkpoint does not provide a q scaling factor. Setting it to k_scale"). This is why `oracle.a.out` reproduced the committed golden `32/32` on prompts 0 and 1 and `26/32` on prompt 2: the oracle is CONFIG-SENSITIVE at fixed pin, and it is DETERMINISTIC at fixed configuration (`ORACLE_LEG 1` and `ORACLE_LEG 2`, one process, identical results, `ORACLE TOKEN MATCH: 180/192`). A ratified distributional gate is therefore INADMISSIBLE here — AGENTS.md permits one only where the oracle's greedy decode is non-deterministic — and what is licensed is re-deriving or re-pinning the golden against a NAMED engine configuration. **PRECEDES [#1431](https://github.com/mudler/vllm.cpp/issues/1431) and needs no GPU lease.** Owned by `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` and listed under `## Owed` in [`nemotron-h-a2q1-fp8-mamba.md`](specs/nemotron-h-a2q1-fp8-mamba.md). Re-found while repairing the A2-Q1 divergence record | bug | | [#1413](https://github.com/mudler/vllm.cpp/issues/1413) | `LTX25-DEVICE-RESIDENCY` | The LTX-2.5 render phase table [#1408](https://github.com/mudler/vllm.cpp/pull/1408) lands for [#1010](https://github.com/mudler/vllm.cpp/issues/1010) is written by the SUCCESS PATH ONLY — `WritePhaseLog` sits after `im.trace.completed = true` at `src/vllm/multimodal/ltx2_video.cpp:4655-4658 @ 4f3c24380` — so a render that is killed, aborted by a lease governor, or still running writes no table at all. That is the population the campaign actually has: [#1375](https://github.com/mudler/vllm.cpp/issues/1375) is `ABORT[92] PROJECTED OVERRUN`, `child exit=-15`, 0 frames; [`ltx25-decode-speed.md`](specs/ltx25-decode-speed.md) rung 1 is `EXIT=137`, 0 frames; rung 2 is `EXIT=1`, 0 frames. **And nothing is emitted WHILE a render runs**: `PhaseLog::Open`/`Close` print nothing and `VLLM_RENDER_PHASE_LOG_STDERR` fires inside `WriteJson`, i.e. on the success path again, so between `ltx2-gen: family=...` and `wrote N frames` a 2.5-hour render is silent and working is byte-identical to hung. **The unit that costs the wall has no counter**: `denoise` is one leaf covering ~2.7 h, against #1375's measured ~162 s per DiT forward at 60 structural forwards (30 steps x 2 CFG legs — `cfg_scale != 1.0` forces the unconditional branch at `ltx2_pipeline.cpp:521-523`). External sampling is NOT the fallback and was tried: #1375 records `phase=OTHER` throughout because `eu-stack` unwinds ZERO frames inside the `rc` worker container, and its own text names an in-process phase marker as the way to attribute the 162 s. Owed: a stderr line on every phase open and close, so the LAST LINE PRINTED names the phase in flight, plus one tick per DiT forward carrying phase, step `k/N`, cumulative forward index, elapsed and seconds-since-previous — on the shipped default, because the failure happened on default settings and `VT_H3_PROGRESS` (`minimax_h3.cpp:776-793`) is the opt-in shape that is exactly why no LTX-2.5 run has one. Precondition for attributing #1375, which caps render resolution. Stage W0-live in [`ltx25-device-residency.md`](specs/ltx25-device-residency.md) | feature | | [#1320](https://github.com/mudler/vllm.cpp/issues/1320) | `ENG-EXPERT-STREAM-DEVICE` | `VT_CPU_QUANT_REPACK` rewrites a Q8_0 weight into the `block_q8_0x4` i8mm interleave at load, and only the CPU `MatmulBTKernel` understands that layout — the CUDA quant dot reads plain `block_q8_0` and never consults `Tensor::repacked`, so a repacked weight uploaded to a device produces WRONG TOKENS rather than a crash. Its sibling `elem_kn_repack` has BOTH a CPU-platform gate in the loader policy (`gguf_keep_quant.cpp`, `CurrentPlatform().device_type() == kCPU`) and a refusal in `ResidentWeight`; `quant_repack` had NEITHER, because it rides `vt::cpu::QuantRepackActive()`, a HOST-CPU Arm i8mm probe that says nothing about where the weight will execute. An aarch64 box doing `--device cuda` — which is exactly `dgx:gpu0`, a Grace-Blackwell GB10 — satisfies it. Currently silent rather than firing, and that is measured: eligibility is Q8_0 with `n % 4 == 0 && k % 32 == 0`, `Qwen3.8-2.4T-A95B UD-Q1_0` has exactly one Q8_0 tensor at 0.01% of parameters, and the instrumented load of that checkpoint recorded `quant_repack = 0`. Found while implementing W0f ([#1299](https://github.com/mudler/vllm.cpp/issues/1299)), which changes the same function and whose G0-CORRECT gate compares CUDA tokens against CPU tokens — precisely the instrument that would have reported this as a W0f defect. FIXED IN FLOW as the tripwire its sibling already has, placed beside it so it covers the staging branch AND W0f's aliasing branch (where the bytes live was never the question; the kernel that reads them is), gated red-first on both arms and mutation-proven. `VT_CPU_QUANT_REPACK=0` is the operator's way past it. NOT fixed here, deliberately: adding the missing CPU-platform gate to `p.quant_repack` itself is a loader-policy semantics change owned by `QUANT-GGUF-KEEPQ-LOADER`, carried under `## Owed` in [`expert-stream-device-slots.md`](specs/expert-stream-device-slots.md) | bug | | [#1426](https://github.com/mudler/vllm.cpp/issues/1426) | `LTX25-GUIDED-VIDEO` | The LTX-2.5 x0 model picks its DiT forward with one ternary on `im.on_device` (`src/vllm/multimodal/ltx2_video.cpp`), and **no test on a box without an accelerator enters the DEVICE branch**, so dropping its `perturbations` argument is undetectable. MEASURED, both mutations built clean at `compile_err=0` and both restored byte-for-byte with the target rebuilt: **D10** (drop `p` from the device branch alone, 1 insertion / 1 deletion) is **GREEN**, `test_ltx2_video` exit 0, 88 cases / 2755 assertions; **D11** (drop it from BOTH branches, 3 insertions / 2 deletions) is **RED**, exit 1, 2 cases / 3 assertions, failing `CHECK(MaxAbsDiffOf(t.video_first_perturbed, t.video_first_cond) > 1e-6 * latent_span)`, the same for `video_first_modality`, and `CHECK(n.video_first_perturbed != n.video_first_cond)`. D11 also has to be written with `(void)p;` to be a RESULT at all, because the naive form leaves the lambda parameter unused and `-Werror=unused-parameter` refuses the build - a mutation that fails to build reads exactly like a passing test. This is a missing RUNNER and not a missing test: `Ltx2VideoEngine::Load` refuses `device != 0` in three places (`ltx2_video.cpp:749-758`, `:779-788`, and the queue-creation check below them) unless a non-CPU platform resolves, registers a backend and does not decline `kLtx2VideoFamily`, and `.github/workflows/ci.yml` has no GPU runner. A SYNTHETIC accelerator is DEFERRED rather than impossible, and the PARTS ALREADY EXIST. `vt::RegisterOp(OpId, DeviceType, void*)` is public (`include/vt/op_provider.h:127`) and is exactly how `src/vt/cpu/cpu_ltx2.cpp:201-205` installs the CPU glue table, so a test can read that pointer back with `vt::GetOp(OpId::kLtx2, kCPU)` and register it for a second device type; `RegisterReferenceTier` loops EVERY OpId (`src/vt/op_provider.cpp:795-807`) and `MaybeInstallReferenceTier` copies `src->fn` verbatim (`:204-225`), which covers the `vt::` ops the staging and the forward need; `Glue()` `VT_CHECK`-refuses BY NAME and never returns nullptr (`src/vllm/model_executor/models/ltx2_device.cpp:81-86`); and the closed 7-value `DeviceType` enum (`include/vt/device.h:16-28`) is handled the way `tests/vllm/multimodal/test_diffusion_device_seam.cpp` ALREADY handles it FOR THIS FAMILY — a SEPARATE EXECUTABLE (`:19-23`, `tests/CMakeLists.txt:317`) that registers a unified-memory `FakeXpuBackend` (`:60-73`), a platform in the XPU and CUDA slots (`:116-121`) and a `device = 1` `kLtx2VideoFamily` load (`:127-133`) — which leaves `test_ltx2_video`'s 88-case baseline byte-for-byte untouched. One design fact goes with that: `Ltx2DeviceKernelsAvailable` reads `vt::OpRegistered`, which EXCLUDES the reference tier deliberately (`op_provider.cpp:677-701`, "a FALLBACK, not a native kernel"), so the glue table needs the direct `RegisterOp` and the tier serves the rest. A fake also DOES prove link B as this row defines it — the x0 model reaching `Ltx2DitForwardDevice(..., p)`, which would turn D10 RED; what it cannot prove is that a GPU runs it, which is the hardware half the leased run below closes. WHAT IT COSTS is that new executable plus a fixture DiT staged through the fake backend (`tests/vllm/multimodal/ltx2_video_fixture.h` already writes a complete reduced model). TWO gates on the route were checked and pass by construction: `vt::RegisterBackend(kXPU, ...)` writes slot 0 (`src/vt/backend.cpp:81-84`) so `:800`'s `TryGetBackend(Device{kXPU, 0})` resolves for `params.device = 1`, and `vt::CreateQueue(Device)` takes its `entry.backend->CreateQueue()` arm (`backend.cpp:146-153`) whose `q.device == device` assertion the fake satisfies. TWO are NOT verified and NOBODY HAS BUILT IT: whether `Ltx2StreamDitToDevice` (`:847`) stages the checkpoint onto that queue, and whether the bf16 render that follows (`:812`) completes through the reference tier. PLAUSIBLE, NOT PROVEN — recorded in those terms because the first draft of this row asserted the opposite three facts as measured. What IS gated: `Ltx2DitForwardDevice` honours all four perturbation fields on the CPU backend, nine mutations RED (`tests/vllm/models/test_ltx2_device.cpp`). Closes on one leased `one_stage --device cuda` run at the model's own guider defaults recording `dit_forwards == 4 * dit_evaluations`. Same shape as [#1131](https://github.com/mudler/vllm.cpp/issues/1131). Listed under `## Owed` in [`ltx25-guided-video.md`](specs/ltx25-guided-video.md) §12.8 link B | bug | diff --git a/.agents/model-matrix.md b/.agents/model-matrix.md index 3c8bd3ae4..45df9fda9 100644 --- a/.agents/model-matrix.md +++ b/.agents/model-matrix.md @@ -282,7 +282,7 @@ Transformers compatibility is capability-driven and excluded from finite counts. | `MODEL-TEXT-mimo-v2-mi-mo-v2-flash-for-causal-lm` | `MiMoV2FlashForCausalLM` | `registry.py:166`; `vllm/model_executor/models/mimo_v2.py::MiMoV2FlashForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; FusedMoE/grouped GEMM; sliding-window attention | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-mimo-v2-mi-mo-v2-for-causal-lm` | `MiMoV2ForCausalLM` | `registry.py:167`; `vllm/model_executor/models/mimo_v2.py::MiMoV2ForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; FusedMoE/grouped GEMM; sliding-window attention | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-nemotron-nemotron-for-causal-lm` | `NemotronForCausalLM` | `registry.py:168`; `vllm/model_executor/models/nemotron.py::NemotronForCausalLM` | causal generation / text | model loader/forward; paged attention/KV | ☐ required | `INVENTORIED` | none | unassigned | -| `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | `NemotronHForCausalLM`, `NemotronHPuzzleForCausalLM` | `registry.py:169-170`; `vllm/model_executor/models/nemotron_h.py::NemotronHForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; FusedMoE/grouped GEMM; Mamba/SSM state; sliding-window attention. **2026-08-12: row spec committed ([#517](https://github.com/mudler/vllm.cpp/issues/517), [spec](specs/nemotron-h-model.md)); BLOCKED on `KERNEL-SSM-MAMBA` [#496](https://github.com/mudler/vllm.cpp/issues/496)** ([spec](specs/mamba2-ssd.md)) — the Mamba2 SSD core is unported. Beyond that kernel this arch owes three further things, none of which exist locally: non-gated `relu²` MoE (every grouped-MoE op we have is SwiGLU-shaped), ModelOpt `MIXED_PRECISION` per-module loading (`quantization/modelopt.py:2280` — NVFP4 W4A16 g16 experts and FP8 W8A8 mamba projections in one checkpoint), and the DeepSeek-style MTP head. Driver checkpoint `nvidia/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-NVFP4` (20.1 GiB, fits one GB10). **2026-08-17: A2-P lands the PAGED forward** ([#810](https://github.com/mudler/vllm.cpp/issues/810), [spec](specs/nemotron-h-a2p-paged-forward.md)): `ForwardNemotronHForCausalLM` selects `NemotronHPagedForward` whenever the runner supplies paged KV and recurrent state, so K/V go into the runner's pages at `attn_meta.slot_mapping` and the conv + SSM rows are gathered from and scattered to `GdnStateCache` at the metadata's state indices — the first time this architecture carries state between decode steps. G-SAFE narrows from three clauses to `num_reqs <= 1` (batching is A2-B). The row stays `INVENTORIED`: the A3 end-to-end token gate against the released checkpoint is PENDING on `dgx.casa`, and `lm_head` (A2-Q2b), the FP8 mamba projections (A2-Q1, [#940](https://github.com/mudler/vllm.cpp/issues/940)), MTP (W5) and GGUF (W7) are all still owed | ☐ required | `INVENTORIED` | none | unassigned | +| `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | `NemotronHForCausalLM`, `NemotronHPuzzleForCausalLM` | `registry.py:169-170`; `vllm/model_executor/models/nemotron_h.py::NemotronHForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; FusedMoE/grouped GEMM; Mamba/SSM state; sliding-window attention. **2026-08-12: row spec committed ([#517](https://github.com/mudler/vllm.cpp/issues/517), [spec](specs/nemotron-h-model.md)); BLOCKED on `KERNEL-SSM-MAMBA` [#496](https://github.com/mudler/vllm.cpp/issues/496)** ([spec](specs/mamba2-ssd.md)) — the Mamba2 SSD core is unported. Beyond that kernel this arch owes three further things, none of which exist locally: non-gated `relu²` MoE (every grouped-MoE op we have is SwiGLU-shaped), ModelOpt `MIXED_PRECISION` per-module loading (`quantization/modelopt.py:2280` — NVFP4 W4A16 g16 experts and FP8 W8A8 mamba projections in one checkpoint), and the DeepSeek-style MTP head. Driver checkpoint `nvidia/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-NVFP4` (20.1 GiB, fits one GB10). **2026-08-17: A2-P lands the PAGED forward** ([#810](https://github.com/mudler/vllm.cpp/issues/810), [spec](specs/nemotron-h-a2p-paged-forward.md)): `ForwardNemotronHForCausalLM` selects `NemotronHPagedForward` whenever the runner supplies paged KV and recurrent state, so K/V go into the runner's pages at `attn_meta.slot_mapping` and the conv + SSM rows are gathered from and scattered to `GdnStateCache` at the metadata's state indices — the first time this architecture carries state between decode steps. G-SAFE narrows from three clauses to `num_reqs <= 1` (batching is A2-B). The row stays `INVENTORIED`: the A3 end-to-end token gate against the released checkpoint is PENDING on `dgx.casa`, and `lm_head` (A2-Q2b), MTP (W5) and GGUF (W7) are all still owed. **2026-08-18: A2-Q1 lands the FP8 W8A8 MAMBA DEVICE ARM** ([#810](https://github.com/mudler/vllm.cpp/issues/810), [spec](specs/nemotron-h-a2q1-fp8-mamba.md)), on the shared seam [#940](https://github.com/mudler/vllm.cpp/issues/940) extracted for it: `NemotronHMamba2MixerDevice` mirrors the host arm op for op and both forwards select it at runtime, so the 23 Mamba2 layers no longer download the normed hidden, dequantize 890e6 fp8 elements to bf16 per token and upload the result. The recurrent rows the gather produced are advanced IN PLACE, and G-SAFE is untouched. The GB10 per-block numeric gate and the GPU-occupancy measurement the unit is accepted on are PENDING a lease | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-olmo-olmo-for-causal-lm` | `OlmoForCausalLM` | `registry.py:171`; `vllm/model_executor/models/olmo.py::OlmoForCausalLM` | causal generation / text | model loader/forward; paged attention/KV | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-olmo2-olmo2-for-causal-lm` | `Olmo2ForCausalLM`, `Olmo3ForCausalLM` | `registry.py:172-173`; `vllm/model_executor/models/olmo2.py::Olmo2ForCausalLM` | causal generation / text | OLMo-2-0425-1B dense LANDED (W0-W4). Impl: `include/vllm/model_executor/models/olmo2.h` + `src/vllm/model_executor/models/{olmo2,olmo2_weights,olmo2_registry}.cpp` (TWO `REGISTER_VLLM_MODEL` for `Olmo2ForCausalLM`+`Olmo3ForCausalLM`, one shared factory; reuses the shared dense glue). ZERO new compute kernel — the two "distinctives" both reduce to WIRING over landed ops: (1) PURE POST-NORM (`norm_after`, `olmo2.py:261-277`) = standalone `vt::RmsNorm` on each sublayer OUTPUT + plain `vt::Add` residual re-join, NO pre-norm (the GLM-4 standalone-output-norm op, pre-norms dropped); (2) FULL-WIDTH QK-norm (`olmo2.py:113-117,160-172`) = two standalone `vt::RmsNorm` over `[T,q_size]`/`[T,kv_size]` before NeoX rope (NOT the per-head `kAttnQkNormRope`). Reuse: SiLU SwiGLU (`kSiluAndMul`), NeoX RoPE (`RopeFromCache`/`RopeNeox`), GQA/MHA paged glue (`dense_attn_block.h`), merged qkv/gate_up loader. Checkpoint is F32 on-disk → loader downcasts f32→bf16 (round-to-nearest-even, `vt::F32ToBF16`) to match vLLM-bf16. UNTIED lm_head (`tie_word_embeddings:false`). Loader 179 tensors, zero missing/unmapped. Tokenizer: ONE shared-TU touch — `tokenizer.cpp` accepts the OLMo-2 Split `behavior=Removed,invert=true` encoding (equiv. to `Isolated` for the full-cover cl100k regex → `kLlama3`); guarded additive branch, no other checkpoint uses it. NO-BOS-verified real ByteLevel gate. Runs EAGER (bf16, no decode graph). Gate: 16/16 vs vLLM 0.25.0 (STRICT 13/16 + near-tie 3/16, max gap 0.094 nats, 0 forward-divergent) — speed pending. **OLMo-3 W5 (batch3) IMPLEMENTED** (guarded additive edits to `olmo2.{h,cpp,weights}`, diff-inert for OLMo-2 — re-run OLMo-2 gate 16/16 UNCHANGED): per-layer interleaved routing off `config.layer_types` — sliding_attention layers use plain NeoX rope (theta 500000) + finite window (masked at the FA kernel, inert for short contexts), full_attention layers use a precomputed YaRN cos/sin cache (get_rope yarn: factor 8, original 8192, mscale=yarn_get_mscale(8)=1.2079=config attention_factor, indexed by real positions); dtype-aware loader (OLMo-3 is BF16 on-disk vs OLMo-2 F32). **BUT the pinned vLLM 0.25.0 oracle CANNOT run `allenai/OLMo-3-1025-7B`** — its transformers version predates OLMo-3's nested per-layer-type rope schema: `olmo2.py:143` does `rope_parameters["rope_theta"]` → `KeyError: 'rope_theta'` (top-level rope_theta not folded into standardized rope_parameters), and forcing it surfaces `Unrecognized keys {'sliding_attention','full_attention'}` → `TypeError: unhashable type: 'dict'`. So there is NO pinned-oracle SACRED bar for OLMo-3 (DEP-blocked, spec D5); our engine LOADS + RUNS it (bootstrap loaded clean). W5 SACRED gate pending an oracle that can construct OLMo-3's rope config. | ✅ [sweep-olmo2](specs/sweep-olmo2.md) | `PARTIAL` | `test_olmo2_paged_engine` 16/16 (dgx, 92 assertions, batch3 re-run UNCHANGED); `test_olmo3_paged_engine` present (oracle-blocked, skips); loader 179 tensors zero-unmapped; registry resolves both arch strings; registration `src/vllm/model_executor/models/olmo2_registry.cpp:126`; test `tests/vllm/models/test_model_registry.cpp:118` | `CLAIM-SWEEP-OLMO2` (Claude Code opus-4-8) | | `MODEL-TEXT-olmo-hybrid-olmo-hybrid-for-causal-lm` | `OlmoHybridForCausalLM` | `registry.py:174`; `vllm/model_executor/models/olmo_hybrid.py::OlmoHybridForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; Mamba/SSM state; GDN/linear-attention state | ☐ required | `INVENTORIED` | none | unassigned | diff --git a/.agents/specs/nemotron-h-a2q1-fp8-mamba.md b/.agents/specs/nemotron-h-a2q1-fp8-mamba.md index 8c4f593f3..6b14ba29c 100644 --- a/.agents/specs/nemotron-h-a2q1-fp8-mamba.md +++ b/.agents/specs/nemotron-h-a2q1-fp8-mamba.md @@ -77,6 +77,24 @@ read or written, so no reconciliation is owed and none may be smuggled in. Recorded explicitly because the temptation runs the other way: a unit touching the conv looks like the place to fix the conv dtype. It is not. +**The persistent conv page IS bf16 today, and the guard is not what holds it +there.** `nemotron_h_registry.cpp:264` allocates the page as +`vt::DType::kBF16`, and `runner.cpp:885-893` sizes the buffer from +`vt::SizeOf(gdn_conv_cache_dtype_)`, so the ALLOCATION fixes the bytes. Only the +transient gather row is f32 (`nemotron_h_device.cpp:1954`, two `DBuf(d, +DType::kF32, …)`), which is `vt::GdnStateGather`'s op contract and matches +upstream's bf16 page. The falsification of "the page was widened to f32" is +therefore correct. + +The comment above the paged guard is not. It says the conv page is "never +widened to f32", and the `VT_CHECK` on the next line admits +`kBF16 || kF16 || kF32` (`nemotron_h_device.cpp:1581-1583`). `runner.cpp:643-644` +admits the same three. A `conv_dtype` mutation to f32 would sail past both +checks; nothing but the allocation would stop it, and a token gate cannot see the +doubled bytes. This is a record defect on `main` — the comment and the guard are +both older than this row — and it is listed under §11 rather than repaired here, +because A2-Q1 owns no line of it. + --- ## 2. Upstream anchors — `file:line` on both sides, at `555967922` @@ -418,26 +436,497 @@ stated removal condition is the device/paged runner path, which is A2-P's. ## 10. Now -**State at this commit:** spec only. No product code, no lifecycle change. - -A2-Q1 is **BLOCKED on #960** and is not claimable until it lands. When it does, a -fresh implementer claims this file, captures a RED per §5.3 first, and lands the -mamba device arm with G-SAFE untouched. A fresh reviewer — never the implementer -— runs the §5.3 mutations. - -**Three things to read before the first edit:** §3, so the Thor dependency is -consumed rather than re-derived; §1.1, so the conv-state dtype is left to A2-P; -and §5.2, so the bands are measured rather than chosen. +**State at this commit:** the device arm is IMPLEMENTED and REACHED, and its +binding measurement is PENDING a lease. + +`NemotronHMamba2MixerDevice` (`nemotron_h_device.cpp`) runs the whole block on +the device on the shared FP8 W8A8 seam, and BOTH forwards select it at runtime: +`NemotronHDeviceForward` (non-paged, discarding the recurrence) and +`NemotronHPagedForward` (production, advancing the gathered recurrent rows in +place). The selection is a runtime op-table query plus a weight-form predicate, +never a preprocessor guard, so a dense NemotronH or a device without the fp8 pair +keeps the host bounce. + +**Three design points differ from §4 and the reasons are recorded here rather +than left to be re-derived:** + +1. **The `Fp8Weight` pair is built ON FIRST DEVICE USE, not by the loader** + (§4.1/§4.2 put it in the loader). Building it in a `ResidentSlot` the weights + own is A2-Q2a's newer idiom and it is strictly better here: it does not double + the 890 MB fp8 tower in host memory at load, it does not upload anything on a + host-only run, and it leaves `rep.host_bytes` — the literal `18888922112` that + `test_nemotron_h_loader.cpp:310` pins — untouched, so §4.2's re-derivation + obligation does not arise at all. The e4m3 staging copy is released as soon as + `ResidentFp8` has uploaded it, so the peak cost of the conversion is one + projection. +2. **The upload IS accounted** (§4.4 expected the report to say it was short). + `dense_fp8::ResidentFp8` still does not call `load_stats::AddDeviceUpload` + — that is [#974](https://github.com/mudler/vllm.cpp/issues/974), unchanged, and + the shared header is not touched — so A2-Q1 accounts what IT uploads at the + site that causes it, exactly as `ResidentWeight` and `ResidentNvfp4` do. The + counter is then also the instrument the residency case reads, because an arm + that re-uploaded the tower every step returns identical numbers to one that + uploads it once. +3. **The paged selection carries an `ssm_dtype == f32` term.** + `vt::GdnStateGather` widens the page into an f32 working buffer by op + contract, and the HOST arm then narrows it back to `ssm_dtype` before the + mixer sees it. On a checkpoint whose `mamba_ssm_cache_dtype` is not f32 the + two arms would round differently and the per-block numeric gate would be + comparing two different computations. The released checkpoint resolves f32. + +**The comparison is NOT bit-comparable by construction, and §5.2's "measure the +band" is therefore binding rather than cautionary.** The host reference is +W8A16 — `DenseFor` dequantizes the fp8 weight to bf16 and leaves the activation +alone, as `DenseBf16` says outright — while the device arm is W8A8 as vLLM is. +The difference between them is the e4m3 activation quantization, and every band +in `tests/vllm/models/test_nemotron_h_mamba_device.cpp` is measured in the run +against a defect the fixture separates. + +### 10.1 Thor (sm_110) RAN the arm, and it answers §3's question + +Measured 2026-08-18 under an `rc` lease on `thor:gpu0`, product tree behaviourally +identical to this row's head (the later commits touch tests, scripts and one +comment only). `cmake --build -j 4` returned 0. The feature table read +**`ENABLED for [110]: 1 ; DISABLED cells: 7`** — only `marlin-nvfp4`, with +`cutlass-fp8` and both `scaledmm-c3x` cells DISABLED. + +**That configuration is the whole point.** §3.1 measured Thor as having half an +fp8 arm: the GEMM present through `kMatmulFp8CublasLt`, the activation quant +trapped in a CUTLASS-gated TU. #991 moved the registration out. The arm running +here, on a build with no CUTLASS fp8 at all, is what closes that question. + +`test_nemotron_h_mamba_device` reported **49 assertions** where a GPU-less box +reports 4, so the device path executed rather than skipping: + +| case | result | +|---|---| +| fresh block vs host reference, T=1 / 8 / 12 | agreed 0.164 / 0.282 / 0.309 against a measured band of 0.5; 128 / 1024 / 1536 elements examined; 3 widths covered | +| the fp8 tower uploads ONCE | first call 61760 B == expected 61760 B, second call 0 B | +| refuses a dense projection, refuses a missing `input_scale` | both threw | +| the carry across two legs | **FAILED, and the instrument was the defect** — see below | + +Neighbouring suites, same run, all `Status: SUCCESS!`: `test_nemotron_h_forward` +16/16 (5716 assertions), `test_nemotron_h_paged_forward` 12/12 (3256), +`test_nemotron_h_loader` 2/2, `test_nemotron_h_moe_device` 2/2 (29), +`test_ops_mamba2_ssd` 12/12 (2095), `test_ops_fp8_cpu` 5/5 (62). + +### 10.2 The carry gate banded a defect smaller than the noise it had to accept + +The case banded the SECOND LEG'S OUTPUT against the separation of a dropped +carry. Thor measured the second leg agreeing to 0.705 while a dropped carry +separated by only 0.205, so the derived band (0.102) sat BELOW the deviation a +FRESH leg already shows (0.164 at T=1, from the case above). + +That is §5.2's lesson arriving from the other direction. The two arms are W8A8 +against W8A16, so a fresh leg already disagrees by the e4m3 activation +quantization and a second leg compounds that with the same disagreement +propagated through the carried state. **A defect whose separation is smaller than +the noise the comparison must accept is not resolvable from that comparison**, and +widening the band until it passes is what §8.1 says to stop for. + +The repair gates what the carry IS: the STATE. A dropped carry hands the next leg +zeros, so the separation between the advanced state and a zeroed one is 1.0 by +construction — about six times the noise floor. The conv window and the SSM state +are banded separately, each against its own zeroed twin. The noise floor is +measured in the run at the same width and printed beside the separation, and the +second leg's output carries an assertion only when the separation exceeds twice +that floor, with the condition printed either way. + +The failure also exposed two fixture defects. `mamba_ssm_cache_dtype` was unset +and resolved bf16 — NOT the configuration the paged forward selects the device arm +for (`ssm_dtype == f32`), so the cheap arm was gating a path production does not +take. And the whole file was a skip on a GPU-less box, so a CPU-runnable case now +pins the op contract the split depends on. + +### 10.2b The repaired carry gate, re-run on hardware, and what it measured + +Thor re-run at the branch head: `test_nemotron_h_mamba_device` **5 cases, 63 +assertions, 0 failed**, and every neighbouring suite green again. The repair of +§10.2 therefore holds on the silicon that exposed the defect, and the numbers it +prints now justify the repair rather than merely passing: + +| quantity | measured | separation | band | margin | +|---|---|---|---|---| +| W8A8-vs-W8A16 noise floor, T=1 | **0.2465** | — | — | the reference for everything below | +| carried conv window | **0.1746** over 576 elements | 1.0 (zeroed) | 0.5 | 2.9x | +| carried SSM state | **0.0614** over 2048 elements | 1.0 (zeroed) | 0.5 | 8.1x | +| second leg's output | 0.7055 | 0.2045 | — | NO assertion, see below | + +**The data now proves the diagnosis that drove the repair.** A dropped carry +separates the second leg's output by 0.2045, while the noise the comparison must +accept is 0.2465 — the defect is genuinely SMALLER than the noise, so +`separation > 2 * noise_floor` is false and the case makes no assertion there, by +design. The state comparison carries the case instead, at 2.9x and 8.1x margins. +Had the original band survived, it would have been asserting on a quantity it +cannot resolve. + +**One diagnostic was defective and is fixed.** The line that reports whether the +second leg is resolvable printed `1` rather than its prose, because doctest +stringifies a `const char*` as a BOOL and the message streamed a `char*` ternary. +That line's whole job is to make "no assertion was made here" a STATED result +rather than a silent hole, so a version of it that cannot say what it means is +the same class of defect as the band it reports on. It now builds a +`std::string`; reproduced against doctest 2.5.2 both ways before the fix. + +### 10.2c The decode-window sampler works, and it quantifies the dilution + +Thor re-run, A3 gate with the arm ON, sampler starting only after +`engine loaded in Ns`: + +``` +RC[a3 on]=0 +[nemotron-h] engine loaded in 500.9s +[nemotron-h] TOKEN MATCH: 96/96 over 3 prompt(s) (full rows=3, short rows=0, mode=decode) +[nemotron-h] STRICT PASS +on: GPU busy in 240 of 564 DECODE samples = 42.55% busy +on: decode window 75.418 s (the engine load is OUTSIDE it) +on: per output token 0.785606 s +reference-tier lines in on: 0 +``` + +**42.55% over the decode alone against 15.33% over load+decode.** The load is +500.9 s and the decode is 75.418 s, so the old window was 87% load — the +dilution §10.3 diagnosed is now measured rather than argued. + +**The same cross-silicon defect was still live on the per-token line and is +fixed.** That run printed `ratio 54.7x` against vLLM's 0.014369 s, which is a +GB10 figure, for a decode measured on Thor. It is exactly the defect the +busy-fraction reporter carried, and fixing one surface while leaving its twin is +how a wrong comparison survives a correction. The ratio is now quoted only when +`ARCH` is the arch it was measured on; elsewhere the rate still prints and the +comparison is withheld by name. Both arms are pinned in +`tests/scripts/test_nemotron_h_a2q1_per_token.py`, and quoting the ratio +unconditionally reds the suite. + +### 10.3 The A/B, on the corrected instrument: the busy fraction ROSE + +Second Thor lease, fresh build and clone, sampler measuring the DECODE window +alone. Same binary, same checkpoint, same golden, differing only by +`VT_NEMOTRON_H_DEVICE_MAMBA`: + +| flag | mamba arm | A3 | exit | decode GPU busy | per output token | decode wall | +|---|---|---|---|---|---|---| +| `1` (default) | device FP8 W8A8 | `96/96 STRICT PASS` | 0 | **240/564 = 42.55%** | **0.785606 s** | 75.4 s | +| `0` | host, dequant to bf16 | `93/96 DIVERGENCE` | 1 | **700/3808 = 18.38%** | **5.633442 s** | 540.8 s | + +**THE BUSY FRACTION ROSE, WHICH IS THIS UNIT'S ACCEPTANCE TEST: 18.38% to +42.55%, +24.17 points, a 2.31x rise.** A decode token costs 7.17x less. Peak host +44070 MiB. `reference-tier lines: 0` on both arms, so neither took the portable +tier. + +**Read on the box it was taken on, and nowhere else.** These are sm_110 figures. +The 6.31% baseline and the 0.014369 s per-token reference are BOTH GB10's, so +neither supports a ratio against these numbers, and the instrument now withholds +both comparisons by name on any other arch. What is established is the ON/OFF +difference on one box, and that is exactly the comparison the hypothesis needed. + +**The `ratio 54.7x` and `ratio 392.1x` strings in that run's log are stale and +must not be quoted.** That job cloned before the per-token arch gate landed, so it +still printed the GB10 comparison unconditionally; the per-token VALUES are sound +measurements of that box, the ratios beside them are not. + +### 10.4 The divergence reproduces, so it is not n=1 on Thor + +Two independent Thor leases, separate builds and separate clones, agree exactly +on both arms: + +| run | arm ON | arm OFF | +|---|---|---| +| `20260818T222352Z` | 96/96 | 93/96 | +| `20260818T232910Z` | 96/96 | 93/96 | + +So the `n=1` caveat is lifted FOR THOR: the host arm's 93/96 is reproducible +there, and the device arm's 96/96 is too. GB10 remains n=1 in the other +direction (the #1157 run, host arm, 96/96). +[#1290](https://github.com/mudler/vllm.cpp/issues/1290) carries this. + +**A GB10 run of the DEVICE arm now exists**, and §10.5 records it. This +paragraph said it did not until 2026-08-20; the discriminator run of 2026-08-19 +superseded that and the sentence outlived it, so the file contradicted itself at +its own head. What GB10 still lacks is the SPEED axis, not a token result. + +**What is still owed.** The occupancy hypothesis is now SUPPORTED on sm_110 and +UNMEASURED on sm_121a, and the two are not interchangeable. Specifically owed: + +- the GB10 OCCUPANCY run, which is the only one that can be read against the + 6.31% baseline and the 0.014369 s per-token reference. Its A3 TOKEN result was + taken on 2026-08-19 (§10.5); the per-token and busy-fraction numbers on `121a` + were not; +- the §5.1 per-block numeric gate against `trace.mixer[l]` on the real checkpoint + — the A3 gate is token-level and cannot see a per-layer defect whose argmax is + unchanged; +- the §5.3 mutations, which belong to the fresh reviewer; +- the three moved tokens on the OFF arm, whose oracle top-2 margin decides + whether [#1290](https://github.com/mudler/vllm.cpp/issues/1290) is a near-tie + sensitivity or a wrong answer. + +`scripts/nemotron-h-a2q1-dgx-gate.sh` is the recipe, and it now measures the +decode window, withholds both GB10 references off `121a`, and refuses a fraction +outright when the load boundary never appears. + +### 10.5 The moved token is UNDETERMINED, it belongs to the device mamba block, and the measurement that decides it is blocked + +**Result: NEEDS_DECISION, undetermined. [#1289](https://github.com/mudler/vllm.cpp/pull/1289) +stays held.** The GB10 three-leg discriminator localized the moved token. It did +not exonerate the arm. + +**An earlier version of this section is withdrawn.** It was headed "three +independent lines say it is a tie" and it read each line as exculpatory. Every +fact it stated is correct and is kept below. The reading was not: line 1 narrows +suspicion ONTO the changed path rather than away from it, line 2 excludes only a +defect nobody proposed, and line 3 is refuted by the golden's own arithmetic and +by the GB10 polarity. The bottom line did not move, because it was already +"undetermined". The argument for it is now the one the evidence supports. + +#### 10.5.1 The discriminator run, and why nothing in this tree can regenerate it + +| field | value | +|---|---| +| run | `20260819T200231Z` | +| evidence | `/mnt/nas_share/rc/a2d1-discriminate/20260819T200231Z/` — the same directory a lease sees as `/workspace/a2d1-discriminate/20260819T200231Z/` | +| box | `dgx:gpu0`, GB10, `sm_121a`, under an `rc` lease | +| when | 2026-08-19, `configure.log` 20:02Z, last leg 21:16Z | +| tree | `/root/src-a2d1d` — a clone of this row's branch PLUS an uncommitted diagnostic, see below | +| toolchain | CUDA `13.0.88`, GCC `13.3.0`, `CUDA target architectures: 121a` | +| features | `fa2`, `cutlass-fp8`, `cutlass-nvfp4`, `marlin-nvfp4` ENABLED (`features.txt`) | +| library | `libvllm 0.0.3+cuda (ABI 21, header 21)` | +| checkpoint | `/workspace/a3/ckpt-stage`, revision `29f2d1746d8f41e316523194b19018707749b1b1` | +| golden | `tests/parity/goldens/nemotron_35_lightning_greedy/oracle.json`, oracle `vllm=0.23.1rc1.dev1511+g555967922` | +| contention | `contention.txt` holds the CSV header only: no other compute application held the device | + +★ **THE COUNTERS BELOW CANNOT BE REGENERATED FROM THIS TREE, AND NO BINARY HASH +WAS TAKEN.** The `[NH-DIAG] ARM step …` line that carries `state_update_rows`, +`chunk_scan_calls`, `conv_update_rows`, `conv_fwd_calls`, `gathers` and +`scatters` came from an UNCOMMITTED edit to `nemotron_h_device.cpp` in +`/root/src-a2d1d`. None of those six names occurs anywhere in `src/`, and the +committed `[NH-DIAG]` facility prints a different line +(`nemotron_h_device.cpp:1764`, `[NH-DIAG] step T=… R=… nd=… np=… idx=[`). The +logs are real and they are readable at the path above. Nobody can rebuild the +instrument that produced them from a checkout, and nobody can prove which source +the binary was built from. Treat the counters as a recorded observation of one +unreproducible build, never as a gate result. Committing that instrument is owed +in §11. + +#### 10.5.2 The three legs + +All three legs ran on that one box, from that one build, in that one lease. + +| leg | log | decode-step recurrent counters | A3 | +|---|---|---|---| +| device arm ON | `a3_on.log` | `state_update_rows=23 chunk_scan_calls=0 conv_update_rows=23 conv_fwd_calls=0 gathers=0 scatters=0` | `95/96` DIVERGENCE | +| device arm OFF | `a3_off.log` | `state_update_rows=0 chunk_scan_calls=23 conv_update_rows=0 conv_fwd_calls=23 gathers=46 scatters=46` | `95/96` DIVERGENCE | +| host mamba | `a3_hostmamba.log` | `state_update_rows=0 chunk_scan_calls=0 conv_update_rows=0 conv_fwd_calls=0 gathers=46 scatters=46` | `96/96` STRICT PASS | + +Both device legs lose the same token, on the same prompt, at the same position: +prompt 2, generation position 32, `11286` where the golden holds `3468`. +Positions 1 to 31 are byte-identical in every leg. The host leg is exact. + +**What that settles.** The moved token belongs to the DEVICE MAMBA BLOCK. The +attention, MoE, norm and sampling paths are common to all three legs and cannot +account for a difference between them. + +#### 10.5.3 What the three lines actually support + +**Line 1 — the two device legs run different recurrent kernels and lose the same +token. This narrows suspicion ONTO the changed path.** The counters exclude the +recurrent reduction order and the gather/scatter state indexing, because neither +is invariant across `state_update_rows=23 gathers=0` and `chunk_scan_calls=23 +gathers=46`. The third leg is what the earlier reading omitted: `a3_hostmamba` +is `96/96`, so the token is inside the device mamba block, and the recurrent +kernel is now excluded INSIDE it. What both failing legs still share, and the +passing leg does not, is the FP8 W8A8 projection path — which is what +[#1289](https://github.com/mudler/vllm.cpp/pull/1289) adds. The elimination +points at the change, not away from it. + +**Line 2 — the polarity flips with silicon. This excludes only a GROSS +systematic defect, which nobody proposed.** The device arm reads `96/96` on +`thor:gpu0` and `95/96` on `dgx:gpu0`; the host arm reads `93/96` on Thor and +`96/96` on GB10. A defect that corrupted every FP8 projection could not be +`96/96` on Thor, and that is the whole of what the flip rules out. A MARGINAL +perturbation whose token flip depends on the rest of the tower survives it +untouched — and that is the mechanism [#1290](https://github.com/mudler/vllm.cpp/issues/1290) +already applies to the HOST arm. Applying it to one arm and refusing it to the +other is not a reading of the evidence; it is a choice about which arm to +protect. + +**Line 3 — the more precise arm is the worse tracker. This is refuted twice.** +The facts hold. Our host arm is W8A16: `nemotron_h.cpp:416-422` selects +`kFp8W8A8Static`, calls `DequantFp8ToBf16`, and says "Weight-only: input_scale +is carried, not applied". The oracle is W8A8: +`FlashInferFP8ScaledMMLinearKernel.input_quant_key()` returns +`kFp8StaticTensorSym` and `apply_scaled_mm` passes `scale_a` +(`vllm/model_executor/kernels/linear/scaled_mm/flashinfer.py:67-85`), +confirmed at run time by the +oracle's own startup line, `Selected FlashInferFP8ScaledMMLinearKernel for +ModelOptFp8LinearMethod`. + +The inference from those facts does not follow. + +1. **The golden IS a W8A8 computation.** The quantity measured is + difference-from-reference, not absolute accuracy. A W8A16 arm differs from a + W8A8 reference by the FULL activation-quantization error that the reference + applied and the arm did not. A W8A8 arm carries a rounding difference against + it instead. "More precise" therefore predicts a LARGER distance from this + golden, which is the ordering observed on Thor. The line reads a confirming + observation as a contradiction. +2. **On GB10 the ordering REVERSES.** Host (W8A16) is `96/96`; device (W8A8) is + `95/96`. The earlier text cited Thor only. Neither ordering survives both + boxes, so neither supports a conclusion about either arm. + +#### 10.5.4 The oracle is CONFIG-SENSITIVE here, not non-deterministic — and that decides the gate FORM + +`/workspace/nhspeed/oracle.a.out` (2026-08-18, `dgx:gpu0`) ran the SAME +configuration twice inside one process. `ORACLE_LEG 1` and `ORACLE_LEG 2` each +returned `matched=32`, `matched=32`, `matched=26`, for `ORACLE TOKEN MATCH: +180/192`. **At a fixed configuration the pinned oracle is deterministic.** The +committed `32/32` golden came from a DIFFERENT, and unrecorded, configuration. + +This is [#926](https://github.com/mudler/vllm.cpp/issues/926), open since +2026-08-15 and not previously linked from this spec. +`tests/parity/goldens/nemotron_35_lightning_greedy/oracle.json` records the +model, the revision, the sampling parameters, and the vLLM, `transformers` and +`flashinfer` versions. It records NO engine configuration, and `af8170154` +committed no generator — that capture ran from `$HOME/venvs/vllm-oracle-next`. +So the `26/32` result compares a known configuration against an unrecorded one, +and the difference cannot be attributed to anything. #926 already names +`enforce_eager`, `gpu_memory_utilization`, `max_model_len` and the batch shape as +the unrecorded terms; the lease logs add one more, and it is the one with a named +accuracy mechanism. `kv_cache_dtype=fp8_e4m3` is AUTO-SELECTED on this checkpoint +(`cache.py:296`, "it may cause accuracy drop without a proper scaling factor") +and the checkpoint does not carry the q scale, so vLLM imputes it: +`kv_cache.py:134`, "Checkpoint does not provide a q scaling factor. Setting it to +k_scale." + +**#926 also already recorded the same fork.** Its rebuilt oracle, at this pin +under its own configuration, diverged on prompt 2 at index 29 to `11286` +` transition` — which is the Thor host arm's continuation, position for position. +A pinned vLLM oracle takes that fork. That is evidence about the POSITION, and it +still says nothing about the device arm's position 32. + +**A distributional gate is INADMISSIBLE on this evidence.** AGENTS.md permits a +ratified distributional gate only when the oracle's greedy decode is +NON-DETERMINISTIC. `oracle.a.out` measured the opposite. What the evidence +licenses is re-deriving or re-pinning the golden against a NAMED engine +configuration, which is #926's own "what done looks like". Nothing in this +section may be read as a case for weakening the token gate. + +#### 10.5.5 The order of work, and why the margin is not first + +The oracle's top-2 margin at generation position 32 decides whether this is a +near-tie sensitivity or a wrong answer. `scripts/nemotron-h-a2q1-neartie-gap.py` +is the instrument. It has never produced a number: five runs on `dgx:gpu0` were +each killed by the host-memory watchdog during engine start-up, before a token +existed. That is [#1431](https://github.com/mudler/vllm.cpp/issues/1431). + +**[#926](https://github.com/mudler/vllm.cpp/issues/926) precedes +[#1431](https://github.com/mudler/vllm.cpp/issues/1431).** It needs no GPU lease. +Until the golden names the configuration it was captured under, a margin measured +against it is a margin against an unknown, and a `26/32` cannot be told from a +defect. Recording the configuration, or re-capturing under a recorded one, is +cheap and it is prior. + +Decoded, the four continuations are one sentence taking a different fork after +"…observed outputs, typically using a": + +| arm | positions 30, 31, 32 | +|---|---| +| vLLM golden | ` combination`, ` of`, ` state` | +| Thor `sm_110`, host arm | ` transition`, ` equation`, ` for` | +| GB10 `sm_121a`, device arm | ` combination`, ` of`, ` transition` | +| rebuilt oracle, #926, own config | ` transition`, ` equation`, ` for` | + +The competing token is the same one, `11286` ` transition`, in every divergent +row. + +### 10.6 REFUTED: the f32 SSM cache is not a too-wide dtype, and it is not mirrored from `transformers` + +Recorded because it was raised as a defect during review, is wrong, and is easy +to raise again from the anchors this spec already carried. + +**The hypothesis.** `nemotron_h_weights.cpp:855` reads `mamba_ssm_cache_dtype` +from the HF config with a `"float32"` default. vLLM's `CacheConfig` defaults the +same name to `"auto"` (`vllm/config/cache.py:135`), and `"auto"` resolves the +temporal state to the conv state dtype, which is the model dtype, bf16 +(`mamba_utils.py:101-108`). So we appear to hold 2x the temporal-state bytes, and +to take the value from `transformers` where vLLM defines something else — which +AGENTS.md forbids, and which a token gate cannot see. + +**Why it is wrong.** vLLM does not stop at `CacheConfig`. It has a per-model +verify hook, and NemotronH is one of the models that has one. +`vllm/model_executor/models/config.py:605-631` at the pin declares +`class NemotronHForCausalLMConfig(VerifyAndUpdateConfig)` with +`DEFAULT_MAMBA_SSM_CACHE_DTYPE = "float32"`, documented in its own docstring as +"Only `float32` is known to have no accuracy issues by default". Its +`update_mamba_ssm_cache_dtype` fires when `cache_config.mamba_ssm_cache_dtype == +"auto"`, reads `getattr(hf_config, "mamba_ssm_cache_dtype", +cls.DEFAULT_MAMBA_SSM_CACHE_DTYPE)`, and writes the result back into +`cache_config`. It is registered for `NemotronHForCausalLM` and +`NemotronHPuzzleForCausalLM` at `config.py:879-880`. + +So the hook reads the key FROM THE HF CONFIG and defaults it to FLOAT32 when the +key is absent. `nemotron_h_weights.cpp:855` mirrors vLLM exactly: same key, same +source, same default. The `"auto"` at `cache.py:135` is the CLI default BEFORE +the hook runs, not the value the model gets. + +**Observed, not only read.** The pinned oracle logs the hook firing on this +checkpoint in five separate lease runs, including the 2026-08-18 run behind +`oracle.a.out`: + +```text +INFO 08-18 20:59:29 [config.py:621] Updating mamba_ssm_cache_dtype to 'float32' for NemotronH model +``` + +The checkpoint's own `config.json` carries `mamba_ssm_cache_dtype: "float32"`, so +the hook takes the checkpoint's value and the default is not even reached. + +**What to fix instead.** Nothing in the dtype. The defect is the CITATION. The +comments at `nemotron_h_registry.cpp:256-263` and `nemotron_h_device.cpp:1586` +cite `mamba_utils.py` alone, which is the half of the chain that makes the value +look unmirrored. Any future reader who checks those anchors reaches the same +wrong conclusion. Add `config.py:605-631` beside them — owed in §11. ## 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). +- [#926](https://github.com/mudler/vllm.cpp/issues/926) — the golden records no + engine configuration and has no committed generator, so it cannot be + re-derived. **This PRECEDES the top-2 margin + ([#1431](https://github.com/mudler/vllm.cpp/issues/1431)) and needs no GPU + lease** (§10.5.4, §10.5.5). Until it closes, every comparison against + `oracle.json` measures a known configuration against an unknown one. +- The `[NH-DIAG] ARM step` recurrent counters that §10.5.1 quotes exist only in + an uncommitted edit on a lease box. Commit the instrument, or the + discriminator cannot be re-run. +- **Upstream anchor drift, for #1289's reviewer, NOT repaired here.** Checked + against the pin `555967922`. Introduced by + [#1289](https://github.com/mudler/vllm.cpp/pull/1289), in §2 above and in + `nemotron_h_device.cpp`: `mamba_mixer2.py:550` is the `in_proj` call at + **554**; `:586` is the `out_proj` call at **585**; `:583` and `:583-585` for + the gated norm point at the `gate` slice, and `self.norm(ssm_output, gate)` is + at **582**. Already on `main` and therefore wider than #1289: + `nemotron_h.py:440` for `self.scaling = self.head_dim**-0.5` is **441** and is + cited three times (`nemotron_h_device.cpp:332`, `:1882`, `nemotron_h.cpp:668`); + `nemotron_h.py:627-631` for `residual is None` holds no such statement — the + four occurrences are at 307, 350, 399 and 521, and `residual = None` is at 618; + `gdn_attn.py:405` names the branch that ASSIGNS `prefill_has_initial_state`, + while the comment describes the `else` at **406-407**. +- **The conv-page guard admits `kF32` while its own comment forbids it** + (§1.1). Pre-existing on `main`; A2-Q1 owns no line of it. +- The `mamba_ssm_cache_dtype` comments at `nemotron_h_registry.cpp:256-263` and + `nemotron_h_device.cpp:1586` cite `mamba_utils.py` without the per-model verify + hook that supplies the value, so the anchors read as if we mirror + `transformers`. §10.6 records the refutation; the anchors still need + `vllm/model_executor/models/config.py:605-631` beside them. ## 12. Outcome diff --git a/.agents/specs/oracle-wheel-in-lease.md b/.agents/specs/oracle-wheel-in-lease.md index 13d939454..6bcbd72c3 100644 --- a/.agents/specs/oracle-wheel-in-lease.md +++ b/.agents/specs/oracle-wheel-in-lease.md @@ -191,6 +191,47 @@ at its 90-minute ceiling and its partial tree was removed. - Stop if a correction needs an edit to an existing `.agents/issue-index.md` row. That file is append-only. +## The durable wheel is NOT installable under the name it was staged with + +Measured 2026-08-19 on `dgx:gpu0`, run `/workspace/a2q1-neartie/20260819T215514Z` +([#1416](https://github.com/mudler/vllm.cpp/issues/1416)). The artifact this spec +calls durable is staged as +`vllm-0.1.dev1+g555967922-FLASHINFER-ONLY-cp312-cp312-linux_aarch64.whl`, and +`pip` refuses it before opening the file: + +``` +ERROR: Invalid wheel filename (wrong number of parts): +'vllm-0.1.dev1+g555967922-FLASHINFER-ONLY-cp312-cp312-linux_aarch64' +``` + +PEP 427 allows five `-`-separated parts, six with a build tag. The +`FLASHINFER-ONLY` marker adds two more, so +`packaging.utils.parse_wheel_filename` raises `InvalidWheelFilename`. The bytes +are unaffected: the wheel's `METADATA` version is `0.1.dev1+g555967922`, exactly +what a conforming name carries. + +**Copy to a conforming name; do not rename the staged artifact**, whose name is +what identifies it to a reader: + +```sh +WHEEL_OK=/tmp/vllm-0.1.dev1+g555967922-cp312-cp312-linux_aarch64.whl +cp "$WHEEL" "$WHEEL_OK" && pip install "$WHEEL_OK" +``` + +**Two facts about WHERE this failure lands, because they are what made it +expensive.** `pip install -q torch==2.13.0` must run first and takes about +thirteen minutes, so `RC[pip wheel]=1` arrives long after the job looks healthy; +and a driver that does not stop there reports the failure only as +`ModuleNotFoundError: No module named 'vllm'` at the measurement step. That is an +infrastructure failure presenting as a verdict about the model, which is the +shape this repository has been caught by before. A lease driver that installs +this wheel therefore asserts the identity and **exits** on failure rather than +continuing. + +Everything else in that job was green, so nothing else here is in doubt: +`RC[apt-get install]=0`, `nvcc` already present at `cuda_13.0.r13.0` with no +install needed, `RC[pip torch]=0`, and `cuda True NVIDIA GB10`. + ## Owed - [#1185](https://github.com/mudler/vllm.cpp/issues/1185) stays open. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 7f9b0bc2e..e836d5032 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -13,6 +13,8 @@ | **Developer/row protocol** | Contribution entry point; `ENG-NOW-DERIVED` #374 @`dbd0d51c` | Entry-point gates retained. #374 W1-W5 DONE; benchmark/runtime/parity `VOID`; row specs now carry `## Now` | n/a | | **NemotronH paged forward** (`MODEL-NEMOTRON-H-ABI-A2P`, [#810](https://github.com/mudler/vllm.cpp/issues/810)) | **No speed number, by the unit's own rule** ([spec](../.agents/specs/nemotron-h-a2p-paged-forward.md) §5) | **A3 host gate PASSES 96/96 `STRICT PASS`.** GB10 read 4/24; cause and fix [#1157](https://github.com/mudler/vllm.cpp/issues/1157), sm_121a re-run pending a lease | CPU gate 12/12. Load 209.0 s, peak RSS 20 142 392 KB | | **Record-anchor ratchet** (`ENG-RECORD-ANCHOR-RATCHET`, #632) | **No number owed:** a record checker. At `8daa67b39`, **832 of 867** in-scope citations (**96.0%**) were already parsed; no symbol test and no report ran. Rot **38** (32 stale, 6 broken), **32 in range** | +| **Record-anchor ratchet** (`ENG-RECORD-ANCHOR-RATCHET`, #632) | **No number owed:** a record checker. At `8daa67b39`, **832 of 867** in-scope citations (**96.0%**) were already parsed; no symbol test and no report ran. Rot **38** (32 stale, 6 broken) over **844** OK, **32 in range** | +| **NemotronH FP8 mamba on device** (`A2-Q1`, [#810](https://github.com/mudler/vllm.cpp/issues/810)) | Thor sm_110, 20.1 GiB checkpoint, c1 greedy, same-binary A/B, n=2 | **Decode GPU busy ROSE 18.38% -> 42.55% (2.31x); per output token 5.6334 -> 0.7856 s (7.17x).** A3 `96/96 STRICT PASS` on; `93/96 DIVERGENCE` off ([#1290](https://github.com/mudler/vllm.cpp/issues/1290)) | Thor figures; the 6.31% and 0.014369 s references are GB10's and NOT comparable. GB10 leg pending. Peak host 44070 MiB | | **LoRA runtime W2** (`LORA-RUNTIME`, #278) | **No number owed:** correctness-only; a grid PENDS the W7 model gate | | **ARCH audit: ABI is text-only** | 4 capabilities (H3 video, Laguna, Kimi-Linear, DeepSeek-V4) reachable only from `examples/`, none registry-backed. No gate asks whether a CONSUMER can reach a capability. Documentation only | | **CUDA-graph break seam W1** (`ENG-CUDAGRAPH-BREAK`, [#1192](https://github.com/mudler/vllm.cpp/issues/1192)) | **No number owed, and none taken:** coverage and correctness row, no throughput gate declared | Capability, not a rate: mid-forward capture re-begin holds on a leased GPU; scoped forward matches eager, 500 logits, 0 differing. Probe committed, recipe and sha256 in the [record](../.agents/benchmark-record.md) | diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index dcbee45e2..db981553e 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -245,6 +245,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 5ec0a74ac..e54821d47 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 85aa52503..fd2a12e8d 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -140,7 +140,7 @@ token-for-token correctness against the pinned oracle. | GLM-4 dense (sandwich norms, partial rope) | Correctness-complete, speed-pending | Token-exact 16/16 (GLM-4-9B-0414); first GLM-family model; partial interleaved RoPE + Gemma2 sandwich norms + biased qkv | | GLM-4.7-Flash (MLA + GLM MoE) | Correctness-complete, speed-pending | Token-exact 8/8 (GLM-4.7-Flash, 31.2B); reuses the DeepSeek-V2 MLA stack; first e2e coverage of the q_lora query branch + noaux_tc sigmoid router with routed-scaling | | Kimi-Linear-48B-A3B (KDA + NoPE-MLA + MoE hybrid) | **RUNNER FOLD LANDS (ROW 7 §21, #122): engine==CLI 128/128 byte-identical; golden 122/128 (near-tie profile); FA2 MLA default-ON; `vllm_complete_tokens` (ABI v13).** Grouped-router top-k block-parallel (byte-identical); no binding speed number: ckpt is tiktoken-only, so no warm-server harness. STRICT stays CLOSED. Server 19.0 tok/s wall (~0.90× vLLM floor) = speed open | paged suite 8/8·206; SACRED post-fold 35B 315/315 + 27B 235/235; thin ABI client (ratchet 8) | -| Nemotron-3.5-Lightning-30B-A3B (Mamba2 + GQA + relu2 MoE) | **Host gate PASSES 96/96 `STRICT PASS`; GB10 fixed ([#1157](https://github.com/mudler/vllm.cpp/issues/1157)), sm_121a re-run pending** | ABI-only driver; G-SAFE `num_reqs <= 1`. The paged forward now honours `device_token_ids`; seam [#1217](https://github.com/mudler/vllm.cpp/issues/1217) | +| Nemotron-3.5-Lightning-30B-A3B (Mamba2 + GQA + relu2 MoE) | **Host and Thor sm_110 gates PASS 96/96 `STRICT PASS`, the latter with the device FP8 mamba arm; sm_121a re-run pending** | ABI-only driver; G-SAFE `num_reqs <= 1`. The 23 FP8 Mamba2 blocks reach the device (A2-Q1, [#810](https://github.com/mudler/vllm.cpp/issues/810)); `lm_head` is the last host arm | | Gemma-3 dense (GeGLU, dual rope, sandwich norms) | Correctness-complete, speed-pending | STRICT token-exact 48/48 greedy (gemma-3-1b-it); first Gemma-family model; GeGLU (gelu_pytorch_tanh) + dual per-layer RoPE theta + Gemma-RMSNorm sandwich norms + sqrt(hidden) embed-scale + query_pre_attn_scalar scaling | | Gemma-2 dense (attn + final logit soft-cap) | Correctness-complete, speed-pending | Near-tie-band 48/48 (gemma-2-2b-it): 44/48 strict on vLLM's greedy + 4/48 at 0.0-nat ties in vLLM's own logits; proves the attention + final logit soft-cap primitives (attn_logit_softcapping 50 + final 30); the inverse of Gemma-3 (both soft-caps, no QK-norm) | | Gemma-1 dense (the original Gemma) | Correctness-complete, speed-pending | STRICT token-exact 48/48 greedy (gemma-2b); two fused norms/layer, head_dim scale, GeGLU + sqrt(hidden) embed-scale, tied lm_head; no soft-cap/QK-norm/sliding. **D1 (2026-07-31): the whole Gemma family (1/2/3/4) folded to the default-ON bf16 merged-QKV descriptor (`MergedQkvEnabled`); re-gated Gemma-2 SACRED 48/48 (global+sliding) + Gemma-4 STRICT 32/32 — its existing gate held** | diff --git a/docs/USAGE.md b/docs/USAGE.md index b662d1800..6a68296a8 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -872,7 +872,7 @@ tokens quietly. | Architecture | Why it refuses | |---|---| | `KimiK3ForConditionalGeneration` | Needs ~1.56 TB (MXFP4); no host here can run it | -| `NemotronHForCausalLM` | **Only BATCHED decode still refuses.** A2-P (#810) narrowed this: `ForwardNemotronHForCausalLM` now selects the paged forward whenever the runner supplies paged KV and recurrent state, so K/V go into the runner's pages and the conv/SSM rows are carried across steps, and `examples/nemotron_h_gen` reaches all of it through `include/vllm.h` alone. What is left is `num_reqs > 1`, refused by name because one request's pages and one request's recurrent state are carried per step and a multi-request step would be decoded as ONE concatenated causal sequence — plausible wrong tokens rather than a failure. Owed to A2-B. **The end-to-end token gate against the pinned oracle has NOT run**, so no claim is made here about what this checkpoint emits; `docs/BENCHMARKS.md` records that as pending rather than as silence. `lm_head` and the FP8 Mamba2 projections still compute on the host, and a GGUF file is refused by name since no GGUF arm exists for it. See *Nemotron-3.5-Lightning-30B: the exact weights, and which arms run* below | +| `NemotronHForCausalLM` | **Only BATCHED decode still refuses.** A2-P (#810) narrowed this: `ForwardNemotronHForCausalLM` now selects the paged forward whenever the runner supplies paged KV and recurrent state, so K/V go into the runner's pages and the conv/SSM rows are carried across steps, and `examples/nemotron_h_gen` reaches all of it through `include/vllm.h` alone. What is left is `num_reqs > 1`, refused by name because one request's pages and one request's recurrent state are carried per step and a multi-request step would be decoded as ONE concatenated causal sequence — plausible wrong tokens rather than a failure. Owed to A2-B. **The end-to-end token gate against the pinned oracle has NOT run**, so no claim is made here about what this checkpoint emits; `docs/BENCHMARKS.md` records that as pending rather than as silence. A2-Q1 puts the 23 FP8 W8A8 Mamba2 blocks on the device, so `lm_head` is the last arm that computes on the host, and a GGUF file is refused by name since no GGUF arm exists for it. See *Nemotron-3.5-Lightning-30B: the exact weights, and which arms run* below | This is a deliberate state, not a bug: registering the architecture is what lets the config parse and weight-name mapping be tested before the forward exists. diff --git a/scripts/agent-preflight.sh b/scripts/agent-preflight.sh index a9e555837..8f6c4a515 100755 --- a/scripts/agent-preflight.sh +++ b/scripts/agent-preflight.sh @@ -135,6 +135,7 @@ SUITES=( test_check_model_checklist test_check_supported_models test_check_env_doc + test_nemotron_h_a2q1_per_token test_checker_text test_check_fusion_consistency test_check_fp4_resident_consistency diff --git a/scripts/nemotron-h-a2q1-dgx-gate.sh b/scripts/nemotron-h-a2q1-dgx-gate.sh new file mode 100755 index 000000000..9a74c24a8 --- /dev/null +++ b/scripts/nemotron-h-a2q1-dgx-gate.sh @@ -0,0 +1,231 @@ +#!/usr/bin/env bash +# A2-Q1 (#810, .agents/specs/nemotron-h-a2q1-fp8-mamba.md) — the GB10 gate for +# the FP8 W8A8 Mamba2 device arm, recorded as a script so the recipe is +# reproducible rather than retyped (AGENTS.md §Gates: "Record the exact build and +# run recipe"). +# +# RUN IT INSIDE A LEASE, NEVER OVER ssh: +# rc run -d dgx:gpu0 --max-runtime 6h -- bash -lc \ +# 'git clone --depth 50 -b row/A2-Q1 https://github.com/mudler/vllm.cpp /root/src \ +# && bash /root/src/scripts/nemotron-h-a2q1-dgx-gate.sh' +# +# THREE ENVIRONMENT FACTS THIS SCRIPT ENCODES RATHER THAN REDISCOVERS: +# * the CUDA lane on this box is `sbsa`, not `arm64`, and an UNANCHORED +# `cuda-toolkit-13*` match selects `cuda-toolkit-13-config-common`, which +# installs cleanly, ships no compiler and returns 0; +# * `nvcc --version` is NOT a sufficient postcondition — a partial CUDA install +# printed every feature line and then failed to link `CUDA::cublasLt`, so the +# toolkit is proved by an actual link; +# * `/workspace` is CIFS and refuses symlinks, so the build lives on local disk +# and only the log is copied out. +set -u -o pipefail + +LOG_ROOT=${LOG_ROOT:-/workspace/a2q1} +CKPT=${CKPT:-/workspace/a3/ckpt-stage} +SRC=${SRC:-/root/src} +# 121a is the GB10; 110 is Thor. The spec gates BOTH hosts, and the FP8 W8A8 arm +# on Thor is the whole point of #960/#991, so the arch is a parameter. +ARCH=${ARCH:-121a} +BUILD=${BUILD:-/root/build-cuda} +CUTLASS=${CUTLASS:-/root/cutlass} +STAMP=$(date -u +%Y%m%dT%H%M%SZ) +RUN=$LOG_ROOT/$STAMP +mkdir -p "$RUN" +echo "A2Q1 run dir: $RUN" + +step() { echo; echo "=== $* ==="; } +# Every gate command runs BARE and echoes its own status. Never pipe a command +# whose exit code matters -- a pipeline reports the LAST stage, which is how a +# failing `mount` once reported rc=0. +rc() { "$@"; local r=$?; echo "RC[$*]=$r"; return $r; } + +step "0. the box, before anything is installed" +rc uname -m +rc id -u +rc nvidia-smi --query-gpu=name,compute_cap,memory.total --format=csv +rc df -h /root /workspace +rc free -m + +step "1. contention -- a timing number measured beside another job is VOID" +nvidia-smi --query-compute-apps=pid,used_memory --format=csv | tee "$RUN/contention.txt" +echo "RC[nvidia-smi compute-apps]=${PIPESTATUS[0]}" + +step "2. toolchain" +export DEBIAN_FRONTEND=noninteractive +rc apt-get update -qq +rc apt-get install -y -qq git cmake ninja-build g++ curl ca-certificates python3 python3-dev +if ! command -v nvcc >/dev/null 2>&1; then + # ANCHORED package names on the sbsa lane. `cuda-nvcc-13-0` is the compiler; + # the libraries this build links are named individually so a metapackage that + # ships nothing cannot satisfy the check. + curl -fsSL -o /tmp/cuda-keyring.deb \ + https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2404/sbsa/cuda-keyring_1.1-1_all.deb + echo "RC[curl keyring]=$?" + rc dpkg -i /tmp/cuda-keyring.deb + rc apt-get update -qq + rc apt-get install -y -qq cuda-nvcc-13-0 cuda-cudart-dev-13-0 libcublas-dev-13-0 \ + cuda-nvrtc-dev-13-0 cuda-nvtx-13-0 cuda-profiler-api-13-0 libcurand-dev-13-0 +fi +export PATH=/usr/local/cuda/bin:$PATH +rc nvcc --version + +step "2b. the toolkit is proved by a LINK, not by --version" +cat > /tmp/probe.cu <<'EOF' +#include +#include +int main() { + cublasLtHandle_t h = nullptr; + const auto s = cublasLtCreate(&h); + std::printf("cublasLtCreate=%d\n", static_cast(s)); + return s == CUBLAS_STATUS_SUCCESS ? 0 : 1; +} +EOF +rc nvcc -arch=sm_$ARCH /tmp/probe.cu -o /tmp/probe -lcublasLt +TOOLCHAIN_OK=$? +if [ "$TOOLCHAIN_OK" -ne 0 ]; then + echo "VOID: the CUDA toolkit does not link cublasLt; every number below would be a lie" + exit 2 +fi +rc /tmp/probe + +step "3. cutlass (the fp8/fp4 fast-path cells resolve from it)" +if [ ! -f "$CUTLASS/include/cutlass/cutlass.h" ]; then + rc git clone --depth 1 --branch v4.5.0 https://github.com/NVIDIA/cutlass "$CUTLASS" +fi +rc test -f "$CUTLASS/include/cutlass/cutlass.h" + +step "4. configure" +rc cmake -S "$SRC" -B "$BUILD" -G Ninja \ + -DCMAKE_BUILD_TYPE=Release \ + -DVLLM_CPP_CUDA=ON \ + -DVLLM_CPP_CUDA_ARCHITECTURES="$ARCH" \ + -DVLLM_CPP_CUTLASS_DIR="$CUTLASS" 2>&1 | tee "$RUN/configure.log" +echo "RC[cmake configure]=${PIPESTATUS[0]}" + +step "4b. the feature table -- a DISABLED or wrong-arch cell VOIDS the run" +grep -E "CUDA feature .*(ENABLED|DISABLED)" "$RUN/configure.log" | tee "$RUN/features.txt" +BAD=$(grep -cE "CUDA feature .*DISABLED" "$RUN/features.txt") +ENABLED=$(grep -cE "ENABLED for \[$ARCH\]" "$RUN/features.txt") +echo "feature cells ENABLED for [$ARCH]: $ENABLED ; DISABLED cells: $BAD" + +step "5. build (-j 4: unconstrained parallelism has OOM-REBOOTED this box)" +rc cmake --build "$BUILD" -j 4 2>&1 | tail -40 | tee "$RUN/build.tail" +echo "RC[cmake build]=${PIPESTATUS[0]}" + +step "6. the focused device gate -- A2-Q1's own cases" +"$BUILD/tests/test_nemotron_h_mamba_device" -s 2>&1 | tee "$RUN/mamba_device.log" +echo "RC[test_nemotron_h_mamba_device]=${PIPESTATUS[0]}" +grep -E "test cases:|assertions:|Status:" "$RUN/mamba_device.log" + +step "6b. the neighbouring suites the arm can break" +for t in test_nemotron_h_forward test_nemotron_h_paged_forward test_nemotron_h_loader \ + test_nemotron_h_moe_device test_ops_mamba2_ssd test_ops_fp8_cpu; do + if [ -x "$BUILD/tests/$t" ]; then + "$BUILD/tests/$t" > "$RUN/$t.log" 2>&1 + echo "RC[$t]=$?" + grep -E "test cases:|assertions:|Status:" "$RUN/$t.log" + else + echo "MISSING BINARY: $t" + fi +done + +step "7. the A3 gate + the GPU busy fraction, device mamba ON" +# The acceptance test of this unit is NOT a ratio: the GPU busy fraction must +# RISE from the 6.31% baseline, so it is sampled on a loop for the whole decode +# and the SAMPLE COUNT is reported beside it. A fraction with no denominator is +# not a measurement. +run_gate() { # $1 = label, $2 = VT_NEMOTRON_H_DEVICE_MAMBA value + local label=$1 flag=$2 + local log="$RUN/a3_$label.log" + : > "$log" + + # ★ THE SAMPLING WINDOW IS THE DECODE, NOT THE WHOLE PROCESS. + # + # The first Thor run sampled from process start, so the window contained the + # multi-minute 20.1 GiB engine load, which is GPU-IDLE. That dilutes both arms + # toward each other and toward zero: it read 15.33% with the arm on and 14.73% + # with it off, a 0.59-point difference that says almost nothing about what the + # DECODE does, and it is not comparable to the 6.31% figure either. Mixing a + # load phase into a decode measurement is the same defect as summing prefill + # and decode in one profile. + # + # So the driver starts FIRST, and the sampler starts only once the driver has + # printed `engine loaded in Ns`. Everything after that line is decode. + VT_NEMOTRON_H_DEVICE_MAMBA=$flag "$BUILD/examples/nemotron-h-gen" \ + --model "$CKPT" \ + --golden "$SRC/tests/parity/goldens/nemotron_35_lightning_greedy/oracle.json" \ + > "$log" 2>&1 & + local pid=$! + local waited=0 + local loaded=0 + while kill -0 "$pid" 2>/dev/null; do + if grep -q "engine loaded in" "$log" 2>/dev/null; then loaded=1; break; fi + sleep 1 + waited=$((waited + 1)) + if [ "$waited" -ge 5400 ]; then break; fi + done + if [ "$loaded" -ne 1 ]; then + # Say so rather than sampling a window whose meaning is unknown. A fraction + # over the wrong window is worse than no fraction, because it still formats + # like a measurement. + echo "$label: the engine never reported a load in ${waited}s -- the decode window is UNKNOWN, so NO busy fraction is sampled" + fi + local t0=$(date +%s.%N) + ( while true; do nvidia-smi --query-gpu=utilization.gpu --format=csv,noheader,nounits; sleep 0.1; done ) \ + > "$RUN/util_$label.txt" 2>/dev/null & + local sampler=$! + wait "$pid" + local r=$? + local t1=$(date +%s.%N) + kill "$sampler" 2>/dev/null + wait "$sampler" 2>/dev/null + + echo "RC[a3 $label]=$r" + grep -E "STRICT|PASS|FAIL|DIVERGENCE|REFUSING|SHORT|mode=|engine loaded in" "$log" | tail -20 + # The busy fraction WITH its denominator, over the DECODE window alone. The + # 6.31% baseline is a GB10 number, so it is quoted ONLY on GB10: printing it + # beside a Thor figure invites a cross-silicon comparison that answers a + # different question, and a number quoted often enough becomes treated as + # measured. + python3 - "$RUN/util_$label.txt" "$label" "$ARCH" "$loaded" <<'PY' +import sys +vals = [int(x) for x in open(sys.argv[1]).read().split() if x.strip().isdigit()] +label, arch, loaded = sys.argv[2], sys.argv[3], sys.argv[4] +if loaded != "1": + print(f"{label}: busy fraction NOT REPORTED -- the decode window was never identified") +elif not vals: + print(f"{label}: NO SAMPLES -- the busy fraction is unmeasured, not 0") +else: + busy = sum(1 for v in vals if v > 0) + pct = 100.0 * busy / len(vals) + note = ("baseline 6.31% on this same GB10 workload" if arch == "121a" + else f"NO baseline for arch {arch}: the 6.31% figure is GB10's, so it " + f"is not comparable here; use the ON/OFF A/B on this box") + print(f"{label}: GPU busy in {busy} of {len(vals)} DECODE samples = {pct:.2f}% busy ({note})") +PY + # The per-output-token time over the SAME decode window, with the terms shown. + python3 "$SRC/scripts/nemotron-h-a2q1-per-token.py" "$log" "$label" "$t0" "$t1" "$ARCH" + # The reference tier is numerically CORRECT, so a pass obtained on it is + # invisible in the numbers and only this line separates them (spec R2). + echo "reference-tier lines in $label: $(grep -c 'reference-tier' "$log")" +} + +if [ ! -d "$CKPT" ]; then + echo "NO CHECKPOINT AT $CKPT -- the A3 gate and the GPU busy fraction are NOT MEASURED on this host." + echo "That is a stated absence, not a pass. Steps 7-9 are skipped." + echo "ALL LOGS: $RUN" + exit 0 +fi + +( while true; do free -m | awk '/^Mem:/{print $3}'; sleep 1; done ) > "$RUN/rss.txt" 2>/dev/null & +MEMPID=$! +run_gate on 1 + +step "8. the same binary with the arm OFF -- the A/B this unit is measured by" +run_gate off 0 +kill "$MEMPID" 2>/dev/null +echo "peak host MiB used during the run: $(sort -n "$RUN/rss.txt" | tail -1)" + +step "9. contention, after" +nvidia-smi --query-compute-apps=pid,used_memory --format=csv +echo "ALL LOGS: $RUN" diff --git a/scripts/nemotron-h-a2q1-neartie-gap.py b/scripts/nemotron-h-a2q1-neartie-gap.py new file mode 100755 index 000000000..76421a1e1 --- /dev/null +++ b/scripts/nemotron-h-a2q1-neartie-gap.py @@ -0,0 +1,239 @@ +#!/usr/bin/env python3 +# A2-Q1 / #1388 -- the oracle's OWN top-2 margin at the position the device +# mamba arm loses, and the per-position margin profile of the whole golden. +# +# ★ THIS HAS NEVER PRODUCED A MEASUREMENT, and the reason is not in this file. +# Five runs on `dgx:gpu0` (2026-08-19/20, logs under /workspace/a2q1-neartie/) +# were each killed by the host-memory watchdog during ENGINE START-UP, before +# any token was generated. That is #1431, and it also makes the 2026-08-18 +# result behind "the oracle DOES run a model in a lease" non-reproducing. Read +# #1431 before spending a lease on this script: the block is the host, and no +# engine knob tried there avoided it. +# +# A sibling of scripts/*-neartie-gap.py, not an edit of one: those teacher-force +# vLLM on OUR generated ids for a 16-prompt corpus and emit .npy gap files a +# gate reads. This one answers a single triage question about ONE recorded +# divergence and prints it, so it is one file per row, read with a glob +# (AGENTS.md "Records"), rather than another writer of a shared surface. +# +# RUN IT INSIDE A LEASE, NEVER OVER ssh. The WHEEL is the one measured on that +# box. The LOAD CONFIGURATION is not: run 4 of the five ran the configuration +# that survived on 2026-08-18 and died anyway, so it is a starting point under +# investigation, not a known-good recipe. For the wheel see +# /workspace/oracle-vllm/README-WHEELS.md +# (vLLM's default FLASH_ATTN carries no sm_12x SASS and cannot run on a GB10; +# and the wheel needs a PEP 427 name before pip will take it, #1416): +# +# rc run -d dgx:gpu0 --max-runtime 3h -- bash /workspace/a2q1-neartie/job.sh +# +# Environment: MODEL (checkpoint dir), GOLDEN (oracle.json), GMU, KVBYTES. +# A2-Q1 / #1388 STEP ONE -- the cheapest question, asked before any numerics work. +# +# The device mamba arm loses exactly ONE token: prompt 2, generation position 32 +# (the LAST position on the LONGEST prompt), got 11286 where the golden says +# 3468. Positions 1..31 are byte-identical, so the two arms reach position 32 +# with the SAME prefix. The question is therefore entirely about the ORACLE: +# how confidently does vLLM itself separate its argmax from the runner-up there? +# +# This TEACHER-FORCES the pinned oracle on the golden prefix and reads its own +# top-K. A margin at the bf16 noise floor makes this a tie-break, not a defect +# ([[near-tie-distributional-gate]]); a large margin makes it a real divergence. +import json, os, sys, time + +OURS = {2: {32: 11286}} # arm token, 1-based generation position, per prompt + +def topk(d, k=8): + return sorted(d.items(), key=lambda kv: -kv[1].logprob)[:k] + +def main(): + from vllm import LLM, SamplingParams + from vllm.inputs import TokensPrompt + import vllm + print("VLLM_FILE", vllm.__file__, flush=True) + print("VLLM_VERSION", vllm.__version__, flush=True) + assert "555967922" in vllm.__version__, f"WRONG COMMIT {vllm.__version__}" + assert "site-packages" in vllm.__file__, f"NOT AN INSTALLED WHEEL {vllm.__file__}" + + golden = json.load(open(os.environ["GOLDEN"])) + print("GOLDEN_MODEL", golden["model"], "rev", golden["revision"], flush=True) + print("GOLDEN_VLLM", golden["vllm"], flush=True) + assert golden["vllm"].endswith("g555967922"), golden["vllm"] + entries = golden["golden"] + T = int(golden["sampling"]["max_tokens"]) + + # ★ THE CONFIGURATION BELOW IS NOT KNOWN TO SURVIVE. It is the one that used + # to survive, plus one variable. Both halves of that sentence were measured. + # + # FIVE runs on `dgx:gpu0`, 2026-08-19/20, logs under /workspace/a2q1-neartie/. + # Every one was killed by the host-memory watchdog inside engine START-UP, + # before a token existed (#1431). The watchdog FLOOR was not constant across + # them, so the kill values are not comparable as a series: + # + # run stamp deviation from the reference config floor killed at + # 1 20260820T002359Z + max_logprobs=64 15000 12597 MB + # 2 20260820T024243Z + enforce_eager=True, len 256, seqs 1 20000 19433 MB + # 3 20260820T025540Z + num_gpu_blocks_override=8 20000 19797 MB + # 4 20260820T030818Z NONE -- the reference config itself 15000 13941 MB + # 5 20260820T031644Z + kv_cache_memory_bytes=4 GiB 15000 14846 MB + # + # RUN 4 IS THE ONE THAT MATTERS, and it falsifies the premise this block used + # to state. /workspace/nhspeed/oracle_only.sh attempt `a` loaded this same + # 20.1 GiB checkpoint on this same box on 2026-08-18, bottomed out at + # minMemAvailable_MB=51528, and generated all three prompts + # (`ORACLE TOKEN MATCH: 180/192`). Run 4 is that configuration, unchanged, and + # it died. So the reference is NOT a safe base to copy; it is a REGRESSION to + # be explained (#1431). The environment is not byte-identical either -- run 4 + # built its venv at /tmp/a2q1-oracle with a fresh `pip install torch==2.13.0` + # and a reinstalled wheel, where 08-18 used /tmp/nhspeed-oracle -- and an + # identity assert on the vLLM COMMIT cannot see a torch or flashinfer delta. + # That delta is an UNEXCLUDED candidate for the regression. + # + # Runs 2 and 3 were killed ~5 GB earlier in the drawdown than the others. That + # they would ALSO have crossed 15000 MB is an INFERENCE, not a measurement: + # the host trace was still falling at ~1.18 GB/s with no arrest when the kill + # landed. Do not quote runs 2 and 3 as 15000 MB results. + # + # WHAT THE FIVE RUNS EXCLUDE. Host memory sits flat through the whole weight + # load, then falls ~76 GB in ~60 s at ~1.2 GB/s just after the mamba page-size + # print. Not `torch.compile`: run 5 hit the AOT cache (`torch.compile took + # 0.30 s`). Not CUDA graph capture: run 2 ran `enforce_eager=True`. Not KV + # SIZING: run 3 overrode the block count and run 5 set an absolute budget. + # What is left is THE FIRST FORWARD. Under `kv_cache_memory_bytes` vLLM runs + # `profile_run()` FIRST (v1/worker/gpu_worker.py:465-468, "still need a profile + # run which compiles the model for max_num_batched_tokens") and only THEN logs + # "Initial free memory ... skipped memory profiling" (:470-482). Run 5's log + # reaches neither string, so the process died INSIDE that forward. This + # supersedes #1185's "after compilation -- profiling forward or graph capture", + # and it is recorded there rather than chased here. + # + # WHY THE ENGINE IS OTHERWISE UNTOUCHED. The question is asked with + # REQUEST-level sampling params, which cost nothing at start-up. + # `max_logprobs` stays at its engine default of 20 (config/model.py:228) and + # the requests below ask for exactly 20 -- ample for a two-way tie, and one + # more engine knob not turned. + # + # THE ONE VARIABLE. `kv_cache_memory_bytes` (config/cache.py:182, honoured at + # v1/worker/gpu_worker.py:465-476) takes an ABSOLUTE KV budget, so it does not + # depend on a utilization fraction -- which matters because `nvidia-smi` + # reports `memory.total = [N/A]` on this GB10. 4 GiB against three sequences of + # at most 45 tokens is an over-provision, not a tuned value: a block here is + # roughly 120 MB (23 mamba layers of conv + f32 SSM state, plus an attention + # page vLLM raises to 4176 tokens so it is >= the mamba page), so ~30 blocks + # for a job that needs one per prompt. `gpu_memory_utilization` is left at the + # reference 0.30 and is NOT the knob being turned. + # + # It did not help. Run 5 IS this configuration, and it died too. The knob is + # kept because it removes KV sizing from the candidate list, not because it + # works. READ #1431 BEFORE SPENDING A LEASE HERE. + kw = dict(model=os.environ["MODEL"], max_model_len=512, max_num_seqs=8, + gpu_memory_utilization=float(os.environ.get("GMU", "0.30")), + max_num_batched_tokens=512, enforce_eager=False, + kv_cache_memory_bytes=int(os.environ.get("KVBYTES", str(4 * 1024**3)))) + mode = "reference-config+kvbytes" + print("MODE", mode, flush=True) + print("ORACLE_KW", {k: v for k, v in kw.items() if k != "model"}, flush=True) + llm = LLM(**kw) + try: + cc = llm.llm_engine.vllm_config.cache_config + print("ORACLE_CACHE block_size=%s num_gpu_blocks=%s override=%s" % + (cc.block_size, getattr(cc, "num_gpu_blocks", None), + getattr(cc, "num_gpu_blocks_override", None)), flush=True) + except Exception as e: + print("ORACLE_CACHE_UNREADABLE:", e, flush=True) + tok = llm.get_tokenizer() + def show(t): + try: + return repr(tok.decode([int(t)])) + except Exception: + return "" + + # ── LEG A: free-running greedy, the gate's own workload. Context only: it + # says whether the oracle at THIS config reproduces its own committed + # golden. It does NOT answer the margin question, because a diverging + # oracle reaches position 32 on a different prefix. + print(f"\n########## [{mode}] LEG A: free-running greedy vs the committed golden ##########", flush=True) + spA = SamplingParams(temperature=0.0, max_tokens=T, ignore_eos=True) + for i, e in enumerate(entries): + out = llm.generate([TokensPrompt(prompt_token_ids=e["prompt_token_ids"])], spA)[0] + got = list(out.outputs[0].token_ids) + exp = e["token_ids"] + n = min(len(got), len(exp)) + mism = [j for j in range(n) if got[j] != exp[j]] + print(f"LEGA[{mode}] prompt {i}: compared={n} matched={n-len(mism)} " + f"first_mismatch={(mism[0]+1) if mism else None} " + f"mismatch_positions_1based={[j+1 for j in mism]}", flush=True) + print(f"LEGA[{mode}] prompt {i} got: {','.join(str(x) for x in got)}", flush=True) + + # ── LEG B: TEACHER-FORCED on the golden prefix. This is the measurement. + # One prefill over prompt + golden[0:T-1] gives, via prompt_logprobs, the + # oracle's own top-K at generation positions 1..T-1, and the single sampled + # step gives position T -- the position that diverges. + print(f"\n########## [{mode}] LEG B: teacher-forced top-K on the GOLDEN prefix ##########", flush=True) + K = 20 # == the engine default cap; raising it is an engine knob, see above + spB = SamplingParams(temperature=0.0, max_tokens=1, logprobs=K, prompt_logprobs=K) + for i, e in enumerate(entries): + P = e["prompt_token_ids"] + g = e["token_ids"] + full = list(P) + list(g[:T-1]) # predict generation position T + out = llm.generate([TokensPrompt(prompt_token_ids=full)], spB)[0] + + # positions 1..T-1 come from prompt_logprobs; index p in `full` holds + # the distribution that PRODUCED full[p], so generation position j + # (1-based) is index len(P)+j-1. + plp = out.prompt_logprobs + if plp is None: + print(f"LEGB[{mode}] p{i}: prompt_logprobs came back None -- the INSTRUMENT " + f"did not measure positions 1..{T-1}; position {T} below still stands", + flush=True) + plp = [None] * len(full) + margins = [] + for j in range(1, T): + d = plp[len(P) + j - 1] or {} + if len(d) < 2: + margins.append(None); continue + tk = topk(d, 2) + margins.append(tk[0][1].logprob - tk[1][1].logprob) + fin = margins + [None] + + # position T: the sampled step's own logprob dict + dT = out.outputs[0].logprobs[0] + tkT = topk(dT, 8) + argT = tkT[0][0] + marginT = tkT[0][1].logprob - tkT[1][1].logprob if len(tkT) > 1 else float("inf") + fin[T-1] = marginT + + print(f"\n--- prompt {i} ({len(P)} prompt tokens) ---", flush=True) + print(f"LEGB[{mode}] p{i} teacher-forced argmax at generation position {T}: " + f"{argT} {show(argT)} golden={g[T-1]} {show(g[T-1])} " + f"match={argT == g[T-1]}", flush=True) + print(f"LEGB[{mode}] p{i} TOP-2 MARGIN at position {T}: {marginT:.6f} nats", flush=True) + print(f"LEGB[{mode}] p{i} top-8 at position {T}:", flush=True) + for r, (t, lp) in enumerate(tkT): + mark = "" + if t == g[T-1]: mark += " <== GOLDEN/vLLM-greedy" + if t == OURS.get(i, {}).get(T): mark += " <== OURS (device mamba arm)" + print(f" rank {r} (vllm_rank={lp.rank}): id={t:<7d} logprob={lp.logprob:+.6f} " + f"p={pow(2.718281828459045, lp.logprob):.6f} {show(t)}{mark}", flush=True) + ours = OURS.get(i, {}).get(T) + if ours is not None: + if ours in dT: + gap = tkT[0][1].logprob - dT[ours].logprob + print(f"LEGB[{mode}] p{i} OUR TOKEN {ours} {show(ours)}: " + f"vllm_rank={dT[ours].rank} logprob={dT[ours].logprob:+.6f} " + f"gap_to_vllm_argmax={gap:.6f} nats", flush=True) + else: + print(f"LEGB[{mode}] p{i} OUR TOKEN {ours} OUTSIDE vLLM top-{K} -- REAL DIVERGENCE", flush=True) + + fmt = ", ".join("None" if m is None else f"{j+1}:{m:.4f}" for j, m in enumerate(fin)) + print(f"LEGB[{mode}] p{i} per-position top-2 margins (nats): {fmt}", flush=True) + fin2 = [m for m in fin if m is not None] + if fin2: + order = sorted(range(len(fin)), key=lambda j: (fin[j] if fin[j] is not None else 1e9)) + print(f"LEGB[{mode}] p{i} TIGHTEST 5 positions: " + f"{[(j+1, round(fin[j], 5)) for j in order[:5]]}", flush=True) + print(f"LEGB[{mode}] p{i} median margin={sorted(fin2)[len(fin2)//2]:.4f} " + f"min={min(fin2):.6f} max={max(fin2):.4f}", flush=True) + print(f"\nDONE_MARKER_NEARTIE mode={mode}", flush=True) + +if __name__ == "__main__": + main() diff --git a/scripts/nemotron-h-a2q1-per-token.py b/scripts/nemotron-h-a2q1-per-token.py new file mode 100755 index 000000000..56baba5b0 --- /dev/null +++ b/scripts/nemotron-h-a2q1-per-token.py @@ -0,0 +1,91 @@ +#!/usr/bin/env python3 +"""Derive the per-output-token time of one A2-Q1 A3 gate run, WITH its terms. + +`examples/nemotron_h_gen` reports neither a rate nor a duration: it prints the +engine load time and a `TOKEN MATCH: m/n` line and nothing else. So the per-token +number this row is measured on has to be derived, and a derived number that hides +its terms is how a rate over an unknown denominator comes to be quoted as if it +were measured. + +THE WINDOW IS THE DECODE, AND THE CALLER BRACKETS IT. `t0`/`t1` are taken AFTER +the driver prints `engine loaded in Ns`, so the 20.1 GiB load is already outside +them and nothing is subtracted here. An earlier version bracketed the whole +process and subtracted the load, which put a multi-minute GPU-idle phase inside +the same window as the decode -- the same defect as summing prefill and decode +into one profile. The load is still read out of the log and printed, because it +is context for the number, not a term in it. + +This prints the window, the load it excludes and the token count it divides by on +separate lines, and it REFUSES rather than printing 0 when a term is missing or +the window is not positive. + +THE vLLM RATIO IS QUOTED ONLY ON THE SILICON IT WAS MEASURED ON. 0.014369 s is a +GB10 figure. Printing it beside a Thor per-token number invites a comparison +across two different pieces of silicon, and the first Thor run did exactly that +-- it printed `ratio 54.7x` for a number that was never measured against vLLM on +that box. This is the same defect the busy-fraction reporter carried, and fixing +one surface while leaving its twin is how a wrong comparison survives. + + nemotron-h-a2q1-per-token.py