diff --git a/.agents/NOW.md b/.agents/NOW.md index 3995dd8ae..2832a8f0a 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -18,7 +18,7 @@ checkpoint on `upstream/main` at `59674cf1d`. | DeepSeek-V4-Flash decode | **Closed: beats ds4 1.144x** (`VT_V4_RESIDENT_W`, byte-exact); phase-2 residency NEG, default-OFF | — | | f32-out GEMV audit | Only laguna + ds4 bf16 tower affected; gate models unaffected | Re-verify ds4 tower same-tool | | Invocation-parity prevention | CI guard + AGENTS.md checklist landing | Merge; build-verify `kGemvHeuristicAlgos` on dgx | -| MiniMax-H3 lane | **RENDER BUG CLOSED** (`H3-RENDER-CLOSE` PR #77): #70/#74 white = t2va on the REF2VA ckpt; the FL2VA GGUF t2va renders COHERENT (adj-cos 0.95) | Follow-up: partition guard + vision tower | +| MiniMax-H3 lane | **RENDER + task/partition guard CLOSED** (#77 render; `H3-TASK-PARTITION-GUARD`: t2va-on-ref2va now RAISES like `_resolve_task`) | Follow-up: vision tower (W3) | | Kimi-Linear-48B (KDA+NoPE-MLA+MoE) | **e2e RUNS** (bf16-resident §13): 13/13·656. Token gate **NEAR-TIE 106/128** | device GDN/MLA islands; 1.59 tok/s; default OFF | | 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/benchmark-record.md b/.agents/benchmark-record.md index 064121e13..5f98481c8 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -14167,3 +14167,48 @@ noise fix from #70 stands but was already known not to be the render fix. | NVFP4 dequant | independent torch dequant of `blocks.0.attn.qkv_proj` + Laguna/Qwen3 already prove `DequantNvfp4ToBf16` byte-exact | sane trained weight (rms 0.089, absmax=ws2·6·maxscale=3.61) | | CUDA kernels at real seq | NEW gate `test_minimax_h3 :: CUDA device forward tracks the host at the REAL render seq (1920)` (RealRatioParams head_dim=128, seq 1920) | **CUDA device == CPU host** (28/28) — no scale-dependent kernel bug (#74 only ran device-vs-host on the CPU backend) | | forward math | RefDiT restatement vs true upstream source, read side by side (block, attention, AdaLN view(m*3,6H), 3D-RoPE, modulate) | identical | + +## MiniMax-H3 TASK/PARTITION GUARD — mirror `_resolve_task`'s raise; the two arms have NO structural discriminator so a stripped file must DECLARE its partition (2026-08-06, `row/H3-TASK-PARTITION-GUARD` PR #84, `ROAD-V1-H3`, CPU-only dev box) + +The #70/#74 follow-up. The white grid was `task=t2va` run on the Ref2VA-partition +checkpoint; upstream `pipeline._resolve_task` RAISES on the mismatch +(`pipeline_minimax_h3.py:374-391`, raise at 387-390), and the recipe documents the split +(`recipes/MiniMaxAI/MiniMax-H3.md:50-51,289`). This row mirrors the raise 1:1 and adds the +community-file fallback. + +**The discriminator finding (grounded both sides).** Upstream derives the partition from +the release config `model_index.json` → `_minimax_h3` → `{partition, tasks}` +(`pipeline_minimax_h3.py:279-282`). Community GGUF/NVFP4 STRIP that block, and there is +NO structural fallback — MEASURED on the two real captured manifests: + +| Arm | Manifest (tensors) | Normalized base names | Key shapes (video/audio/condition/time patch) | +|---|---|---|---| +| Ref2VA | `minimax_h3_nvfp4_manifest.inc` (1051) | 535 after collapsing `{weight,weight_scale,weight_scale_2}` | `[5376,96]`/`[5376,32]`/`[5376,5120]`/`[5376,256]` | +| FL2VA | `minimax_h3_gguf_manifest.inc` (535) | 535 | identical | + +`comm -23`/`-13` of the two normalized name sets is EMPTY both ways; all reference-relevant +shapes match. Ref2VA prepends reference rows through the SAME `video/audio_patch_proj`, so +it adds no reference-specific tensor. A name/shape auto-detector is impossible in principle, +so a stripped file must DECLARE `--partition fl2va|ref2va` (server: `--video-partition`); +`MiniMaxH3PartitionFromFlag` maps it to the recipe's served-task set (fl2va→{t2va,fl2va}, +ref2va→{ref2va}). + +**Guard behavior table (task × partition → pass/refuse):** + +| task \ partition | FL2VA {t2va,fl2va} | Ref2VA {ref2va} | unknown/stripped | +|---|---|---|---| +| t2va | pass | **REFUSE (#77 mismatch)** | REFUSE (declare `--partition`) | +| fl2va | pass | REFUSE | REFUSE | +| ref2va| REFUSE | pass | REFUSE | + +**RED-first proof (reviewer mutation).** New case `test_minimax_h3 :: "the task/partition +guard refuses the #77 mismatch"`, 38 assertions: the #77 combo throws, correct pairings +pass, stripped refuses every task, `--partition` recovers, `MiniMaxH3TaskOfRequest` maps +the three request shapes, and the two real manifests are asserted to reduce to the same +535-name set. Neutralizing the guard body (`return;` at the top of +`MiniMaxH3CheckTaskPartition`) turned the case RED at **10 failed assertions** (t2va-on- +ref2va, the stripped refusals, the dispatch-level `t2va_req` on Ref2VA); restoring it went +GREEN. Suite **67/67** (66 prior + this), 46549 assertions; `test_video_api` 4/4 (server +wiring). Wired at both loading entry points: `MiniMaxH3GenerateT2va` guards every full +render; the pure pipeline-math tests build `declared=false` requests so the guard is inert +there. No numbers changed — this is a correctness/refusal gate, not a perf lever. diff --git a/.agents/specs/minimax-h3.md b/.agents/specs/minimax-h3.md index 28dddd0d2..ad056ae2e 100644 --- a/.agents/specs/minimax-h3.md +++ b/.agents/specs/minimax-h3.md @@ -565,3 +565,62 @@ trailing target rows are the clip; the old code fed unpatchify the full buffer a guard mirroring upstream (community files strip the release config); the encoder vision tower (W3) is still unported, so image/video-conditioned ref2va/fl2va renders are not yet clean (ref2va with a synthetic reference + text-only encoder still grids). + +## 8.7 TASK/PARTITION GUARD — mirror `_resolve_task`'s raise (2026-08-06, `row/H3-TASK-PARTITION-GUARD` PR #84) + +The #70/#74 white grid cost three campaigns because our driver silently accepted +`task=t2va` on the Ref2VA-partition checkpoint. Upstream `pipeline._resolve_task` +RAISES on the mismatch (`pipeline_minimax_h3.py:374-391`, esp. 387-390); the recipe +documents the split (`recipes/MiniMaxAI/MiniMax-H3.md:50-51,289`: "One server loads one +checkpoint partition … must match the served partition"). This row mirrors the raise 1:1. + +**Partition detection — two paths, and the definitive no-discriminator finding.** +Upstream reads the served-task set from the release config +(`pipeline_minimax_h3.py:279-282`): + +``` +release = model_index.get("_minimax_h3") or {} +self.partition = str(release.get("partition", "")) # "fl2va" | "ref2va" +self.supported_tasks = frozenset(release.get("tasks") or ()) +``` + +`MiniMaxH3PartitionFromModelIndex(model_index)` mirrors those exact keys. But community +GGUF/NVFP4 redistributions STRIP that block, and — measured on the two real manifests +this spec already captured — there is **NO structural fallback**: the Ref2VA NVFP4 +(1051 tensors) and FL2VA GGUF (535 tensors) carry the **IDENTICAL DiT**. Normalizing the +NVFP4 `{weight, weight_scale, weight_scale_2}` split, both files reduce to the **SAME 535 +base tensor names AND the SAME shapes** (video_patch_proj `[5376,96]`, audio_patch_proj +`[5376,32]`, condition_proj `[5376,5120]`, time_embedder.proj_in `[5376,256]` on both; +`comm -23`/`-13` of the normalized name sets is empty both ways). Ref2VA conditioning is +achieved by PREPENDING reference rows through the SAME `video/audio_patch_proj` weights, +so it introduces no reference-specific tensor to key on. A name/shape auto-detector is +therefore impossible in principle. When the config is stripped the partition must be +**DECLARED** (`--partition fl2va|ref2va`), never guessed; `MiniMaxH3PartitionFromFlag` +maps it to the recipe's served-task set (fl2va→{t2va,fl2va}, ref2va→{ref2va}). + +**The refuse.** `MiniMaxH3CheckTaskPartition(task, info)` is the raise half of +`_resolve_task`. The task is what the request ENCODES (`MiniMaxH3TaskOfRequest`: +`ref_blocks`→ref2va, `keyframe_frame_indices`→fl2va, else t2va), and +`MiniMaxH3GenerateT2va` calls the pair before denoising. A declared partition refuses a +task it does not serve; an UNKNOWN partition (stripped file, no `--partition`) refuses +EVERY task as ambiguous and names the recipe lines. A default-constructed +`MiniMaxH3PartitionInfo` (`declared=false`) leaves the guard inactive, so the pure +pipeline-math unit tests are unaffected. Wired at both checkpoint-loading entry points: +the driver (`--partition`) and the server (`--video-partition`). + +**Guard behavior table (task × partition → pass/refuse):** + +| task \ partition | FL2VA {t2va,fl2va} | Ref2VA {ref2va} | unknown/stripped | +|---|---|---|---| +| **t2va** | pass | **REFUSE (the #77 mismatch)** | REFUSE (declare `--partition`) | +| **fl2va** | pass | REFUSE | REFUSE | +| **ref2va**| REFUSE | pass | REFUSE | + +**RED-first proof.** New case `test_minimax_h3 :: "the task/partition guard refuses the +#77 mismatch"` (38 assertions): the #77 combo `MiniMaxH3CheckTaskPartition("t2va", +ref2va)` throws; the correct pairings pass; the stripped case refuses every task and +`--partition` recovers it; `MiniMaxH3TaskOfRequest` maps the three request shapes; and it +asserts the two real manifests reduce to the identical 535-name set (proving the +no-discriminator premise in the harness). Neutralizing the guard body (reviewer mutation) +turned the case RED at 10 assertions, restoring it turned it GREEN — the test has teeth. +Suite: 67/67 (66 prior + this), 46549 assertions. `test_video_api` 4/4 (server wiring). diff --git a/.agents/state.md b/.agents/state.md index c15656049..78f5dcde4 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -39742,3 +39742,41 @@ partition/supported_tasks guard mirroring upstream (community files strip the re and the encoder vision tower (W3) for clean image/video-conditioned ref2va/fl2va (ref2va with a synthetic reference + text-only encoder still grids). dgx assets: `~/h3fp4/ckpt/MiniMax-H3- FL2VA-Q3_K_M.gguf`, `~/h3fp4/fl2va_t2va_20/`. Box left clean. + +## MiniMax-H3 task/partition GUARD landed — mirror `_resolve_task`'s raise (`row/H3-TASK-PARTITION-GUARD` PR #84, helper, DRAFT PR, CPU-only) + + +The #77 follow-up: the driver silently accepted `task=t2va` on the Ref2VA-partition +checkpoint (the #70/#74 white grid, three campaigns). Upstream `pipeline._resolve_task` +RAISES on the mismatch (`vllm_omni/diffusion/models/minimax_h3/pipeline_minimax_h3.py:374-391`, +raise at 387-390); the recipe documents the split (`recipes/MiniMaxAI/MiniMax-H3.md:50-51,289`: +one server serves one partition, FL2VA→{t2va,fl2va}, Ref2VA→{ref2va}). Mirrored 1:1. + +DISCRIMINATOR FINDING (definitive, both sides). Upstream reads the served-task set from the +release config `model_index.json` → `_minimax_h3` → `{partition,tasks}` (pipeline:279-282); +`MiniMaxH3PartitionFromModelIndex` mirrors those exact keys. Community GGUF/NVFP4 STRIP that +block, and there is NO structural fallback — MEASURED on the two real captured manifests: the +Ref2VA NVFP4 (`minimax_h3_nvfp4_manifest.inc`, 1051 tensors) and FL2VA GGUF +(`minimax_h3_gguf_manifest.inc`, 535) carry the IDENTICAL DiT. Collapsing the NVFP4 +`{weight,weight_scale,weight_scale_2}` split, both reduce to the SAME 535 base tensor names AND +the SAME shapes (`comm` empty both ways; video_patch_proj `[5376,96]`, condition_proj +`[5376,5120]`, etc. equal). Ref2VA prepends reference rows through the SAME +`video/audio_patch_proj`, adding no reference-specific tensor. So a stripped file must DECLARE +`--partition fl2va|ref2va` (server `--video-partition`); `MiniMaxH3PartitionFromFlag` maps it. + +CODE. New (pure, gated-on-any-machine, `minimax_h3_planner.cpp`): `MiniMaxH3PartitionInfo` +(`declared`/`partition`/`supported_tasks`), `MiniMaxH3PartitionFromModelIndex`, +`MiniMaxH3PartitionFromFlag`, `MiniMaxH3TaskOfRequest` (ref_blocks→ref2va, keyframes→fl2va, +else t2va), and the raise `MiniMaxH3CheckTaskPartition`. `MiniMaxH3GenerateT2va` calls the +pair before denoising; a default `declared=false` request leaves it inert (pipeline-math tests +unaffected). Driver `--partition` + server `--video-partition` wired. An unknown partition +refuses EVERY task (ambiguous) and names the recipe lines. + +GATE. New case `test_minimax_h3 :: "the task/partition guard refuses the #77 mismatch"` (38 +assertions): #77 combo throws, correct pairings pass, stripped refuses+`--partition` recovers, +task-of-request maps all three shapes, and the two real manifests are asserted to the same +535-name set. RED-first (reviewer mutation): neutralizing the guard body → 10 failed +assertions; restored → GREEN. Suite 67/67 (66 prior +1), 46549 assertions; `test_video_api` +4/4 (server wiring). No numbers changed (refusal gate, not a perf lever). Records: spec §8.7, +STATUS/BENCHMARKS H3 rows, benchmark-record, NOW. Pre-existing preflight red +(check-fusion-consistency `minimax_h3_video_vae_device`) is not this row. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index d6bf145c1..7099b0af7 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -310,7 +310,7 @@ built on it rather than keeping the flattering one. | Kimi-Linear-48B-A3B (KDA+MLA+MoE) | Full-model GB10 e2e RUNS (bf16-resident §13), NEAR-TIE 106/128, pool math CLOSES; default OFF | Full model RUNS on GB10 (bf16-resident, RSS peak 1.7 GiB, min-avail 21 GiB, no OOM). Token NEAR-TIE 106/128 (6/8 prompts exact, numerics vs deterministic oracle). 1.59 tok/s. Detail: spec §13 | | 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 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; dequant byte-exact. spec §8.6 | +| 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) | | MXFP4 Qwen3-8B (W4A16 Marlin) | **`KERNEL-MARLIN-DENSE-EXEC` x3 (dense-ON default): c1 1.020, c2/c4/c8 0.962/0.966/0.969, GPU mem 2.63x less** (beats #51 1.005/0.925/0.939/0.953 EVERY axis); #44 3/3, 32B-NVFP4A16 6/6; -Werror test-guard fixes x2 | **VT_MARLIN_DENSE default-ON** (+951us). `FLASH-OCCUPANCY` #75: matched-c8 ncu, occupancy IDENTICAL 8.33%; built vLLM's exact flash recipe, matched reg+instr, STILL +10us, gap is ptxas SASS quality, no lever/flip | | Vulkan vs llama.cpp Vulkan (`BENCH-VK-LLAMA`) | **NOT APPLICABLE: no number measured, claimed or owed.** Vulkan registers 8 of the CPU backend's 83 ops and runs no model. The 2026-08-06 campaign spec is PLAN ONLY. [Detail](../.agents/specs/vulkan-full-support.md) | `VK-E`: build llama.cpp `-DGGML_VULKAN=ON` at pin `237ad9b96` on dgx, `llama-bench` on the same GGUF, record three columns (ours-Vulkan, llama.cpp-Vulkan, ours-CUDA). Blocked by `VK-B` and by `glslc` on the gate box | | SGLang floor arms | Never ran | Both arms of the SGLang comparison | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index ba2162f1f..f39f92cc7 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -159,7 +159,7 @@ model architecture is wired. | Image | ✅ correctness-gated | ✅ | ✅ | ◐ | | Video | ✅ correctness-gated | ✅ | ✅ | ☐ | | Audio | ✅ correctness-gated | ✅ | ◐ | ◐ | -| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va renders a COHERENT prompt-matched scene on GB10 (FL2VA-partition GGUF → h264/AAC mp4); render bug CLOSED (was wrong-partition usage); Marlin W4A16 byte-exact | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | +| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va renders a COHERENT prompt-matched scene on GB10 (FL2VA-partition GGUF → h264/AAC mp4); render bug CLOSED (was wrong-partition usage); task/partition guard mirrors `_resolve_task`; Marlin W4A16 byte-exact | ✅ (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 17f5ecd63..f54f93a25 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -80,7 +80,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 --audio-vae-config // --prompt-embeds (rows of text_dim, little-endian f32) // --out +// [--partition fl2va|ref2va] (REQUIRED for a full render: the +// served checkpoint partition — community GGUF/NVFP4 strip it and +// the FL2VA/Ref2VA DiTs are indistinguishable. t2va/fl2va need +// fl2va, ref2va needs ref2va; recipe:50-51,289) // [--keep-quant] [--steps N] [--frames N] [--height N] [--width N] // [--workdir DIR] [--ffmpeg PATH] [--dry-run] // @@ -170,6 +174,11 @@ int main(int argc, char** argv) { std::string roundtrip_path; // diagnostic: encode->decode a real image std::vector ref_image_paths; std::string ref_video_prefix, ref_audio_path; + // The served checkpoint PARTITION. Community GGUF/NVFP4 files strip the release + // model_index.json `_minimax_h3` block, and the FL2VA/Ref2VA DiTs are structurally + // identical, so it cannot be inferred from the weights — it must be DECLARED. Empty + // => the guard refuses a full render and tells the user to pass it (the #77 catch). + std::string partition_flag; double imgvid_noise_aug = 1.0; int64_t encoder_max_layers = 0; int64_t steps = 0, frames = 0, height = 0, width = 0; @@ -204,6 +213,7 @@ int main(int argc, char** argv) { else if (f == "--ref-image") ref_image_paths.push_back(Need(argc, argv, ++i, f)); else if (f == "--ref-video") ref_video_prefix = Need(argc, argv, ++i, f); else if (f == "--ref-audio") ref_audio_path = Need(argc, argv, ++i, f); + else if (f == "--partition") partition_flag = Need(argc, argv, ++i, f); else if (f == "--noise-aug") imgvid_noise_aug = std::stod(Need(argc, argv, ++i, f)); else if (f == "--encoder-max-layers") encoder_max_layers = std::stoll(Need(argc, argv, ++i, f)); else if (f == "--steps") steps = std::stoll(Need(argc, argv, ++i, f)); @@ -234,7 +244,8 @@ int main(int argc, char** argv) { "[--height N] [--width N] [--device cpu|cuda] [--workdir DIR] [--ffmpeg PATH] " "[--dry-run] [--denoise-only] [--dump-params] " "[--first-frame f.ppm] [--last-frame f.ppm] [--noise-aug A] " - "[--ref-image f.ppm ...] [--ref-video DIR] [--ref-audio f.wav]\n"; + "[--ref-image f.ppm ...] [--ref-video DIR] [--ref-audio f.wav] " + "[--partition fl2va|ref2va]\n"; return 2; } @@ -578,6 +589,11 @@ int main(int argc, char** argv) { request.video_latents_std = video_stats.std_dev; request.audio_latents_mean = audio_stats.mean; request.audio_latents_std = audio_stats.std_dev; + // Declare the served partition so MiniMaxH3GenerateT2va can refuse a task the + // checkpoint does not serve (the #77 catch: t2va on the Ref2VA arm). An empty + // --partition is declared-but-unknown; the guard then names the recipe lines and + // asks for fl2va|ref2va rather than silently rendering the wrong combination. + request.partition = vllm::MiniMaxH3PartitionFromFlag(partition_flag); // --- ref2va VIDEO reference: a CLIP prepended to the sequence. Reads // DIR/frame_%06d.ppm, which is exactly what this example WRITES, so a previous diff --git a/examples/server/main.cpp b/examples/server/main.cpp index 409d80448..3284b22cb 100644 --- a/examples/server/main.cpp +++ b/examples/server/main.cpp @@ -108,6 +108,7 @@ struct Args { std::string video_encoder, video_tokenizer; int video_encoder_max_layers = 50; std::string video_ffmpeg = "ffmpeg", video_device = "cuda"; + std::string video_partition; // served partition (fl2va|ref2va); see the #77 guard bool video_keep_quant = false; int cuda_profile_graph_replays = 0; // trace-only diagnostic build seam. int cuda_profile_graph_batch = 0; // 0 => accepted c16 trace contract. @@ -244,6 +245,8 @@ Args ParseArgs(int argc, char** argv) { a.video_ffmpeg = NextArg(argc, argv, i, argv[0]); } else if (flag == "--video-device") { a.video_device = NextArg(argc, argv, i, argv[0]); + } else if (flag == "--video-partition") { + a.video_partition = NextArg(argc, argv, i, argv[0]); } else if (flag == "--video-keep-quant") { a.video_keep_quant = true; } else if (flag == "--enable-server-dev-mode") { @@ -604,6 +607,10 @@ int main(int argc, char** argv) { std::unique_ptr tokenizer; vt::Queue enc_queue{}; std::string workdir, ffmpeg; + // The served partition (fl2va|ref2va). Resolved ONCE at load from + // --video-partition; the per-request guard in MiniMaxH3GenerateT2va refuses a + // task this partition does not serve (the #77 follow-up). + vllm::MiniMaxH3PartitionInfo partition_info; vt::Device device; std::atomic counter{0}; }; @@ -669,6 +676,7 @@ int main(int argc, char** argv) { } video->workdir = args.video_workdir; video->ffmpeg = args.video_ffmpeg; + video->partition_info = vllm::MiniMaxH3PartitionFromFlag(args.video_partition); if (args.video_device == "cuda") { video->device = vt::GetBackend(vt::DeviceType::kCUDA).CreateQueue().device; } @@ -716,6 +724,8 @@ int main(int argc, char** argv) { } const vllm::MiniMaxH3DitParams& p = video->dit.params; vllm::MiniMaxH3T2vaRequest r; + r.partition = video->partition_info; // #77 guard: MiniMaxH3GenerateT2va + // refuses a task this partition can't serve. const vllm::MiniMaxH3ShapePlan plan = vllm::MiniMaxH3ResolveShape( req.task.empty() ? "t2va" : req.task, req.duration_seconds, req.num_frames, req.height, req.width, 0, 0); diff --git a/include/vllm/model_executor/models/minimax_h3.h b/include/vllm/model_executor/models/minimax_h3.h index cdbe9447a..cea888cb7 100644 --- a/include/vllm/model_executor/models/minimax_h3.h +++ b/include/vllm/model_executor/models/minimax_h3.h @@ -280,6 +280,72 @@ MiniMaxH3ShapePlan MiniMaxH3ResolveShape(const std::string& task, double duratio std::string MiniMaxH3ResolveTask(const std::string& requested, const std::string& partition, bool has_image, const std::vector& supported_tasks); +// --------------------------------------------------------------------------- +// Checkpoint PARTITION / task guard (the #77 follow-up) +// --------------------------------------------------------------------------- +struct MiniMaxH3T2vaRequest; // defined below; MiniMaxH3TaskOfRequest takes it by ref +// MiniMax-H3 ships as TWO independently-served DiT partitions and the task MUST +// match the partition upstream loads (recipes/MiniMaxAI/MiniMax-H3.md:50-51,289; +// pipeline_minimax_h3.py:374-391 raises on a mismatch): +// +// FL2VA serves {t2va, fl2va} +// Ref2VA serves {ref2va} +// +// The #70/#74 white render was exactly this: t2va run on the Ref2VA NVFP4 +// checkpoint — an out-of-distribution task/partition combination upstream +// rejects. Our driver silently accepted it. This mirrors the raise. +// +// PARTITION DETECTION. The release safetensors carries the served-task set in +// `model_index.json` -> `_minimax_h3` -> {"partition","tasks"} +// (pipeline_minimax_h3.py:279-282). Community GGUF/NVFP4 redistributions STRIP +// that block, and there is NO structural fallback: the FL2VA and Ref2VA DiTs are +// byte-structurally identical — the two real manifests carry the SAME 535 base +// tensor names AND the SAME shapes (ref2va prepends reference rows through the +// SAME video/audio_patch_proj, adding no tensor), so nothing in a stripped file +// discriminates the partition. When the block is gone the partition must be +// DECLARED (an explicit --partition), never guessed. +struct MiniMaxH3PartitionInfo { + // false => no partition metadata was supplied. The guard is then a no-op: it + // gates the entry points that LOAD a checkpoint (driver/server), not the pure + // pipeline-math unit tests that build a request by hand. Every resolver below + // sets it true. + bool declared = false; + // The served partition ("fl2va" | "ref2va"), or "" when a stripped file gave no + // way to know it. Mirrors `self.partition` (pipeline_minimax_h3.py:281). + std::string partition; + // The served-task set. Mirrors `self.supported_tasks` (pipeline:282). Empty when + // the partition is unknown — then EVERY task is ambiguous and refused. + std::vector supported_tasks; +}; + +// Read `_minimax_h3.{partition,tasks}` from a parsed model_index.json, mirroring +// pipeline_minimax_h3.py:279-282 exactly. Always sets `declared` true; a missing +// `_minimax_h3` block (or missing keys) yields {true,"",{}}, which the guard +// treats as an unknown partition and refuses rather than guesses. +MiniMaxH3PartitionInfo MiniMaxH3PartitionFromModelIndex(const nlohmann::json& model_index); + +// Build partition info from an explicit --partition override, grounded in the +// recipe's one-server-one-partition split (recipes/MiniMaxAI/MiniMax-H3.md:50-51): +// "fl2va" serves {t2va, fl2va}, "ref2va" serves {ref2va}. An empty string yields +// the unknown/declared state (the guard then refuses every task); any other +// non-empty value is refused as an invalid partition name. +MiniMaxH3PartitionInfo MiniMaxH3PartitionFromFlag(const std::string& partition); + +// The task a built request ENCODES: non-empty `ref_blocks` => "ref2va", else +// non-empty `keyframe_frame_indices` => "fl2va", else "t2va". The two conditioning +// fields are mutually exclusive (packed layout enforces it), so this is the same +// key the pipeline dispatch already branches on. +std::string MiniMaxH3TaskOfRequest(const MiniMaxH3T2vaRequest& request); + +// The task/partition guard — the raise half of `_resolve_task` +// (pipeline_minimax_h3.py:387-390). Throws when `task` is not served by the +// partition ("checkpoint partition 'ref2va' supports {ref2va}, got task='t2va'"). +// When the partition is UNKNOWN (stripped community file, no --partition) EVERY +// task is ambiguous — t2va/fl2va are valid only on FL2VA, ref2va only on Ref2VA, +// and the two DiTs are indistinguishable — so it refuses and names the recipe +// lines. A no-op when `info.declared` is false. +void MiniMaxH3CheckTaskPartition(const std::string& task, const MiniMaxH3PartitionInfo& info); + // --------------------------------------------------------------------------- // Flow-matching scheduler (scheduling_minimax_h3_euler_ancestral.py) // --------------------------------------------------------------------------- @@ -1380,6 +1446,15 @@ struct MiniMaxH3T2vaRequest { int64_t num_frames = 0; int64_t num_steps = kMiniMaxH3DefaultSteps; + // The served checkpoint partition, for the task/partition guard mirrored from + // upstream `_resolve_task` (pipeline_minimax_h3.py:374-391). MiniMaxH3GenerateT2va + // refuses a task the partition does not serve (t2va/fl2va need FL2VA, ref2va needs + // Ref2VA). Default-constructed (`declared=false`) leaves the guard INACTIVE, so a + // direct caller building a request by hand is unaffected; the checkpoint-loading + // entry points (driver/server) fill it via MiniMaxH3PartitionFromFlag / + // MiniMaxH3PartitionFromModelIndex. See the #77 follow-up. + MiniMaxH3PartitionInfo partition; + // --- fl2va KEYFRAME CONDITIONING (empty => plain t2va) --- // Which generated frames the supplied keyframes pin. Upstream allows exactly // {}, {0}, {-1} or {0, -1}: first frame, last frame, or both. diff --git a/src/vllm/model_executor/models/minimax_h3_pipeline.cpp b/src/vllm/model_executor/models/minimax_h3_pipeline.cpp index 9e6372db4..e60b9d829 100644 --- a/src/vllm/model_executor/models/minimax_h3_pipeline.cpp +++ b/src/vllm/model_executor/models/minimax_h3_pipeline.cpp @@ -362,6 +362,15 @@ MiniMaxH3T2vaResult MiniMaxH3GenerateT2va(vt::Device device, const MiniMaxH3T2va const std::vector& initial_audio_rows, vt::DType compute_dtype, const MiniMaxH3DitDeviceWeights* prestaged) { + // The task/partition guard — the raise half of `_resolve_task` + // (pipeline_minimax_h3.py:374-391), which the #70/#74 white render bypassed by + // running t2va on the Ref2VA NVFP4 checkpoint. The task is what the request + // ENCODES (ref_blocks => ref2va, keyframes => fl2va, else t2va); a checkpoint that + // declared its partition (driver --partition / server model_index.json) refuses a + // task it does not serve. A request that never set `partition` (declared=false) + // leaves the guard inactive, so the pipeline-math tests are unaffected. + MiniMaxH3CheckTaskPartition(MiniMaxH3TaskOfRequest(request), request.partition); + const MiniMaxH3DenoiseResult denoised = MiniMaxH3DenoiseT2va(device, request, dit_params, dit_weights, prompt_embeds, initial_video_rows, initial_audio_rows, compute_dtype, prestaged); diff --git a/src/vllm/model_executor/models/minimax_h3_planner.cpp b/src/vllm/model_executor/models/minimax_h3_planner.cpp index 7afc353c5..71475ab0d 100644 --- a/src/vllm/model_executor/models/minimax_h3_planner.cpp +++ b/src/vllm/model_executor/models/minimax_h3_planner.cpp @@ -152,6 +152,115 @@ std::string MiniMaxH3ResolveTask(const std::string& requested, const std::string return task; } +// --------------------------------------------------------------------------- +// Checkpoint PARTITION / task guard (the #77 follow-up) +// --------------------------------------------------------------------------- +namespace { + +// "{a, b, c}" for an error message. +std::string JoinBraced(const std::vector& items) { + std::string out = "{"; + for (size_t i = 0; i < items.size(); ++i) { + if (i != 0) out += ", "; + out += items[i]; + } + out += "}"; + return out; +} + +} // namespace + +// pipeline_minimax_h3.py:279-282 — +// release = model_index.get("_minimax_h3") or {} +// self.partition = str(release.get("partition", "")) +// self.supported_tasks = frozenset(release.get("tasks") or ()) +MiniMaxH3PartitionInfo MiniMaxH3PartitionFromModelIndex(const nlohmann::json& model_index) { + MiniMaxH3PartitionInfo info; + info.declared = true; // a real model_index.json was read; absence of the block + // below leaves partition unknown (and the guard refuses). + if (!model_index.is_object()) return info; + const auto release_it = model_index.find("_minimax_h3"); + if (release_it == model_index.end() || !release_it->is_object()) return info; + const nlohmann::json& release = *release_it; + const auto part_it = release.find("partition"); + if (part_it != release.end() && part_it->is_string()) { + info.partition = ToLower(part_it->get()); + } + const auto tasks_it = release.find("tasks"); + if (tasks_it != release.end() && tasks_it->is_array()) { + for (const nlohmann::json& t : *tasks_it) { + if (t.is_string()) info.supported_tasks.push_back(ToLower(t.get())); + } + } + return info; +} + +// The recipe's one-server-one-partition split (recipes/MiniMaxAI/MiniMax-H3.md:50-51): +// FL2VA serves t2va + fl2va, Ref2VA serves ref2va. Mirrors the served-task set the +// release model_index.json would have carried, for the community files that strip it. +MiniMaxH3PartitionInfo MiniMaxH3PartitionFromFlag(const std::string& partition) { + MiniMaxH3PartitionInfo info; + info.declared = true; + info.partition = ToLower(partition); + if (info.partition.empty()) { + return info; // declared-but-unknown: the guard refuses every task. + } + if (info.partition == "fl2va") { + info.supported_tasks = {"t2va", "fl2va"}; + } else if (info.partition == "ref2va") { + info.supported_tasks = {"ref2va"}; + } else { + VT_CHECK(false, + "minimax_h3: unknown --partition '" + partition + + "' — expected fl2va or ref2va (recipes/MiniMaxAI/MiniMax-H3.md:50-51). " + "FL2VA serves t2va+fl2va, Ref2VA serves ref2va."); + } + return info; +} + +// The request fields are mutually exclusive (packed layout enforces it via +// BuildMiniMaxH3PackedSequenceRef2va vs the keyframe path); ref2va prepends whole +// reference BLOCKS, fl2va pins keyframe FRAMES of the output. +std::string MiniMaxH3TaskOfRequest(const MiniMaxH3T2vaRequest& request) { + if (!request.ref_blocks.empty()) return "ref2va"; + if (!request.keyframe_frame_indices.empty()) return "fl2va"; + return "t2va"; +} + +// The raise half of `_resolve_task` (pipeline_minimax_h3.py:387-390): +// if task not in self.supported_tasks: +// raise ValueError(f"checkpoint partition {self.partition!r} supports " +// f"{sorted(self.supported_tasks)}, got task={task!r}") +// Extended for the stripped community files that carry no served-task set: an +// unknown partition makes EVERY task ambiguous (the FL2VA/Ref2VA DiTs are +// indistinguishable), so it refuses and names the recipe lines instead of guessing. +void MiniMaxH3CheckTaskPartition(const std::string& task, const MiniMaxH3PartitionInfo& info) { + if (!info.declared) return; // no checkpoint metadata supplied → guard inactive. + const std::string t = ToLower(task); + if (info.supported_tasks.empty()) { + VT_CHECK(false, + "minimax_h3: cannot run task '" + t + + "' — this checkpoint does not declare its partition. Community " + "GGUF/NVFP4 files strip model_index.json `_minimax_h3`, and the " + "FL2VA and Ref2VA DiTs are byte-structurally identical (same tensor " + "names AND shapes), so the partition cannot be inferred from the " + "weights. Pass --partition fl2va|ref2va. One server serves one " + "partition: FL2VA serves t2va+fl2va, Ref2VA serves ref2va " + "(recipes/MiniMaxAI/MiniMax-H3.md:50-51,289; " + "pipeline_minimax_h3.py:374-391)."); + } + const bool served = + std::find(info.supported_tasks.begin(), info.supported_tasks.end(), t) != + info.supported_tasks.end(); + VT_CHECK(served, + "minimax_h3: checkpoint partition '" + info.partition + "' supports " + + JoinBraced(info.supported_tasks) + ", got task='" + t + + "'. t2va/fl2va require the FL2VA partition, ref2va requires Ref2VA; " + "one server serves one partition " + "(recipes/MiniMaxAI/MiniMax-H3.md:50-51,289; " + "pipeline_minimax_h3.py:387-390)."); +} + // _resolve_shape (pipeline_minimax_h3.py:393-434). MiniMaxH3ShapePlan MiniMaxH3ResolveShape(const std::string& task, double duration_seconds, int64_t requested_frames, int64_t height, int64_t width, diff --git a/tests/vllm/models/test_minimax_h3.cpp b/tests/vllm/models/test_minimax_h3.cpp index ae995e51e..31eebebb9 100644 --- a/tests/vllm/models/test_minimax_h3.cpp +++ b/tests/vllm/models/test_minimax_h3.cpp @@ -1299,6 +1299,118 @@ TEST_CASE("minimax_h3: request planning matches upstream") { CHECK_THROWS(vllm::MiniMaxH3ResolveTask("ref2va", "FL2VA", false, fl2va_partition)); } +// The #77 follow-up: the task/partition GUARD. The #70/#74 white render cost three +// campaigns because the driver silently accepted t2va on the Ref2VA NVFP4 +// checkpoint — a task/partition mismatch upstream `_resolve_task` RAISES on +// (pipeline_minimax_h3.py:374-391; recipes/MiniMaxAI/MiniMax-H3.md:50-51,289). This +// gates the mirror: the mismatch now FAILS LOUDLY, the correct pairings pass, and +// the community-file (stripped-config) path refuses ambiguity rather than guessing. +TEST_CASE("minimax_h3: the task/partition guard refuses the #77 mismatch") { + using vllm::MiniMaxH3CheckTaskPartition; + using vllm::MiniMaxH3PartitionFromFlag; + using vllm::MiniMaxH3PartitionFromModelIndex; + + // --- Partition detection path 1: the release model_index.json `_minimax_h3` + // block (pipeline_minimax_h3.py:279-282). Synthetic manifests of BOTH shapes. + const nlohmann::json fl2va_index = { + {"_minimax_h3", {{"partition", "fl2va"}, {"tasks", {"t2va", "fl2va"}}}}}; + const nlohmann::json ref2va_index = { + {"_minimax_h3", {{"partition", "ref2va"}, {"tasks", {"ref2va"}}}}}; + const vllm::MiniMaxH3PartitionInfo fl2va = MiniMaxH3PartitionFromModelIndex(fl2va_index); + const vllm::MiniMaxH3PartitionInfo ref2va = MiniMaxH3PartitionFromModelIndex(ref2va_index); + CHECK(fl2va.declared); + CHECK(fl2va.partition == "fl2va"); + CHECK(fl2va.supported_tasks == std::vector{"t2va", "fl2va"}); + CHECK(ref2va.partition == "ref2va"); + CHECK(ref2va.supported_tasks == std::vector{"ref2va"}); + + // The GUARD BEHAVIOR TABLE (task x partition -> pass/refuse). The whole point of + // the row: the top-left cell (t2va on Ref2VA) is the #77 failure mode. + CHECK_THROWS(MiniMaxH3CheckTaskPartition("t2va", ref2va)); // <-- #77: FAIL LOUDLY + CHECK_THROWS(MiniMaxH3CheckTaskPartition("fl2va", ref2va)); // fl2va needs FL2VA + CHECK_NOTHROW(MiniMaxH3CheckTaskPartition("ref2va", ref2va)); + CHECK_NOTHROW(MiniMaxH3CheckTaskPartition("t2va", fl2va)); + CHECK_NOTHROW(MiniMaxH3CheckTaskPartition("fl2va", fl2va)); + CHECK_THROWS(MiniMaxH3CheckTaskPartition("ref2va", fl2va)); // ref2va needs Ref2VA + // Case-insensitive, mirroring `str(requested).lower()` (pipeline:386). + CHECK_NOTHROW(MiniMaxH3CheckTaskPartition("T2VA", fl2va)); + CHECK_THROWS(MiniMaxH3CheckTaskPartition("T2VA", ref2va)); + + // --- Partition detection path 2: an explicit --partition flag (community + // GGUF/NVFP4 strip the config). The recipe's one-server-one-partition split. + const vllm::MiniMaxH3PartitionInfo fl2va_flag = MiniMaxH3PartitionFromFlag("fl2va"); + const vllm::MiniMaxH3PartitionInfo ref2va_flag = MiniMaxH3PartitionFromFlag("ref2va"); + CHECK(fl2va_flag.supported_tasks == std::vector{"t2va", "fl2va"}); + CHECK(ref2va_flag.supported_tasks == std::vector{"ref2va"}); + CHECK_THROWS(MiniMaxH3CheckTaskPartition("t2va", ref2va_flag)); // same #77 catch + CHECK_NOTHROW(MiniMaxH3CheckTaskPartition("fl2va", fl2va_flag)); + CHECK_THROWS(MiniMaxH3PartitionFromFlag("bogus")); // invalid partition name + + // --- The ambiguous / stripped case: NO config block AND NO --partition. The + // partition is unknown, and (see the structural check below) unknowable from the + // weights, so the guard REFUSES every task rather than guessing. + const vllm::MiniMaxH3PartitionInfo stripped = MiniMaxH3PartitionFromModelIndex(nlohmann::json::object()); + CHECK(stripped.declared); + CHECK(stripped.partition.empty()); + CHECK(stripped.supported_tasks.empty()); + CHECK_THROWS(MiniMaxH3CheckTaskPartition("t2va", stripped)); + CHECK_THROWS(MiniMaxH3CheckTaskPartition("fl2va", stripped)); + CHECK_THROWS(MiniMaxH3CheckTaskPartition("ref2va", stripped)); + const vllm::MiniMaxH3PartitionInfo empty_flag = MiniMaxH3PartitionFromFlag(""); + CHECK(empty_flag.declared); + CHECK(empty_flag.supported_tasks.empty()); + CHECK_THROWS(MiniMaxH3CheckTaskPartition("t2va", empty_flag)); // --partition required + + // --- A DEFAULT-CONSTRUCTED (declared=false) request leaves the guard INACTIVE: + // the pure pipeline-math tests build requests by hand and must be unaffected. + const vllm::MiniMaxH3PartitionInfo undeclared; + CHECK_FALSE(undeclared.declared); + CHECK_NOTHROW(MiniMaxH3CheckTaskPartition("t2va", undeclared)); + CHECK_NOTHROW(MiniMaxH3CheckTaskPartition("ref2va", undeclared)); + + // --- The dispatch key: MiniMaxH3TaskOfRequest, the same discriminator + // MiniMaxH3GenerateT2va feeds the guard. ref_blocks => ref2va, keyframes => fl2va, + // else t2va. + vllm::MiniMaxH3T2vaRequest t2va_req; // no conditioning + CHECK(vllm::MiniMaxH3TaskOfRequest(t2va_req) == "t2va"); + vllm::MiniMaxH3T2vaRequest fl2va_req; + fl2va_req.keyframe_frame_indices = {0}; + CHECK(vllm::MiniMaxH3TaskOfRequest(fl2va_req) == "fl2va"); + vllm::MiniMaxH3T2vaRequest ref2va_req; + ref2va_req.ref_blocks.push_back(vllm::MiniMaxH3RefBlock{}); + CHECK(vllm::MiniMaxH3TaskOfRequest(ref2va_req) == "ref2va"); + + // The exact #77 combination as the dispatch sees it: a t2va request pointed at a + // Ref2VA checkpoint. This is what MiniMaxH3GenerateT2va now refuses. + t2va_req.partition = ref2va; + CHECK_THROWS(MiniMaxH3CheckTaskPartition(vllm::MiniMaxH3TaskOfRequest(t2va_req), t2va_req.partition)); + ref2va_req.partition = ref2va; // the correct pairing passes + CHECK_NOTHROW(MiniMaxH3CheckTaskPartition(vllm::MiniMaxH3TaskOfRequest(ref2va_req), ref2va_req.partition)); + + // --- WHY the stripped case must refuse rather than auto-detect: the two REAL + // captured manifests (ref2va NVFP4 = 1051 tensors, FL2VA GGUF = 535) carry the + // IDENTICAL DiT — same base tensor names AND shapes. NVFP4 splits each quantized + // weight into {weight, weight_scale, weight_scale_2}; normalizing that away, the + // two name sets are equal, so no tensor-name/shape discriminator exists. + auto normalize_nvfp4 = [](std::string n) -> std::string { + for (const char* suffix : {".weight_scale_2", ".weight_scale"}) { + const std::string s = suffix; + if (n.size() >= s.size() && n.compare(n.size() - s.size(), s.size(), s) == 0) { + return n.substr(0, n.size() - s.size()) + ".weight"; + } + } + return n; + }; + std::set nvfp4_names, gguf_names; + for (const vllm_test::H3Nvfp4Tensor& t : vllm_test::kH3Nvfp4Tensors) { + nvfp4_names.insert(normalize_nvfp4(t.name)); + } + for (const vllm_test::H3GgufTensor& t : vllm_test::kH3GgufTensors) gguf_names.insert(t.name); + CHECK(nvfp4_names.size() == 535); + CHECK(gguf_names.size() == 535); + CHECK(nvfp4_names == gguf_names); // no structural discriminator => --partition is required +} + TEST_CASE("minimax_h3: the denoise loop advances targets and pins condition rows") { // The loop itself has no upstream golden (upstream's own loop test needs the // checkpoint), so this gates its INVARIANTS, which are what the CFG-distilled