From 2a448ab493473bc6310c4ba5aab07023c2c4ae83 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Fri, 7 Aug 2026 07:13:42 +0000 Subject: [PATCH] feat(minimax-h3): stream the ORIGINAL bf16 DiT to the device, one tensor at a time MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Branch: row/H3-BF16-SHARDED-STREAM (helper; CPU-only, no GPU job, no download), STACKED on row/H3-BF16-SHARDED-DIT. Split from that row so each PR stays inside the 900-line review cap; the seam is checkpoint-vs-device, not an arbitrary cut. row/H3-BF16-SHARDED-DIT made the 13-shard 66.3 GB release INDEXABLE, but its only loader was host-f32 — ~132 GB on the real release. The box has 122 GiB of UNIFIED memory (host and device draw on ONE pool), so "load to host, then stage" holds the model TWICE against that budget, and the non-streaming NVFP4 loader was already OOM-KILLED at anon-rss 125 GB on HALF this size. Without a streaming path the original release is not loadable at all. - StreamMiniMaxH3ShardedToDeviceBf16(queue, ckpt, out_params) in minimax_h3_device.cpp, beside its GGUF and NVFP4 twins so it reuses BindStreamedDitViews. Manifest first (names+shapes, no payload) so the geometry is known and a wrong name map is caught before a byte is allocated, then ONE tensor at a time. - Peak host memory, precisely: BF16-on-disk -> bf16 device slot (essentially the whole 66.3 GB) and F32 -> f32 island are both uploaded DIRECTLY out of the read-only mmap, with no host buffer at any point. Only a dtype MISMATCH costs one tensor's conversion buffer, freed before the next iteration. Each source range goes to MaybeReleaseSourcePages the moment its copy returns, so the page cache does not accumulate against the pool the weights live in. rope.inv_freq stays HOST-resident: BuildRopeCosSin runs before any kernel, so a device pointer there segfaults on the first forward. - MiniMaxH3ShardStreamStats (mirroring Nvfp4W4A16Stats) makes the path OBSERVABLE. This codebase has shipped a never-executing guarded kernel under a green suite before; counters are how that is prevented. - examples/minimax_h3_gen: --dit --device cuda streams and prints the stats line; --device cpu keeps the host reference loader. Gates (CPU, test_minimax_h3 73/73 cases / 55203 assertions; clean Release build of libvllm.a, test_minimax_h3 and minimax-h3-gen), over a 3-shard synthetic set whose dtypes exercise all four (on-disk x device) combinations: (1) streamed == non-streamed — all 46 weight views BIT-EXACT (memcmp == 0) vs StageMiniMaxH3DitWeights(kBF16), dtypes included (12 fp32 islands), and both device forwards return IDENTICAL logits (video and audio max|diff| both exactly 0.0); rope.inv_freq asserted HOST-resident. (2) the loader RAN — the counters are ASSERTED, not printed: observed shards=3 tensors=46 direct=37 converted=9 bytes=444504 host_peak=8192, so BOTH upload paths were taken, every bound view is owned by this loader's staging, and host_peak_bytes (one tensor's buffer) is under 1/4 of the bytes uploaded — the peak cannot scale with the model. Honest residuals: the real 66.3 GB load and its measured peak RSS are UNVERIFIED, as is CUDA memcpy straight from a file-backed mmap (valid pageable-source usage, every copy followed by a synchronize, but unexercised on device) — this row ran CPU-only. NO bf16-vs-quantized render or speed number is claimed. The quality A/B is now RUNNABLE; it has not been run. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context) --- .agents/NOW.md | 2 +- .agents/model-matrix.md | 2 +- .agents/parity-ledger.md | 1 + .agents/roadmap_v1.md | 2 +- .agents/specs/minimax-h3.md | 54 +++++ .agents/state.md | 55 +++++ docs/BENCHMARKS.md | 2 +- docs/FEATURES.md | 4 +- docs/STATUS.md | 2 +- examples/minimax_h3_gen/main.cpp | 28 ++- .../vllm/model_executor/models/minimax_h3.h | 54 ++++- .../models/minimax_h3_device.cpp | 136 +++++++++++ tests/vllm/models/test_minimax_h3.cpp | 219 +++++++++++++++++- 13 files changed, 542 insertions(+), 19 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index 590a3c32d..304c222c8 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -17,7 +17,7 @@ Working head: `row/backend-rocm-w0` (#41). Prior: benchmark checkpoint | Laguna NVFP4 / DeepSeek-V4 decode | **Both CLOSED, byte-exact, default-ON**: 1.03x vLLM, 1.144x ds4 | Laguna vLLM K-run when convenient | | f32-out GEMV audit | Only laguna + ds4 bf16 tower affected; gate models unaffected | Re-verify ds4 tower same-tool | | Invocation-parity prevention | CI guard + checklist landing | Merge; build-verify `kGemvHeuristicAlgos` on dgx | -| MiniMax-H3 lane | **fl2va COHERENT; ref2va grid DIAGNOSED (#95): NO loader bug; bf16 13-shard DiT INDEXES** | residual = community-NVFP4 quant fidelity §8.12; no bf16 render yet | +| MiniMax-H3 lane | **fl2va COHERENT; ref2va grid DIAGNOSED (#95): NO loader bug; bf16 13-shard DiT STREAMS** | residual = community-NVFP4 quant fidelity §8.12; no bf16 render yet | | Kimi-Linear-48B (KDA+NoPE-MLA+MoE) | bf16 knobs **106→120/128**, NOT STRICT (§14, `row/KIMI-LINEAR-STRICT-SPEED`); default OFF | residual = device islands; 1.30 tok/s | | 35B fresh grid | **BOUND** @`1ea26427`: 0.93-1.03x, c16 0.93x. INTAKE + Option A both NEGATIVE | Lever left: prefill glue (#61) | | Qwen3.5-4B revalidation | 0.9971x @`59674cf1` (#35); TTFT/PSS pass, TPOT/ITL open | `docs/bench-evidence/` | diff --git a/.agents/model-matrix.md b/.agents/model-matrix.md index 3380c571d..b8df5e10b 100644 --- a/.agents/model-matrix.md +++ b/.agents/model-matrix.md @@ -81,7 +81,7 @@ Engaged architectures (the 45 non-`INVENTORIED` rows): | ✅ | `Glm4MoeLiteForCausalLM` | GLM-4.7-Flash (31.2B MLA + GLM MoE) | SACRED gate 8/8 vs vLLM 0.25.0 (STRICT token-exact 1/8 + near-tie-band 7/8, 69/128 tokens strictly exact, max teacher-forced gap 0.0 nats, 0 forward-divergent; vLLM K=5 self-deterministic → STRICT bar); FIRST e2e coverage of the q_lora query branch AND the noaux_tc sigmoid router (closes the MLA campaign's two gaps, C2); speed pending | `MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm` | | 🚧 | `KimiLinearForCausalLM` | Kimi-Linear-48B-A3B | **FULL-MODEL GB10 e2e RUNS — NEAR-TIE 106/128 (2026-08-06, `row/MODEL-KIMI-LINEAR-BF16`):** the bf16-resident path CLEARS the f32-loader block — the full 48.9B model now runs e2e on one GB10. dgx CUDA build (`-Werror` clean, 14 GDN AOT symbols nm-linked, `test_kimi_linear_forward` 13/13·656 in the CUDA binary); `kimi-linear-gen --gpu` greedy-decodes the §12 8-prompt battery x16 vs `greedy_ids.npy`. MEMORY: load 117.6s, host RSS PEAK **1.7 GiB** (stage-then-ReleaseHost), device peak 98.5 GiB, min-avail **21.6 GiB** (above the 15 GiB floor, matches the ~25 GiB pool-math headroom), NO OOM/reboot. TOKEN gate **NEAR-TIE 106/128 (82.8%)** — prompts 0,1,3,4,5,6 are 16/16 token-exact, p2/p7 diverge at punctuation/word near-ties; 96 consecutive exact tokens across 6 prompts prove the WIRING (a wiring bug can't). Root cause (honest): the f32 residual stream + host-f64 islands are MORE precise than vLLM's bf16 device kernels, so they flip the argmax where vLLM's deterministic bf16 top-1 has a small margin. STRICT path = the named W7-speed residuals (device GDN/MLA islands -> bf16 stream matching vLLM's rounding). 1.59 tok/s (recompute+island rate). `VT_KIMI_DEVICE_COMPUTE` STAYS OFF (parity-enablers: near-tie != token-exact). Row STAYS 🚧. **bf16-RESIDENT loader/forward IMPLEMENTED + CPU-gated (2026-08-06, `row/MODEL-KIMI-LINEAR-BF16`):** the §13 design is coded — `LoadKimiLinearResidentBf16Weights`/`StageKimiResidentBf16`/`BuildKimiResidentFromHost` (`kimi_linear_weights.cpp`; `LoadBf16Direct` -> `OwnedTensor`, per-tensor stage-to-`d_dev` + `ReleaseHost`, tiny vectors host f32), `KimiLinearResidentWeights` (`kimi_linear.h`), bf16 device forward `DeviceForwardBodyBf16` + `Gemm Bf16` cast-act at ~20 GEMM sites with the two host-fallback islands EXTRACTED+shared (`kimi_linear_device.cpp`), `ForwardDevice` resident-path dispatch (`kimi_linear.cpp`), and the `kimi-linear-gen` e2e harness. CPU **13/13·656** (12/12·614 f32 path UNTOUCHED + NEW tiny-config bf16-vs-f32 gate). PENDING: dgx CUDA build + full-model GB10 e2e vs the STRICT golden. Row STAYS 🚧. **bf16-RESIDENT brick POOL-MATH+DESIGN (2026-08-06, `row/MODEL-KIMI-LINEAR-BF16`):** pool math CLOSES (91.5 GiB bf16 device-resident + ~2.4 GiB act/norms/ctx ≈ 94 GiB, ~25 GiB headroom); design grounded §13 (Laguna `GemmBf16` cast-act + `OwnedTensor::d_dev`, `LoadBf16Direct`, f32 `MaterializeHost` kept for the unit gate). Impl (loader/forward rewrite + gate + e2e) pending. Row STAYS 🚧. **§8 GOLDEN CAPTURED — STRICT (2026-08-06, `row/MODEL-KIMI-LINEAR-E2E`):** the §8 SACRED oracle golden is captured on GB10 (0.25.0-stage, util 0.82, moe=triton, min 15 GiB avail, NO reboot), **8/8 prompts DETERMINISTIC over K=3 → STRICT gate**, committed at `tests/parity/goldens/kimi_linear_greedy/`. Full our-engine e2e BLOCKED on OUR f32 loader (materializes ~183 GiB > 119 pool), the bf16-residency residual; row STAYS 🚧. **W7 GPU-VERIFY (2026-08-06, branch `row/MODEL-KIMI-LINEAR-GPU`):** the device compute runs **12/12·614 GREEN on GB10 sm_121a CUDA build**, BOTH arms (`VT_KIMI_DEVICE_COMPUTE=1` + host-ref); prod stack (CUTLASS-NVFP4 GEMM + FA2 ENABLED + Triton-AOT GDN, 14 cubins nm-verified); f32 device==W2 ref, no divergence, no DeepSeek-class trap. Oracle gateability re-confirmed (0.25.0-stage registers `KimiLinearForCausalLM`). e2e §8 SACRED golden STILL disk-blocked (91.5 GiB checkpoint absent, dgx root 100% full, 34G free). Row STAYS 🚧. **W7 DBuf-resident device COMPUTE landed, CPU-gated** (`CLAIM-KIMI-LINEAR-W7`): the real device compute (`ForwardDeviceCompute`, `kimi_linear_device.cpp`) composes the whole 27-layer KDA/NoPE-MLA + 256-expert-MoE hybrid over pooled f32 `DBuf`s through the SHARED `vt::` ops (embed/`FusedChain` add+RMSNorm/`MatmulBT` projections/`CausalConv1dFwd` convs/`L2Norm`/`RmsNormGated`/`MoeRouterTopK` sigmoid-`noaux_tc`/`MoeSiluMul`/`MoeCombine`/lm_head), returning DEVICE-RESIDENT logits; 2 documented HOST-FALLBACK islands (the KDA per-k-channel gated-delta recurrence + its exp/softplus decay gate — `vt::GdnDecode` carries only a per-HEAD scalar decay; the NoPE-MLA softmax core — the paged `mla::ForwardMlaAttentionBlock` device path is born-on-runner) are the W7-speed residuals. CPU-gated vs the W2 host reference (the CPU backend runs the SAME `vt::` dispatch): `test_kimi_linear_forward` **12/12·614** (per-op KDA/NoPE-MLA/MoE/dense device==ref within f32-accumulation tolerance; the whole `ForwardDeviceCompute` == ref logits + greedy-token-identical + device-resident). Runner opt-in via `VT_KIMI_DEVICE_COMPUTE=1` (default OFF keeps the CPU-verified W6 host-ref compose). GPU numerics (bf16 activations, GDN Triton-AOT cubins, paged het-KV, grouped-MoE slabs) + the e2e SACRED golden stay a NAMED pending (box down) — row STAYS 🚧. ON TOP OF **W6 DEVICE forward SEAM** (`CLAIM-KIMI-LINEAR-W6`): the born-on-the-runner `ForwardDevice` (the DEFAULT `gather_logits` runner path) no longer refuses — it composes the `[rows,vocab]` logits via the CPU reference and hands them back DEVICE-RESIDENT (a pooled `DBuf`, wrapped like deepseek_v2 `WrapDeviceLogits`; `on_device()==true` on CPU+CUDA) so the on-GPU sampler consumes them with NO host download. Kimi-Linear now ROUTES device-resident (`check-runner-routing-consistency` reclassifies it, refuse-skipped stubs 2→1, NO allowlist; `check-fusion-consistency` green); `test_kimi_linear_forward` **7/7·300** (adds the `ForwardDevice`==host-ref device-resident gate). The DBuf-resident device COMPUTE (KDA via the GDN family, NoPE-MLA via `mla::ForwardMlaAttentionBlock`, DeepSeek-V2 grouped-MoE over the paged het-KV; full plan in `kimi_linear.cpp`) is the GPU-verify-pending W7 residual. ON TOP OF **W2-W6 CPU REFERENCE forward** (`CLAIM-KIMI-LINEAR-W2`): the real host `KimiLinearModel::Forward` composes the whole 27-layer hybrid from the landed primitives (KDA layer via `vllm::kimi_kda` refs + the gated-delta recurrence; NoPE-MLA materialized-MHA ref; sigmoid `noaux_tc` MoE + shared expert; dense SwiGLU); loader now materializes host float weights; `test_kimi_linear_forward` 6/6·246 (per-op gates + finite whole forward + greedy decode). ON TOP OF **W1 scaffolding** (registry + `ParseKimiLinearParams` 20 KDA + 7 NoPE-MLA + index-verified name-map + het-KV spec). e2e-gateable (FITS one GB10, 0.77× pool). RESIDUAL = the DEVICE born-on-runner forward (KDA kernel/absorbed-MLA/grouped-MoE slabs) + the W0/W7 e2e SACRED golden. Row → `ACTIVE` (device SEAM wired; the DBuf device compute + e2e SACRED golden pending) | `MODEL-TEXT-kimi-linear-kimi-linear-for-causal-lm` | | 📋 | `KimiK3ForConditionalGeneration` | Kimi K3 (2.8T MoE + MoonViT-V2, DERIVE-AND-SHIP) | **W2/W5 CPU scaffolding landed** (registry stub + nested text/vision/quant config descent + text-backbone structural name-map + REFUSE-by-name forward + MXFP4-refuse loader; clean CPU build, scaffold gate 6/6). text backbone IS `KimiLinearForCausalLM` (KDA+MLA+MoE hybrid, HEAVY reuse); **does NOT fit GB10 (~1.56 TB MXFP4, ~12×)** and NOT in the pinned oracle ⇒ no on-box golden — DERIVED, proxy-gated on Kimi-Linear-48B; forward + MXFP4 + KDA delta + MoonViT-V2 not implemented (NOT-YET-BUILDABLE) | `MODEL-MM-kimi-k3-kimi-k3-for-conditional-generation` | -| 🚧 | `MiniMaxH3DiTModel` | MiniMax-H3 (33.1B omni-modal video+audio DiT, DERIVE-AND-SHIP) | **W1/W2 landed**: packed layout (fl2va + ref2va, fp64 position grid BIT-EXACT), latent packing, euler-ancestral eta0 scheduler, and the full DiT forward all parity-gated against the UPSTREAM vLLM-Omni modules executed at reduced dimensions (**max abs diff 1.6e-7**, 10/10 cases / 2539 assertions). NOT autoregressive (no KV cache, no sampler, no logits) and **e2e HW-BLOCKED** (~354 GB checkpoint, ~133 GB/rank on 4x B300 vs 119 GiB unified); bf16 production stream + request planning + the ComfyUI-GGUF arm also landed (535 REAL tensors resolve onto our contract, geometry from shapes alone). **HW verdict CORRECTED: quantized arms FIT (~41 GB in 119 GiB)**, so e2e + speed are reachable; encoder/VAEs/audio VAE DONE (4.2e-9 vs the checkpoint's remote code); NVFP4 layout GATED as identical to ours (speed path is loader wiring); BOTH VAE DECODERS done (audio 4.2e-9, video ViT3D 8.9e-8); video tiling + 3D-CNN encoder (conditioning only) pending; encoder TEXT tower done (1.2e-7); **serving `/v1/videos` DONE and the DEVICE-RESIDENT forward (W2b, f32) LANDED + GPU-VERIFIED on Thor sm_110 at video 1.49e-7 / audio 8.94e-8**; bf16 stream + fusion folds + the FP4 path (needs sm_121a) + a real-checkpoint run pending. **2026-08-05: the AUDIO-VAE ENCODER is ported** (DAC analysis stack + `pre_block` AttnProjection + `mean_proj`, gated stage by stage vs the checkpoint's own remote code at 2.98e-8 / 1.64e-7 / 1.86e-8) with its own checkpoint loader gated on the real 1087-tensor manifest — so **ref2va AUDIO and VIDEO+AUDIO references are now WIRED** (audio rows move by 0.51 / 0.71; a different waveform still moves them by 7.1e-4). Both VAEs are now complete in both directions. **bf16 13-SHARD RELEASE INDEXES 2026-08-07 (`row/H3-BF16-SHARDED-DIT`)**: `MiniMaxH3ShardedCheckpoint` resolves the ORIGINAL 66.3 GB release through its own `model.safetensors.index.json` (a tensor named in the index but missing from its shard throws BY NAME), `EnumerateMiniMaxH3ShardedTensors` feeds the shared shapes-only geometry parser, and `LoadMiniMaxH3DitFromShards` is the host-f32 reference loader. Gated CPU-only at 72/72/54497 (post-rebase): every tensor resolves to the shard the index named AND to the bytes written there, the derived geometry equals the single-file path field for field, and a SPARSE 13-shard release with the REAL 535 tensors at REAL shapes (66.3 GB declared, 144 KB on disk) derives the SHIPPED geometry (50/5376/56/128/14336/24/32/1x2x2/5120). The DEVICE streamer is the stacked follow-up `row/H3-BF16-SHARDED-STREAM`. This UNBLOCKS the quantization-quality question; no bf16-vs-quant render or speed number is claimed. Spec §8.13. **W-FP4a LANDED (CPU) 2026-08-06 (`row/H3-FP4-SPEED`)**: the device DiT forward now routes the NVFP4 projections through the shared Marlin W4A16 dispatcher (fp4 kept packed; no new quant code), fp4-vs-bf16 wiring gate GREEN (62/62·30039). **W-FP4a GB10 leg LANDED 2026-08-06 (`row/H3-FP4-GPU-E2E`, PR #64):** on sm_121a the Marlin W4A16 path RAN for all 11 projections (`dense_gemms==11` default — VT_MARLIN_DENSE is default-ON → vLLM's own DENSE Marlin GEMM, not the grouped route; `marlin_gemms==11` under VT_MARLIN_DENSE=0; `fallback_gemms==0`), fp4-vs-bf16 BYTE-EXACT (max\|diff\|=0), and the fp4 arm is a MEMORY win not a diffusion-forward speed win (per-forward bf16/fp4 3.47× @seq64 → 0.79–0.83× @seq4224–7040; ~16 vs ~66 GB device). Real-checkpoint fp4-resident t2va e2e RUNS (real 18.75 GB NVFP4 DiT + VAEs + GGUF Qwen3-VL-32B encoder → valid mp4/wav; DiT s/step 5.45/20.0/209 s @512/768/REF-209f) but frames are a non-scene patch-grid at 12/20/50 steps → OPEN render bug (device VAE/denoise). vLLM-Omni has no quantized H3 arm (BF16-only) so any comparison is HW/loader-forced-indirect — spec §8 | `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit` | +| 🚧 | `MiniMaxH3DiTModel` | MiniMax-H3 (33.1B omni-modal video+audio DiT, DERIVE-AND-SHIP) | **W1/W2 landed**: packed layout (fl2va + ref2va, fp64 position grid BIT-EXACT), latent packing, euler-ancestral eta0 scheduler, and the full DiT forward all parity-gated against the UPSTREAM vLLM-Omni modules executed at reduced dimensions (**max abs diff 1.6e-7**, 10/10 cases / 2539 assertions). NOT autoregressive (no KV cache, no sampler, no logits) and **e2e HW-BLOCKED** (~354 GB checkpoint, ~133 GB/rank on 4x B300 vs 119 GiB unified); bf16 production stream + request planning + the ComfyUI-GGUF arm also landed (535 REAL tensors resolve onto our contract, geometry from shapes alone). **HW verdict CORRECTED: quantized arms FIT (~41 GB in 119 GiB)**, so e2e + speed are reachable; encoder/VAEs/audio VAE DONE (4.2e-9 vs the checkpoint's remote code); NVFP4 layout GATED as identical to ours (speed path is loader wiring); BOTH VAE DECODERS done (audio 4.2e-9, video ViT3D 8.9e-8); video tiling + 3D-CNN encoder (conditioning only) pending; encoder TEXT tower done (1.2e-7); **serving `/v1/videos` DONE and the DEVICE-RESIDENT forward (W2b, f32) LANDED + GPU-VERIFIED on Thor sm_110 at video 1.49e-7 / audio 8.94e-8**; bf16 stream + fusion folds + the FP4 path (needs sm_121a) + a real-checkpoint run pending. **2026-08-05: the AUDIO-VAE ENCODER is ported** (DAC analysis stack + `pre_block` AttnProjection + `mean_proj`, gated stage by stage vs the checkpoint's own remote code at 2.98e-8 / 1.64e-7 / 1.86e-8) with its own checkpoint loader gated on the real 1087-tensor manifest — so **ref2va AUDIO and VIDEO+AUDIO references are now WIRED** (audio rows move by 0.51 / 0.71; a different waveform still moves them by 7.1e-4). Both VAEs are now complete in both directions. **bf16 13-SHARD RELEASE INDEXES 2026-08-07 (`row/H3-BF16-SHARDED-DIT`)**: `MiniMaxH3ShardedCheckpoint` resolves the ORIGINAL 66.3 GB release through its own `model.safetensors.index.json` (a tensor named in the index but missing from its shard throws BY NAME), `EnumerateMiniMaxH3ShardedTensors` feeds the shared shapes-only geometry parser, and `LoadMiniMaxH3DitFromShards` is the host-f32 reference loader. Gated CPU-only at 72/72/54497 (post-rebase): every tensor resolves to the shard the index named AND to the bytes written there, the derived geometry equals the single-file path field for field, and a SPARSE 13-shard release with the REAL 535 tensors at REAL shapes (66.3 GB declared, 144 KB on disk) derives the SHIPPED geometry (50/5376/56/128/14336/24/32/1x2x2/5120). **STREAMS 2026-08-07 (`row/H3-BF16-SHARDED-STREAM`)**: `StreamMiniMaxH3ShardedToDeviceBf16` uploads it one tensor at a time — a BF16 tensor bound for a bf16 device slot goes straight from the mmap with ZERO host buffer, so peak host is bounded by ONE tensor (observed `host_peak=8192`, `direct=37 converted=9`); bit-exact vs the non-streamed `StageMiniMaxH3DitWeights` reference over all 46 views with identical logits, rope.inv_freq host-resident, 73/73/55203. Spec §8.14. This UNBLOCKS the quantization-quality question; no bf16-vs-quant render or speed number is claimed. Spec §8.13. **W-FP4a LANDED (CPU) 2026-08-06 (`row/H3-FP4-SPEED`)**: the device DiT forward now routes the NVFP4 projections through the shared Marlin W4A16 dispatcher (fp4 kept packed; no new quant code), fp4-vs-bf16 wiring gate GREEN (62/62·30039). **W-FP4a GB10 leg LANDED 2026-08-06 (`row/H3-FP4-GPU-E2E`, PR #64):** on sm_121a the Marlin W4A16 path RAN for all 11 projections (`dense_gemms==11` default — VT_MARLIN_DENSE is default-ON → vLLM's own DENSE Marlin GEMM, not the grouped route; `marlin_gemms==11` under VT_MARLIN_DENSE=0; `fallback_gemms==0`), fp4-vs-bf16 BYTE-EXACT (max\|diff\|=0), and the fp4 arm is a MEMORY win not a diffusion-forward speed win (per-forward bf16/fp4 3.47× @seq64 → 0.79–0.83× @seq4224–7040; ~16 vs ~66 GB device). Real-checkpoint fp4-resident t2va e2e RUNS (real 18.75 GB NVFP4 DiT + VAEs + GGUF Qwen3-VL-32B encoder → valid mp4/wav; DiT s/step 5.45/20.0/209 s @512/768/REF-209f) but frames are a non-scene patch-grid at 12/20/50 steps → OPEN render bug (device VAE/denoise). vLLM-Omni has no quantized H3 arm (BF16-only) so any comparison is HW/loader-forced-indirect — spec §8 | `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit` | | ✅ | `LagunaForCausalLM` | Poolside Laguna-S-2.1 (118B/8B MoE) | **LONG-CTX DECODE LEVERS LANDED + MEASURED (2026-08-03, `CLAIM-LAGUNA-LONGCTX-LEVERS`): window-bounded SWA reads (`VT_LAGUNA_SWA_WINDOW`, default-ON, BYTE-EXACT) bound the four `DecodeAttnGqa*` kernels' read to the ~512 sliding window (vLLM `laguna.py:412`) — GB10 A/B token-IDENTICAL `=1` vs `=0` at 520-token context (truncation active), MEASURED −0.30 ms/step at ~2k (~0 at ≤512, grows linearly). bf16 paged KV (`VT_LAGUNA_KV_BF16`, default-OFF opt-in) a distributional near-tie left UNRATIFIED. See BENCHMARKS `CLAIM-LAGUNA-LONGCTX-LEVERS`.** — **NVFP4 W4A4 ARM RAN on GB10 (N4, 2026-08-01, `CLAIM-LAGUNA-NVFP4-N4`): the additive safetensors NVFP4 arm (N1a/N1b/N2/N3 — `Nvfp4Weight` expert fields + `LoadLagunaForCausalLMWeights` + `LqGemmNvfp4Fp4` per-expert TRUE-W4A4 + `LagunaFfnBlock` `fp4` branch + `laguna_gen` dir-autodetect; CPU-gated `test_laguna_nvfp4_loader` 3/3·61, GGUF path byte-identical) generates COHERENTLY on the real 67 GiB `poolside/Laguna-S-2.1-NVFP4`. vs the vLLM MARLIN golden (vLLM's exact prompt ids injected): FIRST 2 TOKENS MATCH exactly, then near-tie divergence (our TRUE-W4A4 fp4-activations vs the MARLIN golden's W4A16 bf16-activations — different precision, EXPECTED; shares golden vocab). SPEED (N5, trace-driven, 2026-08-01): 0.16 → ~4.5 tok/s (~28× THIS SESSION), now ~4× from vLLM 18.8. **Lever #2** (nsys found the bf16 tower running host `MatmulNK` on the CUDA queue): route it to the GPU (`LqGemm` bf16 → `CastBf16` + `MatmulBT`, weight stays bf16) → 6.34 → 0.39 s/tok (16×). **Lever #1** (nsys found the emulation expert GEMM at 92%, GPU 87% busy): the engine's native sm120a fp4 tensor-core MMA (`MatmulNvfp4Fp4Native`) reads the SAME linear scales — it was gated OFF behind `VT_NVFP4_FP4_NATIVE`; default it ON in the driver → 0.39 → ~0.20-0.24 s/tok (~2×). Both coherent + near-tie (byte-identical ids to emulation; first token matches golden). Two GB10 memory fixes landed to run (shard-release + context-before-load). OPEN #234 (remaining ~4×): grouped W4A4 MoE (top_k×3 launches → 3), `ResidentNvfp4`, decode CUDA-graph + on-GPU sampling (the host-orchestration tail). Spec `.agents/specs/laguna-nvfp4-arm-2026-07-31.md` §N4/§N5. The GGUF-Q4_K track (below) is the separate keep-quant vehicle.** Prior **FASTER DECODE (W9, 2026-07-31, `CLAIM-LAGUNA-W9-GROUPED`): the 30 un-grouped per-expert keep-quant GEMV launches/step (top_k × {gate,up,down} `LqGemmRowSlice`) fold onto the SHARED `vt::MatmulBTQuantGrouped` op — per token, Pk experts' gate/up/down each collapse to ONE grouped launch over the already-stacked `[E*N,H]` tower (no loader change). Same-binary A/B on real UD-Q4_K_XL (GB10, `--gpu`, drop_caches cold, 24 tok): grouped (`VT_LAGUNA_GROUPED_MOE=1`, default) == per-expert (`=0`) BYTE-IDENTICAL (md5 `754728c6`, both == W6 golden) + decode 0.18 → 0.13 s/tok (1.38×). Routes through the shared vt op (fold policy). Cumulative with W8: decode 0.66 → 0.13 s/tok (5.1×; 1.5 → 7.7 tok/s; 18× → 3.6× vs llama.cpp 27.8). Next lever: device-resident decode (#1). See spec §W9.** Prior **FASTER DECODE (W8, 2026-07-31, `CLAIM-LAGUNA-W8-EMBED`): `LagunaEmbed` no longer converts the whole 1.23 GB embed table to f32 every token (it gathered T rows out of the whole [Vsz,H] table via `ReadF32` — ~311M host element-converts/token, the DOMINANT decode cost the W7 profile under-filed as "#5"); now gathers only the T needed rows directly (BIT-IDENTICAL — same per-element conversion, same rows). GATED on the real 3-shard UD-Q4_K_XL GGUF (GB10, `--gpu`, W6 cached, drop_caches cold, 24 tok): TOKEN-IDENTICAL to the W5/W6 golden (`22345 83 350 785 …`, coherent " Paris.") + decode 0.66 → 0.17 s/tok = 3.9× (1.5 → 5.9 tok/s; 18× → 4.7× vs llama.cpp 27.8). See `.agents/specs/laguna-s21-w7-speed-2026-07-31.md` §W8. Next: grouped-expert GEMM (=A3) then device-resident decode.** Prior **DECODE-SPEED ATTRIBUTED (W7 profile-only, 2026-07-31, `CLAIM-LAGUNA-W7-SPEED`): `nsys` of the W6 decode (real UD-Q4_K_XL GGUF, GB10) attributes the 0.66 s/tok (~1.5 tok/s vs llama.cpp 27.8 on identical bytes, ~15-18x) to HOST-ORCHESTRATION, not kernel compute — GPU active only 32.7% of the step, 67.3% host/idle; 22,115 `cudaStreamSynchronize` (~2,764/step, zero GPU overlap) from the ~1,795 per-GEMM `DrainQueue` in `LagunaForwardGgufCached` + scalar host glue; 39.4% of GPU time is `QuantizeQ8K` activation-quant (per-GEMM), weight GEMVs un-grouped at ~22% of the 240 GB/s peak (llama.cpp ~76%); no H2D/D2H (unified memory). Ranked levers (all in-tree from ds4): device-resident decode 1.5->~5-7 tok/s, grouped-expert GEMM (`MatmulBTQuantGrouped`) +1.5-2x + dedupes the activation-quant, decode CUDA-graph, tuned MMVQ; + free host cleanups (`LagunaEmbed` copies the whole 1.23 GB embed table/token, per-token RoPE-cache rebuild). Honest reachable ~13-20 tok/s, 27.8 a stretch. NO code changed. See `.agents/specs/laguna-s21-w7-speed-2026-07-31.md`. Prior RUNNABLE + FAST DECODE (W6, 2026-07-31): a per-layer K/V cache + single-token incremental decode replaces W5's O(n²) STATELESS recompute — TOKEN-IDENTICAL (byte-equal ids, md5 `754728c6…` match, == the W5 golden) and 5.05× faster per token: decode 3.33 → 0.66 s/tok on the real UD-Q4_K_XL GGUF (GB10, `--gpu`, keep-quant), same " Paris.…" text. `LagunaKvCache` (mirrors `DeepseekV4KvCache`, MLA-latent → GQA multi-head K/V; caches post-QK-RMSNorm/post-RoPE K + raw V at f32 — bit-exact since RoPE/QK-norm are position-only and attention is causal), MIXED attention per-layer: 12 GLOBAL layers grow unbounded + 36 SLIDING-WINDOW-512 layers EVICT rows beyond the 512 window (gemma2/3 `is_sliding`); `LagunaForwardGgufCached` + shared `LagunaAttention`/`LagunaFfnBlock` helpers used by BOTH forwards (identical float ops; recompute ids unchanged after refactor), `examples/laguna_gen --stateless` A/B flag. No cache bug (bit-exact first run). Next speed = grouped-expert GEMM + device-resident decode (both in-tree from ds4). See `.agents/specs/laguna-s21-w6-2026-07-31.md`. Prior RUNNABLE (W5, 2026-07-31): our engine greedy-generates COHERENT text on the REAL 3-shard UD-Q4_K_XL GGUF (GB10 keep-quant) — "The capital of France is" → " Paris. …", first token "Paris." matches the llama.cpp-Poolside reference. Multi-shard GGUF reader + keep-quant tower (`LoadLagunaFromGgufShards`) + `LagunaForwardGguf` (ds4 keep-quant Gemm/GemmRowSlice) + `examples/laguna_gen`; load 20.6s, peak 71 GiB, 3.27 s/tok stateless recompute (speed=W6).** Prior W3: **W3 REAL forward + 3 new ops landed** (`laguna_ops.cpp`: per-head softplus attn out-gate + ungrouped sigmoid-noaux router + dual per-layer RoPE cos/sin builders; `LagunaModel::Forward` now a REAL runnable host-reference composition — variable-Q-head GQA + dual RoPE + sliding-window mask + softplus gate + dense L0 / ungrouped-MoE L1..47 + untied lm_head — replacing the W1/W2 `VT_CHECK(false)` stub; CPU `-Werror` full-library build clean; `test_laguna_scaffold` **8/8·166** incl. softplus math, router selection+tie-break RED-first, dual-RoPE cos/sin bit-match, variable-Q-head shapes, forward composition on synthetic weights; `test_model_registry` 24/24). W1 oracle DECISION: vLLM native `laguna.py` in pin ⇒ config constructs; dual-oracle = vLLM-NVFP4/-FP8 (fits GB10, BF16 235 GiB does NOT) + llama.cpp-Q4_K token-exact. DEFERRED to W4 (needs 73 GB checkpoint): GGUF keep-quant tower materialization + device/paged production forward + strict dual-oracle greedy gate. ~85-90% reuse (ds4-MoE + gemma-sliding + olmo3-dual-rope + landed Q4_K keep-quant); NEW = the 3 landed host ops + name-map + variable-Q-head device runner. **W4 (2026-07-31, `CLAIM-LAGUNA-W4`, in progress):** the UD-Q4_K_XL GGUF (73.4 GiB, 3 shards) FETCHED to dgx + its metadata/tensor-map READ AUTHORITATIVELY (814 tensors, arch `laguna`, `expert_gating_func=2` sigmoid, `leading_dense_block_count=1`, `expert_weights_scale=2.5`). Three CPU-verified FIDELITY corrections the W1-W3 scaffold got wrong, each grounded in the real GGUF + llama.cpp: (1) **per-head QK-RMSNorm** (`attn_q_norm`/`attn_k_norm` F32[128]) added to params+forward — the scope MISSED it (surfaces only in the tensor map); (2) **dual-RoPE mscale** now uses llama.cpp's `yarn_attn_factor·(1+0.1·ln(factor))` off the GGUF-authoritative `factor=32`/`yarn_attn_factor=1.0` (256K-ctx build, NOT HF's factor-128/1.4852 1M-ctx scalar) — resolves the numerics-delicate residual; (3) **separate** `ffn_gate_exps`/`ffn_up_exps` (Q4_K) + `ffn_down_exps` (Q5_K) + Q8_0 shared/attn (the scaffold assumed merged gate_up). GGUF keep-quant tower materialization (`Mw`/`Sew` mirror of ds4) + keep-quant `ForwardGguf` (vt::MatmulBT/GemmRowSlice) + the real-model greedy run vs the llama.cpp-laguna same-quant oracle remain the W5 close (73 GB single-GB10, host-orchestrated) | `MODEL-TEXT-laguna-laguna-for-causal-lm` | | 🚫 | `DeepseekV3ForCausalLM` / `DeepseekV32ForCausalLM` | DeepSeek-V3 / V3.2 | HW-blocked (671B, ~642 GiB fp8 vs 119 GiB unified memory); V3.2 additionally DEP-blocked (DSA indexer) | `MODEL-TEXT-deepseek-v2-deepseek-v3-for-causal-lm` | | 🚫 | `GlmMoeDsaForCausalLM` | GLM-5 (DSA) | HW-blocked (1404 GiB bf16) and DEP-blocked (GLM-5.x is DeepSeek-V3.2 verbatim) | `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` | diff --git a/.agents/parity-ledger.md b/.agents/parity-ledger.md index 70e42f1d5..066855ffe 100644 --- a/.agents/parity-ledger.md +++ b/.agents/parity-ledger.md @@ -917,3 +917,4 @@ Columns: | 2026-08-06 (startup-latency axis becomes measurable; extends `SERVE-GATE-ONLINE`; no new row; isolated worktree `.claude/worktrees/startup-axis` branch `feat/startup-latency-axis` off `362a3c99`; CPU-only, NO GPU; `benchmark_binding=false`) | **What it does.** Makes cold launch-to-first-`/health` a recordable axis on both arms. `scripts/dgx-online-serving.sh` gains `--startup-only` (3 interleaved ours/vLLM legs under one `/tmp/gpu` lock, page cache dropped per leg, GPU idle proven before/after, no timed client, `server` target only); `wait_ready` moves from a 5 s cadence / `seq 1 360` to a 0.2 s cadence on a deadline preserving the identical 1800 s budget; `start_server` stamps launch and ready immediately around the spawn; new `online_gate.py record-startup` + `summarize-startup`. The 5 s cadence was ~12% of a ~40 s startup, which is why this axis was never reportable despite being in the gate protocol. | **Not a vLLM mirror** — measurement tooling, recorded as such (`porting-inventory.md` §9). vLLM has no startup-latency test to port. The measured definition is deliberately stack-inclusive: whatever each engine really does before `/health` answers (vLLM: torch import, engine init, flashinfer JIT, graph capture; ours: weight load, graph capture). | **PASS (harness correctness; NO throughput number owed or claimed).** RED-first: the new suite failed on `cannot import name 'record_startup'` before the implementation existed. `tests/tools/test_online_gate_startup.py` **19/19**; run together with `test_online_gate_client`, `test_online_gate_summary`, `test_online_gate_trace`, `test_drop_file_cache` -> **66/66**. `shellcheck` + `bash -n` clean. The existing `--execute` purity contract CAUGHT a real regression (a second `for repetition in 1 2 3; do` header broke its split anchor) — fixed by renaming the loop variable. Cold-gate: `record-startup` validates and embeds the leg's cache-drop report, so a warm leg cannot produce an artifact. **This change also REPAIRS main**: squash `b95543c4` (#77) swept the half-written test file onto main without its implementation, leaving the tool suite red. **RESIDUAL (honest): NO ours-vs-vLLM startup number exists.** dgx.casa was at 100% disk with no CUDA `server` build on 2026-08-06; the 27B 3-repetition run is owed. | | 2026-08-07 (startup-latency FIRST NUMBERS, provisional; extends `SERVE-GATE-ONLINE`; no new row; `benchmark_binding=false`) | **What it does.** Runs the `--startup-only` series landed the day before: Qwen3.6-27B-NVFP4, GB10, 3 interleaved ours/vLLM repetitions under one `/tmp/gpu` lock, page cache dropped per leg, GPU idle proven before and after each. | Reference = vLLM oracle 0.25.0 in the same production server config the throughput grid launches (`--gpu-memory-utilization 0.6`, matched `--max-num-seqs`/`--max-num-batched-tokens`, prefix caching off). Attribution taken from vLLM's OWN log, not inferred. | **MEASURED, PROVISIONAL, NOT BINDING.** ours 37.94/36.51/35.88 s (median **36.51**), vLLM 460.36/221.51/217.86 s (warm median **221.51**) => **6.07x**. Ours ±3%; vLLM warm legs within 1.7%. vLLM r1's 460 s is one-time FlashInfer autotune+compile (`saved 64 configs`, `init engine 259.42 s`) vs r2 (`loaded 64 configs`, `26.95 s`); even warm, init is only ~27 s of ~221 s, so the gap is process start + imports + weight load. Our cold-autotune start = 69.29 s (+33 s over warm). **Two reasons it is not binding:** (a) a concurrent build session overlapped r2/r3 of both arms, including BOTH warm-cache vLLM legs, biasing vLLM slow and inflating the ratio; (b) the uncontended repeat was destroyed when the box HARD-REBOOTED mid-leg during vLLM's cold-autotune start (previous boot's journal ends with no shutdown sequence) - a NEW trigger for the known GB10 unified-memory reboot hazard, since the warm-cache config ran six times without incident. Owed: one uncontended 3-rep series on a quiet box. | | 2026-08-07 (`row/H3-BF16-SHARDED-DIT`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; CPU-only, no GPU and no download; lifecycle unchanged) | **MiniMax-H3 — the ORIGINAL bf16 release (13 safetensors shards, 66.3 GB) is now INDEXABLE.** Every H3 render so far used a QUANTIZED DiT and H3 is unusually quantization-sensitive (Q3_K_M -> Q4_K_M alone turned a murky lattice into a photoreal close-up; ComfyUI PR 15298 blames the partial split-half RoPE's channel-wise magnitude outliers), but the full-precision question was unaskable because every DiT loader took a SINGLE file. Adds (a) `MiniMaxH3ShardedCheckpoint::Open(dir)` (`src/vllm/model_executor/models/minimax_h3_sharded.cpp`), which resolves tensors through the checkpoint's own `model.safetensors.index.json` weight map (never by scanning) with one index over every shard, mirroring the in-tree multi-shard template `LoadMiniMaxH3EncoderWeights(const std::vector&, ...)`, and throws BY NAME when the index names a tensor its shard does not hold; (b) `EnumerateMiniMaxH3ShardedTensors`, the shapes-only manifest the geometry parser consumes; (c) `LoadMiniMaxH3DitFromShards`, the host-f32 reference loader; (d) `MiniMaxH3IsFp32IslandTensor`, single-sourcing the upstream fp32-ISLAND split the three existing streamers each hand-rolled; (e) `--dit ` in `examples/minimax_h3_gen` for both `--dump-params` and the run path, every existing `--dit` form unchanged. The DEVICE streamer is the stacked follow-up `row/H3-BF16-SHARDED-STREAM`, split out to stay inside the 900-line PR cap. | vLLM-Omni `vllm_omni/diffusion/models/minimax_h3/minimax_h3_transformer.py:85-101` (MINIMAX_H3_FP32_PARAM_NAMES / _BUFFER_NAMES, the island split) and `:906-922` (the parameter set); the shard-index container convention is HF safetensors' own `model.safetensors.index.json` weight_map, already consumed in-tree by `LoadSafetensorsIndex` and the multi-shard encoder/VAE loaders. No vLLM behavior changed; H3 remains BEYOND-PIN (vllm-omni, not the pinned vLLM repo). | **LANDED + CPU-GATED (loader brick; `benchmark_binding=false` — no throughput owed, and NO bf16-vs-quant render or speed number is claimed).** Re-gated AFTER the rebase onto `f34e0d17`: `test_minimax_h3` 72/72 cases / 54497 assertions, clean Release build of `libvllm.a`, `test_minimax_h3` and `minimax-h3-gen`. Two gates: (1) index+name mapping over a synthetic 4-shard set — every tensor resolves to the shard the index named AND to the bytes written there, a tensor missing from its shard throws WITH ITS NAME, and the derived geometry equals the single-file path field for field; (2) a SPARSE 13-shard release declaring the REAL 535 tensors at REAL shapes (66.3 GB declared, 144 KB on disk) derives the SHIPPED geometry 50/5376/56/128/14336/24/32/1x2x2/5120, and `minimax-h3-gen --dit --dump-params` prints all 20 fields on it. Also FIXES a real latent defect this row's sanitizer lane exposed: `MiniMaxH3ReadSafetensorF32` read 16-bit payloads through `reinterpret_cast`, which is UB on a safetensors file whose JSON header leaves the payload odd-aligned (the format does not require padding); now a byte-wise `memcpy`. RED-first proven: reverting it reproduces UBSan's `load of misaligned address` at the same line and exits 1. Honest residuals: no device load of the real 66.3 GB release, no measured peak RSS, and no bf16-vs-quantized render/speed comparison — the quality question is UNBLOCKED, not answered. | +| 2026-08-07 (`row/H3-BF16-SHARDED-STREAM`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; stacked on `row/H3-BF16-SHARDED-DIT`; CPU-only, no GPU and no download; lifecycle unchanged) | **MiniMax-H3 — the ORIGINAL bf16 release (13 shards, 66.3 GB) now STREAMS to the device.** §8.13 made the checkpoint indexable but its only loader was host-f32 (~132 GB on the real release); on a 122 GiB UNIFIED pool that holds the model TWICE, and the non-streaming NVFP4 loader was already OOM-killed at anon-rss 125 GB on HALF this size, so the real release was not loadable at all. Adds `StreamMiniMaxH3ShardedToDeviceBf16` (`minimax_h3_device.cpp`, sharing `BindStreamedDitViews` with the GGUF and NVFP4 streamers): manifest first, then ONE tensor at a time, with a BF16-on-disk tensor bound for a bf16 device slot — essentially the whole 66.3 GB — uploaded DIRECTLY out of the read-only mmap with NO host buffer, each source range released via `MaybeReleaseSourcePages`, and `rope.inv_freq` kept HOST-resident. Adds `MiniMaxH3ShardStreamStats` (mirroring `Nvfp4W4A16Stats`) so the path is observable, and `--dit --device cuda` in `examples/minimax_h3_gen`. | vLLM-Omni `vllm_omni/diffusion/models/minimax_h3/minimax_h3_transformer.py:85-101` (MINIMAX_H3_FP32_PARAM_NAMES / _BUFFER_NAMES, the fp32-island split the stream honours). The streaming SHAPE is our own in-tree convention (`StreamMiniMaxH3Nvfp4ToDeviceBf16`), which exists because upstream never has to load this checkpoint on one unified-memory device; recorded as a deviation in porting-inventory §9 terms. No vLLM behavior changed; H3 remains BEYOND-PIN. | **LANDED + CPU-GATED (loader brick; `benchmark_binding=false` — no throughput owed, and NO bf16-vs-quant render or speed number is claimed).** `test_minimax_h3` 73/73 cases / 55203 assertions, clean Release build of `libvllm.a`, `test_minimax_h3` and `minimax-h3-gen`. Two gates: (1) streamed == non-streamed — all 46 weight views BIT-EXACT (`memcmp == 0`) vs `StageMiniMaxH3DitWeights(kBF16)`, dtypes included (12 fp32 islands), both device forwards IDENTICAL (video and audio max|diff| == 0.0), `rope.inv_freq` host-resident; (2) the loader RAN — counters ASSERTED, observed `shards=3 tensors=46 direct=37 converted=9 bytes=444504 host_peak=8192`, i.e. BOTH upload paths taken, every view owned by this loader, and `host_peak_bytes` bounded by one tensor (< 1/4 of bytes uploaded) so the peak cannot scale with the model. Honest residuals: the real 66.3 GB load, its measured peak RSS, and CUDA memcpy from a file-backed mmap are all UNVERIFIED (CPU-only row); the bf16-vs-quant A/B is now RUNNABLE and has not been run. | diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index 42049f3b8..dcef39ff4 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -78,7 +78,7 @@ models we already ship + benchmark. Full seam map + M0–M5 W-plan: | 13 | `ROAD-V1-D4` | **KV persistent state to disk, and external KV-cache provider interoperability with LMCache** (user-directed 2026-07-22: "let's do the KV persistent state to disk support, and LMCache support too", under the standing same-featureset-as-vLLM-and-better bar) | [`KV-OFFLOAD`](engine-matrix.md), [`KV-EXTERNAL-CACHE`](engine-matrix.md), [`KV-CONNECTORS`](engine-matrix.md), [coverage view §2](feature-matrix.md#2-kv-cache--memory), [LMCache quickstart](https://docs.lmcache.ai/getting_started/quickstart.html) | spike ACCEPTED [kv-persistence-lmcache.md](specs/kv-persistence-lmcache.md) — 60 enumerated features across `vllm/v1/kv_offload/`, the `KVConnectorBase_V1` ABI and the LMCache integration, each with a DONE/PARTIAL/MISSING verdict read out of our source. **The two halves of the user's ask are NOT the same kind of work.** Disk persistence is a faithful MIRROR job and is tractable: vLLM's `fs` tier is ~101 lines of `open`/`write`/`readv` with nothing Python-specific in the byte path, one raw file per block, temp-file + atomic rename under `O_DIRECT`. LMCache is NOT: the vLLM-facing glue is vendored in-tree (~2396 lines) but every file of it imports the EXTERNAL PyPI package at module scope, and the storage engine, wire protocol, config schema and CUDA-IPC handoff all live outside the tree with no upstream test that runs without it — so it is scoped as an interop STUDY with a go/no-go, never a from-scratch client. **REOPENED 2026-07-23 ([LMCache client wire analysis](specs/lmcache-cpp-client-connector.md)), and the "no specified wire protocol" half of that verdict is REFUTED by reading the LMCache package: vLLM connects to a RUNNING LMCache over TWO fully-specified portable wires — the `lm://` remote-store (plain TCP + fixed `struct` header + raw KV bytes, no ZMQ/msgpack/pickle/CUDA-IPC) and the MP server (ZMQ + `msgspec.msgpack` + CUDA-IPC, the user's "zmq" recollection). A from-scratch C++ client is FEASIBLE with ZERO `lmcache` in-process; both wires sidestep the hash blocker because LMCache keys on its own blake3 token hash. Recommend the `lm://` mode first. Residual risk is that LMCache is an unpinned moving target — an interop feature with a version-sync cost, not a mechanical core port. LMCACHE-CLIENT W1 LANDED 2026-07-23 (`CLAIM-LMCACHE-CPP-CLIENT`, `KV-EXTERNAL-CACHE` `SPIKE`→`ACTIVE`): the pure-CPU `lm://` wire codec — fixed-`struct` `ClientMetaMessage`/`ServerMetaMessage` framing, the `CacheEngineKey` string, the blake3 rolling token hash (vendored BLAKE3 1.5.5), and the `KV_2LTD` `[2,L,T,D]` repack — is BYTE/BIT-EXACT vs fixtures from the real Python codec (`test_lmcache_codec` 6/6, 2074 assertions), blake3 verified byte-identical on x86-64 + aarch64, and INERT (no call site; the connector is client-W3). LMCACHE-CLIENT W2 LANDED 2026-07-23 — the go/no-go PASSED: a blocking POSIX-socket `LMCacheRemoteClient` (PUT/GET/EXIST/HEALTH/LIST + `KV_2LTD` repack + `VT_LMCACHE_*` config) round-trips a REAL `lmcache.v1.server` (`8570aad`, run headless from source in a throwaway venv — torch imported before lmcache to dodge a torch circular import, the compiled `c_ops` ext stubbed as unused by the lm:// CPU store) byte-identical (`test_lmcache_client` 36/36), with BIDIRECTIONAL interop proven against LMCache's OWN Python protocol codec; the always-on CI gate is a same-binary C++ mock-server round-trip (45/45, no Python). STILL `ACTIVE`, not DONE. Resume at client-W3 (wire as a `KVConnector` over the parent W5 seam, then key-agreement + DGX every-axis gates).** **Blocking correction found in OUR source:** `NONE_HASH` is seeded from `std::random_device` with no escape hatch, so every block hash differs across processes and a content-addressed disk tier would score 0% hits on restart — we are WORSE than vLLM here, which at least exposes `PYTHONHASHSEED`. This also FALSIFIES the caching spike's §B2 claim that we are deterministic by construction. **Two upstream weaknesses recorded as beyond-parity targets:** the `fs` tier's `config.json` is written and never read (its only identity check is a path digest that omits checkpoint content, weight quantization, rope config and `sliding_window` — a silent-wrong-output hazard we will not copy), and the disk tier has no capacity accounting and no eviction. Three matrix rows `INVENTORIED` -> `SPIKE`; `SharedStorageConnector` found RENAMED to `ExampleConnector` and `P2pNcclConnector` found DELETED at the pin, both stale in the prior record **W1-W3 IMPLEMENTED 2026-07-22, CPU-only.** W1 deterministic block hashes: `init_none_hash` now resolves explicit arg > `$VLLM_PREFIX_CACHING_HASH_SEED` > `$PYTHONHASHSEED` > a fixed built-in default, so hashes are identical across processes with ZERO configuration — the blocking correction is CLOSED, and we now BEAT upstream on this axis rather than trailing it (upstream is random-by-default and documents `PYTHONHASHSEED` as the operator's problem). Proven by comparing hash chains emitted by SEPARATELY LAUNCHED processes, with a negative control confirming the opt-in `=random` mode genuinely disagrees. W2 CPU primary tier: `CachePolicy` (LRU + ARC) with the `ref_cnt == -1` tri-state and the ATOMIC evict, `CPUOffloadingManager` incl. the `prepare_store -> nullopt` skip control path, pinned backing store, and a side-queue event-polled device/host transfer worker. W3 disk `fs` tier: one raw file per block, temp-file + atomic rename publish, self-healing unlink, dual-queue read/write pool. **BOTH recorded upstream weaknesses are now EXCEEDED rather than merely noted** — the identity block is a VERIFIED header read on EVERY open that REFUSES on mismatch across 27 fields (tested per field, with a positive control), and the tier carries a byte budget with policy eviction honoured across restarts. `O_DIRECT` deliberately NOT ported (a header+payload file breaks its alignment requirement); the GIL-releasing batch-lookup C extension is unconditionally unnecessary without a GIL. `KV-OFFLOAD` `SPIKE` -> `PARTIAL`. **W4 IMPLEMENTED 2026-07-23:** the TIERING MANAGER (ONE manager over CPU primary + disk secondary — disk→CPU promotion RETRY→flush→HIT, cascade demotion, reset drains the secondary first and never resets it so a persisted cache survives) and the CONNECTOR/SCHEDULER HALF (`OffloadingConnector`, the semantics of `KVConnectorBase_V1`'s scheduler hooks — nullopt third state, `block_hashes` striding, load-before-compute — wired OPT-IN + DEFAULT-OFF into the scheduler). First measured offload speedup: a restarted-prefix workload through the REAL scheduler saved 32/48 prefill tokens (2/3 blocks HIT from disk), promoted bytes byte-identical to the cold store; identity refusal holds through a promotion. Ported the SEMANTICS not the Python plugin ABI (compile-time wiring); the full abstract ABI is W5. **W5 IMPLEMENTED 2026-07-23, CPU-only:** the connector seam is now a first-class C++ ABI — the abstract `KVConnector` base carrying the full scheduler + worker method set of `KVConnectorBase_V1` (the scheduler methods load-bearing, the worker hooks defaulted no-ops for our synchronous runner, documented), a compile-time `KVConnectorFactory` + `REGISTER_KV_CONNECTOR` (the C++ analogue of vLLM's `importlib` module path), and a `KVTransferConfig` selection surface (default `kv_connector` empty == no connector == zero behaviour change, `kv_role` validation, `fail`-default load policy). The W4 disk connector was refactored ONTO this base behaviour-identically — the restart-hit e2e reproduces byte-for-byte and a config-selected owning connector shortcuts prefill by the identical 32/48. `KV-CONNECTORS` `SPIKE`→`ACTIVE`. This closes the seam so LMCache client-W3 is 'implement the abstract `KVConnector` with the landed W2 `lm://` client'. **LMCACHE-CLIENT W3 LANDED 2026-07-23 — the `lm://` client wired as a `KVConnector` over the W5 seam (`LMCacheConnector`, `REGISTER_KV_CONNECTOR("LMCacheConnector", …)`, default OFF), the FIRST time the whole chain engine -> connector -> W2 client -> a running lm:// server -> back runs.** Scheduler side computes rolling-blake3 chunk hashes and `Exist`-probes the remote store for the longest cached prefix (synchronous `(n, false)`, mirroring `lmcache_connector.py:230-259`); worker `StoreChunk`/`LoadChunk` drive the W2 client with foreign-block REFUSAL. **Gate ACHIEVED = the connector-level round-trip: STORE a prefix -> a fresh "restarted" connector LOOKS UP + shortcuts prefill through the REAL scheduler (32/48 tokens saved) -> LOAD byte-identical; foreign-key REFUSAL; default-off INERT** (`test_lmcache_connector` 5 cases / 50 assertions vs an in-process mock; store->load ALSO GREEN vs a REAL `lmcache.v1.server` 8570aad, 16 assertions, `VT_LMCACHE_LIVE_*`). **LMCACHE-CLIENT W4 LANDED 2026-07-23 — REAL peer KEY-AGREEMENT + a peer->us interop LOAD, both PROVEN — the interop-correctness milestone is COMPLETE; `KV-EXTERNAL-CACHE` stays `ACTIVE` for the DGX full-model output-invariance + throughput arm.** The actual `lm://` key derivation is `ChunkedTokenDatabase` (NOT the blake3 MP hasher): chunk_size 256, a rolling prefix-hash over `(prefix_int, tuple(tokens), extra=())` keyed by vLLM's OWN hash (portable `sha256_cbor`), folded to uint64 each step, `NONE_HASH=fold8(sha256_cbor(str(PYTHONHASHSEED)))`. Mirrored byte-exact (`chunked_token_database.{h,cpp}`, reusing `CborValue`+`sha256_cbor`) and wired as connector `key_mode=kVllmSha256Cbor` (chunk 256) alongside W3's kept-green blake3 path. Key-agreement GREEN: `test_lmcache_key_agreement` 4/85 == the REAL lmcache `ChunkedTokenDatabase.process_tokens()` BYTE-FOR-BYTE (fixtures dumped from the unmodified real driver + vLLM's pinned `sha256_cbor`/`init_none_hash`), sample `meta-llama/Llama-3.1-8B@1@0@33d6862800fff40c@bfloat16`. Peer->us LOAD GREEN over the wire: a REAL lmcache `ChunkedTokenDatabase` derives a key + PUTs KV to a REAL `lmcache.v1.server`, our C++ re-derives the SAME key and GETs the 512 B byte-identical (`run_key_interop.sh`). ASan+UBSan clean. Text-only (mm-hash extra_keys deferred). **LMCACHE-CLIENT W5 LANDED 2026-07-24 — the LAST open arm, connector-ON full-model OUTPUT-INVARIANCE + throughput in a REAL generation loop, is CLOSED (spec gates 4/6 met).** The worker side is now wired into the engine: `GPUModelRunner::execute_model` calls `ConnectorLoadExternalKv` before the forward (writes the external-prefix KV into the allocated GPU blocks, load-before-compute) and `ConnectorStorePromptKv` after (stores each newly-complete prompt block), and `LoadedEngine` builds the connector from an `EngineParams` `KVTransferConfig` and wires it to BOTH the scheduler and the runner. **OUTPUT-INVARIANCE PROVEN on a real OPT-125m loop vs a live `lmcache.v1.server`: connector-ON generated tokens are BIT-IDENTICAL to connector-OFF (cold full prefill) — first-divergence index -1 — on BOTH (a) a store->restart->load cycle in one process AND (b) a genuinely cold second process that only hits the server; prefill saved on the hit = 48 tokens (3×16-token blocks).** `tests/vllm/models/test_lmcache_output_invariance.cpp` PASSES both modes via `scripts/lmcache/run_output_invariance.sh`. THROUGHPUT reported HONESTLY: on a 125M model the wall-clock delta is noise-dominated (fixed TCP/copy overhead ~ tiny compute saved), so NO binding speedup is claimed — a real speed number is owed by an every-axis grid on a larger model + long shared-prefix corpus (docs/BENCHMARKS.md). No-regression: OPT SACRED UNCHANGED default-off (6/6, 96/96, 63/63); connector unit tests green (codec 6/6, client 3/3, connector 5/5, key-agreement 4/4, kv_offload_connector 11/11); ASan+UBSan clean on the connector path; CUDA `-Werror` 0 warnings. Additive + default-off inert (all worker/loader changes are behind a null-connector guard)| `PARTIAL` | **W7 the one genuine beyond-parity item (imperative named per-sequence save/restore), which now has the verified header it depends on; and a binding every-axis LMCache throughput grid on a larger model vs vLLM's `--kv-transfer-config`.** W5 (the abstract ABI) is DONE; LMCache client W1 (codec) + W2 (client) + W3 (connector round-trip) + W4 (peer key-agreement + interop load) + W5 (full-model output-invariance) are DONE. **The benchmark blocker is CLEARED:** the caching spike's W1 prefix-cache counters landed earlier, so the W4 offload arm proved its hits | | 13a | `ROAD-V1-D4-APC` | **Prompt / prefix caching to full vLLM parity, then beyond (user-directed 2026-07-22: "same featureset of vLLM and better")** — the headline user-facing caching feature, previously mentioned only once in this roadmap despite being a shipped, default-ON behaviour for dense models | [`KV-PREFIX-CACHE`](engine-matrix.md), [`KV-BLOCK-POOL`](engine-matrix.md), [`KV-HYBRID-COORD`](engine-matrix.md), [`KV-MAMBA-ALIGN`](engine-matrix.md), [`KV-EVENTS`](engine-matrix.md), [`KV-PREFIX-MATCH-UNIT`](engine-matrix.md), [`ENG-CASCADE-ATTN`](engine-matrix.md), [coverage view §2](feature-matrix.md#2-kv-cache--memory) | umbrella spike ACCEPTED [prefix-prompt-caching-parity.md](specs/prefix-prompt-caching-parity.md) — enumerates the complete pinned-vLLM caching surface (38 features) with a per-feature DONE/PARTIAL/MISSING verdict grounded in our source. **The ported core is deeper than the record claimed** (chain hashing, block pool, all three coordinators, the full hybrid fixed-point intersection, four single-type managers); the real gaps are narrower and different: block-hash extra keys are a no-op stub, there are NO prefix-cache statistics at any level, KV events are inert, `cache_salt` and 3 of 4 hash algos are absent, and `reset_prefix_cache` is implemented but unreachable. Three matrix rows corrected, two of them in our favour. `ENG-CASCADE-ATTN` DISPOSITIONED as not owed (default-off, absent from the MRV2 runner we port, unreachable on Blackwell). llama.cpp comparison completed: its "prompt cache" is session/slot state serialization, strictly weaker than APC on every reuse axis, and vLLM already covers disk persistence via the `kv_offload` fs tier — the ONE genuine capability neither vLLM nor we have is an imperative named per-sequence save/restore **W1 IMPLEMENTED 2026-07-22: prefix-cache statistics exist for the first time.** `BaseCacheStats`/`PrefixCacheStats`/`CachingMetrics` ported 1:1 from `vllm/v1/metrics/stats.py:35-142`, recorded in `get_computed_blocks`, flagged by `reset_prefix_cache`, taken-and-swapped per step and folded into a 1000-request sliding window exposed on `Scheduler`/`EngineCore`/`LLMEngine`. Per the standing parity-enabler rule `log_stats` is DEFAULTED ON (mirroring upstream's `disable_log_stats=False`), so no benchmark arm is void for want of a counter. `Request::num_preemptions` un-deferred to feed the mutually-exclusive `preempted_*` triple. **FIRST MEASURED HIT RATE: 0.75** (1920 of 2560 queried tokens over 16 requests sharing a 128-token prefix), with a caching-OFF 0.0 negative control — the first demonstration in this project that APC actually serves cached tokens. The hard blocker on [`BACKEND-GATE-CUDA-SGLANG-PREFIX`](backend-matrix.md) is CLOSED **W2 DONE 2026-07-27 (`CLAIM-ROADMAP-D4APC`, CPU-gated on dgx GB10, NOT pushed):** `generate_block_hash_extra_keys` ported 1:1 (`kv_cache_utils.py:451-591`) — mm hash + LoRA name + `cache_salt`, fixed order lora->mm->salt (prompt_embeds deferred: no prompt-embeds path); `cache_salt`/`lora_name` carried on `Request`/`EngineCoreRequest`, set before the first hash in `FromEngineCoreRequest` (fixes a latent ordering bug). RED-first no-false-share PROVEN: with the stub a differently-salted request false-hits the prior tenant's 48 cached tokens (`n1==48`), with extra keys `n1==0`. Ported extra-key/ordering cases + hash- and manager-level no-false-share (`test_kv_cache_utils.cpp` 29/29, `test_kv_cache_manager.cpp` 10/10). **This unblocks the MM + LoRA cache consumers.** **W3 DONE 2026-07-27 (`CLAIM-ROADMAP-D4APC-W3`, dgx GB10, NOT pushed) — the FIRST-EVER cache-ON model gate:** on `Qwen/Qwen3-4B` (dense, full-attention, APC-default-ON — the vehicle the prior "vehicle-blocked" note missed) a shared-prefix workload runs APC-ON and APC-OFF through the full paged engine, gating token-identity + hits + prefill drop. **NO engine code changed** (`git diff --stat` = tests+scripts+goldens) ⇒ pure GATE over the already-shipped default-ON path; binary byte-identical ⇒ SACRED unaffected. RESULT (`test_qwen3_apc_e2e` 2/2, 84/84 asserts): APC-ON hits **2240/2777 (rate 0.807)**, APC-OFF 0; APC-ON == APC-OFF EXACT on 5/6 (the 1 diff a vLLM-confirmed 0.125-nat near-tie, RCA'd = attention-kernel-path near-tie flip, not a cache bug); **== vLLM-APC-ON** teacher-forced (APC-OFF 6/6 max gap 0.0 nats = exact argmax, APC-ON 6/6 max gap 0.125 nats, 0 outside top-20); **TTFT drop 70.1→39.9 ms = 1.76×** on a cache hit. Existing 4B SACRED gate 16/16 GREEN (no regression). Oracle vLLM 0.25.0 (0.26 venv broken — editable source disk-reclaimed; 4B byte-stable across the pin). | `DONE` (headline) | **Row DONE for the default dense APC path (block hashing incl. extra_keys, pool, coordinators, stats, scheduling, cache-ON e2e all gated).** Named NON-BLOCKING tails tracked in their own rows / future items: W4 KV events (`KV-EVENTS` — event GENERATION + `msgpack` PAYLOAD DONE 2026-07-27 `CLAIM-ROADMAP-D4-KV-EVENTS`, `SPIKE`→`ACTIVE`, byte-exact vs `msgspec`; live ZMQ transport + engine batch wiring DEFERRED), W5 partial-block primitive (upstream dead-code), W6 Mamba-`align` hybrid cache-on (`KV-MAMBA-ALIGN`, SPIKE — feeds `BACKEND-GATE-CUDA-SGLANG-PREFIX`), W7 `reset_prefix_cache` dev-endpoint + `--prefix-caching-hash-algo` + `skip_reading_prefix_cache`, W8/W9 the beyond-vLLM named session save/restore. The every-axis cache-on grid vs vLLM/SGLang is a separate perf follow-on under `ROAD-V1-A`. No `/metrics` route yet (`SERVE-METRICS`), so the hit rate is read from the engine API. **`--prefix-match-unit` (0.26-new fine-grained matching unit) W0 spike + W1 resolver LANDED 2026-07-28 (`CLAIM-PREFIX-MATCH-UNIT`, `KV-PREFIX-MATCH-UNIT` PARTIAL): `resolve_kv_cache_block_sizes` computes `hash_block_size = prefix_match_unit if set else gcd(group_block_sizes)`, RED-first unit-gated; config/CLI/ABI field (W2) + scheduler threading of `hash_block_size != block_size` (W3, needs the `KV-BLOCK-POOL` align path) + benchmark (W4) deferred.** | | 14 | `ROAD-V1-D5` | LoRA, local KV/weight offload, expert streaming, wider model zoo | [engine matrix](engine-matrix.md), [model matrix](model-matrix.md) | corrected expert-streaming spike accepted (`ENG-EXPERT-STREAM` READY): bank-only safetensors→Marlin bank, fixed contiguous cache slots matching Marlin dense strides, logical→slot remap after explicit router D2H, chunked C` working everywhere `--dit ` did; gated CPU-only (72/72, 54497 post-rebase) on index/name mapping and on the REAL 535-tensor geometry read from a sparse 13-shard release. The DEVICE streamer the real 66.3 GB load needs is the stacked follow-up `row/H3-BF16-SHARDED-STREAM`. This unblocks the bf16-vs-quantized quality A/B; no render or speed number is claimed. Spec §8.13. **W-FP4a LANDED (CPU) 2026-08-06 (`row/H3-FP4-SPEED`)**: the NVFP4 DiT projections now keep FP4 PACKED and route through the shared `dense_nvfp4::MatmulNvfp4W4A16D` (Marlin W4A16 — vLLM's own forced-a16 selection; SAME kernel as Laguna/dense-Qwen3 NVFP4; no new quant code); fp4-vs-bf16 WIRING gate GREEN (62/62·30039, W4A16 dispatcher runs all 11 quantized GEMMs). **GB10 leg LANDED 2026-08-06 (`row/H3-FP4-GPU-E2E`, PR #64):** Marlin W4A16 RAN on sm_121a (`dense_gemms==11` default / `marlin_gemms==11` VT_MARLIN_DENSE=0, `fallback_gemms==0`), fp4-vs-bf16 BYTE-EXACT; fp4 is a MEMORY win (~16 vs ~66 GB), ~0.79–0.83× the bf16 arm per diffusion forward (compute-bound large M; 3.47× faster at small decode-like M). Real-checkpoint fp4-resident t2va e2e RUNS end-to-end (real 18.75 GB NVFP4 DiT + VAEs + GGUF Qwen3-VL-32B encoder → valid mp4/wav; DiT s/step 5.45/20.0/209 s @512/768/REF-768×1344-209f) but frames are a non-scene patch-grid at 12/20/50 steps → OPEN render-coherence bug (device VAE decode / denoise), separate from the fp4 speed work. vLLM-Omni serves NO quantized H3 (BF16-only) -> HW/loader-forced-indirect (4×B300 209f render 86.964 s vs 1×GB10 209 s/forward). | +| H3 | `ROAD-V1-H3` | **DIFFUSION generation — a new capability class.** MiniMax-H3 (`MiniMaxH3DiTModel`): omni-modal video+audio generation via a 50-step flow-matching denoise loop, ported from vLLM-Omni. Not autoregressive: no KV cache, sampler or logits. | [`MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`](model-matrix.md) | [minimax-h3 spike](specs/minimax-h3.md) | `PARTIAL` | **W0-W2 landed 2026-08-03**: packed layout (fp64 grid bit-exact), latent packing, scheduler and the full DiT forward parity-gated vs the upstream vLLM-Omni modules at reduced dims (max abs diff 1.6e-7, 10/10 cases). **W2b device-resident forward LANDED (f32) and GPU-VERIFIED 2026-08-03** — the whole DiT graph runs with activations resident in device memory, gated vs the same upstream goldens on a Thor sm_110 GPU at video 1.49e-7 / audio 8.94e-8. Only 3 H3 kernels were needed; the port reuses the tuned shared ops. Next gate: bf16 stream + `vt::FusedChain` glue folds, then the FP4 path — which needs sm_121a, since sm_110 resolves every fp4/cutlass feature DISABLED. **HW verdict CORRECTED 2026-08-03: e2e is NOT blocked** — quantized H3 checkpoints fit (GGUF ~41 GB working set; NVFP4 likewise) and the ComfyUI-GGUF arm's 535-tensor manifest already resolves onto our contract, so e2e + a speed comparison are reachable. W7 `/v1/videos` still needs a NEW MP4/AV-encoder dependency decision. **bf16 13-SHARD RELEASE INDEXES 2026-08-07 (`row/H3-BF16-SHARDED-DIT`)**: the ORIGINAL 66.3 GB bf16 DiT (13 safetensors shards) is now resolvable through its own `model.safetensors.index.json`, with a host-f32 reference loader and `--dit ` working everywhere `--dit ` did; gated CPU-only (72/72, 54497 post-rebase) on index/name mapping and on the REAL 535-tensor geometry read from a sparse 13-shard release. The DEVICE streamer landed 2026-08-07 (`row/H3-BF16-SHARDED-STREAM`, spec §8.14): one tensor at a time, zero host buffer for the bulk, bit-exact vs the non-streamed reference (73/73, 55203). This unblocks the bf16-vs-quantized quality A/B; no render or speed number is claimed. Spec §8.13. **W-FP4a LANDED (CPU) 2026-08-06 (`row/H3-FP4-SPEED`)**: the NVFP4 DiT projections now keep FP4 PACKED and route through the shared `dense_nvfp4::MatmulNvfp4W4A16D` (Marlin W4A16 — vLLM's own forced-a16 selection; SAME kernel as Laguna/dense-Qwen3 NVFP4; no new quant code); fp4-vs-bf16 WIRING gate GREEN (62/62·30039, W4A16 dispatcher runs all 11 quantized GEMMs). **GB10 leg LANDED 2026-08-06 (`row/H3-FP4-GPU-E2E`, PR #64):** Marlin W4A16 RAN on sm_121a (`dense_gemms==11` default / `marlin_gemms==11` VT_MARLIN_DENSE=0, `fallback_gemms==0`), fp4-vs-bf16 BYTE-EXACT; fp4 is a MEMORY win (~16 vs ~66 GB), ~0.79–0.83× the bf16 arm per diffusion forward (compute-bound large M; 3.47× faster at small decode-like M). Real-checkpoint fp4-resident t2va e2e RUNS end-to-end (real 18.75 GB NVFP4 DiT + VAEs + GGUF Qwen3-VL-32B encoder → valid mp4/wav; DiT s/step 5.45/20.0/209 s @512/768/REF-768×1344-209f) but frames are a non-scene patch-grid at 12/20/50 steps → OPEN render-coherence bug (device VAE decode / denoise), separate from the fp4 speed work. vLLM-Omni serves NO quantized H3 (BF16-only) -> HW/loader-forced-indirect (4×B300 209f render 86.964 s vs 1×GB10 209 s/forward). | | 15 | `ROAD-V1-D6` | **llama.cpp device breadth folded into scope (user-directed 2026-08-05):** the 11 ggml backends vLLM has no platform for — cann, musa, opencl, openvino, rpc, webgpu, zdnn, zendnn, hexagon, blas, virtgpu — inventoried as `BACKEND-GGML-*`. **SPIKES FIRST:** no implementation before each row's `.agents/specs/.md` clears the spike contract, per the standing directive. vLLM stays the mirror source; llama.cpp is the breadth reference. | [backend matrix](backend-matrix.md) | ☐ per-row spike required | `INVENTORIED` | first spike accepted | An area row cannot enter `READY` without a real spike under `specs/`, and cannot diff --git a/.agents/specs/minimax-h3.md b/.agents/specs/minimax-h3.md index fd04c7584..7bb3514a8 100644 --- a/.agents/specs/minimax-h3.md +++ b/.agents/specs/minimax-h3.md @@ -946,3 +946,57 @@ the 900-line review cap. loader ships in this row - see §8.14), no measured peak RSS, and no bf16-vs-quantized RENDER or SPEED comparison. The bf16-vs-quant quality question is UNBLOCKED, not answered. + +## 8.14 The ORIGINAL bf16 release — the DEVICE STREAMER (2026-08-07, `row/H3-BF16-SHARDED-STREAM`, CPU-only) + +Stacked on §8.13, which landed the multi-shard checkpoint and its host-f32 reference +loader. Split as its own row so each PR stays inside the 900-line review cap; the seam +is checkpoint-vs-device. + +**Why it must stream.** The reference loader materializes the whole DiT as host f32, +~132 GB on the real release. The box has 122 GiB of UNIFIED memory — host and device +draw on ONE pool — so "load to host, then stage" holds the model TWICE against that +budget, and the non-streaming NVFP4 loader was already OOM-KILLED at anon-rss 125 GB on +HALF this size. Without a streaming path the 66.3 GB release is not loadable at all. + +**What landed.** +- `StreamMiniMaxH3ShardedToDeviceBf16(queue, ckpt, out_params)` (`minimax_h3_device.cpp`, + next to its GGUF and NVFP4 twins so it shares `BindStreamedDitViews`). Manifest first + (names+shapes, no payload) so the geometry is known and a wrong name map is caught + before anything is allocated; then ONE tensor at a time. +- PEAK HOST MEMORY, precisely. Two of the three cases cost NOTHING: BF16-on-disk -> + bf16 device slot (essentially the whole 66.3 GB) and F32 -> f32 island are both + uploaded DIRECTLY out of the read-only mmap, with no host buffer at any point. Only a + dtype MISMATCH (a BF16 island widened to f32, an F32 body rounded to bf16, an F16 + shard) costs one tensor's conversion buffer, freed before the next iteration. + `host_peak_bytes` reports the largest such buffer, so the gate asserts the bound + rather than trusting the comment. +- Each source range goes to `MaybeReleaseSourcePages` the moment its copy returns, so + the page cache does not accumulate against the same pool the weights live in. + `rope.inv_freq` stays HOST-resident (`BuildRopeCosSin` runs before any kernel, so a + device pointer there segfaults on the first forward). +- `MiniMaxH3ShardStreamStats`, mirroring `Nvfp4W4A16Stats`, makes the path OBSERVABLE. + This codebase has shipped a never-executing guarded kernel under a green suite before; + counters are how that is prevented. +- `examples/minimax_h3_gen`: `--dit --device cuda` streams and prints the stats + line; `--device cpu` keeps §8.13's host reference loader. + +**Gates (CPU, `test_minimax_h3` 73/73, 55203 assertions).** +1. *Streamed == non-streamed*: over a 3-shard synthetic set whose dtypes exercise all + four (on-disk x device) combinations, every one of the 46 weight views is BIT-EXACT + (`memcmp == 0`) against `StageMiniMaxH3DitWeights(kBF16)` over the same checkpoint, + dtypes included (12 fp32 islands, the rest bf16), and both device forwards return + IDENTICAL logits (video and audio max|diff| both exactly 0.0). `rope.inv_freq` is + asserted HOST-resident. +2. *The loader RAN*: the counters are asserted, not merely printed. Observed + `shards=3 tensors=46 direct=37 converted=9 bytes=444504 host_peak=8192` — BOTH upload + paths taken, every bound view pointing at an allocation this loader owns, and + `host_peak_bytes` (one tensor's buffer) under 1/4 of the bytes uploaded, i.e. the peak + cannot scale with the model. + +**Not claimed here (needs the real weights / a GPU):** the 66.3 GB load itself and its +measured peak RSS, CUDA `cudaMemcpy` straight from a file-backed mmap (valid +pageable-source usage, and every copy is followed by a synchronize, but unexercised on +device), and any bf16-vs-quantized RENDER or SPEED comparison. The quality A/B is now +runnable; it has not been run. + diff --git a/.agents/state.md b/.agents/state.md index 2353b467b..933bfe2a0 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -40965,3 +40965,58 @@ question is UNBLOCKED, not answered. Next: `row/H3-BF16-SHARDED-STREAM` (device streamer), then the operator runs the real 13-shard bf16 DiT. + +## 2026-08-07T10:00 - H3: the ORIGINAL bf16 DiT now STREAMS to the device - one tensor at a time, zero host buffer for the bulk (row/H3-BF16-SHARDED-STREAM, CPU-only) + + +`row/H3-BF16-SHARDED-STREAM` (helper; CPU-only, no GPU job, no download), STACKED on +`row/H3-BF16-SHARDED-DIT`, which landed the multi-shard checkpoint and its host-f32 +reference loader. Split from that row so each PR stays inside the 900-line review cap; +the seam is checkpoint-vs-device, not an arbitrary cut. + +**Why a streamer and not the reference loader.** The reference loader materializes the +whole DiT as host f32 - ~132 GB on the real release. The box has 122 GiB of UNIFIED +memory, so host and device draw on ONE pool and "load to host, then stage" holds the +model TWICE against that budget; the non-streaming NVFP4 loader was already OOM-KILLED +at anon-rss 125 GB on HALF this size. So the real 66.3 GB release is not loadable at +all without a streaming path. + +**What landed.** +- `StreamMiniMaxH3ShardedToDeviceBf16(queue, ckpt, out_params)` in + `minimax_h3_device.cpp`, beside its GGUF and NVFP4 twins so it reuses + `BindStreamedDitViews`. Manifest first (names+shapes, no payload), so the geometry is + known and a wrong name map is caught before a byte is allocated; then ONE tensor at a + time. A BF16-on-disk tensor bound for a bf16 device slot - essentially the whole + 66.3 GB - is uploaded DIRECTLY out of the read-only mmap with NO host buffer at all; + same for F32 -> f32 island. Only a dtype MISMATCH (BF16 island widened, F32 body + rounded, an F16 shard) costs a host buffer, and it is freed before the next tensor. + Each source range goes to `MaybeReleaseSourcePages` the moment its copy returns, so + the page cache does not accumulate against the pool the weights live in. + `rope.inv_freq` stays HOST-resident (`BuildRopeCosSin` runs before any kernel; a + device pointer there segfaults on the first forward). +- `MiniMaxH3ShardStreamStats` (mirroring `Nvfp4W4A16Stats`) makes the path OBSERVABLE: + shards opened, tensors streamed, direct vs converted uploads, host-resident count, + bytes uploaded, and the largest host conversion buffer alive at once. +- `examples/minimax_h3_gen`: `--dit --device cuda` now streams and prints the + stats line; `--device cpu` keeps the host reference loader. + +**Gates (CPU, `test_minimax_h3` 73/73 cases / 55203 assertions; clean Release build of +`libvllm.a`, `test_minimax_h3`, `minimax-h3-gen`).** Over a 3-shard synthetic set whose +dtypes exercise all four (on-disk x device) combinations: every one of the 46 weight +views is BIT-EXACT (`memcmp == 0`) against `StageMiniMaxH3DitWeights(kBF16)` over the +same checkpoint, dtypes included (12 fp32 islands), and the two device forwards return +IDENTICAL logits (video and audio max|diff| both exactly 0.0). The counters are asserted +on, not just printed - observed `shards=3 tensors=46 direct=37 converted=9 bytes=444504 +host_peak=8192`: BOTH upload paths were taken, every bound view is owned by this +loader's staging, and `host_peak_bytes` (8192 = one tensor's buffer) is under 1/4 of the +bytes uploaded, so the peak cannot scale with the model. `rope.inv_freq` asserted +HOST-resident. + +**Honest residuals.** The real 66.3 GB load and its measured peak RSS are UNVERIFIED +here, as is CUDA `cudaMemcpy` straight from a file-backed mmap (valid pageable-source +usage, and every copy is followed by a synchronize, but unexercised on device): this row +ran CPU-only, with no GPU job and no download. NO bf16-vs-quantized RENDER or SPEED +number is claimed - the quality A/B is now runnable, and has not been run. + +Next: the operator runs the real 13-shard bf16 DiT on GB10 and answers the quantization +quality question. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 59cb9b08d..58221c66b 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -311,7 +311,7 @@ built on it rather than keeping the flattering one. | Qwen3-dense decode CUDA-graph | Token-exact pass, ~4.3% e2e directional | Steady-state per-step tok/s | | Kimi-Linear-48B-A3B (KDA+MLA+MoE) | e2e RUNS (bf16-resident §13); bf16-regime knobs 106→120/128 (7/8 exact), NOT STRICT; default OFF | bf16 residual+island-inputs → 120/128 best (control/each-alone 106; output-bf16 & f32-accum NEGATIVE); 1 near-tie left. 1.30 tok/s (O(n²)); vLLM HW-can't-serve bf16 on 1 GB10. Residual = device islands. §14 | | vLLM 0.26 re-benchmark | Pending | Re-run the binding grids on the advanced pin | -| MiniMax-H3 FP4 speed (W-FP4a) | **Measured GB10 (`row/H3-FP4-GPU-E2E`).** Marlin W4A16 byte-exact vs bf16; fp4 a memory win, 0.8x bf16/forward. Real-ckpt fp4-resident e2e RUNS (mp4/wav) | fp4 speed CLOSED. bf16-vs-quant A/B UNBLOCKED (the 13-shard bf16 DiT is now indexable) but NOT MEASURED: no bf16 render exists. Detail: benchmark-record + spec §8 | +| MiniMax-H3 FP4 speed (W-FP4a) | **Measured GB10 (`row/H3-FP4-GPU-E2E`).** Marlin W4A16 byte-exact vs bf16; fp4 a memory win, 0.8x bf16/forward. Real-ckpt fp4-resident e2e RUNS (mp4/wav) | fp4 speed CLOSED. bf16-vs-quant A/B UNBLOCKED (the 13-shard bf16 DiT now STREAMS to device) but NOT MEASURED: no bf16 render exists. Detail: benchmark-record + spec §8 | | MiniMax-H3 render coherence (`row/H3-RENDER-CLOSE` #77) | **CLOSED: a COHERENT scene on GB10.** #70/#74 white was wrong-PARTITION usage (t2va on the ref2va ckpt); t2va on the FL2VA GGUF renders a prompt-matched orange cat (adj-cos 0.95 vs 0.06, no patch-grid) | Verified first: t2va inputs byte-exact vs upstream; CUDA device==host at seq 1920. Follow-up `H3-TASK-PARTITION-GUARD`: the task/partition mismatch now RAISES 1:1 with `_resolve_task` (spec §8.6-8.7) | | MiniMax-H3 image conditioning (`row/H3-CONDITIONED-E2E`, `row/H3-VISION-SCATTER`, `row/H3-REF2VA-ASSEMBLY`) | **fl2va COHERENT; ref2va assembly bug FIXED+gated.** vision→cond scatter gated; ref2va block-dim double-division fixed + RED-first gated (128 vs 512) + a permanent ref2va DiT-forward rung (§8.10) | grid RE-ATTRIBUTED: with the fix ref2va grids in fp4 AND bf16, and t2va with no refs on the ref2va NVFP4 also grids while FL2VA-GGUF renders, so it is the **NVFP4 checkpoint/loader**, NOT assembly/fp4 (§8.10) | | MiniMax-H3 Thor render speed (sm_110, no FA2) | **34.6 s/step** at 864x480/124f/50 steps on Q4_K_M, **16.6x** off 574.5 (render ~28 min, was ~8 h). Landed: warp-per-query, chunked warp reduce-scatter (1.76x), bf16 `mma.sync` (9.82x) | Shared-memory K/V tiling (23% SLOWER) and register Q-blocking (-0.8%) both measured and REVERTED: memory traffic is not the bound (one head's K+V is 3.9 MB against 32 MB of L2) | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 6d93848e9..78c142886 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -133,7 +133,7 @@ they sit outside the gated list above. |---|---|---|---| | Voxtral audio (`VoxtralForConditionalGeneration`) | Voxtral-Mini-3B-2507 | near-tie-robust 16/16 vs vLLM 0.25.0 | decode 0.97x (beats vLLM); encoder TTFT ~17x, pending | | Whisper audio encoder | openai/whisper-small; whisper-large-v3 (Voxtral cfg) | encoder tower 77/77; large-v3 tower 203/203 | pending | -| MiniMax-H3 DiT (`MiniMaxH3DiTModel`, vllm-omni lane) | MiniMax-H3 (33.1B video+audio) | portable 72/72; t2va+fl2va COHERENT; ref2va NVFP4 grid = the community checkpoint's own quant fidelity, NO loader bug (§8.12); loads GGUF + NVFP4, INDEXES the bf16 13-shard release | FP4/Marlin landed; ref2va NVFP4 render blocked on checkpoint quant (needs official modelopt NVFP4), speed pending; no bf16 render yet | +| MiniMax-H3 DiT (`MiniMaxH3DiTModel`, vllm-omni lane) | MiniMax-H3 (33.1B video+audio) | portable 72/72; t2va+fl2va COHERENT; ref2va NVFP4 grid = the community checkpoint's own quant fidelity, NO loader bug (§8.12); loads GGUF + NVFP4, STREAMS the bf16 13-shard release to device | FP4/Marlin landed; ref2va NVFP4 render blocked on checkpoint quant (needs official modelopt NVFP4), speed pending; no bf16 render yet | | MTP speculator | Qwen3.6-27B, Qwen3.6-35B-A3B | token-identical to vLLM `mtp` at c1 | ~4% faster c1; +16% output tput (MoE) | | DFlash block-diffusion | Qwen3 (DFlash draft) | near-tie e2e 27/27 vs vLLM | 2.9x over spec-off, 1.003x vs vLLM DFlash-on | | DeepSeek-V4 MTP | DeepSeek-V4-Flash (nextn head) | lossless 5/5; real-model weight-blocked | pending | @@ -161,7 +161,7 @@ model architecture is wired. | Image | ✅ correctness-gated | ✅ | ✅ | ◐ | | Video | ✅ correctness-gated | ✅ | ✅ | ☐ | | Audio | ✅ correctness-gated | ✅ | ◐ | ◐ | -| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va+fl2va COHERENT on GB10; ref2va NVFP4 grid = the community checkpoint's own quant fidelity, NO loader bug (§8.12); DiT loads GGUF or NVFP4, and indexes the bf16 13-shard release | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | +| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va+fl2va COHERENT on GB10; ref2va NVFP4 grid = the community checkpoint's own quant fidelity, NO loader bug (§8.12); DiT loads GGUF or NVFP4, and streams the bf16 13-shard release to device | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | | Multimodal over the OpenAI server | ☐ | ✅ | ✅ | ◐ | Image, video and audio are correct through the CLI and library. Serving them diff --git a/docs/STATUS.md b/docs/STATUS.md index b15fc4e11..92396f8df 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -85,7 +85,7 @@ token-for-token correctness against the pinned oracle. | OLMo-3 dense (dual rope, interleaved sliding window) | Implemented, oracle-blocked | Loads + runs in our engine (dual rope: plain sliding + YaRN full-attn, per-layer sliding window); no SACRED gate: vLLM 0.25.0 oracle cannot run OLMo-3-1025-7B (`KeyError: 'rope_theta'`; transformers 5.13.1 nests `rope_parameters` per layer-type, no flat `rope_theta`; run-verified W0 2026-07-26) | | Laguna-S-2.1 MoE (`LagunaForCausalLM`, 118B/8B) | **BINDING 2026-08-04: 87% of vLLM (37.55 vs 43.10, SAME-TOOL nsys both engines); the whole +3.1 ms/step is the bf16 M=1 GEMV bucket (2/3 o_proj, ~196-204 vs 139 us/call, identical `gemvx` kernel); attention/MoE/glue tied or ours-ahead. Invocation match (bf16-out `cublasGemmEx`) A/B'd = WASH, ruled out; ROOT CAUSE FOUND 2026-08-04 (`VT_LAGUNA_RESIDENT_BF16W`): the bf16 projections read UNIFIED/ATS host memory, not `cudaMalloc`'d device memory — staging them device-resident (byte-exact ids) gives 38.8→44.6 tok/s (o_proj 194→131, lm_head 2410→1620 us/call), parity+ vs vLLM 43.1; **default-ON** (flip smoke-verified: canonical byte-exact ids, 44.6 clean-median). Earlier ceiling/diffuse verdicts below were cross-tool artifacts.** **REAL vLLM BAR ESTABLISHED (2026-07-31, `CLAIM-LAGUNA-VLLM-NVFP4`): FIRST-EVER vLLM Laguna run** — prior numbers (incl. the correctness oracle) were all llama.cpp, never vLLM. vLLM on official `poolside/Laguna-S-2.1-NVFP4` (single GB10, greedy, eager, MARLIN backend forced via `VLLM_TEST_FORCE_FP8_MARLIN=1` because the auto-default `FLASHINFER_CUTLASS` needs an absent `nvcc`): **~18.8 tok/s** (64-tok steady) — a LOWER bound. Our GGUF-Q4_K engine = 7.7 tok/s (vLLM ~2.4×); llama.cpp GGUF = 27.8 (still fastest at batch-1). llama.cpp is now a labeled SECONDARY "beat best-in-class GGUF" note; vLLM-NVFP4 is the headline bar. TRUE apples-to-apple still owes OUR NVFP4 Laguna forward arm (same tensor-core path as 27B/35B) — bring-up W-plan SPEC'D in `.agents/specs/laguna-nvfp4-arm-2026-07-31.md` (~85% reuse of the 35B NVFP4 W4A4 MoE infra + a name-map; bf16 attn/dense + fp4 experts; N1-N5 bricks, DGX-gated). **N1-scaffold LANDED (2026-07-31):** additive `LagunaMoeWeights.experts_{gate,up,down}_fp4` + `shared_{gate,up,down}_fp4` (`Nvfp4Weight`, mirror qwen3_5), dead until the N1 loader; CPU build clean + `test_laguna_scaffold` 8/8·167 unchanged. **N1b loader IMPLEMENTED (2026-07-31, build-verified):** `LoadLagunaForCausalLMWeights` (`laguna_weights.cpp`) replaces the `VT_CHECK(false)` stub — resolver + per-layer `LoadBf16Direct` (attn/dense/norms/embed/lm_head/router/shared-expert) + F32 `e_score_correction_bias` + `LnLoadCtNvfp4Raw` W4A4 experts. Name-map + dtypes VERIFIED against the real `poolside/Laguna-S-2.1-NVFP4` index (router `mlp.gate` BF16, bias F32, experts W4A4, shared-expert BF16). **N1b RUN-VERIFIED (2026-07-31):** loader round-trips a synthetic NVFP4 checkpoint byte-identically (`test_laguna_nvfp4_loader` 2/2·29; full detail in the benchmark record). **N2 FORWARD-BRANCH LANDED + CPU-GATED (2026-07-31):** `LqGemmNvfp4Fp4` (per-expert TRUE-W4A4: `ScaledFp4Quant(input_global_scale_inv)`→`MatmulNvfp4Fp4(alpha)`, unified-memory pattern like `LqGemm`) + `LagunaFfnBlock` branches on `fp4=!experts_gate_fp4.empty()` (routed experts fp4; keep-quant grouped fast-path gated off `!fp4`; bf16 attn/dense/router/shared-expert/lm_head unchanged) + both `LagunaForwardGguf{,Cached}` guards relaxed to `has_gguf_weights||has_nvfp4_weights`. **CORRECTION:** routed experts are W4A4 ⇒ per-expert `MatmulNvfp4Fp4`, NOT the grouped W4A16 `MoeGroupedGemmNvfp4` (grouped W4A4 deferred to N5 speed). `test_laguna_nvfp4_loader` 3/3·61 (added a forward run-gate: fp4 MoE branch runs through the real `LagunaForwardGguf` → finite+deterministic logits + routed-experts-consumed); `test_laguna_scaffold` 8/8 unchanged (GGUF byte-identical). **N3 DRIVER LANDED + CPU-SMOKE-VERIFIED (2026-07-31):** `examples/laguna_gen` auto-detects a safetensors DIRECTORY (→ NVFP4: `LoadHfConfig(config.json)` + `LoadLagunaForCausalLMWeights` + `LagunaForwardGguf{,Cached}`) vs a `.gguf` FILE (→ keep-quant), sharing the greedy loop; `--token-ids` bypass the tokenizer for the id-vs-golden gate. Verified on a synthetic NVFP4 dir with a REAL config.json (exercises the `LoadHfConfig`→`ParseLagunaParams` seam the loader test bypassed) → `has_nvfp4=1`, KV-cache decode runs finite. **N4 RAN on GB10 (2026-08-01) — the arm works end-to-end; correctness coherent+near-tie, speed 120× off.** git-archived `84fab587` → clean CUDA build (`121a`) → `laguna-gen --gpu` on the real 67 GiB `ckpt` with vLLM's exact prompt ids injected (`2,785,9626,377,15360,395`, captured via the HF tokenizer). Two GB10 memory fixes landed to run: release the mmap'd shards after the loader's memcpy-copy (114→67 GiB RSS), and create the CUDA context BEFORE the load (the 67 GiB reclaimable page cache otherwise starves `cudaStreamCreate`). **Correctness:** ours `22345 83 350 71070 395 340 9626 372 1703 …` vs golden `22345 83 290 350 674 330 5541 966 340 9626 377 15360 …` — **first 2 tokens match vLLM exactly**, then near-tie divergence; coherent ("France is" = 9626/377/15360; shares golden vocab). EXPECTED: our TRUE-W4A4 (fp4 activations) vs the MARLIN golden's W4A16 (bf16 activations) — different precision, not a bug. **Speed: 6.34 s/tok (0.16 tok/s), prefill 17.3s — ~120× slower than vLLM 18.8.** ROOT CAUSE (source-confirmed): `LqGemmNvfp4Fp4` uses the generic `vt::MatmulNvfp4Fp4` = the hand-written EMULATION CUDA kernel, NOT the cutlass sm120a fp4 tensor-core path the 27B/35B W4A4 use (`MatmulNvfp4Fp4DirectD`); + per-expert loop + per-GEMM host sync + no device residency. **nsys (2026-08-01) trace-confirmed + refined:** only 2 GPU kernels — `MatmulNvfp4Fp4Naive` = 99.3% of GPU time + fp4-quant 0.7%; GPU busy only ~18% of wall. NO bf16 GEMM on the GPU ⇒ `LqGemm`'s bf16 branch runs the host `MatmulNK` reference on the CUDA queue (attention/dense/router/shared/lm_head are CPU-bound, ~4.8 s/tok) — a second lever the source scan missed. **N5 LEVER #2 LANDED (2026-08-01) — 16× decode.** Routed the bf16 tower (attention/dense/router/shared/lm_head) off the host `MatmulNK` onto the GPU (`LqGemm` bf16 branch: `vt::CastBf16` the small activation + `vt::MatmulBT` bf16×bf16→f32, weight stays bf16 — no per-token `ReadF32` of `lm_head [100352,H]`): **decode 6.34 → 0.39 s/tok (16.3×; 0.16 → 2.56 tok/s), prefill 17.3 → 2.24s**; coherence preserved (near-tie). CPU path unchanged (run-gate byte-identical). **N5 LEVER #1 LANDED (2026-08-01) — native fp4 tensor-core, another ~2×.** The engine's native sm120a fp4 tensor-core MMA (`MatmulNvfp4Fp4Native`, `mma.sync kind::mxf4nvf4`) reads the same linear scale layout `LqGemmNvfp4Fp4` produces — it was gated OFF behind `VT_NVFP4_FP4_NATIVE`; the Laguna driver now defaults it ON (scoped; 27B/35B untouched). **decode 0.39 → ~0.20-0.24 s/tok (~2×; ~4.2-5.0 tok/s)**; coherent (byte-identical ids to the emulation path — numerically equivalent), first token matches the golden. **Cumulative N5: 0.16 → ~4.5 tok/s (~28×), now ~4× from vLLM 18.8.** **Device-resident MoE block LANDED + MEASURED (2026-08-01, `LagunaMoeResidentFp4`, `VT_LAGUNA_RESIDENT_MOE` default-ON):** the whole token's routed experts as ONE async device chain (fp4-quant→GEMM gate/up, `MoeSiluMul`, →down stacked, ONE `MoeCombine`), draining once vs ~Pk×3 syncs. **Speed EAGER-NEUTRAL (0.20 s/tok)** — empirically confirms the ds4 precedent (per-op syncs overlap GPU compute; wall is GPU-serial-bound; the graph is the payoff). **CORRECTNESS WIN: golden-token match 2 → 13** (the device `MoeSiluMul`/`MoeCombine` mirror vLLM's fused MoE faithfully). Lands default-ON (better correctness, no speed cost, graph prerequisite). **CORRECTED CEILING (from the measured state): a perfect decode graph caps at ~5.9 tok/s** (GPU already ~87% busy at 0.20 s/tok), still 3.3× short of vLLM 18.8 — the graph is necessary but NOT sufficient; the remaining 3.3× is KERNEL EFFICIENCY (native fp4 MMA ~302µs/M=1 expert GEMM vs vLLM's tuned cutlass sm120a fp4 + fused norm/quant/silu). Parity = TWO campaigns: (A) device-resident+graph → ~5.9; (B) cutlass DirectD experts + fused ops + M=1-tuned GEMV → the rest. **CAMPAIGN-B FIRST BRICK LANDED (2026-08-01): coalesced M=1 fp4 GEMV** (`MatmulNvfp4Fp4Gemv`, one warp/column, coalesced weight-row reads, `VT_NVFP4_FP4_GEMV` default-ON) — same-binary A/B: **decode 0.20 → 0.15 s/tok (1.33×; → ~6.7 tok/s), prefill 1.14 → 0.86s**, coherent+near-tie. **Cumulative this session: 0.16 → ~6.7 tok/s (~42×), now ~2.8× from vLLM 18.8.** (ILP variant `kCpw=4` measured SLOWER — 0.21 s/tok, occupancy loss > activation-reuse gain — reverted to `kCpw=1`; kernel kept templated as a re-measurable knob.) **ncu of the GEMV (sudo): sm__throughput 35-71%, DRAM n/a — COMPUTE/LATENCY-bound, not BW-bound.** Corrects the earlier "~6× BW → ~16-17 tok/s" estimate: the next GEMV lever is HARDWARE fp4 dequant (`cvt.e2m1x2`), not more bandwidth. Parity (18.8) is a multi-brick campaign (decode graph + fused norm/quant + hardware-dequant GEMV), not one more kernel. **B0 hw-fp8 SCALE-decode: MEASURED NEGATIVE, reverted (2026-08-01, `ab7a1c1e`).** Replacing the GEMV's per-byte software fp8-e4m3 group-scale decode (`F8E4M3ToF32Dev`/`ldexpf`) with hardware `cvt.rn.f16.e4m3` (`__nv_fp8_e4m3`→float) is bit-exact (ids byte-identical on the real ckpt) but paging-immune ncu shows it NEUTRAL-to-slightly-WORSE (grid768 41.2 vs 41.9µs tie; mean 53.6 vs 49.4µs) — GPU `ldexpf` is a cheap exponent-bit add, not a libcall. NOTE this is the fp8 SCALE decode, NOT the fp4-e2m1 WEIGHT dequant (the `kE2M1` `__constant__` LUT); the LUT→arithmetic/`cvt.e2m1x2` weight-dequant is a SEPARATE still-open lever (spec brick B1). Also: end-to-end wall-clock is unusable for kernel A/B here (67 GiB unified reload swings TPOT 0.16↔1.08 s/tok run-to-run) — kernel-duration ncu is the only honest anchor. **★ B2 SCOPED + DE-RISKED (2026-08-01, zero-DGX) — the real 18.8 lever:** vLLM's 18.8 bar is MARLIN W4A16 (`VLLM_TEST_FORCE_FP8_MARLIN=1`), which is LOW-M-optimized (decode-correct, unlike a tensor-core W4A4 GEMM that wastes M=1 tile rows). The engine already ships the EXACT kernel `vt::MoeGroupedGemmNvfp4Marlin` (1:1 lift of vLLM `moe_wna16_marlin_gemm`) + shared `MarlinRepackExpertWeight`, and qwen3_5 (27B/35B) already routes its NVFP4 experts through it (default-ON `VT_NVFP4_MARLIN`, 16/16-vs-oracle, +22% gate/+80% decode) via `BuildMoeMarlinResident`. So B2 = mirror that for `LagunaMoeWeights.experts_*_fp4` (a `BuildLagunaMoeMarlinResident` reusing the shared repack + route `LagunaFfnBlock`'s fp4 branch to the Marlin grouped GEMM, GEMV kept as the `=0` escape hatch) — pure reuse, no new kernel, matches vLLM's exact W4A16 numerics. **B2 IMPLEMENTED (2026-08-01, `3c49ef37`) — COMPILES CLEAN on GB10 sm_121a, runtime bug pending.** `LagunaMoeResidentMarlin` + `BuildLagunaMoeMarlinResident` (laguna.cpp, `#ifdef VT_MARLIN_NVFP4`) reconstruct the MoE Marlin path over the SHARED `dense_nvfp4::Dev`/`DBuf`/`ResidentNvfp4` + shared `vt::cuda` Marlin repack/align ops + `vt::MoeGroupedGemmNvfp4Marlin`; SACRED 27B/35B path BYTE-UNTOUCHED; gated `VT_LAGUNA_MARLIN_MOE=1` **default-OFF** (zero regression to the default GEMV path). Compiles clean on the full CUDA build. RUN: loads OK (48 layers, 256 experts) but the FIRST FORWARD device-faults silently on the Marlin path — a layout/param bug (suspects: `MoeCombine` bf16-in/f32-out dtype, the down-GEMM reusing the gate/up align, or the fp4-original free omitted → mem ~doubles). NEXT: `compute-sanitizer` localize → fix → near-tie vs the vLLM-Marlin golden + kernel-duration ncu → flip default-ON. Default path unaffected. **UPDATE (`22d6e146`): added the qwen3_5-style fp4-original free after repack** (device transients + host bytes; peak was ~3× the expert tower → past the 119 GiB pool → null-alloc → silent fault the likely cause); compiles clean. The runtime gate stayed INCONCLUSIVE this session (contended/orphaned processes on the shared box, no captured ids) — rerun on a clean uncontended session, compute-sanitizer if it still faults. **★★ B2 VALIDATED on GB10 (2026-08-01, with the mem-free fix): RUN_EXIT=0, coherent, first 13 generated tokens MATCH the vLLM-Marlin golden EXACTLY** (`22345 83 290 350 674 330 5541 966 340 9626 377 15360 81` — the best Laguna-NVFP4 correctness yet, W4A16 matching vLLM's config). **Steady-state decode 0.10 s/tok = ~10 tok/s** (steps 10-17 all 0.10; the TPOT-0.56 average is warmup-polluted — the DevicePool warms over ~9 decode steps then reuses). vs the GEMV path's 6.7 tok/s = **~1.5× faster; the gap to vLLM 18.8 closes from ~3× to ~1.9×.** Memory flat (7.9 GiB host RSS — the fp4-original free worked; it also fixed the first-forward fault). Still `VT_LAGUNA_MARLIN_MOE=1` default-OFF. TO DONE: move the lazy Marlin-resident build (216s first-forward, 48L×256E repack) to model-LOAD time → clean warm A/B + ncu → flip default-ON → matrix/roadmap. Remaining ~1.9×: vLLM graphs its decode (ours still eager) — decode CUDA-graph is the next lever. **REPRODUCED 3× (reproduction gate MET): GB10 runs deterministic — first 18-20 tokens byte-identical, steady-state 0.10 s/tok confirmed each — so the ~10 tok/s + golden-match is gated, not a single sample.** **#234 item (1) — load-time resident-build LANDED (`LagunaBuildMarlinResidents`, called from the example after load; mirrors vLLM process_weights_after_loading): builds all 48L×256E Marlin residents at LOAD so the repack is not a first-token TTFT spike. Fixed an anon-namespace linkage bug (public fn was defined with internal linkage → moved outside the anon namespace); BUILD CLEAN + links on GB10 sm_121a, default-OFF. Runtime prewarm-fires-at-load timing UNVERIFIED this session (repeated ssh-drops ate the run capture); the forward's lazy build is the validated fallback so it cannot regress. Owed: one clean run to confirm the build moved to load + then flip default-ON.** **★★ DONE (2026-08-01): Marlin is now the UNCONDITIONAL DEFAULT (`LagunaMarlinMoeEnabled` default-ON; `=0` is a code-level A/B opt-out no user needs) — "it just works" with NO env. Confirmed in a no-env GB10 run captured via tmux: `MARLIN residents built at load in 238.4s`, prefill 14.78s (build moved OUT of first-forward), golden-matching ids, steady-state 0.10 s/tok = ~10 tok/s (4th reproduction), RSS ~5-8 GiB. So a default Laguna-NVFP4 load on GB10 gets vLLM's own W4A16 Marlin decode (~10 tok/s, ~1.9× from vLLM 18.8) with zero flags. The 238s load-time repack is a one-time cost (mirrors vLLM process_weights_after_loading); optimizing its 48×256 per-expert sync count is a follow-up. Residual to 18.8 = decode CUDA-graph (deferred; user refocusing on DeepSeek next).** Post-lever-1 nsys: the remaining ~4× is HOST-SYNC-bound — 22,115 `cudaStreamSynchronize` (78.6% of API time, ~2,760/token, the per-GEMM `DrainQueue`), GPU kernels fast. Remaining levers: grouped W4A4 MoE (design input: `vt::MoeGroupedGemmNvfp4` is W4A16, so true-W4A4 grouped needs a new fp4×fp4 op or the `use_a16` mode + expert-stacking — needs a spike), device-resident decode (RECOMMENDED — the current forward is host-style so every GEMM drains; keep activations on-device, drain once/step; reuse qwen3_5's `Dev`/`Nvfp4Dev`/`ResidentNvfp4`/device-SwiGLU machinery; kills the 22k syncs; converges with the pending GGUF #228 and lifts both quant paths), decode CUDA-graph. Binding number needs a clean 2-3× re-run. See `docs/BENCHMARKS.md` + the spec N5 plan. See `docs/BENCHMARKS.md` `CLAIM-LAGUNA-VLLM-NVFP4`. Prior W7 nsys attribution: host-orchestration-bound, levers ranked (spec `laguna-s21-w7-speed-2026-07-31.md`, ledger `CLAIM-LAGUNA-W7-SPEED`). Prior RUNNABLE + FAST DECODE (W6, 2026-07-31): a per-layer K/V cache + single-token incremental decode replaces W5's O(n²) STATELESS full-recompute — TOKEN-IDENTICAL (byte-equal ids, md5 match, == the W5 golden) and 5.05× faster per token: decode 3.33 → 0.66 s/tok on the real 3-shard UD-Q4_K_XL GGUF (GB10, `--gpu`, keep-quant), same "The capital of France is" → " Paris.\n\nThe user is seeking a detailed explanation of the concept of \"cultural capital\"…". `LagunaKvCache` (mirrors `DeepseekV4KvCache`, MLA-latent → GQA multi-head K/V) caches post-QK-RMSNorm/post-RoPE K + raw V at f32 (bit-exact by construction: RoPE/QK-norm are position-only and attention is causal). MIXED attention handled per-layer: 12 GLOBAL layers grow the cache unbounded (full causal); 36 SLIDING-WINDOW-512 layers EVICT the oldest rows beyond the 512 window (gemma2/3 `is_sliding`), capping their K/V. `LagunaForwardGgufCached` + shared `LagunaAttention`/`LagunaFfnBlock` helpers used by BOTH forwards (identical float ops — the recompute path's ids are unchanged after the refactor); `examples/laguna_gen --stateless` forces the W5 recompute for the A/B gate. No cache bug: bit-exact on the first run. Next speed: grouped-expert GEMM + device-resident decode (both in-tree from ds4). See `.agents/specs/laguna-s21-w6-2026-07-31.md`. Prior RUNNABLE (W5, 2026-07-31): our engine greedy-generates COHERENT text on the REAL 3-shard UD-Q4_K_XL GGUF (GB10, keep-quant). `laguna-gen` "The capital of France is" → " Paris.\n\nThe user is seeking a detailed explanation of the concept of \"cultural capital\" as developed by French soci…" — the FIRST token is "Paris.", matching the llama.cpp-Poolside reference on the identical bytes. Multi-shard GGUF reader (LagunaGgufCtx routes each of 814 tensors to its shard; shard-1 = header only) + keep-quant tower (attn/dense/shared/experts/lm_head stay Q8_0/Q4_K/Q5_K COMPRESSED, consumed via `vt::MatmulBT`; norms/router/bias/embed → f32) + `LagunaForwardGguf` (the f32 composition with the ~9 GEMM sites swapped to keep-quant Gemm/GemmRowSlice, ds4 precedent) + `examples/laguna_gen`. Real GGUF metadata verified: dual-RoPE freq_base 500000/10000, dims 64/128, YaRN factor 32, sigmoid ungrouped-noaux router (scale 2.5), per-layer Q-head [48 global/72 sliding], per-head softplus out-gate, QK-RMSNorm. Load 20.6s, peak 71 GiB (fits 119 pool). Prior W4 IN PROGRESS (2026-07-31): 73.4 GiB UD-Q4_K_XL GGUF FETCHED + read authoritatively (814 tensors); 3 CPU-verified fidelity corrections grounded in the real GGUF + llama.cpp — per-head QK-RMSNorm (`attn_q/k_norm`, the scope MISSED it), GGUF-authoritative dual-RoPE mscale (llama.cpp `yarn_attn_factor·(1+0.1·ln(factor))`, factor 32 not HF 128), separate `ffn_gate/up_exps`. Keep-quant tower materialization + `ForwardGguf` + the real-model greedy run vs llama.cpp-laguna same-quant oracle = W5 close. Prior: W3 REAL host-reference forward + 3 new ops (`laguna_ops.cpp`, CPU `-Werror` clean, `test_laguna_scaffold` unit-gated)** | Poolside Laguna: 48 layers (12 global + 36 sliding-window-512), 256 routed top-10 + 1 shared expert, per-head **softplus attention output gate**, sigmoid `noaux_tc` router, dual per-layer RoPE (YaRN full-attn / plain sliding), GQA 8 KV / 128 head-dim, 1M ctx. **W3 (2026-07-31):** the 3 genuinely-NEW small host ops landed in `laguna_ops.cpp` — per-head softplus attn out-gate (`LagunaSoftplusHeadGate`), ungrouped sigmoid-noaux router (`LagunaUngroupedRouterTopK`, ds3 noaux_tc MINUS the group step + tie-break razor), dual per-layer RoPE cos/sin builders (`BuildLaguna{FullYarn,Sliding}CosSin`, reusing the pinned YaRN inv_freq over the partial-64 dims); `LagunaModel::Forward` is now a REAL runnable host-reference composition (variable-Q-head GQA + dual RoPE + sliding-window mask + softplus gate + dense L0 / ungrouped-MoE L1..47 + untied lm_head) replacing the `VT_CHECK(false)` stub; `test_laguna_scaffold` **8/8·166** (softplus math, router selection+tie-break RED-first, dual-RoPE bit-match, variable-Q-head shapes, forward composition on synthetic weights), `test_model_registry` 24/24. **W2 (2026-07-30):** registered, `ParseLagunaParams`, GGUF `blk.N.*` name-map + UD-Q4_K_XL quant-mix (Q4_K/Q5_K/Q6_K/Q8_0 ALL already decoded → ZERO new kernel). **W1 oracle DECISION:** vLLM NATIVE `laguna.py` (in pin → config constructs); dual-oracle = vLLM-NVFP4/-FP8 (fits GB10 119 GiB; BF16 235 GiB does NOT) + llama.cpp-Q4_K token-exact. ~85–90% reuse (ds4-MoE + Gemma-sliding + OLMo-3-dual-rope + Q4_K keep-quant, ALREADY landed). DEFERRED (W4): GGUF keep-quant tower materialization + device/paged production forward (loaders still LOUDLY throw) + strict dual-oracle greedy gate on a fetched checkpoint + `poolside_v1` parser. See `.agents/specs/laguna-s21-w3-2026-07-31.md` (+ W1/W2 `laguna-s21-w1w2-2026-07-30.md`, W0 `laguna-s21-scope-2026-07-30.md`). **Decode attention-glue fusion LANDED (2026-08-02, `CLAIM-LAGUNA-GLUE-FUSED`, default-ON `VT_LAGUNA_GLUE_FUSED`, `=0` A/B):** BYTE-EXACT L1 (softplus out-gate → `DecodeAttnCombineKernel` store) + L4 (residual-Add+RMSNorm pairs → the shared `vt::FusedChain(kFusedAddRmsNormStd)` seam) on the resident decode-graph — same-binary A/B ids byte-identical (159/159 @160), paging-immune nsys steady decode **−4.2% GPU-busy (28.90→27.69 ms/step), −120 graph nodes/step (−10%)**, wall drop_caches-tied (no regression). C shared-into-MoeCombine SKIPPED (Laguna's bf16 `MoeCombine` → not byte-exact); L2 qk-norm+RoPE preamble DEFERRED (needs a device-position kernel variant). See BENCHMARKS.md `CLAIM-LAGUNA-GLUE-FUSED`. **On-device greedy sample LANDED (2026-08-02, `CLAIM-LAGUNA-ONDEV-SAMPLE`, default-ON `VT_LAGUNA_ONDEV_SAMPLE`, `=0` A/B):** the resident decode graph used to Synchronize, return the whole `[100352]` logits, and argmax on the HOST between replays (+ host embed-gather of the next token) — the off-framework "born-on-host" seam the decode-framework-routing audit flagged. Now BOTH run ON-DEVICE inside the captured graph: `vt::GreedyArgmax` (lowest-index tie = the exact host winner) → 1-elem device token buffer, + a new capture-safe `embed_gather` kernel gathers the next input embedding from it (the stock `vt::Embedding` is NOT capture-safe: per-call event-sync + D2H ring). BYTE-EXACT (160-id stream identical `=0`/`=1` on `~/laguna-xs-nvfp4`) + faster: paired drop_caches decode wall **+0.28% median** (8/8 reps ≥0; removes ~150 us/step host argmax) at GPU-busy parity (nsys 2-length 27.44→27.42 ms/step). Aligns Laguna decode with vLLM on-device sampling. **Lever 2 (lm_head GEMV DRAM eff) MEASURED, NOT landed:** `[M=1,100352,2048]` bf16 = **170 GB/s (2.41 ms)** = ~91% of the cuBLAS M=1×large-N reference (~187 GB/s / 2.2 ms) — at the M=1 practical floor (the 273 GB/s ceiling is streaming-only, unreachable for a once-read GEMV); ≤0.7%-of-step headroom needs a reduction reorder (near-tie re-gate) ⇒ not chased, per prior "lm_head optimal". See BENCHMARKS.md `CLAIM-LAGUNA-ONDEV-SAMPLE`. **MoE add_rms_norm fold LANDED (2026-08-02, `CLAIM-LAGUNA-MOE-ADDNORM`, default-ON `VT_LAGUNA_MOE_ADDNORM_FUSED`, `=0` A/B):** the glue-fused MoE tail ran its residual update as TWO graph nodes — `vt::Add(hidden,routed)` [`AddKernel`] + `FusedChain(kFusedAddRmsNormStd)` [shared-add+RMSNorm, `RmsNormRowKernel`] — now ONE `fused_add2_rmsnorm` device node/MoE-layer (`hidden=(hidden+routed)+shared; hn=rms_norm(hidden)*w`). BYTE-EXACT (IEEE add commutes + the identical 256-thread shared-tree norm reduction; 160-id stream byte-identical `=0`/`=1` on `~/laguna-xs-nvfp4`) + faster: **−39 `AddKernel` graph nodes/step** (2.63ms→0 over 69 steps), paging-immune nsys 2-length **~−46 us/tok GPU (27339→27293)**, nsys wall **+0.4% (34.00→34.14 tok/s @70-tok)**. Small (byte-exact node-count trim on the graph-captured, GPU-bound decode; the dominant ~72% cost is the bf16 projection GEMVs — see the Lever-B negative in BENCHMARKS.md). See BENCHMARKS.md `CLAIM-LAGUNA-MOE-ADDNORM`. **Shared expert kept fp4 LANDED (2026-08-03, `CLAIM-LAGUNA-SHARED-FP4`, default-ON `VT_LAGUNA_SHARED_FP4`, `=0` A/B):** the XS-NVFP4 shared expert was DEQUANTIZED to bf16 at load (`LnLoadSharedExpertBf16`) → the M=1 decode GEMV read 4× the DRAM bytes of vLLM (which keeps it fp4). Now kept fp4-resident and routed through the SAME Marlin W4A16 single-expert (num_experts=1) grouped GEMM the routed experts win on (`dense_nvfp4::GateUpFusedMarlinD`+`MatmulNvfp4MarlinD`); the decode GEMV drops to router-ONLY (`moe.router`), shared gate/up/down go fp4. ADDITIVE new `laguna_shared_fp4.cpp` re-reads the on-disk fp4 from the gen driver before shard release (does NOT touch SACRED `laguna_weights.cpp`); bf16 shared KEPT for the T>1 prefill. NEAR-TIE (fp4≠bf16): coherent, first-20 ids == documented golden, byte-identical to bf16 for ~85 tokens then diverges; **DISTRIBUTIONAL GATE PASS 40/40** (ours' first-40 ids ∈ vLLM's 8-run greedy candidate set; vLLM XS-greedy is bf16-non-det, 8 unique of 8). FASTER: paging-immune nsys 2-length **GPU 27.24→26.53 ms/step (−2.6%)**, wall drop_caches **35.8→36.3 tok/s (+1.4%, fp4 wins all 3 reps)**; shared-expert kernel bucket ~1.68→~0.90 ms/step (halved); vs vLLM ~43 tok/s 83.3%→84.4%; RSS 22.2→22.1 GiB (freed the decode-only fused router-shared projection). Modest by design — XS's shared expert is small (`shared_expert_intermediate_size==moe_intermediate_size==512`). Default-ON per parity (matches vLLM's fp4 shared). See BENCHMARKS.md `CLAIM-LAGUNA-SHARED-FP4`. **qk-norm+RoPE preamble fusion LANDED (2026-08-03, `CLAIM-LAGUNA-PREAMBLE-FUSED`, default-ON `VT_LAGUNA_PREAMBLE_FUSED`, `=0` A/B):** closes the `CLAIM-LAGUNA-GLUE-FUSED` L2 deferral — the decode graph ran the per-layer attention preamble as FOUR under-occupied M=1 nodes (`rms_norm_seq(q)`+`rms_norm_seq(k)`+`rope_from_cache_g(q)`+`rope_from_cache_g(k)`); now ONE capture-safe `fused_qk_norm_rope_g` node/layer (`FusedQkNormRopeGKernel`, one block/head, reads the decode position from DEVICE `*pos_buf`, handles the per-layer dual-RoPE 64/128 + `Hq` 48/64). BYTE-EXACT BY CONSTRUCTION: it replicates the composed path's f32 MEMORY round-trip (Phase A 256-thread Σx² == `RmsNormSeqKernel`; Phase B the same `(x*inv)*w` store; `__syncthreads`; Phase C the `RopeFromCacheGKernel` rope read back) — an earlier register-only recompute was numerically-equivalent but diverged at a token-110 near-tie via compiler fma-contraction; the memory boundary forces bit-identity. 160-id stream byte-identical `=0`/`=1` on `~/laguna-xs-nvfp4` (determinism verified `=0`×3/`=1`×3 each run-to-run identical). FASTER: preamble norm+rope kernels **160→40 launches/tok, 326→154 us/tok (−0.17 ms/step)**; all decode-scaling kernels 26.53→26.37 ms/step; wall drop_caches **36.42→36.64 tok/s (+0.6%, fused wins all 3 paired reps)**; vs vLLM ~43 84.7%→85.2%. Modest (preamble ~1.2% of the 26.5 ms/step decode; the dominant cost stays the bf16 projection GEMVs at cuBLAS parity) — a byte-exact graph-node/launch trim (the glue-fusion residual mechanism). Default-ON per parity. See BENCHMARKS.md `CLAIM-LAGUNA-PREAMBLE-FUSED`. **W7 two-front pass LANDED (2026-08-03, `CLAIM-LAGUNA-W7-DECODE`):** FRONT 1 — the example driver logged `[gen] step N …(RSS)` EVERY decode step, and the RSS arg calls `CurResidentGiB()` (a `/proc/self/status` read) + an unbuffered stderr write in the GPU-idle gap between replays; guarded behind `VT_LAGUNA_STEP_LOG` (default OFF) + added a `decode_wall` line (TRUE end-to-end throughput incl. per-step gaps) next to the gap-free `decode_hp`. Since the fprintf sat OUTSIDE the `s0→s1` timer, `decode_hp` was ALREADY honest; with the log off `decode_wall == decode_hp` (within 0.001 tok/s, every LOG_OFF rep) and the recovered host tax is only ~0.1% (drop_caches noise floor). CONCLUSION: the ~86% gap to vLLM 43 is genuine device compute, NOT a harness artifact. FRONT 2 — `VT_LAGUNA_MOE_ONECAST` (default ON): a MoE layer cast the same `hn[1,H]` f32→bf16 THREE times (router GEMV + routed Marlin + shared Marlin); now cast ONCE into a persistent buffer and reuse (`CastHnBf16`/`GemmBf16Pre` + optional pre-cast param on both `…Into` helpers). BYTE-EXACT (deterministic truncation; `=1` vs `=0` byte-identical 300-tok ids); `CastBf16` **200→122 nodes/step (−78 = 2×39 MoE layers)**, GPU-busy parity within nsys noise, decode_hp +0.29%. Combined (onecast on + log off) **36.97 tok/s = 86.0% of vLLM-NVFP4 43** (from 36.64/85.2%). See BENCHMARKS.md `CLAIM-LAGUNA-W7-DECODE`. **Tail-fold follow-up LANDED (2026-08-03, `CLAIM-LAGUNA-TAIL-FUSED`, default-ON `VT_LAGUNA_TAIL_FUSED`, `=0` A/B):** a fresh node-ranking of the baseline decode graph found the routed-MoE `CastF32` as the one clean byte-exact fold left; it folds into the trailing `fused_add2_rmsnorm` via a new bf16-x1 sibling kernel (`AddAdd2RmsNormStdBf16Kernel` — `MoeCombine` writes bf16 straight to a persistent buffer, widened in-kernel by `__bfloat162float`). BYTE-EXACT (`=1` vs `=0` byte-identical 160-tok ids), `CastF32` **78→39 nodes/step**, total graph nodes **919→880**, GPU-busy parity; decode_hp a WASH (median +0.14% / mean −0.04%, at the drop_caches noise floor). Lands on the deterministic node-count basis (like onecast/preamble/addnorm), NOT a wall win; combined headline UNCHANGED **36.97 tok/s = 86.0%**. The ranking confirms the byte-exact decode-tail fold tier is now essentially EXHAUSTED (residual tail = already-folded norms + attention compute + cuBLAS-adjacent router/topk + ported-Marlin `MoeAlign`/`SiluAndMul`/`MoeCombine`); the gap to vLLM 43 is genuine device compute at the practical ceiling. See BENCHMARKS.md `CLAIM-LAGUNA-TAIL-FUSED`. **KERNEL-EFFICIENCY tier (2026-08-03, `VT_LAGUNA_FAST_NORM` default ON + f32 ext of `VT_RMSNORM_DECODE_FAST`):** the fold tier was exhausted but the residual-stream norm KERNELS were still under-occupied — `ncu` on the shipped `<<<1,256>>>` `AddAdd2RmsNormStdBf16`/`RmsNormRow` decode norms: `launch__waves_per_multiprocessor≈0.00`, `sm__throughput≈0.06%` (one 256-thread block on 1 SM of ~100+, latency-bound). Porting the PROVEN bit-identical `RmsNormRowFastKernel` structure (1024-thread float4 memory passes; 256-strided-partial + tree reduction reproduced byte-for-byte) to the f32 kernels cut each **286→~155 µs/tok (1.85×)**, **byte-exact** (160-tok ids identical `=1`vs`=0`; the f32 fix vs the bf16 sibling: store `v` not `v²` and square in the reduction so nvcc emits shipped's `acc += v*v` **fma** — a pre-squared f32 `v²` is not exact and flipped an XS near-tie at tok 108). **−0.81% decode-step GPU time** (paging-immune 70-vs-20 2-length diff, 26192→25980 µs/step); wall-clock ON/OFF overlap (noise floor). Residual: the byte-exact 256-strided reduction can't reach vLLM's per-kernel norm floor (~2.4× vLLM) without breaking byte-exactness → that remainder is byte-exactness-BLOCKED. See BENCHMARKS.md `CLAIM-LAGUNA-FAST-NORM`. **Router top-k warp-shuffle LANDED (2026-08-03, `CLAIM-LAGUNA-TOPK-SHFL`, default-ON `VT_LAGUNA_TOPK_SHFL`, `=0` A/B): BYTE-EXACT** — an nsys 2-length rank of the remaining small kernels (past the at-parity `gemvx` projection GEMVs ~69% of step + Marlin MoE) put the router `SigmoidTopKKernel` top (415 µs/step); `ncu` showed it `<<<1,256>>>` at `waves≈0.000`/`sm≈0.2%` — pure latency (8 serially-dependent rounds × a ~10-sync `sh[256]` argmax tree). New `SigmoidTopKShflKernel` reduces each round by warp-shuffle argmax (2 syncs/round; argmax over the total order is associative ⇒ SAME winner) → **`SigmoidTopK` 414.6→248.8 µs/step (1.67×)**, decode-step GPU **−0.57%** (26.018→25.869 ms/step), 37.39→37.49 tok/s decode_hp (**87.2% of vLLM-NVFP4 43**); 160-id stream byte-identical `=1`vs`=0`. **NOT landed — norm warp-shuffle (`VT_LAGUNA_NORM_SHFL`):** a near-tie register-accumulate+shuffle reduce for the Laguna `AddAdd2RmsNormStd{,Bf16}Fast` norms PASSED the distributional gate (coherent, in-set 38/40 = baseline, one near-tie fork at pos 37) and was −19.3% per-kernel (`AddAdd2RmsNormStdBf16` 150.3→121.3 µs/step) BUT washed at whole-step (0.6% of step; +0.02% within noise) — a near-tie fork isn't justified by a below-noise gain, so it was dropped. The small-kernel norm tail is at its occupancy floor; the decode step is dominated by the at-parity projection GEMVs. See BENCHMARKS.md `CLAIM-LAGUNA-TOPK-SHFL`. **Shared-expert 2-stream overlap LANDED (2026-08-03, `CLAIM-LAGUNA-SHARED-AUX`, default-ON `VT_LAGUNA_SHARED_AUX`, `=0` A/B):** mirror of vLLM's `MULTI_STREAM_OVERLAPPED` — in `LagunaGraph::RunChain` the fp4-shared arm's shared expert is EARLY-forked onto a second CUDA stream from the post-attn hidden `hn` BEFORE the router GEMV (aux reads `hn` f32 + does its own byte-identical cast; scratch from `AuxPool`), overlapping router+`sigmoid_topk`+routed grouped GEMM, joined before the combine — the SAME machinery the 35B ships default-ON (ENG-MOE-SHARED-AUX, runs inside the captured graph). This is the EARLY fork the prior fused-`router_shared_gu` attempt (`89e0d074`, −0.35% wash) could not reach. Capture-safe (aux stream+2 events in the ctor; gstate-0 warm-run builds residents + warms `AuxPool`). **BYTE-EXACT** (`=1`vs`=0` byte-identical 63-tok ids). REAL concurrency: nsys `--cuda-graph-trace=node` 20↔70 sum-vs-union → OVERLAP **2.34 ms/step** (SUM/UNION 1.092) vs `=0`'s 0.0004 ms; net GPU-busy wall **26.213→25.467 ms/step (−2.9%, 38.15→39.27 tok/s)**, wall @200 37.08→37.93 (+2.3%). Net device in " + << std::chrono::duration(std::chrono::steady_clock::now() - t0).count() + << " s (" << st.tensors_streamed << " tensors, " << st.direct_uploads + << " direct, " << st.converted_uploads << " converted, peak host buffer " + << (st.host_peak_bytes / (1024.0 * 1024.0)) << " MiB)\n"; + } else { + // Host f32 reference path. ~132 GB on the real release — usable only on a + // reduced checkpoint; a real run wants --device cuda. + dit = vllm::LoadMiniMaxH3DitFromShards(ckpt); + } } else { const vllm::SafetensorsFile f = vllm::SafetensorsFile::Open(dit_path); if (device_name == "cuda") { diff --git a/include/vllm/model_executor/models/minimax_h3.h b/include/vllm/model_executor/models/minimax_h3.h index 20d22d7b7..2dbfb8f42 100644 --- a/include/vllm/model_executor/models/minimax_h3.h +++ b/include/vllm/model_executor/models/minimax_h3.h @@ -1499,12 +1499,58 @@ std::vector EnumerateMiniMaxH3ShardedTensors( // The REFERENCE (non-streaming) multi-shard loader: materialize every tensor as // host f32 and bind the forward's views, exactly as LoadMiniMaxH3DitFromNvfp4 -// does for the single-file NVFP4 arm. It is the CPU path for small checkpoints -// and the comparison baseline a device streamer is gated against; on the REAL -// 66.3 GB release it would need ~132 GB of host f32 and must NOT be used — that -// release needs a streaming device loader, which this row does not yet ship. +// does for the single-file NVFP4 arm. It is the comparison baseline the streamer +// is gated against and the CPU path for small checkpoints; on the REAL 66.3 GB +// release it would need ~132 GB of host f32 and must not be used — that is what +// StreamMiniMaxH3ShardedToDeviceBf16 exists for. MiniMaxH3GgufDit LoadMiniMaxH3DitFromShards(const MiniMaxH3ShardedCheckpoint& ckpt); +// ★ Stream the ORIGINAL bf16 DiT from its shards STRAIGHT ONTO THE DEVICE, one +// tensor at a time — the multi-shard twin of StreamMiniMaxH3Nvfp4ToDeviceBf16. +// +// It MUST stream. 66.3 GB cannot be materialized on the host and then staged: the +// box has 122 GiB of UNIFIED memory (host and device share ONE pool), and the +// non-streaming NVFP4 loader was already OOM-KILLED at anon-rss 125 GB doing +// exactly that. Peak host anonymous memory here is ONE tensor's conversion buffer +// (freed before the next), and for the real release it is very nearly ZERO: a +// BF16-on-disk tensor bound for a bf16 device slot is uploaded DIRECTLY out of +// the read-only mmap with no host buffer at all, and each source range is +// released (MaybeReleaseSourcePages) the moment its copy returns, so the page +// cache does not accumulate against the same pool the weights live in either. +// +// The fp32 ISLANDS (see MiniMaxH3IsFp32IslandTensor) stay f32; everything else is +// staged bf16. `rope.inv_freq` stays HOST-resident on the returned struct. +MiniMaxH3DitDeviceWeights StreamMiniMaxH3ShardedToDeviceBf16( + vt::Queue& queue, const MiniMaxH3ShardedCheckpoint& ckpt, + MiniMaxH3DitParams* out_params = nullptr); + +// --- "this loader actually RAN" counters ----------------------------------- +// A green suite over a path that silently fell back to another loader is a +// failure mode this codebase has hit before (a kernel guarded to `num_reqs == 1` +// shipped never-executing while the suite stayed green), so the streamer is made +// OBSERVABLE and the gate asserts on it. Mirrors dense_nvfp4_gemm.h's +// Nvfp4W4A16Stats. +struct MiniMaxH3ShardStreamStats { + uint64_t shards_opened = 0; // shards the checkpoint resolved to + uint64_t tensors_streamed = 0; // tensors uploaded to the device + uint64_t direct_uploads = 0; // uploaded straight from the mmap, NO host copy + uint64_t converted_uploads = 0; // needed one host dtype conversion first + uint64_t host_resident = 0; // kept on the HOST on purpose (rope.inv_freq) + uint64_t bytes_uploaded = 0; // total device bytes staged + uint64_t host_peak_bytes = 0; // largest host conversion buffer alive at once +}; + +inline MiniMaxH3ShardStreamStats& MutableMiniMaxH3ShardStreamStats() { + static MiniMaxH3ShardStreamStats s; + return s; +} +inline MiniMaxH3ShardStreamStats GetMiniMaxH3ShardStreamStats() { + return MutableMiniMaxH3ShardStreamStats(); +} +inline void ResetMiniMaxH3ShardStreamStats() { + MutableMiniMaxH3ShardStreamStats() = MiniMaxH3ShardStreamStats{}; +} + MiniMaxH3DitDeviceWeights StageMiniMaxH3DitWeights(vt::Queue& queue, const MiniMaxH3DitParams& params, const MiniMaxH3DitWeights& host, diff --git a/src/vllm/model_executor/models/minimax_h3_device.cpp b/src/vllm/model_executor/models/minimax_h3_device.cpp index d4eeb832e..a0e2a0b34 100644 --- a/src/vllm/model_executor/models/minimax_h3_device.cpp +++ b/src/vllm/model_executor/models/minimax_h3_device.cpp @@ -30,6 +30,7 @@ // own accumulation orders. f32 is what upstream torch RMSNorm does, so the device // path is arguably the closer mirror. It is gated against the SAME upstream // goldens at the SAME tolerance as the CPU forward. +#include #include #include #include @@ -1434,4 +1435,139 @@ MiniMaxH3DitDeviceWeights StreamMiniMaxH3Nvfp4ToDeviceFp4(vt::Queue& queue, return staged; } +// ★ The ORIGINAL bf16 release (13 shards, 66.3 GB), streamed STRAIGHT ONTO THE +// DEVICE — the multi-shard twin of StreamMiniMaxH3Nvfp4ToDeviceBf16, and the +// first loader that can open the full-precision DiT at all. +// +// WHY IT MUST STREAM. The box has 122 GiB of UNIFIED memory: host and device +// share ONE pool, so "load to host, then stage to device" holds the model TWICE +// against the same budget. The non-streaming NVFP4 loader was already OOM-KILLED +// at anon-rss 125 GB doing exactly that, and this checkpoint is twice its size. +// So: convert ONE tensor, upload it, let the host buffer die before the next. +// +// PEAK HOST MEMORY, precisely. Two of the three cases cost NOTHING: +// * BF16 on disk -> bf16 device slot: uploaded DIRECTLY out of the read-only +// mmap. No host buffer exists at any point. This is the case for essentially +// the entire 66.3 GB (every norm, every projection). +// * F32 on disk -> f32 island: same, a direct mmap upload. +// * anything else (a BF16 island widened to f32, an F32 tensor rounded to +// bf16, an F16 shard): ONE tensor's conversion buffer, freed at the end of +// its iteration. `host_peak_bytes` reports the largest such buffer, so the +// gate asserts on the bound rather than trusting this comment. +// Each source range is handed to MaybeReleaseSourcePages the moment its copy +// returns (the same windowed release every other loader uses, default ON), so +// the page cache does not accumulate against the pool the weights live in. +MiniMaxH3DitDeviceWeights StreamMiniMaxH3ShardedToDeviceBf16( + vt::Queue& queue, const MiniMaxH3ShardedCheckpoint& ckpt, + MiniMaxH3DitParams* out_params) { + vt::Backend& backend = vt::GetBackend(queue.device.type); + const bool trace = std::getenv("VT_H3_PROGRESS") != nullptr; + + // Manifest first: names + shapes only, no payload, so the geometry is known + // (and a wrong name map is caught) before a single byte is allocated. + const std::vector manifest = EnumerateMiniMaxH3ShardedTensors(ckpt); + const MiniMaxH3DitParams params = ParseMiniMaxH3DitParamsFromGgufManifest(manifest); + if (out_params != nullptr) *out_params = params; + + MiniMaxH3ShardStreamStats& stats = MutableMiniMaxH3ShardStreamStats(); + stats = MiniMaxH3ShardStreamStats{}; + stats.shards_opened = static_cast(ckpt.ShardCount()); + + MiniMaxH3DitDeviceWeights staged; + std::map views; + size_t done = 0; + for (const MiniMaxH3TensorSpec& spec : manifest) { + const StTensor& t = ckpt.Get(spec.name); + + // rope.inv_freq is consumed on the HOST (BuildRopeCosSin runs before any + // kernel); binding a device pointer here segfaults on the first forward. + if (spec.name == "rope.inv_freq") { + staged.rope_inv_freq_host = MiniMaxH3ReadSafetensorF32(t); + staged.weights.rope_inv_freq = vt::Tensor::Contiguous( + staged.rope_inv_freq_host.data(), DType::kF32, vt::Device{}, + {static_cast(staged.rope_inv_freq_host.size())}); + MaybeReleaseSourcePages(t.data, t.nbytes); + ++stats.host_resident; + continue; + } + + int64_t numel = 1; + for (int64_t d : spec.shape) numel *= d; + VT_CHECK(numel > 0, "minimax_h3 sharded stream: tensor '" + spec.name + + "' has an empty logical shape"); + + const bool island = MiniMaxH3IsFp32IslandTensor(spec.name); + const DType want = island ? DType::kF32 : DType::kBF16; + const bool already = (island && t.dtype == "F32") || (!island && t.dtype == "BF16"); + + std::vector f32; + std::vector bf16; + const void* src = nullptr; + size_t bytes = 0; + if (already) { + // The whole point: the on-disk bytes ARE the device bytes, so there is no + // host copy to peak on. + VT_CHECK(t.nbytes == static_cast(numel) * (island ? 4u : 2u), + "minimax_h3 sharded stream: tensor '" + spec.name + + "' byte span does not match its shape"); + src = t.data; + bytes = t.nbytes; + ++stats.direct_uploads; + } else { + // Whatever the shard stored it as (BF16 island, F32 body, F16 either way), + // land on the dtype the forward needs — one tensor at a time. + f32 = MiniMaxH3ReadSafetensorF32(t); + VT_CHECK(static_cast(f32.size()) == numel, + "minimax_h3 sharded stream: tensor '" + spec.name + + "' read produced the wrong element count"); + if (island) { + src = f32.data(); + bytes = f32.size() * sizeof(float); + } else { + bf16.resize(f32.size()); + for (size_t i = 0; i < f32.size(); ++i) { + uint32_t bits; + std::memcpy(&bits, &f32[i], sizeof(bits)); + // round-to-nearest-even, the same rounding vt uses on a bf16 store + const uint32_t rounded = bits + 0x7FFFu + ((bits >> 16) & 1u); + bf16[i] = static_cast(rounded >> 16); + } + // The f32 staging buffer dies HERE, not at the end of the iteration: it + // is the larger of the two, and holding both is the only way this branch + // could double its own peak. + const size_t f32_bytes = f32.size() * sizeof(float); + f32.clear(); + f32.shrink_to_fit(); + stats.host_peak_bytes = std::max( + stats.host_peak_bytes, static_cast(f32_bytes + bf16.size() * 2)); + src = bf16.data(); + bytes = bf16.size() * sizeof(uint16_t); + } + stats.host_peak_bytes = std::max(stats.host_peak_bytes, bytes); + ++stats.converted_uploads; + } + + void* pdev = backend.Alloc(bytes); + std::shared_ptr owner(pdev, [&backend](void* q) { backend.Free(q); }); + backend.Copy(queue, pdev, src, bytes); + backend.Synchronize(queue); // the host buffer dies at the end of this iteration + views[spec.name] = dense_attn::MakeTensor(pdev, want, queue.device, spec.shape); + staged.storage.push_back(std::move(owner)); + // Copied-then-dead: drop the shard pages this tensor occupied. On a unified + // box the page cache competes with the model for the same pool. + MaybeReleaseSourcePages(t.data, t.nbytes); + + ++stats.tensors_streamed; + stats.bytes_uploaded += bytes; + if (trace && (++done % 50 == 0 || done == manifest.size())) { + std::fprintf(stderr, "[h3] shard-streamed %zu/%zu tensors (last: %s)\n", done, + manifest.size(), spec.name.c_str()); + std::fflush(stderr); + } + } + + BindStreamedDitViews(views, params, &staged.weights); + return staged; +} + } // namespace vllm diff --git a/tests/vllm/models/test_minimax_h3.cpp b/tests/vllm/models/test_minimax_h3.cpp index 3f7171bf3..c8d406307 100644 --- a/tests/vllm/models/test_minimax_h3.cpp +++ b/tests/vllm/models/test_minimax_h3.cpp @@ -4180,6 +4180,56 @@ std::set ShardedGateF32Names() { "rope.inv_freq", "blocks.0.norm1.weight"}; } +// Every weight view in binder order, so a comparison covers the WHOLE contract +// rather than a spot check. +std::vector> AllDitViews( + const vllm::MiniMaxH3DitWeights& w) { + std::vector> out; + auto add = [&out](const std::string& name, const vt::Tensor& t) { + out.emplace_back(name, &t); + }; + add("video_patch_proj.weight", w.video_patch_proj_w); + add("video_patch_proj.bias", w.video_patch_proj_b); + add("audio_patch_proj.weight", w.audio_patch_proj_w); + add("audio_patch_proj.bias", w.audio_patch_proj_b); + add("condition_proj.weight", w.condition_proj_w); + add("condition_proj.bias", w.condition_proj_b); + add("time_embedder.proj_in.weight", w.time_proj_in_w); + add("time_embedder.proj_in.bias", w.time_proj_in_b); + add("time_embedder.proj_out.weight", w.time_proj_out_w); + add("time_embedder.proj_out.bias", w.time_proj_out_b); + auto add_block = [&](const std::string& prefix, const vllm::MiniMaxH3DitBlockWeights& b, + bool adaln) { + add(prefix + ".norm1.weight", b.norm1); + add(prefix + ".norm2.weight", b.norm2); + add(prefix + ".attn.qkv_proj.weight", b.qkv_proj); + add(prefix + ".attn.q_norm.weight", b.q_norm); + add(prefix + ".attn.k_norm.weight", b.k_norm); + add(prefix + ".attn.out_proj.weight", b.out_proj); + add(prefix + ".mlp.fc1.weight", b.fc1); + add(prefix + ".mlp.fc2.weight", b.fc2); + if (adaln) { + add(prefix + ".adaln_proj.linear.weight", b.adaln_w); + add(prefix + ".adaln_proj.linear.bias", b.adaln_b); + } + }; + for (size_t i = 0; i < w.refiner.size(); ++i) { + add_block("token_refiner.blocks." + std::to_string(i), w.refiner[i], false); + } + add("token_refiner.final_norm.weight", w.refiner_final_norm); + for (size_t i = 0; i < w.blocks.size(); ++i) { + add_block("blocks." + std::to_string(i), w.blocks[i], true); + } + add("final_layer.norm.weight", w.final_norm); + add("final_layer.adaln_proj.linear.weight", w.final_adaln_w); + add("final_layer.adaln_proj.linear.bias", w.final_adaln_b); + add("final_layer.video_out.weight", w.video_out_w); + add("final_layer.video_out.bias", w.video_out_b); + add("final_layer.audio_out.weight", w.audio_out_w); + add("final_layer.audio_out.bias", w.audio_out_b); + return out; +} + } // namespace TEST_CASE("minimax_h3: the multi-shard index resolves every tensor to its own shard") { @@ -4293,11 +4343,176 @@ TEST_CASE("minimax_h3: the multi-shard index resolves every tensor to its own sh CHECK_THROWS(vllm::MiniMaxH3ShardedCheckpoint::Open("/tmp")); } +TEST_CASE("minimax_h3: the sharded bf16 DiT STREAMS to the device, matching the reference") { + const MiniMaxH3DitParams want = ShardedGateParams(); + const std::vector entries = + BuildMiniMaxH3DitEntries(want, /*quantize=*/false, /*plain_bf16=*/true, + ShardedGateF32Names()); + const std::string dir = "/tmp/minimax_h3_sharded_stream"; + const size_t kShards = 3; + WriteMiniMaxH3ShardedDit(entries, dir, kShards); + const vllm::MiniMaxH3ShardedCheckpoint ckpt = vllm::MiniMaxH3ShardedCheckpoint::Open(dir); + + // The NON-STREAMED reference: materialize to host f32, then stage with the + // existing (load-everything-then-upload) stager. This is the arm the streamer + // must reproduce, and it is the arm that CANNOT be used on the real 66.3 GB + // release — it holds the model twice against a 122 GiB UNIFIED pool. + vt::Queue q = vt::GetBackend(vt::DeviceType::kCPU).CreateQueue(); + const vllm::MiniMaxH3GgufDit reference = vllm::LoadMiniMaxH3DitFromShards(ckpt); + const vllm::MiniMaxH3DitDeviceWeights ref_staged = + vllm::StageMiniMaxH3DitWeights(q, reference.params, reference.weights, vt::DType::kBF16); + + vllm::ResetMiniMaxH3ShardStreamStats(); + MiniMaxH3DitParams streamed_params; + const vllm::MiniMaxH3DitDeviceWeights streamed = + vllm::StreamMiniMaxH3ShardedToDeviceBf16(q, ckpt, &streamed_params); + const vllm::MiniMaxH3ShardStreamStats stats = vllm::GetMiniMaxH3ShardStreamStats(); + + CHECK(streamed_params.num_layers == reference.params.num_layers); + CHECK(streamed_params.hidden_size == reference.params.hidden_size); + CHECK(streamed_params.text_dim == reference.params.text_dim); + + // ★ THE LOADER RAN, and it produced DEVICE tensors. A green suite over a path + // that silently fell back to another loader is a failure mode this codebase has + // already shipped once, so the positive signal is asserted, not assumed. + INFO("shard-stream stats: shards=" << stats.shards_opened << " tensors=" + << stats.tensors_streamed << " direct=" << stats.direct_uploads << " converted=" + << stats.converted_uploads << " bytes=" << stats.bytes_uploaded << " host_peak=" + << stats.host_peak_bytes); + CHECK(stats.shards_opened == kShards); + CHECK(stats.tensors_streamed == entries.size() - 1); // rope.inv_freq stays host + CHECK(stats.host_resident == 1); + CHECK(stats.direct_uploads + stats.converted_uploads == stats.tensors_streamed); + CHECK(stats.direct_uploads > 0); // the zero-host-copy path really is taken + CHECK(stats.converted_uploads > 0); // and so is the converting one + CHECK(stats.bytes_uploaded > 0); + + // ★ PEAK HOST MEMORY CANNOT BALLOON. Only ONE tensor's conversion buffer is ever + // alive, so the peak is bounded by the largest single tensor — NOT by the model. + // On the real release that bound is ~1 GB against 66.3 GB of weights, and the + // BF16->bf16 bulk costs nothing at all. + uint64_t largest = 0; + for (const H3StEntry& e : entries) { + largest = std::max(largest, static_cast(e.bytes.size()) * 2); + } + CHECK(stats.host_peak_bytes <= largest); + CHECK(stats.host_peak_bytes < stats.bytes_uploaded / 4); + + // Every streamed view points at an allocation this loader OWNS (a device buffer), + // never into the mmap or into a host reference. + std::set owned; + for (const std::shared_ptr& s : streamed.storage) owned.insert(s.get()); + CHECK(owned.size() == stats.tensors_streamed); + + // ★ STREAMED == NON-STREAMED, tensor for tensor. Both round f32->bf16 with the + // same round-to-nearest-even rule and keep the same fp32 islands, so this is + // BIT-EXACT, not a tolerance. + const auto ref_views = AllDitViews(ref_staged.weights); + const auto got_views = AllDitViews(streamed.weights); + REQUIRE(ref_views.size() == got_views.size()); + size_t compared = 0, island_count = 0; + for (size_t i = 0; i < ref_views.size(); ++i) { + const std::string& name = ref_views[i].first; + const vt::Tensor& a = *ref_views[i].second; + const vt::Tensor& b = *got_views[i].second; + INFO("weight " << name); + REQUIRE(got_views[i].first == name); + REQUIRE(b.data != nullptr); + CHECK(owned.count(b.data) == 1); // it came from THIS loader's device staging + REQUIRE(a.rank == b.rank); + for (int r = 0; r < a.rank; ++r) CHECK(a.shape[r] == b.shape[r]); + // The dtype policy: fp32 ISLANDS stay f32, everything else is bf16. + const bool island = vllm::MiniMaxH3IsFp32IslandTensor(name); + CHECK(b.dtype == (island ? vt::DType::kF32 : vt::DType::kBF16)); + CHECK(a.dtype == b.dtype); + island_count += island ? 1 : 0; + const size_t bytes = static_cast(a.Numel()) * vt::SizeOf(a.dtype); + CHECK(std::memcmp(a.data, b.data, bytes) == 0); + ++compared; + } + CHECK(compared == ref_views.size()); + CHECK(island_count == 12); // both patch projections, the time embedder, both heads + + // ★ rope.inv_freq stays on the HOST. The forward builds the cos/sin cache from it + // BEFORE any kernel runs, so a device pointer here segfaults rather than + // misbehaving. + CHECK(!streamed.rope_inv_freq_host.empty()); + CHECK(streamed.weights.rope_inv_freq.data == streamed.rope_inv_freq_host.data()); + CHECK(streamed.weights.rope_inv_freq.dtype == vt::DType::kF32); + CHECK(owned.count(streamed.weights.rope_inv_freq.data) == 0); + REQUIRE(static_cast(streamed.rope_inv_freq_host.size()) == want.rope_inv_freq_len); + for (int64_t i = 0; i < want.rope_inv_freq_len; ++i) { + CHECK(streamed.rope_inv_freq_host[static_cast(i)] == + ref_staged.weights.rope_inv_freq.Ptr()[i]); + } + + // And it must RUN: the streamed weights drive a real device forward, and the + // non-streamed reference produces the SAME velocity prediction. + const MiniMaxH3PackedSequence packed = + BuildMiniMaxH3PackedSequence(4, 2, 4, 4, 2, 2, /*include_keyframe_cond=*/false, {}, 0); + const int64_t seq = packed.seq_len; + const int64_t num_img = static_cast(packed.img_pos.size()); + const int64_t num_audio = static_cast(packed.audio_pos.size()); + const int64_t num_text = static_cast(packed.text_pos.size()); + const int64_t video_width = want.video_row_width(); + const std::vector x(static_cast(seq * video_width), 0.25f); + const std::vector audio_x(static_cast(seq * want.audio_latents_dim), 0.1f); + const std::vector prompt(static_cast(num_text * want.text_dim), 0.2f); + const std::vector unique_ts = {0.4f}; + const std::vector inverse(static_cast(seq), 0); + const std::vector refiner_cu = {0, static_cast(num_text), + static_cast(num_text)}; + MiniMaxH3DitInputs in; + in.seq_len = seq; + in.x = x.data(); + in.audio_x = audio_x.data(); + in.img_position_ids = packed.img_position_ids.data(); + in.unique_timesteps = unique_ts.data(); + in.num_unique_timesteps = 1; + in.inverse_indices = inverse.data(); + in.token_tags = packed.token_tags.data(); + in.prompt_embeds = prompt.data(); + in.img_pos = packed.img_pos.data(); + in.num_img_pos = num_img; + in.audio_pos = packed.audio_pos.data(); + in.num_audio_pos = num_audio; + in.text_pos = packed.text_pos.data(); + in.num_text_pos = num_text; + in.infer_out_pos = packed.img_pos.data(); + in.num_infer_out_pos = num_img; + in.update_mask = packed.update_mask.data(); + in.cu_seqlens = packed.cu_seqlens.data(); + in.num_cu_seqlens = static_cast(packed.cu_seqlens.size()); + in.refiner_cu_seqlens = refiner_cu.data(); + in.num_refiner_cu_seqlens = static_cast(refiner_cu.size()); + + const MiniMaxH3DitOutputs got = + MiniMaxH3DitForwardDevice(q, streamed_params, streamed.weights, in, vt::DType::kBF16); + const MiniMaxH3DitOutputs ref = + MiniMaxH3DitForwardDevice(q, reference.params, ref_staged.weights, in, vt::DType::kBF16); + REQUIRE(got.video_logits.size() == static_cast(num_img * video_width)); + REQUIRE(got.video_logits.size() == ref.video_logits.size()); + REQUIRE(got.audio_logits.size() == ref.audio_logits.size()); + for (float v : got.video_logits) REQUIRE(std::isfinite(v)); + for (float v : got.audio_logits) REQUIRE(std::isfinite(v)); + const double video_delta = + MaxAbsDiff(got.video_logits, ref.video_logits.data(), got.video_logits.size()); + const double audio_delta = + MaxAbsDiff(got.audio_logits, ref.audio_logits.data(), got.audio_logits.size()); + INFO("streamed-vs-reference forward: video max|diff| = " << video_delta + << ", audio max|diff| = " << audio_delta); + // Bit-identical weights through the same graph: the outputs match EXACTLY. + CHECK(video_delta == 0.0); + CHECK(audio_delta == 0.0); + + RemoveShardedDit(dir, kShards); +} + TEST_CASE("minimax_h3: a 13-shard 66 GB bf16 release derives the SHIPPED geometry") { // The real release's SHAPE at its real size, with the payload as a sparse hole: // the headers are byte-for-byte what the 66.3 GB checkpoint declares, so the - // manifest path — the one `--dump-params ` uses, and the one any device - // loader must derive its geometry from before allocating — is gated on the REAL + // manifest path — the one `--dump-params ` uses, and the one the streamer + // derives its geometry from before allocating anything — is gated on the REAL // names and shapes without downloading or storing a weight byte. MiniMaxH3DitParams shipped; // the defaults ARE the shipped H3 geometry const std::vector specs =