From 016a07df89ff241defa506ff8275ceb4ed24aa04 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Fri, 7 Aug 2026 07:09:55 +0000 Subject: [PATCH] =?UTF-8?q?feat(minimax-h3):=20the=20ORIGINAL=20bf16=20DiT?= =?UTF-8?q?=20release=20is=20indexable=20=E2=80=94=2013=20shards,=2066.3?= =?UTF-8?q?=20GB?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Branch: row/H3-BF16-SHARDED-DIT (helper; CPU-only, no GPU job, no download). Written 2026-08-06 off 075b9f21, REBASED onto f34e0d17 and re-gated before landing; every number below is the POST-REBASE re-run, not the earlier report. 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-covered silhouette into a photoreal close-up (ComfyUI PR 15298 attributes it to the partial split-half RoPE producing channel-wise magnitude outliers that corrupt even INT8). "What does FULL PRECISION look like?" could not be asked, because every DiT loader took a SINGLE file while the bf16 release ships 13 safetensors shards totalling 66.3 GB. This is the CHECKPOINT half. The device streamer the real 66.3 GB load needs is the stacked follow-up row/H3-BF16-SHARDED-STREAM, split out so each PR stays inside the 900-line review cap. - MiniMaxH3ShardedCheckpoint::Open(dir) (new minimax_h3_sharded.cpp) resolves tensors through the checkpoint's own model.safetensors.index.json weight map, never by scanning, with one index over every shard — the shape of the in-tree template LoadMiniMaxH3EncoderWeights(const std::vector&, ...). A tensor the index NAMES but whose shard does not contain it throws BY NAME; skipping it would read as zeros and render. - EnumerateMiniMaxH3ShardedTensors builds the same names+shapes manifest the GGUF and NVFP4 arms build, so ParseMiniMaxH3DitParamsFromGgufManifest derives the geometry from SHAPES ALONE on a sharded checkpoint too. - LoadMiniMaxH3DitFromShards is the host-f32 REFERENCE loader (comparison baseline and the CPU path for reduced checkpoints; ~132 GB on the real release, so not for real runs), and MiniMaxH3IsFp32IslandTensor single-sources the upstream fp32-ISLAND split the three existing streamers each hand-rolled. - examples/minimax_h3_gen accepts --dit everywhere it accepted a DiT file, for both --dump-params and the run path; every existing --dit form unchanged. Gates (CPU, re-run post-rebase: test_minimax_h3 72/72 cases / 54497 assertions; clean Release build of libvllm.a, test_minimax_h3 and minimax-h3-gen): (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 missing-from-shard tensor throws WITH ITS NAME, and the geometry equals the single-file path field for field; (2) real geometry without the weights — a 13-shard release declaring the REAL 535 tensors at REAL shapes with a SPARSE payload (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 repairs two gates that f34e0d17 (#95) left RED on main and that block any PR branched from it: docs/ENVIRONMENT.md now documents VT_H3_ACT_DUMP and VT_H3_ACT_CALL (check-env-doc), and the §9 spec table that ran into the next table without a separator now has one (check-agent-record). Also fixes a real latent defect this row's ASan+UBSan 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, and LoadMiniMaxH3DitFromShards is the first caller to hand it such a file. Now a byte-wise memcpy load. RED-first proven locally on the same sanitizer build: reverting it reproduces CI's `load of misaligned address ... requires 2 byte alignment` at the same line and exits 1; with it, test_minimax_h3 is 73/73 / 55203 under -fsanitize=address,undefined with ZERO findings. The synthetic shard writer stays deliberately UNPADDED so the case remains covered. Honest residuals: NO device load of the real 66.3 GB release ships here, no measured peak RSS, and no bf16-vs-quantized render or speed number. The bf16-vs-quant quality question is UNBLOCKED, not answered. 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 | 58 +++ .agents/state.md | 69 +++ CMakeLists.txt | 1 + docs/BENCHMARKS.md | 2 +- docs/ENVIRONMENT.md | 2 + docs/FEATURES.md | 4 +- docs/STATUS.md | 2 +- examples/minimax_h3_gen/main.cpp | 31 +- .../vllm/model_executor/models/minimax_h3.h | 81 ++++ scripts/check-public-doc-tables.py | 8 +- .../models/minimax_h3_device.cpp | 21 +- .../models/minimax_h3_sharded.cpp | 222 ++++++++++ .../models/minimax_h3_vae_loader.cpp | 26 +- tests/vllm/models/test_minimax_h3.cpp | 419 +++++++++++++++++- 18 files changed, 900 insertions(+), 53 deletions(-) create mode 100644 src/vllm/model_executor/models/minimax_h3_sharded.cpp diff --git a/.agents/NOW.md b/.agents/NOW.md index 342823425..590a3c32d 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 NVFP4 grid DIAGNOSED (#95): NO loader bug** | weights/islands/RoPE all quant-noise-close to coherent GGUF; residual = community-NVFP4 quant fidelity §8.12 | +| 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 | | 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 9ef0052e6..3380c571d 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. **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). 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` | | ✅ | `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 9d9b0f337..70e42f1d5 100644 --- a/.agents/parity-ledger.md +++ b/.agents/parity-ledger.md @@ -916,3 +916,4 @@ Columns: | 2026-08-06 (`row/H3-FP4-SPEED`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; lifecycle unchanged) | **MiniMax-H3 W-FP4a — fp4-RESIDENT NVFP4 routing for the device DiT forward (NO new quant code).** Until now both NVFP4 loaders dequantized the packed FP4 projections to bf16 and the device forward ran `vt::MatmulBT`, so the sm_121a FP4 tensor-core route never ran for H3. Adds `Nvfp4Weight` carriers to `MiniMaxH3DitBlockWeights`/`MiniMaxH3DitWeights`, a fp4-resident streamer `StreamMiniMaxH3Nvfp4ToDeviceFp4` (keeps the compressed-tensors triple host-resident; the shared dispatcher uploads+repacks lazily then frees the fp4 originals, so peak device memory is ~1/4 of the bf16 arm), and a `LinearDev` dispatch that routes a non-Empty fp4 projection through `dense_nvfp4::MatmulNvfp4W4A16D`. | The routing is vLLM's OWN forced-Marlin-for-a16 selection: the checkpoint is weight-only NVFP4 (no `input_activations`, `IsTrueW4A4()==false`), so `kernels/linear/__init__.py:879-881` forces the Marlin W4A16 kernel, mirrored by `include/vllm/model_executor/models/dense_nvfp4_gemm.h:12-22,505-549` (`MatmulNvfp4W4A16D` -> single-expert `vt::MoeGroupedGemmNvfp4Marlin`, the SAME kernel Laguna routed-experts + dense Qwen3-32B NVFP4 use). Not cutlass-fp4/W4A4 (needs fp4 activations, private to `qwen3_5.cpp`). fc1 is already merged `[gate;up]` -> one W4A16 GEMM + `vt::SiluAndMul`. | **CPU-GATED (wiring), verified**: `test_minimax_h3` 62/62 cases / 30039 assertions, 0 failed. The synthetic-NVFP4 case streams the fp4 twin, asserts the loader kept the projections PACKED (fp4 slot set / bf16 slot Empty, and the inverse for the bf16 loader), runs fp4 + bf16 device forwards on the SAME file, asserts the W4A16 dispatcher executed ALL 11 quantized GEMMs (the `Nvfp4W4A16Stats` this-path-ran counter), and bounds fp4-vs-bf16 <= 2e-3. On CPU the dispatcher has no Marlin op so it falls to the bf16 arm's own dequant+matmul (hence a WIRING gate here); the Marlin kernel numerics are CUDA-gated independently by `test_ops_nvfp4_matmul` / `test_linear_method` (2e-3 f32-out / 8e-3 bf16-out vs a bf16 reference). `benchmark_binding=false`. PENDING: GB10 CUDA build + the fp4-vs-bf16 numeric delta and steady per-step timing (disk/build window); real-checkpoint t2va e2e DISK-BLOCKED (~41 GB working set). Comparability: vLLM-Omni serves NO quantized H3 (BF16-only in practice; source-audited `a4ea67a2`, spec §8.3) -> HW/loader-forced-indirect. | | 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. | diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index 301a2c231..42049f3b8 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 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 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). | | 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 9c1372694..fd04c7584 100644 --- a/.agents/specs/minimax-h3.md +++ b/.agents/specs/minimax-h3.md @@ -884,7 +884,65 @@ carries. Before it, no reference modality was reachable over HTTP at all. | Dependencies | Row `SERVE-VIDEOS-OAI` (§9), stacked. Code: `MiniMaxH3Encode{KeyframeCondRows,ReferenceVideo,ReferenceAudio}`, `MiniMaxH3ReadWav`, `DecodeDataUri`. Runtime: `--video-vae` for an image or video reference, `--audio-vae` for an audio reference (both encoder halves, loaded lazily and once). No new download, no GPU. | | Work breakdown | (1) `input_reference` parsing (path or `data:` URL) -> fl2va, with the geometry refusal; (2) the `metadata` map + the video/audio reference keys; (3) the combination rule in the parser; (4) the `examples/server` runner branches; (5) both test files; (6) docs + record. | | Risks/decisions | `input_reference` -> fl2va, NOT ref2va: OpenAI documents it as the frame the video starts from; ref2va would silently change what the API promises. The two extra modalities go in `metadata` rather than new top-level fields, so a strict client's schema validation still passes. Combination legality is enforced in the PARSER, not left to the pipeline, so a supplied reference is never silently dropped. | + | OpenAI | Lands on | Notes | +|---|---|---| | `model` | `VideoRequest::model` | Recorded + echoed; an unserved name is a job `warning`, never a rejection (a Sora client cannot know the local model's name) | | `size` | `width`, `height` | `"x"`, whole positive pixels, one `x`/`X` | | `seconds` | `duration_seconds` | Number OR numeric string — OpenAI types it as a string enum ("4"/"8"/"12") | + +## 8.13 The ORIGINAL bf16 release — the multi-shard CHECKPOINT (2026-08-07, `row/H3-BF16-SHARDED-DIT`, CPU-only) + +**Why.** 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-covered +silhouette into a photoreal close-up (ComfyUI PR 15298 attributes it to the partial +split-half RoPE producing channel-wise magnitude outliers that corrupt even INT8). +"What does FULL PRECISION look like?" was unanswerable because every DiT loader took +a SINGLE file (`LoadMiniMaxH3DitFromGguf`/`...Bf16`/`StreamMiniMaxH3DitToDeviceBf16` +one GGUF; `LoadMiniMaxH3DitFromNvfp4`/`StreamMiniMaxH3Nvfp4To*` one safetensors), +while the bf16 release ships **13 safetensors shards totalling 66.3 GB**. + +This section is the CHECKPOINT half. The device streamer that makes the real 66.3 GB +release loadable on a GPU is §8.14, split out as a stacked row so each PR stays inside +the 900-line review cap. + +**What landed.** +- `MiniMaxH3ShardedCheckpoint::Open(dir)` (`minimax_h3_sharded.cpp`) resolves shards + through the checkpoint's own `model.safetensors.index.json` weight map (the + `diffusion_pytorch_model.*` spelling is accepted too). Nothing is discovered by + scanning. Shape mirrors the in-tree multi-shard template + `LoadMiniMaxH3EncoderWeights(const std::vector&, ...)`: one index + over every shard. A tensor the index NAMES but whose shard does not contain it + throws BY NAME (skipping it would read as zeros and render). +- `EnumerateMiniMaxH3ShardedTensors` produces the same names+shapes manifest the GGUF + and NVFP4 arms build, so `ParseMiniMaxH3DitParamsFromGgufManifest` derives the + geometry from SHAPES ALONE here too, and a sharded checkpoint and a single-file one + holding the same tensors produce IDENTICAL params. +- `LoadMiniMaxH3DitFromShards` is the host-f32 REFERENCE loader (the comparison + baseline and the CPU path for reduced checkpoints; ~132 GB on the real release, so + not for real runs). +- The fp32 ISLAND split is single-sourced as `MiniMaxH3IsFp32IslandTensor` and the + three pre-existing streamers now call it. It is load-bearing: `vt::MatmulBT` rejects + an (f32 activation, bf16 weight) pair, so a tensor on the wrong side fails at the + first island GEMM. +- `examples/minimax_h3_gen`: `--dit ` accepts a shard directory everywhere a DiT + file was accepted, for both `--dump-params` and the run path. Every existing `--dit` + form is unchanged. + +**Gates (CPU, re-run after the rebase onto `f34e0d17`: `test_minimax_h3` 72/72 cases, +54497 assertions).** +1. *Index + name mapping, no weights*: a synthetic 4-shard set; every tensor resolves + to the shard the index named AND to the bytes written there; a tensor named in the + index but missing from its shard THROWS with the tensor name in the message; the + derived geometry equals the SINGLE-FILE path over the same tensors, field for field. +2. *Real geometry without the weights*: a 13-shard release whose headers declare the + REAL 535 tensors at their REAL shapes with the payload as a SPARSE hole (61.73 GiB + = 66.3 GB declared, 144 KB on disk) parses to the SHIPPED geometry: num_layers 50, + hidden 5376, heads 56, head_dim 128, ffn 14336, latents 24, audio_latents 32, patch + 1x2x2, text_dim 5120 - the same numbers the working GGUF arm derives from shapes + alone. `minimax-h3-gen --dit --dump-params` prints all 20 fields on it. + +**Not claimed here.** No device load of the real 66.3 GB release (no streaming device +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. diff --git a/.agents/state.md b/.agents/state.md index 81a59e1f9..2353b467b 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -40896,3 +40896,72 @@ Records: spec §14, STATUS/BENCHMARKS/FEATURES Kimi rows, benchmark-record, NOW. fails to create. `sudo` on the host is the working path. No container was holding the GPU (all exited), so nothing was stopped and nothing needed restoring. + +## 2026-08-07T09:55 - H3: the ORIGINAL bf16 release (13 shards, 66.3 GB) is INDEXABLE - multi-shard DiT checkpoint + host reference loader (row/H3-BF16-SHARDED-DIT, CPU-only) + + +`row/H3-BF16-SHARDED-DIT` (helper; CPU-only, no GPU job, no download - the box's GPU +was busy with an attention A/B). Written 2026-08-06 off `075b9f21`, REBASED onto +`f34e0d17` and re-gated on 2026-08-07 before landing; the gate numbers below are the +post-rebase ones, re-run here, not the pre-rebase report. + +**Why.** 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-covered +silhouette into a photoreal close-up (ComfyUI PR 15298 attributes it to the partial +split-half RoPE creating channel-wise magnitude outliers that corrupt even INT8). The +obvious next question - what does FULL PRECISION look like? - could not be asked, +because every DiT loader took a SINGLE file while the bf16 release ships 13 +safetensors shards totalling 66.3 GB. + +**What landed (this row is the CHECKPOINT half; the device streamer is the stacked +follow-up `row/H3-BF16-SHARDED-STREAM`, split out to stay inside the PR-size cap).** +- `MiniMaxH3ShardedCheckpoint::Open(dir)` in the new + `src/vllm/model_executor/models/minimax_h3_sharded.cpp`: resolves every tensor + through the checkpoint's own `model.safetensors.index.json` weight map (the + diffusers `diffusion_pytorch_model.*` spelling is accepted too), one index over all + shards, mirroring the in-tree template `LoadMiniMaxH3EncoderWeights(const + std::vector&, ...)`. A tensor the index NAMES but whose shard does + not contain it throws BY NAME - skipping it would read as zeros and render. +- `EnumerateMiniMaxH3ShardedTensors` builds the same names+shapes manifest the GGUF + and NVFP4 arms build, so `ParseMiniMaxH3DitParamsFromGgufManifest` derives the + geometry from SHAPES ALONE on a sharded checkpoint too. +- `LoadMiniMaxH3DitFromShards` (host-f32 REFERENCE loader - the comparison baseline + and the CPU path for reduced checkpoints; it needs ~132 GB on the real release and + must not be used there), and `MiniMaxH3IsFp32IslandTensor` single-sourcing the + fp32-ISLAND split that the three existing streamers each hand-rolled. +- `--dit ` in `examples/minimax_h3_gen` wherever a single DiT file was accepted, + for both `--dump-params` and the run path; every existing `--dit` form unchanged. + +**Gates (CPU, re-run post-rebase: `test_minimax_h3` 72/72 cases / 54497 assertions, +clean Release build of `libvllm.a`, `test_minimax_h3`, `minimax-h3-gen`).** +(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 named in the index but +missing from its shard throws with its NAME in the message; the derived geometry +equals the SINGLE-FILE path over the same tensors, field for field. (2) real geometry +without the weights - a 13-shard release whose headers declare the REAL 535 tensors at +their REAL shapes with the payload as a SPARSE hole (61.73 GiB = 66.3 GB declared, +144 KB on disk) parses to the SHIPPED geometry (num_layers 50, hidden 5376, heads 56, +head_dim 128, ffn 14336, latents 24, audio_latents 32, patch 1x2x2, text_dim 5120), +and `minimax-h3-gen --dit --dump-params` prints all 20 fields on it. + +**A real defect this row's CI found, fixed here.** `MiniMaxH3ReadSafetensorF32` +(`minimax_h3_vae_loader.cpp:88`) read 16-bit payloads as +`reinterpret_cast(tensor.data)[i]`. safetensors puts the payload +straight after a JSON header of ARBITRARY length, so a tensor's first byte is only +2-byte aligned if the writer happened to pad, and the format does not require it - the +cast is UB on any file with an odd header. It had never fired because every checkpoint +reaching that function so far was padded; `LoadMiniMaxH3DitFromShards` made an unpadded +one reachable and the ASan+UBSan lane caught it immediately (`load of misaligned address +... requires 2 byte alignment`). Now a byte-wise `memcpy` load, which has no alignment +precondition. RED-first proven locally: reverting the fix on the same sanitizer build +reproduces CI's message at the same line and exits 1; with it, `test_minimax_h3` is +73/73 / 55203 under `-fsanitize=address,undefined` with ZERO findings. The synthetic +shard writer is deliberately left UNPADDED so this stays covered. + +**Honest residuals.** NOT claimed here: any device load of the real 66.3 GB release +(this row ships no streaming device loader - see the stacked follow-up), its measured +peak RSS, and any bf16-vs-quantized RENDER or SPEED number. The bf16-vs-quant quality +question is UNBLOCKED, not answered. + +Next: `row/H3-BF16-SHARDED-STREAM` (device streamer), then the operator runs the real +13-shard bf16 DiT. diff --git a/CMakeLists.txt b/CMakeLists.txt index 30adf4a6e..aa5b94204 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -567,6 +567,7 @@ add_library(vllm STATIC src/vllm/model_executor/models/minimax_h3_planner.cpp src/vllm/model_executor/models/minimax_h3_gguf.cpp src/vllm/model_executor/models/minimax_h3_nvfp4.cpp + src/vllm/model_executor/models/minimax_h3_sharded.cpp src/vllm/model_executor/models/minimax_h3_audio_vae.cpp src/vllm/model_executor/models/minimax_h3_video_vae.cpp src/vllm/model_executor/models/minimax_h3_video_vae_device.cpp diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 487c8594e..59cb9b08d 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. 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 is now indexable) 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/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 6a938ebc6..74c79bd9d 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -110,6 +110,8 @@ Read-only observability; none change output. | `VT_H3_VAE_PROBE` | unset | `=1` runs a video-VAE receptive-field probe after the normal decode: it perturbs ONE interior spatial latent cell (across all channels and temporal frames), re-decodes, and prints a per-16px-block RMS-change map (`[h3-vae-probe]`) over output frame 0. If only the perturbed cell's block moves, the ViT3D decoder is not mixing tokens spatially. Byte-identical to production when unset (no second decode) | | `VT_H3_DUMP_DIR` | unset | Directory into which the MiniMax-H3 denoise loop writes the initial and final video latent rows (`init_video_rows.f32`, `final_video_rows.f32`) and the pipeline writes the exact VAE-input latent (`vae_input_video_latent.f32`), all raw little-endian f32. Lets two runs (e.g. 12 vs 50 steps, conditioned vs not) be byte/stat-compared, and the video VAE decode be replayed on a KNOWN latent, without re-running the denoise. Byte-identical to production when unset (no file is opened) | | `VT_H3_DUMP_INPUTS` | unset | Directory into which the MiniMax-H3 denoise loop writes EVERY DiT input at step 0 as raw little-endian binary plus a `manifest.txt` — the packed layout (`input_ids`/`image_mask`/`audio_mask`/`img_pos`/`audio_pos`/`text_pos`/`update_mask`/`cu_seqlens`/`document_id`), the fp64 position grid (`img_position_ids.f64`), the per-token modality tags (`token_tags.i64`), the per-token pre-unique timesteps and their `unique_timesteps`/`inverse_indices`/`combined_indices` AdaLN selection, both sigma schedules, and the raw `prompt_embeds`; the `minimax-h3-gen` driver additionally writes `prompt_token_ids.i32`. Lets the REAL-scale DiT inputs be diffed EXACTLY against upstream `pipeline_minimax_h3.py` (the render-coherence S1 surface the reduced-dim ladder never fed real values into). Byte-identical to production when unset (no file is opened) | +| `VT_H3_ACT_DUMP` | unset | Path to which ONE MiniMax-H3 device DiT forward writes a per-STAGE activation FINGERPRINT: mean/rms/absmax/finite plus a fixed set of positional sample values, for every stage and for the input-independent weight classes (islands, biases, output heads). Two weight arms running the SAME graph on the SAME inputs (e.g. an NVFP4-bf16 stream vs a GGUF-bf16 control) can then be diffed layer by layer: a JUMP at a stage names the guilty tensor class, and a scramble/transpose is caught by the positional samples even when rms matches. This is the instrument that REFUTED a #94 load-path defect (spec §8.12). Byte-identical to production when unset (no file is opened) | +| `VT_H3_ACT_CALL` | 0 | Which forward `VT_H3_ACT_DUMP` captures, as a 0-based call index within the process. Only that one forward dumps, so a single small render (`--denoise-only --steps 1`) yields exactly one clean fingerprint file instead of 50 overwrites. Read only when `VT_H3_ACT_DUMP` is set | ## Kernel-internal knobs (deferred) diff --git a/docs/FEATURES.md b/docs/FEATURES.md index c188e198e..6d93848e9 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 69/69; t2va+fl2va COHERENT; ref2va NVFP4 grid DIAGNOSED (§8.12): NO loader bug (weights/islands/RoPE all quant-noise-close to the coherent GGUF); residual = community-NVFP4 quant fidelity, not a loader fix | FP4/Marlin landed; ref2va NVFP4 render blocked on checkpoint quant (needs official modelopt NVFP4), speed 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 | | 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 DIAGNOSED (§8.12): NO loader bug (weights/islands/RoPE all quant-noise-close to the coherent GGUF); residual = community-NVFP4 quant fidelity | ✅ (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 indexes the bf16 13-shard release | ✅ (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 79c209d6c..b15fc4e11 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 +// minimax-h3-gen --dit +// # a DIRECTORY holding the original bf16 release's shards plus +// # model.safetensors.index.json is accepted wherever a single +// # DiT file is; every existing --dit form is unchanged. // --video-vae --video-vae-config // --audio-vae --audio-vae-config // --prompt-embeds (rows of text_dim, little-endian f32) @@ -276,7 +279,7 @@ int main(int argc, char** argv) { (dit_path.empty() || (need_vaes && (video_vae_path.empty() || audio_vae_path.empty())) || (need_vaes && out_path.empty()) || (need_cond && embeds_path.empty() && (encoder_path.empty() || prompt.empty())))) { - std::cerr << "usage: minimax-h3-gen --dit --video-vae --audio-vae " + std::cerr << "usage: minimax-h3-gen --dit --video-vae --audio-vae " "--prompt-embeds --out [--video-vae-config ] " "[--audio-vae-config ] [--keep-quant] [--steps N] [--frames N] " "[--height N] [--width N] [--device cpu|cuda] [--workdir DIR] [--ffmpeg PATH] " @@ -297,6 +300,16 @@ int main(int argc, char** argv) { if (EndsWith(dit_path, ".gguf")) { const vllm::GgufFile gf = vllm::GgufFile::Open(dit_path); pr = vllm::ParseMiniMaxH3DitParamsFromGgufManifest(vllm::EnumerateMiniMaxH3GgufTensors(gf)); + } else if (vllm::MiniMaxH3ShardedCheckpoint::IsShardedDir(dit_path)) { + // A DIRECTORY of shards + index: the original bf16 release. Manifest only, + // so this answers "do the 13 shards agree on the geometry the GGUF arm + // derives?" on a 66.3 GB checkpoint without reading a single weight byte. + const vllm::MiniMaxH3ShardedCheckpoint ckpt = + vllm::MiniMaxH3ShardedCheckpoint::Open(dit_path); + std::cerr << " " << ckpt.ShardCount() << " shard(s), " << ckpt.Names().size() + << " tensors, index " << ckpt.IndexPath() << "\n"; + pr = vllm::ParseMiniMaxH3DitParamsFromGgufManifest( + vllm::EnumerateMiniMaxH3ShardedTensors(ckpt)); } else { const vllm::SafetensorsFile sf = vllm::SafetensorsFile::Open(dit_path); std::vector manifest; @@ -546,6 +559,20 @@ int main(int argc, char** argv) { dit = dequant_bf16 ? vllm::LoadMiniMaxH3DitFromGgufBf16(f) : vllm::LoadMiniMaxH3DitFromGguf(f, keep_quant); } + } else if (vllm::MiniMaxH3ShardedCheckpoint::IsShardedDir(dit_path)) { + // The ORIGINAL bf16 release: a DIRECTORY of shards plus + // model.safetensors.index.json. Every other --dit form is a single file and + // keeps working unchanged; this is the only one that can open the + // full-precision DiT at all, which is what makes the quantization-quality + // question askable. + const vllm::MiniMaxH3ShardedCheckpoint ckpt = + vllm::MiniMaxH3ShardedCheckpoint::Open(dit_path); + std::cerr << " " << ckpt.ShardCount() << " shard(s), " << ckpt.Names().size() + << " tensors (index " << ckpt.IndexPath() << ")\n"; + // Host f32 reference path. ~132 GB on the real release, so it is usable + // only on a reduced checkpoint; opening the real release on a device needs + // a streaming loader, which this row does not yet ship. + 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 ae5686ff5..20d22d7b7 100644 --- a/include/vllm/model_executor/models/minimax_h3.h +++ b/include/vllm/model_executor/models/minimax_h3.h @@ -1424,6 +1424,87 @@ MiniMaxH3DitDeviceWeights StreamMiniMaxH3Nvfp4ToDeviceFp4(vt::Queue& queue, const SafetensorsFile& file, MiniMaxH3DitParams* out_params = nullptr); +// --------------------------------------------------------------------------- +// The ORIGINAL bf16 release: 13 safetensors shards, 66.3 GB +// (minimax_h3_sharded.cpp + the streamer in minimax_h3_device.cpp) +// --------------------------------------------------------------------------- +// +// 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 frame; ComfyUI PR 15298 traces it to the partial split-half RoPE +// producing channel-wise magnitude outliers that corrupt even INT8). Answering +// "what does FULL PRECISION look like?" needs the original release, and every +// DiT loader before this one took a SINGLE file. + +// The upstream `MINIMAX_H3_FP32_PARAM_NAMES` / `_BUFFER_NAMES` split +// (minimax_h3_transformer.py:85-101): both patch projections, both time-embedder +// projections, both output heads and `rope.inv_freq` stay FP32 even in a bf16 +// stream. This is LOAD-BEARING, not a precision nicety — their activations are +// f32 and `vt::MatmulBT` REJECTS a mixed (f32 activation, bf16 weight) pair, so a +// tensor on the wrong side of this line fails loudly at the first island GEMM. +// Single-sourced here because all four staging paths must agree on it. +bool MiniMaxH3IsFp32IslandTensor(const std::string& name); + +// A MULTI-SHARD safetensors checkpoint, resolved through its +// `model.safetensors.index.json` weight map. The index is USED, never guessed +// around: a tensor the index names but whose shard does not contain it throws BY +// NAME instead of being silently skipped (a skipped weight reads as zeros later, +// which is a plausible-looking render rather than an error). +// +// Shape mirrors `LoadMiniMaxH3EncoderWeights(const std::vector&, +// ...)` — one index over every shard, so a tensor is found wherever it lives. +// Every shard is mmap'd read-only; nothing is materialized at Open() time, which +// is why this is safe on a checkpoint far larger than RAM. +class MiniMaxH3ShardedCheckpoint { + public: + // `dir` holds the shards and their index. `model.safetensors.index.json` is the + // name the H3 release ships; `diffusion_pytorch_model.safetensors.index.json` + // (the diffusers spelling) is accepted as well. Throws naming `dir` when + // neither exists. + static MiniMaxH3ShardedCheckpoint Open(const std::string& dir); + // Whether `path` is a directory holding one of those indexes — the test a + // caller with a single `--dit` flag uses to tell a directory from a file. + static bool IsShardedDir(const std::string& path); + + MiniMaxH3ShardedCheckpoint(); + ~MiniMaxH3ShardedCheckpoint(); + MiniMaxH3ShardedCheckpoint(MiniMaxH3ShardedCheckpoint&&) noexcept; + MiniMaxH3ShardedCheckpoint& operator=(MiniMaxH3ShardedCheckpoint&&) noexcept; + MiniMaxH3ShardedCheckpoint(const MiniMaxH3ShardedCheckpoint&) = delete; + MiniMaxH3ShardedCheckpoint& operator=(const MiniMaxH3ShardedCheckpoint&) = delete; + + // Every tensor the index names, in index order. + const std::vector& Names() const; + bool Has(const std::string& name) const; + // Throws BY NAME when the index does not name `name`. + const StTensor& Get(const std::string& name) const; + // The shard FILENAME `name` was resolved to — the gateable answer to "did this + // tensor come out of the right shard?". + const std::string& ShardOf(const std::string& name) const; + const std::vector& ShardFiles() const; + const std::string& IndexPath() const; + size_t ShardCount() const; + + private: + struct Impl; + std::unique_ptr impl_; +}; + +// Names + shapes (+ the fp32-island flag) over every shard, the same manifest the +// GGUF and NVFP4 arms build — so `ParseMiniMaxH3DitParamsFromGgufManifest` derives +// the geometry from SHAPES ALONE here too, and a sharded checkpoint and a +// single-file one that hold the same tensors produce IDENTICAL params. +std::vector EnumerateMiniMaxH3ShardedTensors( + const MiniMaxH3ShardedCheckpoint& ckpt); + +// 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. +MiniMaxH3GgufDit LoadMiniMaxH3DitFromShards(const MiniMaxH3ShardedCheckpoint& ckpt); + MiniMaxH3DitDeviceWeights StageMiniMaxH3DitWeights(vt::Queue& queue, const MiniMaxH3DitParams& params, const MiniMaxH3DitWeights& host, diff --git a/scripts/check-public-doc-tables.py b/scripts/check-public-doc-tables.py index 1dba5782c..ef7087779 100755 --- a/scripts/check-public-doc-tables.py +++ b/scripts/check-public-doc-tables.py @@ -331,7 +331,13 @@ def features_errors(text: str) -> list[str]: # exactly on it: a ratchet pinned to the byte turns every concurrently # merged row's one-line status edit into a spurious failure. Still strictly # DOWN from 284062, the only direction this number may move. - "chars": 283470, + # + # 283455 since 2026-08-07 (measured 283433): the MiniMax-H3 row had to carry + # a new claim (the ORIGINAL bf16 13-shard DiT release is now indexable), and + # it was paid for inside the same cell rather than out of the page - the + # ref2va activation-diff narrative collapsed to its binding result, with the + # full guilty-class audit kept in .agents/specs/minimax-h3.md 8.12. Net -6. + "chars": 283455, "h2_sections": 11, "long_paragraphs": 89, "oversized_cells": 47, diff --git a/src/vllm/model_executor/models/minimax_h3_device.cpp b/src/vllm/model_executor/models/minimax_h3_device.cpp index 7634429c4..d4eeb832e 100644 --- a/src/vllm/model_executor/models/minimax_h3_device.cpp +++ b/src/vllm/model_executor/models/minimax_h3_device.cpp @@ -1097,12 +1097,9 @@ MiniMaxH3DitDeviceWeights StreamMiniMaxH3DitToDeviceBf16(vt::Queue& queue, const // projections, both time-embedder projections and both output heads stay f32 even // in a bf16 stream. Their ACTIVATIONS are f32 too, and vt::MatmulBT rejects a // mixed (f32 act, bf16 weight) pair — so getting this split wrong is not a - // precision nuance, it fails loudly at the first island GEMM. - auto is_fp32_island = [](const std::string& n) { - return n.rfind("video_patch_proj.", 0) == 0 || n.rfind("audio_patch_proj.", 0) == 0 || - n.rfind("time_embedder.", 0) == 0 || n.rfind("final_layer.video_out.", 0) == 0 || - n.rfind("final_layer.audio_out.", 0) == 0 || n == "rope.inv_freq"; - }; + // precision nuance, it fails loudly at the first island GEMM. Single-sourced in + // MiniMaxH3IsFp32IslandTensor so every staging path agrees on it. + auto is_fp32_island = [](const std::string& n) { return MiniMaxH3IsFp32IslandTensor(n); }; int64_t done = 0; for (const MiniMaxH3TensorSpec& spec : manifest) { const GgufTensorInfo& info = file.Get(spec.name); @@ -1187,11 +1184,7 @@ MiniMaxH3DitDeviceWeights StreamMiniMaxH3Nvfp4ToDeviceBf16(vt::Queue& queue, // Same fp32 ISLANDS as the GGUF stream: vt::MatmulBT rejects a mixed // (f32 activation, bf16 weight) pair, so this split is load-bearing, not a // precision nicety. - auto is_fp32_island = [](const std::string& n) { - return n.rfind("video_patch_proj.", 0) == 0 || n.rfind("audio_patch_proj.", 0) == 0 || - n.rfind("time_embedder.", 0) == 0 || n.rfind("final_layer.video_out.", 0) == 0 || - n.rfind("final_layer.audio_out.", 0) == 0 || n == "rope.inv_freq"; - }; + auto is_fp32_island = [](const std::string& n) { return MiniMaxH3IsFp32IslandTensor(n); }; // Pass 1: LOGICAL shapes only (no payload), so geometry is known before any // allocation. A packed [out, in/2] U8 weight is logically [out, in]. @@ -1327,11 +1320,7 @@ MiniMaxH3DitDeviceWeights StreamMiniMaxH3Nvfp4ToDeviceFp4(vt::Queue& queue, return (n.size() > 12 && n.compare(n.size() - 12, 12, "weight_scale") == 0) || (n.size() > 14 && n.compare(n.size() - 14, 14, "weight_scale_2") == 0); }; - auto is_fp32_island = [](const std::string& n) { - return n.rfind("video_patch_proj.", 0) == 0 || n.rfind("audio_patch_proj.", 0) == 0 || - n.rfind("time_embedder.", 0) == 0 || n.rfind("final_layer.video_out.", 0) == 0 || - n.rfind("final_layer.audio_out.", 0) == 0 || n == "rope.inv_freq"; - }; + auto is_fp32_island = [](const std::string& n) { return MiniMaxH3IsFp32IslandTensor(n); }; // Pass 1: logical shapes (U8 packed [out, in/2] is logically [out, in]). std::vector manifest; diff --git a/src/vllm/model_executor/models/minimax_h3_sharded.cpp b/src/vllm/model_executor/models/minimax_h3_sharded.cpp new file mode 100644 index 000000000..b1eb46f67 --- /dev/null +++ b/src/vllm/model_executor/models/minimax_h3_sharded.cpp @@ -0,0 +1,222 @@ +// MiniMax-H3 — the ORIGINAL bf16 DiT release: 13 safetensors shards, 66.3 GB. +// +// Every DiT loader before this one took a SINGLE file (one GGUF, or one NVFP4 +// safetensors), which made the full-precision checkpoint the one thing we could +// not load — and therefore made "is quantization costing us render quality?" +// unanswerable. H3 is unusually quantization-sensitive (Q3_K_M -> Q4_K_M alone +// turned a murky lattice-covered silhouette into a photoreal close-up; ComfyUI PR +// 15298 traces it to the partial split-half RoPE producing channel-wise magnitude +// outliers that corrupt even INT8), so the comparison is worth the loader. +// +// This file owns the SHARD RESOLUTION half: the checkpoint's own +// `model.safetensors.index.json` weight map is used, never guessed around, and a +// tensor the index names but whose shard does not contain it throws BY NAME +// rather than being skipped (a skipped weight reads as zeros later, which is a +// plausible-looking render rather than an error). It mirrors the in-tree +// multi-shard template `LoadMiniMaxH3EncoderWeights(const +// std::vector&, ...)` in minimax_h3_vae_loader.cpp: one index +// over every shard, so a tensor is found wherever it lives. +// +// The STREAMING stager that a real 66.3 GB run uses lives next to its GGUF and +// NVFP4 twins in minimax_h3_device.cpp, because it shares their view binder. +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "vllm/model_executor/model_loader/safetensors_reader.h" +#include "vllm/model_executor/models/minimax_h3.h" +#include "vt/dtype.h" // VT_CHECK + +namespace vllm { +namespace { + +// The two spellings a multi-shard release uses. The H3 DiT ships the first; the +// second is the diffusers convention for a pipeline sub-model. Nothing is +// discovered by SCANNING the directory — the index file is what maps a tensor to +// a shard, and guessing by filename is exactly the mistake that silently loads a +// stale or partial shard set. +const char* const kIndexNames[] = { + "model.safetensors.index.json", + "diffusion_pytorch_model.safetensors.index.json", +}; + +bool IsDir(const std::string& path) { + struct stat st {}; + return ::stat(path.c_str(), &st) == 0 && S_ISDIR(st.st_mode); +} + +bool IsFile(const std::string& path) { + struct stat st {}; + return ::stat(path.c_str(), &st) == 0 && S_ISREG(st.st_mode); +} + +std::string StripTrailingSlash(const std::string& dir) { + std::string out = dir; + while (out.size() > 1 && out.back() == '/') out.pop_back(); + return out; +} + +// The index path inside `dir`, or "" when the directory holds neither spelling. +std::string FindIndexPath(const std::string& dir) { + for (const char* name : kIndexNames) { + const std::string candidate = StripTrailingSlash(dir) + "/" + name; + if (IsFile(candidate)) return candidate; + } + return std::string(); +} + +} // namespace + +// The fp32 ISLAND split, single-sourced (minimax_h3_transformer.py:85-101). All +// four staging paths (GGUF stream, NVFP4 bf16 stream, NVFP4 fp4 stream, and the +// sharded bf16 stream) must agree on it: `vt::MatmulBT` rejects a mixed (f32 +// activation, bf16 weight) pair, so a tensor on the wrong side fails loudly at +// the first island GEMM rather than drifting numerically. +bool MiniMaxH3IsFp32IslandTensor(const std::string& n) { + return n.rfind("video_patch_proj.", 0) == 0 || n.rfind("audio_patch_proj.", 0) == 0 || + n.rfind("time_embedder.", 0) == 0 || n.rfind("final_layer.video_out.", 0) == 0 || + n.rfind("final_layer.audio_out.", 0) == 0 || n == "rope.inv_freq"; +} + +struct MiniMaxH3ShardedCheckpoint::Impl { + std::string dir; + std::string index_path; + std::vector shard_files; // index-first-seen order, deduplicated + std::vector shards; // parallel to shard_files + std::vector names; // every tensor the index names + std::map shard_of; // name -> slot in shard_files + std::map tensors; // name -> its entry in that shard +}; + +MiniMaxH3ShardedCheckpoint::MiniMaxH3ShardedCheckpoint() : impl_(std::make_unique()) {} +MiniMaxH3ShardedCheckpoint::~MiniMaxH3ShardedCheckpoint() = default; +MiniMaxH3ShardedCheckpoint::MiniMaxH3ShardedCheckpoint(MiniMaxH3ShardedCheckpoint&&) noexcept = + default; +MiniMaxH3ShardedCheckpoint& MiniMaxH3ShardedCheckpoint::operator=( + MiniMaxH3ShardedCheckpoint&&) noexcept = default; + +bool MiniMaxH3ShardedCheckpoint::IsShardedDir(const std::string& path) { + return IsDir(path) && !FindIndexPath(path).empty(); +} + +MiniMaxH3ShardedCheckpoint MiniMaxH3ShardedCheckpoint::Open(const std::string& dir) { + VT_CHECK(IsDir(dir), "minimax_h3 sharded: '" + dir + "' is not a directory"); + const std::string index_path = FindIndexPath(dir); + VT_CHECK(!index_path.empty(), "minimax_h3 sharded: '" + StripTrailingSlash(dir) + + "' holds no model.safetensors.index.json"); + + MiniMaxH3ShardedCheckpoint out; + Impl& impl = *out.impl_; + impl.dir = StripTrailingSlash(dir); + impl.index_path = index_path; + + // The index IS the map. LoadSafetensorsIndex already rejects a shard value that + // is not a plain filename, so a hostile index cannot escape the directory. + const std::map weight_map = LoadSafetensorsIndex(index_path); + VT_CHECK(!weight_map.empty(), + "minimax_h3 sharded: '" + index_path + "' has an empty weight_map"); + + std::map slot_of_file; + for (const auto& entry : weight_map) { + if (slot_of_file.count(entry.second) != 0) continue; + slot_of_file.emplace(entry.second, impl.shard_files.size()); + impl.shard_files.push_back(entry.second); + } + + // Open every shard ONCE, before any tensor pointer is taken: the StTensor + // addresses below point into these objects, so the vector must be final first. + impl.shards.reserve(impl.shard_files.size()); + for (const std::string& file : impl.shard_files) { + impl.shards.push_back(SafetensorsFile::Open(impl.dir + "/" + file)); + } + + // One index over every shard. A name the index promises but whose shard does + // not contain it is a HARD error, reported with the tensor AND the shard. + std::vector> present; + present.reserve(impl.shards.size()); + for (const SafetensorsFile& shard : impl.shards) { + present.emplace_back(shard.Names().begin(), shard.Names().end()); + } + impl.names.reserve(weight_map.size()); + for (const auto& entry : weight_map) { + const std::string& name = entry.first; + const size_t slot = slot_of_file.at(entry.second); + VT_CHECK(present[slot].count(name) != 0, + "minimax_h3 sharded: the index names tensor '" + name + "' in shard '" + + entry.second + "', but that shard does not contain it"); + impl.names.push_back(name); + impl.shard_of.emplace(name, slot); + impl.tensors.emplace(name, &impl.shards[slot].Get(name)); + } + return out; +} + +const std::vector& MiniMaxH3ShardedCheckpoint::Names() const { return impl_->names; } + +bool MiniMaxH3ShardedCheckpoint::Has(const std::string& name) const { + return impl_->tensors.count(name) != 0; +} + +const StTensor& MiniMaxH3ShardedCheckpoint::Get(const std::string& name) const { + const auto it = impl_->tensors.find(name); + VT_CHECK(it != impl_->tensors.end(), + "minimax_h3 sharded: no tensor named '" + name + "' in " + impl_->index_path); + return *it->second; +} + +const std::string& MiniMaxH3ShardedCheckpoint::ShardOf(const std::string& name) const { + const auto it = impl_->shard_of.find(name); + VT_CHECK(it != impl_->shard_of.end(), + "minimax_h3 sharded: no tensor named '" + name + "' in " + impl_->index_path); + return impl_->shard_files[it->second]; +} + +const std::vector& MiniMaxH3ShardedCheckpoint::ShardFiles() const { + return impl_->shard_files; +} + +const std::string& MiniMaxH3ShardedCheckpoint::IndexPath() const { return impl_->index_path; } + +size_t MiniMaxH3ShardedCheckpoint::ShardCount() const { return impl_->shards.size(); } + +std::vector EnumerateMiniMaxH3ShardedTensors( + const MiniMaxH3ShardedCheckpoint& ckpt) { + std::vector out; + out.reserve(ckpt.Names().size()); + for (const std::string& name : ckpt.Names()) { + const StTensor& t = ckpt.Get(name); + // The bf16 release stores plain tensors; a packed quantized weight here would + // mean the caller pointed a bf16 loader at a quantized checkpoint, and its + // logical shape would be half its stored one — so refuse rather than derive + // a silently halved geometry. + VT_CHECK(t.dtype == "F32" || t.dtype == "BF16" || t.dtype == "F16", + "minimax_h3 sharded: tensor '" + name + "' has unsupported dtype '" + t.dtype + + "' (expected F32/BF16/F16)"); + MiniMaxH3TensorSpec spec; + spec.name = name; + spec.shape = t.shape; + spec.fp32 = MiniMaxH3IsFp32IslandTensor(name); + out.push_back(std::move(spec)); + } + return out; +} + +MiniMaxH3GgufDit LoadMiniMaxH3DitFromShards(const MiniMaxH3ShardedCheckpoint& ckpt) { + MiniMaxH3GgufDit out; + const std::vector manifest = EnumerateMiniMaxH3ShardedTensors(ckpt); + out.params = ParseMiniMaxH3DitParamsFromGgufManifest(manifest); + for (const MiniMaxH3TensorSpec& spec : manifest) { + out.storage[spec.name] = MiniMaxH3ReadSafetensorF32(ckpt.Get(spec.name)); + out.shapes[spec.name] = spec.shape; + } + BindMiniMaxH3DitViews(&out); + return out; +} + +} // namespace vllm diff --git a/src/vllm/model_executor/models/minimax_h3_vae_loader.cpp b/src/vllm/model_executor/models/minimax_h3_vae_loader.cpp index 2f66f0cb5..42e3e7a16 100644 --- a/src/vllm/model_executor/models/minimax_h3_vae_loader.cpp +++ b/src/vllm/model_executor/models/minimax_h3_vae_loader.cpp @@ -81,16 +81,24 @@ std::vector MiniMaxH3ReadSafetensorF32(const StTensor& tensor) { VT_CHECK(tensor.nbytes == static_cast(numel) * 4, "minimax_h3: F32 tensor span does not match its shape"); std::memcpy(out.data(), tensor.data, tensor.nbytes); - } else if (tensor.dtype == "BF16") { + } else if (tensor.dtype == "BF16" || tensor.dtype == "F16") { VT_CHECK(tensor.nbytes == static_cast(numel) * 2, - "minimax_h3: BF16 tensor span does not match its shape"); - const uint16_t* src = reinterpret_cast(tensor.data); - for (int64_t i = 0; i < numel; ++i) out[static_cast(i)] = Bf16ToF32(src[i]); - } else if (tensor.dtype == "F16") { - VT_CHECK(tensor.nbytes == static_cast(numel) * 2, - "minimax_h3: F16 tensor span does not match its shape"); - const uint16_t* src = reinterpret_cast(tensor.data); - for (int64_t i = 0; i < numel; ++i) out[static_cast(i)] = F16ToF32(src[i]); + "minimax_h3: 16-bit tensor span does not match its shape"); + // Byte-wise load, NOT `reinterpret_cast(tensor.data)[i]`. + // safetensors puts the payload immediately after a JSON header of ARBITRARY + // length, so a tensor's first byte is only 2-byte aligned if the writer + // happened to pad; the format does not require it. The cast was UB on such a + // file and UBSan caught it ("load of misaligned address ... requires 2 byte + // alignment") the first time a checkpoint with an odd header reached this + // path. memcpy has no alignment precondition and compiles to the same load + // where the address does happen to be aligned. + const auto* bytes = static_cast(tensor.data); + const bool bf16 = (tensor.dtype == "BF16"); + for (int64_t i = 0; i < numel; ++i) { + uint16_t bits; + std::memcpy(&bits, bytes + static_cast(i) * 2, sizeof(bits)); + out[static_cast(i)] = bf16 ? Bf16ToF32(bits) : F16ToF32(bits); + } } else { VT_CHECK(false, "minimax_h3: unsupported tensor dtype (expected F32/BF16/F16)"); } diff --git a/tests/vllm/models/test_minimax_h3.cpp b/tests/vllm/models/test_minimax_h3.cpp index f5361a5a0..3f7171bf3 100644 --- a/tests/vllm/models/test_minimax_h3.cpp +++ b/tests/vllm/models/test_minimax_h3.cpp @@ -25,6 +25,8 @@ #include #include #include +#include +#include #include #include #include @@ -386,32 +388,62 @@ std::unique_ptr BuildDitForwardCase(const MiniMaxH3DitParams& p) return c; } -// Serialize a synthetic compressed-tensors NVFP4 (W4A16) MiniMax-H3 DiT file at -// the geometry `want`, exactly as the real `lilcheaty/MiniMax-H3-NVFP4` file -// stores it: quantized projections as U8 packed [out, in/2] + E4M3 group-16 -// weight_scale + F32 weight_scale_2; islands (patch/time/output/norms) plain F32. -// Factored out of the CPU "NVFP4 checkpoint loads" case so the CUDA speed case -// can build the SAME file at real geometry without duplicating 100 lines. Both -// callers set num_layers == token_refiner_num_layers == 1, which is what makes -// the quantized-GEMM count exactly 11 (refiner 4 + block 5 + condition + final). -void WriteMiniMaxH3Nvfp4File(const MiniMaxH3DitParams& want, const std::string& path) { - struct Entry { - std::string name; - std::string dtype; - std::vector shape; - std::string bytes; - }; +// The synthetic checkpoint writers. `WriteMiniMaxH3Nvfp4File`'s callers set +// num_layers == token_refiner_num_layers == 1, which is what makes the +// quantized-GEMM count exactly 11 (refiner 4 + block 5 + condition + final). +// +// One serialized safetensors entry. Shared by the single-file NVFP4 writer and the +// MULTI-SHARD bf16 writer below, so both emit the SAME tensor set from ONE list -- +// two copies of the ~30-name DiT layout would drift the moment a tensor is added. +struct H3StEntry { + std::string name; + std::string dtype; + std::vector shape; + std::string bytes; +}; + +std::string PackF32(const std::vector& v) { + return std::string(reinterpret_cast(v.data()), v.size() * sizeof(float)); +} + +// Round-to-nearest-even, the rule vt uses on a bf16 store. +std::string PackBf16(const std::vector& v) { + std::string out(v.size() * sizeof(uint16_t), '\0'); + for (size_t i = 0; i < v.size(); ++i) { + uint32_t bits; + std::memcpy(&bits, &v[i], sizeof(bits)); + const uint32_t rounded = bits + 0x7FFFu + ((bits >> 16) & 1u); + const uint16_t half = static_cast(rounded >> 16); + std::memcpy(&out[i * sizeof(uint16_t)], &half, sizeof(half)); + } + return out; +} + +// Build the WHOLE DiT tensor set at geometry `want`. +// quantize -- projections become the NVFP4 triple (the single-file NVFP4 arm); +// otherwise every tensor is written plain. +// plain_bf16 -- plain tensors are stored BF16, which is what the original bf16 +// release does, unless their name is in `f32_names`. It is false +// with `quantize`, keeping the NVFP4 file byte-for-byte what it was. +std::vector BuildMiniMaxH3DitEntries(const MiniMaxH3DitParams& want, bool quantize, + bool plain_bf16, + const std::set& f32_names) { + using Entry = H3StEntry; std::vector entries; auto add_plain = [&](const std::string& name, const std::vector& shape) { int64_t numel = 1; for (int64_t d : shape) numel *= d; const std::vector values = MakeParam("nvfp4." + name, numel, 0.1); - entries.push_back({name, "F32", shape, - std::string(reinterpret_cast(values.data()), - values.size() * sizeof(float))}); + const bool as_f32 = !plain_bf16 || f32_names.count(name) != 0; + entries.push_back({name, as_f32 ? "F32" : "BF16", shape, + as_f32 ? PackF32(values) : PackBf16(values)}); }; auto add_quant = [&](const std::string& name, int64_t out_dim, int64_t in_dim) { + if (!quantize) { + add_plain(name, {out_dim, in_dim}); + return; + } REQUIRE(in_dim % 16 == 0); std::string packed(static_cast(out_dim * (in_dim / 2)), '\0'); for (size_t i = 0; i < packed.size(); ++i) { @@ -468,7 +500,12 @@ void WriteMiniMaxH3Nvfp4File(const MiniMaxH3DitParams& want, const std::string& add_plain("final_layer.video_out.bias", {video_width}); add_plain("final_layer.audio_out.weight", {want.audio_latents_dim, want.hidden_size}); add_plain("final_layer.audio_out.bias", {want.audio_latents_dim}); + return entries; +} +// Serialize `entries` as ONE .safetensors file. +void WriteSafetensorsFromEntries(const std::vector& entries, const std::string& path) { + using Entry = H3StEntry; std::string header = "{"; size_t offset = 0; bool first = true; @@ -494,6 +531,134 @@ void WriteMiniMaxH3Nvfp4File(const MiniMaxH3DitParams& want, const std::string& std::fclose(fh); } +// The synthetic single-file NVFP4 DiT (`lilcheaty/MiniMax-H3-NVFP4`'s layout): +// quantized projections as U8 packed [out, in/2] + E4M3 group-16 weight_scale + +// F32 weight_scale_2; islands (patch/time/output/norms) plain F32. +void WriteMiniMaxH3Nvfp4File(const MiniMaxH3DitParams& want, const std::string& path) { + WriteSafetensorsFromEntries( + BuildMiniMaxH3DitEntries(want, /*quantize=*/true, /*plain_bf16=*/false, {}), path); +} + +// --- the ORIGINAL bf16 release's shape: N shards + model.safetensors.index.json -- + +std::string ShardFileName(size_t i, size_t n) { + char buf[64]; + std::snprintf(buf, sizeof(buf), "model-%05zu-of-%05zu.safetensors", i + 1, n); + return std::string(buf); +} + +// Write `entries` as a MULTI-SHARD checkpoint under `dir`: contiguous chunks +// across `num_shards` files plus the index. Returns the name -> shard-file map +// the index promises, so the test can assert every tensor resolved to the shard +// it was actually written into. +// +// `omit_payload` names tensors the INDEX still lists but that are deliberately +// left OUT of their shard — a checkpoint that is corrupt in the one way that +// otherwise fails SILENTLY (a skipped weight reads as zeros and renders). +std::map WriteMiniMaxH3ShardedDit( + const std::vector& entries, const std::string& dir, size_t num_shards, + const std::set& omit_payload = {}) { + REQUIRE(num_shards > 0); + REQUIRE(entries.size() >= num_shards); + ::mkdir(dir.c_str(), 0755); + + std::map weight_map; + std::vector> per_shard(num_shards); + for (size_t i = 0; i < entries.size(); ++i) { + const size_t shard = i * num_shards / entries.size(); + weight_map[entries[i].name] = ShardFileName(shard, num_shards); + if (omit_payload.count(entries[i].name) == 0) per_shard[shard].push_back(entries[i]); + } + for (size_t s = 0; s < num_shards; ++s) { + WriteSafetensorsFromEntries(per_shard[s], dir + "/" + ShardFileName(s, num_shards)); + } + + nlohmann::json index; + index["metadata"] = {{"total_size", 0}}; + index["weight_map"] = weight_map; + FILE* fh = std::fopen((dir + "/model.safetensors.index.json").c_str(), "wb"); + REQUIRE(fh != nullptr); + const std::string text = index.dump(); + std::fwrite(text.data(), 1, text.size(), fh); + std::fclose(fh); + return weight_map; +} + +// The REAL release's SHAPE without its 66.3 GB of payload: headers declare every +// tensor at its true size and the payload is a SPARSE hole (ftruncate), so the +// manifest a loader reads is byte-for-byte the real one while the files cost a +// few KB of disk. Only names/shapes/dtypes are ever read from it — never a weight +// byte — which is exactly what --dump-params does on the real directory. +// Returns the total DECLARED payload bytes. +uint64_t WriteMiniMaxH3SparseShardedRelease(const std::vector& specs, + const std::string& dir, size_t num_shards) { + REQUIRE(num_shards > 0); + ::mkdir(dir.c_str(), 0755); + std::vector> per_shard(num_shards); + std::map weight_map; + for (size_t i = 0; i < specs.size(); ++i) { + const size_t shard = i * num_shards / specs.size(); + per_shard[shard].push_back(&specs[i]); + weight_map[specs[i].name] = ShardFileName(shard, num_shards); + } + + uint64_t declared = 0; + for (size_t s = 0; s < num_shards; ++s) { + std::string header = "{"; + uint64_t offset = 0; + bool first = true; + for (const vllm::MiniMaxH3TensorSpec* spec : per_shard[s]) { + // The upstream dtype policy: fp32 ISLANDS stay F32, everything else is BF16 + // — which is what makes the whole DiT ~66.3 GB. + const uint64_t width = spec->fp32 ? 4u : 2u; + uint64_t numel = 1; + for (int64_t d : spec->shape) numel *= static_cast(d); + const uint64_t bytes = numel * width; + if (!first) header += ","; + first = false; + header += "\"" + spec->name + "\":{\"dtype\":\"" + (spec->fp32 ? "F32" : "BF16") + + "\",\"shape\":["; + for (size_t i = 0; i < spec->shape.size(); ++i) { + if (i) header += ","; + header += std::to_string(spec->shape[i]); + } + header += "],\"data_offsets\":[" + std::to_string(offset) + "," + + std::to_string(offset + bytes) + "]}"; + offset += bytes; + } + header += "}"; + const std::string path = dir + "/" + ShardFileName(s, num_shards); + FILE* fh = std::fopen(path.c_str(), "wb"); + REQUIRE(fh != nullptr); + const uint64_t n = header.size(); + std::fwrite(&n, sizeof(n), 1, fh); + std::fwrite(header.data(), 1, header.size(), fh); + std::fflush(fh); + // The payload is a HOLE: declared in full, allocated not at all. + REQUIRE(::ftruncate(fileno(fh), static_cast(sizeof(n) + header.size() + offset)) == 0); + std::fclose(fh); + declared += offset; + } + + nlohmann::json index; + index["metadata"] = {{"total_size", declared}}; + index["weight_map"] = weight_map; + FILE* fh = std::fopen((dir + "/model.safetensors.index.json").c_str(), "wb"); + REQUIRE(fh != nullptr); + const std::string text = index.dump(); + std::fwrite(text.data(), 1, text.size(), fh); + std::fclose(fh); + return declared; +} + +void RemoveShardedDit(const std::string& dir, size_t num_shards) { + for (size_t s = 0; s < num_shards; ++s) { + std::remove((dir + "/" + ShardFileName(s, num_shards)).c_str()); + } + std::remove((dir + "/model.safetensors.index.json").c_str()); + ::rmdir(dir.c_str()); +} + } // namespace TEST_CASE("minimax_h3: the deterministic weight stream matches the generator") { @@ -3970,6 +4135,224 @@ TEST_CASE("minimax_h3: an NVFP4 checkpoint loads into a runnable DiT") { std::remove(path.c_str()); } +// --------------------------------------------------------------------------- +// The ORIGINAL bf16 release: 13 safetensors shards, 66.3 GB +// --------------------------------------------------------------------------- +// Every render so far used a QUANTIZED DiT, and Q3_K_M -> Q4_K_M alone turned a +// murky lattice-covered silhouette into a photoreal close-up — so "what does FULL +// PRECISION look like?" is the next question, and until this loader existed we +// could not ask it: every DiT loader took a SINGLE file. + +namespace { + +// The reduced geometry the sharded gates run at. Mirrors the NVFP4 case's dims so +// the two arms are comparable, with text_dim a multiple of 16. +MiniMaxH3DitParams ShardedGateParams() { + MiniMaxH3DitParams want; + want.num_layers = 2; + want.token_refiner_num_layers = 1; + want.hidden_size = 64; + want.num_attention_heads = 4; + want.attention_head_dim = 16; + want.ffn_hidden_size = 128; + want.latents_dim = 8; + want.audio_latents_dim = 6; + want.text_dim = 32; + want.timestep_input_dim = 16; + want.time_embed_hidden_size = 64; + want.time_embed_dim = 32; + want.adaln_out_features = 18 * want.hidden_size; + want.final_adaln_out_features = 2 * want.hidden_size; + want.rope_inv_freq_len = 2; + return want; +} + +// The dtype MIX the loader must survive. The real release stores the model bf16; +// these four names are pinned F32 so ONE synthetic checkpoint exercises all four +// (on-disk dtype x device dtype) combinations: +// BF16 -> bf16 device slot direct mmap upload, NO host buffer (the bulk) +// F32 -> f32 island direct mmap upload, NO host buffer (time_embedder) +// BF16 -> f32 island widened on the host (patch/out heads) +// F32 -> bf16 device slot rounded on the host (blocks.0.norm1) +std::set ShardedGateF32Names() { + return {"time_embedder.proj_in.weight", "time_embedder.proj_in.bias", + "time_embedder.proj_out.weight", "time_embedder.proj_out.bias", + "rope.inv_freq", "blocks.0.norm1.weight"}; +} + +} // namespace + +TEST_CASE("minimax_h3: the multi-shard index resolves every tensor to its own shard") { + const MiniMaxH3DitParams want = ShardedGateParams(); + const std::vector entries = + BuildMiniMaxH3DitEntries(want, /*quantize=*/false, /*plain_bf16=*/true, + ShardedGateF32Names()); + const std::string dir = "/tmp/minimax_h3_sharded_index"; + const size_t kShards = 4; + const std::map promised = + WriteMiniMaxH3ShardedDit(entries, dir, kShards); + + const vllm::MiniMaxH3ShardedCheckpoint ckpt = vllm::MiniMaxH3ShardedCheckpoint::Open(dir); + CHECK(ckpt.ShardCount() == kShards); + CHECK(ckpt.Names().size() == entries.size()); + CHECK(ckpt.IndexPath() == dir + "/model.safetensors.index.json"); + + // ★ EVERY tensor resolves to the shard the index named — and to the tensor that + // was actually written into it, not a same-named one elsewhere. Resolving a + // tensor to the WRONG shard is the failure that yields a loaded-but-wrong model. + size_t checked = 0; + for (const auto& entry : promised) { + REQUIRE(ckpt.Has(entry.first)); + CHECK(ckpt.ShardOf(entry.first) == entry.second); + const vllm::StTensor& t = ckpt.Get(entry.first); + const auto it = std::find_if(entries.begin(), entries.end(), + [&](const H3StEntry& e) { return e.name == entry.first; }); + REQUIRE(it != entries.end()); + CHECK(t.dtype == it->dtype); + CHECK(t.shape == it->shape); + REQUIRE(t.nbytes == it->bytes.size()); + CHECK(std::memcmp(t.data, it->bytes.data(), t.nbytes) == 0); + ++checked; + } + CHECK(checked == entries.size()); + CHECK(!ckpt.Has("blocks.99.mlp.fc1.weight")); + + // The shards really are several files, each holding part of the model. + CHECK(ckpt.ShardFiles().size() == kShards); + std::set distinct; + for (const auto& entry : promised) distinct.insert(entry.second); + CHECK(distinct.size() == kShards); + + // GEOMETRY from the shards must equal the SINGLE-FILE path over the same tensors: + // one file, same entries, same derived params — so sharding is a container + // question and never a model question. + const std::vector manifest = + vllm::EnumerateMiniMaxH3ShardedTensors(ckpt); + const MiniMaxH3DitParams sharded = vllm::ParseMiniMaxH3DitParamsFromGgufManifest(manifest); + const std::string single = "/tmp/minimax_h3_sharded_single.safetensors"; + WriteSafetensorsFromEntries(entries, single); + const vllm::SafetensorsFile sf = vllm::SafetensorsFile::Open(single); + const vllm::MiniMaxH3GgufDit single_loaded = vllm::LoadMiniMaxH3DitFromNvfp4(sf); + const MiniMaxH3DitParams& one = single_loaded.params; + CHECK(sharded.num_layers == one.num_layers); + CHECK(sharded.token_refiner_num_layers == one.token_refiner_num_layers); + CHECK(sharded.hidden_size == one.hidden_size); + CHECK(sharded.num_attention_heads == one.num_attention_heads); + CHECK(sharded.attention_head_dim == one.attention_head_dim); + CHECK(sharded.ffn_hidden_size == one.ffn_hidden_size); + CHECK(sharded.latents_dim == one.latents_dim); + CHECK(sharded.audio_latents_dim == one.audio_latents_dim); + CHECK(sharded.patch_size_t == one.patch_size_t); + CHECK(sharded.patch_size_h == one.patch_size_h); + CHECK(sharded.patch_size_w == one.patch_size_w); + CHECK(sharded.text_dim == one.text_dim); + CHECK(sharded.timestep_input_dim == one.timestep_input_dim); + CHECK(sharded.time_embed_hidden_size == one.time_embed_hidden_size); + CHECK(sharded.time_embed_dim == one.time_embed_dim); + CHECK(sharded.adaln_out_features == one.adaln_out_features); + CHECK(sharded.final_adaln_out_features == one.final_adaln_out_features); + CHECK(sharded.rope_inv_freq_len == one.rope_inv_freq_len); + CHECK(sharded.video_row_width() == one.video_row_width()); + CHECK(sharded.rope_rot_dim() == one.rope_rot_dim()); + // ...and it is the geometry that was asked for. + CHECK(sharded.num_layers == want.num_layers); + CHECK(sharded.hidden_size == want.hidden_size); + CHECK(sharded.text_dim == want.text_dim); + + // The manifest also carries the fp32-ISLAND flag, and it must agree with the + // upstream-derived enumeration name for name (see the island test below). + for (const vllm::MiniMaxH3TensorSpec& spec : manifest) { + CHECK(spec.fp32 == vllm::MiniMaxH3IsFp32IslandTensor(spec.name)); + } + + std::remove(single.c_str()); + RemoveShardedDit(dir, kShards); + + // ★ A tensor the index NAMES but its shard does not contain must throw BY NAME. + // Silently skipping it would leave that weight reading as zeros — a plausible + // but wrong render rather than an error. + const std::string broken_dir = "/tmp/minimax_h3_sharded_broken"; + const std::string missing = "blocks.1.mlp.fc2.weight"; + WriteMiniMaxH3ShardedDit(entries, broken_dir, kShards, {missing}); + bool threw = false; + try { + const vllm::MiniMaxH3ShardedCheckpoint bad = + vllm::MiniMaxH3ShardedCheckpoint::Open(broken_dir); + (void)bad; + } catch (const std::exception& e) { + threw = true; + const std::string what = e.what(); + INFO("missing-tensor error: " << what); + CHECK(what.find(missing) != std::string::npos); + } + CHECK(threw); + RemoveShardedDit(broken_dir, kShards); + + // A directory with no index at all is refused, not half-loaded. + CHECK(!vllm::MiniMaxH3ShardedCheckpoint::IsShardedDir("/tmp")); + CHECK_THROWS(vllm::MiniMaxH3ShardedCheckpoint::Open("/tmp")); +} + +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 + // names and shapes without downloading or storing a weight byte. + MiniMaxH3DitParams shipped; // the defaults ARE the shipped H3 geometry + const std::vector specs = + vllm::EnumerateMiniMaxH3DitTensors(shipped); + const std::string dir = "/tmp/minimax_h3_sharded_release"; + const size_t kShards = 13; // what the release actually ships + const uint64_t declared = WriteMiniMaxH3SparseShardedRelease(specs, dir, kShards); + INFO("declared payload = " << (declared / (1024.0 * 1024.0 * 1024.0)) << " GiB"); + CHECK(declared > 55ull * 1024 * 1024 * 1024); // ~66.3 GB of bf16 weights + CHECK(declared < 70ull * 1024 * 1024 * 1024); + + const vllm::MiniMaxH3ShardedCheckpoint ckpt = vllm::MiniMaxH3ShardedCheckpoint::Open(dir); + CHECK(ckpt.ShardCount() == kShards); + CHECK(ckpt.Names().size() == specs.size()); + + const MiniMaxH3DitParams p = + vllm::ParseMiniMaxH3DitParamsFromGgufManifest(vllm::EnumerateMiniMaxH3ShardedTensors(ckpt)); + // ★ The same 20 fields --dump-params prints, and the same values the WORKING + // GGUF arm derives from shapes alone. A mismatch means the name mapping is wrong. + CHECK(p.num_layers == 50); + CHECK(p.token_refiner_num_layers == 2); + CHECK(p.hidden_size == 5376); + CHECK(p.num_attention_heads == 56); + CHECK(p.attention_head_dim == 128); + CHECK(p.ffn_hidden_size == 14336); + CHECK(p.latents_dim == 24); + CHECK(p.audio_latents_dim == 32); + CHECK(p.patch_size_t == 1); + CHECK(p.patch_size_h == 2); + CHECK(p.patch_size_w == 2); + CHECK(p.text_dim == 5120); + CHECK(p.timestep_input_dim == 256); + CHECK(p.time_embed_hidden_size == 5376); + CHECK(p.time_embed_dim == 2688); + CHECK(p.adaln_out_features == 18 * 5376); + CHECK(p.final_adaln_out_features == 2 * 5376); + CHECK(p.rope_inv_freq_len == 16); + CHECK(p.video_row_width() == 96); + CHECK(p.rope_rot_dim() == 96); + + // Every tensor of the release resolves, and the fp32-ISLAND split matches the + // upstream-derived enumeration name for name — the split vt::MatmulBT enforces + // at the first island GEMM. + size_t islands = 0; + for (const vllm::MiniMaxH3TensorSpec& spec : specs) { + REQUIRE(ckpt.Has(spec.name)); + CHECK(ckpt.Get(spec.name).shape == spec.shape); + CHECK(spec.fp32 == vllm::MiniMaxH3IsFp32IslandTensor(spec.name)); + islands += spec.fp32 ? 1 : 0; + } + CHECK(islands == 13); // 12 island weights/biases + rope.inv_freq + + RemoveShardedDit(dir, kShards); +} + TEST_CASE("minimax_h3: the NVFP4 fp4 forward runs Marlin W4A16 on CUDA (speed)") { // The GB10 leg the CPU wiring gate cannot reach (spec 8.4a): on the CUDA backend // the W4A16 dispatcher hits dense_nvfp4::MatmulNvfp4MarlinD, so the fp4 path RAN