diff --git a/.agents/NOW.md b/.agents/NOW.md index 746fb989..9c4b6421 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 + task/partition guard CLOSED** (#77 render; `H3-TASK-PARTITION-GUARD`: t2va-on-ref2va now RAISES like `_resolve_task`) | Follow-up: vision tower (W3) | +| MiniMax-H3 lane | **vision tower RUNS on real weights + fl2va COHERENT** (`H3-CONDITIONED-E2E` PR#86: `visual.*` loader+probe; fl2va matched cat; ref2va grids) | scatter feats→embeds + DeepStack→text tower | | 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 bc6b641d..4a53c825 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -14239,3 +14239,74 @@ GREEN. Suite **67/67** (66 prior + this), 46549 assertions; `test_video_api` 4/4 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. + +## MiniMax-H3 ENCODER VISION TOWER — record reconciled + real-weights loader; the encoder GGUF DOES carry the vision tensors (2026-08-07, `row/H3-CONDITIONED-E2E`, `ROAD-V1-H3`, dgx GB10 sm_121a) + +**Record reconciliation.** #26/W3 recorded the vision tower "COMPLETE … only the MM +processor remains"; #77 recorded it "still unported". Reading the code, both are true of +different halves: the tower MATH exists as a CPU scalar reference +(`minimax_h3_encoder.cpp` `MiniMaxH3VisionTowerForward:572` + surround) gated ONLY at +reduced dims with SYNTHETIC weights (`test_minimax_h3.cpp:4041`), but it was NEVER wired to +real weights — `LoadMiniMaxH3EncoderFromGguf` (`minimax_h3_encoder_gguf.cpp:47`) loads the +TEXT tower only and skips every `visual.*` tensor; the device forward is text-only; the +driver/server call only the text path. So: math CPU-gated synthetic, zero real-weights +wiring. + +**Encoder-arm decision (measured, no download).** The on-box encoder +`~/h3fp4/ckpt/qwen3vl-32B-MiniMax-H3-Q4_K_M.gguf` (14 GiB) carries the FULL vision tower — +`visual.blocks.{0..26}` (Q4_K/Q5_K), `patch_embed.proj` (F16 `[16,16,6,1152]`), `pos_embed` +(F16 `[2304,1152]`=48²), `merger.*`, and 3 `deepstack_merger_list.*`. Geometry: hidden 1152 +/ 16 heads / depth 27 / intermediate 4304 / out 5120 / patch 16 / temporal 2 / merge 2 / +pos 2304 (== the Qwen3.6-27B vision config, state.md :23310). ComfyUI quant reshapes +non-256-aligned rows (in=1152) to ne0=256; dequant of the flat buffer preserves the +row-major `[out,in]` order the tower reads. `deepstack_visual_indexes` is NOT in the +weights-only GGUF and must be supplied from the H3 vision_config (residual for a bit-correct +DeepStack inject). + +**Landing this row:** GGUF `visual.*`→`Qwen3VLVisionWeights` loader (`LoadQwen3VLVisionFromGguf`, +mirrors safetensors `LoadQwen3VLVisionWeights` `qwen3_vl.cpp:417`) + a CPU/real-weights +vision-tower forward gate + driver `--prompt-image` probe. HONEST residual (recorded, not +claimed done): the vision-ENRICHED DiT render (DeepStack scatter into the DEVICE text tower +actually changing the frames) needs the exact deepstack indexes + a device-text +DeepStack/merge extension + an on-box GPU render; fl2va/ref2va frame-sanity verdicts recorded +here as they are produced. + +### GB10 VERIFICATION (2026-08-07, dgx sm_121a, flock, worker parked) + +**Real-weights vision-tower probe RAN (deliverable 1).** `minimax-h3-gen --device cuda +--encoder qwen3vl-32B-MiniMax-H3-Q4_K_M.gguf --prompt-image <512x512 cat frame>`: the loader +read the real `visual.*` tower (depth 27 / hidden 1152 / 16 heads / out 5120 / 3 DeepStack +mergers); the stock Qwen3VLImageProcessor produced grid_thw [1,32,32] = 1024 tokens → 256 +merged; `Qwen3VLVisionForward` on the GGUF-loaded weights returned [256, 20480] = merged + +3 deepstack concatenated, ALL FINITE and non-degenerate: merged rms 1.45 / maxabs 29.1 (the +expected Qwen MASSIVE-ACTIVATION signature), deepstack rms 0.54/1.04/0.51. The vision tower now +runs on real weights. + +**fl2va e2e — COHERENT (deliverable 2, frame-sanity PASS).** FL2VA-partition GGUF +(`MiniMax-H3-FL2VA-Q3_K_M.gguf`, `--dequant-bf16`, ~66 GB device, free 117→23 G, no OOM) + +`--first-frame` (real coherent orange-cat frame, VAE-keyframe conditioning) + `--partition +fl2va`, 512x512/22f/12steps. Valid h264 512x512 + AAC mp4, 22 frames. VISUAL sanity (frames +0/6/12/21 inspected): every frame is a coherent photorealistic ORANGE CAT sitting on a WOODEN +TABLE, matching the conditioning first-frame's pose/table/gray-background, consistent across +the whole clip — no patch-grid, no white latent. The render visibly CONTINUES/MATCHES the +conditioning image. (This uses the VAE-KEYFRAME path, not the vision tower.) + +**ref2va e2e — STILL GRIDS (deliverable 2, frame-sanity FAIL, recorded honestly).** Ref2VA +NVFP4 (`minimax_h3_ref2va_nvfp4_full`, `--fp4-resident`, ~16 GB, free 117→63 G) + `--ref-image` +(one 256x256 real cat image, VAE-reference conditioning) + `--partition ref2va` + text +prompt_embeds, 256x256/22f/12steps. Valid mp4, 22 frames, but EVERY frame (0/10/21) is a +multicolour PATCH GRID — the #70/#74 degenerate render — NOT a coherent scene. First attempt +(5x512 `--ref-video`) was ABANDONED: the reference-video VAE encode is a single-thread CPU +3D-CNN scalar path (GPU 0%, 98% CPU) and did not finish in ~15 min — a known perf limit, not a +bug. **HONEST verdict vs the mission's expectation:** landing the vision-tower LOADER does NOT +fix the ref2va grid, because the tower is wired as a loader+probe and is NOT yet scattered into +the DiT render-conditioning path — this render did not use it. fl2va (same partition family, +same session, VAE-keyframe path) is COHERENT, so the DiT/VAE/partition are sound; the ref2va +grid is specific to the ref2va conditioning assembly (VAE-reference rows + un-enriched text +prompt). Testing whether the vision-enriched prompt fixes ref2va needs the render-conditioning +scatter (merge merged-features into prompt_embeds + DeepStack inject into the DEVICE text +tower) — the tracked residual. Suite note: the full `test_minimax_h3` run SIGSEGVs at an +UNRELATED CUDA case (line 3503, "an NVFP4 checkpoint loads into a runnable DiT") that PASSES in +isolation (585 assertions) and runs BEFORE the new case — a pre-existing cross-test CUDA +resource-accumulation flake, not this change; the new loader gate passes standalone (59 +assertions). diff --git a/.agents/specs/minimax-h3.md b/.agents/specs/minimax-h3.md index ad056ae2..397c245b 100644 --- a/.agents/specs/minimax-h3.md +++ b/.agents/specs/minimax-h3.md @@ -400,7 +400,7 @@ vLLM-Omni H3 modules at `vllm_omni/diffusion/models/minimax_h3/`; serving in | WebSocket `/v1/video/chat/stream`, `/v1/realtime/video` | `api_server.py:1593,1610` | — | **MISSING** (streaming/realtime) | | Request schema (prompt, size/w/h, num_frames, fps, seed, steps, refs) | `protocol/videos.py:97-249` | request contract (W7) | **PARTIAL** (core fields; frame-interp/lora/generate_sound absent) | | H3 knobs via `extra_params.{task,duration,flow_shift,audio_flow_shift}` | `pipeline:1034,403,1157-1158` | planner reads task/duration/shift | **DONE** | -| Modalities in: text/image/video/audio | `pipeline:1036-1104` | t2va (text) done; fl2va/ref2va image/video/audio WIRED (W6/ref2va) | **PARTIAL** (encoder vision tower still open) | +| Modalities in: text/image/video/audio | `pipeline:1036-1104` | t2va (text) done; fl2va COHERENT on real weights (VAE-keyframe); vision tower now LOADS real `visual.*` + runs (probe); ref2va still grids | **PARTIAL** (vision-tower→DiT-conditioning scatter is the residual; §8.8) | | Output: joint video+audio, 24 fps, 32 kHz stereo | `pipeline:106-111,1187` | frames + WAV + MP4 mux (W7) | **DONE** | | Scheduler: euler-ancestral rectified flow (single) | `scheduling_...euler_ancestral.py`; `time_request.py:34-61` | `MiniMaxH3EulerEta0Step` / `MiniMaxH3TimeShiftSigmas` | **DONE** | | CFG: distilled, no CFG (guidance params accepted+ignored; `cfg_parallel_size==1`) | `pipeline:250,275-276` | no CFG branch | **DONE** (matches) | @@ -624,3 +624,88 @@ 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). + +## 8.8 ENCODER VISION TOWER — record reconciliation + real-weights wiring (2026-08-06, `row/H3-CONDITIONED-E2E`) + +**The contradictory record, reconciled (file:line).** Two prior lanes disagreed. The +#26/W3 lane recorded the vision tower as **"W3 COMPLETE … the FULL vision tower at +6.0e-8 … only the MM processor remains"** (this spec lines 101, 162-163); the #77 +residual recorded **"the encoder vision tower (W3) is still unported"** (lines 565-566). +Reading the actual code resolves it — **both describe different halves and both are +literally true of what they describe**: + +- The vision-tower **MATH exists** as a CPU scalar f32 reference in + `minimax_h3_encoder.cpp`: `MiniMaxH3VisionBlockForward` (:311), the surround + `MiniMaxH3VisionPosEmbedInterpolate` (:430) / `MiniMaxH3VisionRotary` (:500) / + `PatchMerger` (:545) / `MiniMaxH3VisionTowerForward` (:572). It is gated ONLY in + `tests/vllm/models/test_minimax_h3.cpp` (:3942, :4041) at **reduced dims with SYNTHETIC + weights** (`MakeParam`), block 6.0e-8 / tower ≤1e-4 vs a self-restated oracle. +- It is **NEVER wired to real weights.** `LoadMiniMaxH3EncoderFromGguf` + (`minimax_h3_encoder_gguf.cpp:47`) loads the **TEXT tower only** — it iterates + `model.layers.N.*` + `model.embed_tokens.weight` and **skips every `visual.*` tensor** + (the comment at :51-52 even names `visual.*` as present-but-unloaded). The device + encoder `MiniMaxH3EncoderTextForwardDevice` (`minimax_h3_encoder_device.cpp:103`) runs + text only and takes **no deepstack / no visual-mask** argument (the HOST reference + `MiniMaxH3EncoderTextForward` does, :113-118). The driver + (`examples/minimax_h3_gen/main.cpp:476-547`) and server + (`examples/server/main.cpp:659-716`) call only the text path. +- **So the reconciled truth:** the tower math is CPU-gated at reduced dims with synthetic + weights; there is **zero real-weights wiring** — no GGUF `visual.*` loader, no image→patch + MM processor on the H3 path, no device vision forward, no merge/DeepStack injection into + the encoded prompt. The #26 "only the MM processor remains" understated the gap (loader, + real-weights forward, and the merge/inject were ALSO absent); the #77 "still unported" was + right in the sense that matters (nothing real ran through it). + +**The encoder ARM already carries the vision weights (no download).** The on-box encoder +`~/h3fp4/ckpt/qwen3vl-32B-MiniMax-H3-Q4_K_M.gguf` (14 GiB, the ComfyUI-format text-tower +GGUF that already serves text conditioning) **DOES carry the full vision tower**: measured +`visual.blocks.{0..26}` (27, Q4_K/Q5_K), `visual.patch_embed.proj` (F16 `[16,16,6,1152]` = +Conv3d as a linear over `patch_elems`=1536), `visual.pos_embed.weight` (F16 `[2304,1152]` = +48² grid), `visual.merger.*`, and **`visual.deepstack_merger_list.{0,1,2}`** (3 DeepStack +mergers). Names map 1:1 to `MiniMaxH3VisionTowerForward` / the reuse target +`multimodal::Qwen3VLVisionWeights`. So the encoder-arm decision is settled: **reuse the +in-place encoder GGUF; no new download** (disk floor 15 GiB / ~23 GiB free honoured). + +**Vision geometry (from the checkpoint + state.md :23310-23318, the Qwen3.6-27B vision +config which shares this tower):** hidden **1152**, **16 heads** (head_dim 72), depth **27**, +intermediate **4304**, out_hidden **5120** (== encoder text dim), patch **16**, temporal **2**, +merge **2**, num_position_embeddings **2304**, gelu-tanh blocks / exact-erf merger. H3 differs +from the 27B only by having **3 real DeepStack mergers** (the 27B's are empty). The one +config value NOT recoverable from the ComfyUI GGUF (weights-only, no arch metadata) is +`deepstack_visual_indexes` — the WHICH-layers taps — needed for a bit-correct DeepStack +inject; it is inferred + flagged as the residual for a fully-correct conditioned render. + +**The reuse path (mission: "stock Qwen3VLProcessor + our existing front end").** The image +MM processor already exists and is gated: `multimodal::Qwen3VLImageProcessor::ProcessImage` +(`qwen3vl_processor.h`, patch 16 / temporal 2 / merge 2 / 0.5 normalize → pixel_values + +grid_thw), `ExpandImagePlaceholders`, and the device tower +`multimodal::Qwen3VLVisionForward` (`qwen3_vl_vision.cpp`) with `PrepareVisionDeviceWeights`. +The only genuinely-new code is the **GGUF `visual.*` → `Qwen3VLVisionWeights` loader** +(`LoadQwen3VLVisionFromGguf`), mirroring the safetensors `LoadQwen3VLVisionWeights` +(`qwen3_vl.cpp:417`) but dequantizing the Q4_K/Q5_K blocks (ComfyUI reshapes non-256-aligned +rows to ne0=256; dequant preserves the flat row-major order the tower reads as `[out,in]`) +and converting the F16 patch/pos tensors. + +**This row's status (honest):** loader + real-image processor reuse + the real-weights +vision-tower forward gate LAND here (see §8.4-style status in STATUS/BENCHMARKS). The full +vision-ENRICHED DiT render (DeepStack scatter into the DEVICE text tower changing the frames) +depends additionally on the exact `deepstack_visual_indexes` and a device-text DeepStack/merge +extension; its e2e render verdict is recorded honestly in the benchmark record. + +**GB10 VERIFIED (2026-08-07, dgx sm_121a).** +- **Vision-tower probe RAN on real weights:** `--prompt-image` loaded the real `visual.*` + tower (27 blocks / 3 DeepStack mergers), processed a 512×512 image → grid [1,32,32], and + `Qwen3VLVisionForward` returned [256, 20480] all FINITE + non-degenerate (merged rms 1.45 / + maxabs 29.1 — the expected Qwen massive-activation). Deliverable-1 core DONE. +- **fl2va e2e COHERENT:** FL2VA GGUF (`--dequant-bf16`) + a real first-frame (VAE-keyframe) + + `--partition fl2va`, 512×512/22f/12steps → all 22 frames a coherent photorealistic orange + cat on a wooden table matching the conditioning frame (no grid). Frame-sanity PASS. +- **ref2va STILL GRIDS (honest):** Ref2VA NVFP4 (`--fp4-resident`) + a real `--ref-image` + + `--partition ref2va` → every frame a multicolour patch grid. Landing the tower LOADER does + NOT fix it: the tower is a loader+probe, NOT yet scattered into the DiT render-conditioning, + so this render never used it. fl2va (same session, VAE-keyframe) is coherent ⇒ DiT/VAE/ + partition are sound; the ref2va grid is specific to the ref2va conditioning assembly. The + render-conditioning scatter (merge features into prompt_embeds + DeepStack inject into the + DEVICE text tower) is the tracked residual that would let the vision-enriched-prompt + hypothesis be tested. The `--ref-video` VAE encode is a slow single-thread CPU 3D-CNN path + (separate perf limit). diff --git a/.agents/state.md b/.agents/state.md index 65c15e73..a7aacc9e 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -39869,3 +39869,87 @@ shipped; CMakeLists NOTE + benchmark-record (#82) + spec CLOSED capture the clos levers. Box: experiment left artifacts under `~/mxfp4-ptxas/`; box OOM-rebooted mid-confirmatory-run (another agent's 25 GiB vLLM on the 119 GiB unified pool, not this <1 GiB microbench) - run-1 data is on disk and decisive. +## 2026-08-07T01:45 - MiniMax-H3 encoder VISION TOWER — record reconciled + real-weights loader wired (row/H3-CONDITIONED-E2E, helper, DRAFT PR) + + +**Record reconciliation (the mission's first ask — do not re-port what exists).** Two prior +lanes contradicted: #26/W3 recorded the vision tower "W3 COMPLETE … FULL vision tower at +6.0e-8 … only the MM processor remains" (spec :101,:162-163); #77 recorded "encoder vision +tower still unported" (:565-566). Read the actual code — BOTH are true of DIFFERENT halves: +- The tower MATH exists as a CPU scalar f32 reference: `MiniMaxH3VisionBlockForward` + (`minimax_h3_encoder.cpp:311`), `MiniMaxH3VisionPosEmbedInterpolate` (:430), + `MiniMaxH3VisionRotary` (:500), `PatchMerger` (:545), `MiniMaxH3VisionTowerForward` (:572). + Gated ONLY in `test_minimax_h3.cpp` (:3942/:4041) at REDUCED DIMS with SYNTHETIC weights. +- It is NEVER wired to real weights: `LoadMiniMaxH3EncoderFromGguf` + (`minimax_h3_encoder_gguf.cpp:47`) loads the TEXT tower only (`model.layers.*`+embed) and + SKIPS every `visual.*` tensor; `MiniMaxH3EncoderTextForwardDevice` + (`minimax_h3_encoder_device.cpp:103`) is text-only with NO deepstack arg; the driver + (`examples/minimax_h3_gen/main.cpp:476-547`) + server run text only. So: math CPU-gated + synthetic, ZERO real-weights wiring (no `visual.*` loader, no image→patch path on H3, no + device vision forward, no merge/DeepStack inject). #26 understated (loader/forward/inject + all absent), #77 was right in what matters. `MiniMaxH3VisionTowerForward` reads every weight + as a FLAT f32 buffer keyed by name (dims from config), so a loader need only produce correct + flat dequant. + +**Encoder-arm decision — reuse in place, NO download (disk floor 15G / ~23G free honoured).** +The on-box encoder `~/h3fp4/ckpt/qwen3vl-32B-MiniMax-H3-Q4_K_M.gguf` (14 GiB, ComfyUI-format, +already serving TEXT conditioning) DOES carry the full vision tower — MEASURED via a minimal +GGUF parser: `visual.blocks.{0..26}` (27; Q4_K/Q5_K), `visual.patch_embed.proj.weight` (F16 +`[16,16,6,1152]`), `visual.pos_embed.weight` (F16 `[2304,1152]`=48²), `visual.merger.*`, and +`visual.deepstack_merger_list.{0,1,2}` (3 real DeepStack mergers). Names map 1:1 to +`Qwen3VLVisionWeights`. Geometry (checkpoint + state.md :23310-23318, the Qwen3.6-27B vision +config that shares this tower): hidden 1152 / 16 heads (hd 72) / depth 27 / intermediate 4304 / +out_hidden 5120 / patch 16 / temporal 2 / merge 2 / pos 2304. The ONLY value not in the +weights-only GGUF is `deepstack_visual_indexes` (which text layers the DeepStack taps inject +into) — inferred, flagged as the residual for a bit-correct conditioned render. + +**Reuse path (mission-mandated).** The image processor already exists + is gated: +`multimodal::Qwen3VLImageProcessor::ProcessImage` (patch16/temporal2/merge2/0.5-norm → +pixel_values+grid_thw), `ExpandImagePlaceholders`; the device tower +`multimodal::Qwen3VLVisionForward` + `PrepareVisionDeviceWeights` (`qwen3_vl_vision.cpp`). +Only NEW code is the GGUF `visual.*`→`Qwen3VLVisionWeights` loader, mirroring the safetensors +`LoadQwen3VLVisionWeights` (`qwen3_vl.cpp:417`) but dequantizing Q4_K/Q5_K (ComfyUI reshapes +non-256-aligned rows to ne0=256; dequant preserves the flat row-major order the tower reads +as `[out,in]`) + converting F16 patch/pos. + +**IN FLIGHT this row:** the GGUF vision loader + a CPU/real-weights vision-tower gate + driver +`--prompt-image` probe. HONEST residual: the full vision-ENRICHED DiT render (DeepStack scatter +into the DEVICE text tower changing the frames) needs the exact deepstack indexes + a +device-text DeepStack/merge extension + an on-box GPU render; recorded in the benchmark record. +Records: spec §8.8, STATUS/BENCHMARKS/FEATURES H3 rows, benchmark-record, NOW. Pre-existing +preflight red (check-fusion-consistency `minimax_h3_video_vae_device`) is not this row. +## 2026-08-07T03:20 - MiniMax-H3 vision tower GB10-VERIFIED on real weights; fl2va COHERENT, ref2va STILL GRIDS (row/H3-CONDITIONED-E2E, helper, PR #86) + + +Landed on dgx (git-archive of HEAD over `~/h3fp4/src`, golden md5 unchanged, incremental CUDA +build, flock, worker parked). New loader `LoadQwen3VLVisionFromGguf` + CPU gate + driver +`--prompt-image` probe. + +**Deliverable 1 (vision tower → real weights) DONE + VERIFIED.** CPU gate `test_minimax_h3 :: +"the encoder GGUF visual.* loader dequantizes the vision tower"` passes standalone (59 +assertions). Real-weights probe on dgx: `--prompt-image` loaded the real `visual.*` tower (27 +blocks / hidden 1152 / 16 heads / out 5120 / 3 DeepStack mergers), stock Qwen3VLImageProcessor +→ grid [1,32,32]=1024 tokens→256 merged, `Qwen3VLVisionForward` → [256,20480] ALL FINITE + +non-degenerate (merged rms 1.45 / maxabs 29.1 = Qwen massive-activation; 3 deepstack finite). + +**Deliverable 2 renders (frame-sanity, visually inspected).** +- **fl2va COHERENT (PASS):** FL2VA GGUF `--dequant-bf16` (~66 GB, free 117→23 G, no OOM) + + real `--first-frame` (VAE-keyframe) + `--partition fl2va`, 512×512/22f/12steps. All 22 frames + a coherent photorealistic orange cat on a wooden table matching the conditioning frame — no + grid. (VAE-keyframe path, not the vision tower.) +- **ref2va STILL GRIDS (FAIL, honest vs the mission's expectation):** Ref2VA NVFP4 + `--fp4-resident` (~16 GB, free→63 G) + real `--ref-image` + `--partition ref2va`, + 256×256/22f → every frame a multicolour patch grid. Landing the tower LOADER does NOT fix it: + it is a loader+probe, NOT yet scattered into the DiT render-conditioning, so this render never + used it. fl2va (same session) is coherent ⇒ DiT/VAE/partition sound; the ref2va grid is + specific to the ref2va conditioning assembly. The `--ref-video` VAE encode is a slow + single-thread CPU 3D-CNN path (abandoned the 5×512 attempt at 98% CPU / 0% GPU). + +**RESIDUAL (the true deliverable-2 fix):** scatter the vision-tower merged features into +prompt_embeds + DeepStack inject into the DEVICE text tower (`MiniMaxH3EncoderTextForwardDevice` +has no deepstack arg yet) + confirm the exact `deepstack_visual_indexes` (inferred {8,16,24}) — +then re-test whether the vision-enriched prompt fixes ref2va. Suite note: full `test_minimax_h3` +SIGSEGVs at an UNRELATED CUDA case (line 3503) that PASSES in isolation (585) and runs before +the new case — pre-existing cross-test CUDA resource-accumulation flake, not this change. +Box left clean (no procs, GPU idle, locks released, worker parked, my artifacts pruned, ckpts +kept). Records: spec §8.8 + §8.2, STATUS/BENCHMARKS/FEATURES, benchmark-record, NOW. diff --git a/CMakeLists.txt b/CMakeLists.txt index 4b59576b..f89ac7f5 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -500,6 +500,7 @@ add_library(vllm STATIC src/vllm/model_executor/models/minimax_h3_vae_loader.cpp src/vllm/model_executor/models/minimax_h3_encoder_gguf.cpp src/vllm/model_executor/models/minimax_h3_encoder_device.cpp + src/vllm/model_executor/models/minimax_h3_vision_gguf.cpp src/vllm/entrypoints/openai/video_api.cpp src/vllm/model_executor/models/kimi_kda.cpp src/vllm/model_executor/models/kimi_linear_registry.cpp diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 52467c24..61e22f98 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -311,6 +311,7 @@ built on it rather than keeping the flattering one. | 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. 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`) | **Vision tower RUNS on REAL weights (GB10) + fl2va COHERENT.** GGUF `visual.*`→`Qwen3VLVisionWeights` loader + Qwen3VLProcessor reuse; tower forward finite. fl2va (VAE-keyframe) renders a matched cat; ref2va STILL grids | ref2va needs the vision→DiT-conditioning scatter (residual); frame-sanity + probe in benchmark record + spec §8.8 | | 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-PTXAS` #82: cuModule A/B ties our+vLLM PTX across ptxas 13.0/13.2/driver-JIT (~144us); +10us is engine CONTEXT not codegen, no ptxas lever/flip (retires #75) | | 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 | | Memory footprint vs declared workload (`ROAD-V1-MEM`, #83) | **Never measured, and not measurable today**: there is no auto-sizing to compare against, because the KV pool is a hand-typed `--num-blocks`, so "what the run actually needed" has no number | Once M1's `MemoryBudget` lands: predicted-vs-actual bytes per allocation class, then peak footprint ours-auto vs vLLM at its 0.9 default on the same model and config | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index a57d3c36..8a91694a 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 path 66/66 (DiT geometry ladder + CUDA-vs-host at the REAL render seq 1920); t2va renders a COHERENT prompt-matched scene on GB10 (render bug CLOSED: #70/#74 was wrong-partition usage, not a code bug) | FP4/Marlin routing landed, GB10 speed pending | +| MiniMax-H3 DiT (`MiniMaxH3DiTModel`, vllm-omni lane) | MiniMax-H3 (33.1B video+audio) | portable path 66/66 (DiT geometry ladder + CUDA-vs-host at seq 1920); t2va + fl2va render COHERENT prompt/frame-matched scenes on GB10; encoder vision tower now RUNS on real `visual.*` weights (probe); ref2va still grids | FP4/Marlin landed; ref2va conditioning + speed pending | | 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 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) | ☐ | ☐ | +| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va + fl2va render COHERENT scenes on GB10 (fl2va first-frame VAE-keyframe); guard mirrors `_resolve_task`; vision tower RUNS on real weights (probe, §8.8); ref2va still grids | ✅ (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 ae07bc8a..4bc94483 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 #include +#include +#include #include #include #include @@ -51,8 +53,10 @@ #include "vllm/model_executor/model_loader/gguf_reader.h" #include "vllm/model_executor/model_loader/safetensors_reader.h" #include "vllm/model_executor/models/minimax_h3.h" +#include "vllm/multimodal/qwen3vl_processor.h" #include "vllm/tokenizer/tokenizer.h" #include "vt/backend.h" +#include "vt/dtype.h" namespace { @@ -118,6 +122,31 @@ std::vector ReadPpmAsChw(const std::string& path, int64_t* out_h, int64_t return chw; } +// A binary PPM (P6) reader that returns HWC uint8 [0,255] -- the layout the shared +// Qwen3-VL image processor expects (it does its own rescale + 0.5/0.5 normalize). +std::vector ReadPpmAsHwcU8(const std::string& path, int64_t* out_h, int64_t* out_w) { + std::ifstream in(path, std::ios::binary); + if (!in) throw std::runtime_error("cannot open " + path); + std::string magic; + in >> magic; + if (magic != "P6") throw std::runtime_error(path + ": not a binary PPM (P6)"); + auto next_int = [&]() { + int v = 0; + while (in >> std::ws, in.peek() == '#') { std::string skip; std::getline(in, skip); } + in >> v; + return v; + }; + const int w = next_int(), h = next_int(), maxv = next_int(); + if (w <= 0 || h <= 0 || maxv <= 0) throw std::runtime_error(path + ": bad PPM header"); + in.get(); // the single whitespace byte before the payload + std::vector rgb(static_cast(w) * h * 3); + in.read(reinterpret_cast(rgb.data()), static_cast(rgb.size())); + if (!in) throw std::runtime_error(path + ": truncated PPM payload"); + if (out_h != nullptr) *out_h = h; + if (out_w != nullptr) *out_w = w; + return rgb; +} + // The WAV reader lives in the LIBRARY (MiniMaxH3ReadWav), next to the writer and // unit-gated with it; this only opens the file. Returns CHANNEL-MAJOR samples in // [-1, 1], mono repeated up to kMiniMaxH3AudioChannels, and REFUSES a sample rate @@ -172,6 +201,7 @@ int main(int argc, char** argv) { std::string first_frame_path, last_frame_path; std::string decode_latent_path; // diagnostic: decode a dumped VAE-input latent std::string roundtrip_path; // diagnostic: encode->decode a real image + std::string prompt_image_path; // diagnostic: run an image through the vision tower std::vector ref_image_paths; std::string ref_video_prefix, ref_audio_path; // The served checkpoint PARTITION. Community GGUF/NVFP4 files strip the release @@ -203,6 +233,7 @@ int main(int argc, char** argv) { else if (f == "--dump-params") dump_params = true; else if (f == "--decode-latent") decode_latent_path = Need(argc, argv, ++i, f); else if (f == "--roundtrip") roundtrip_path = Need(argc, argv, ++i, f); + else if (f == "--prompt-image") prompt_image_path = Need(argc, argv, ++i, f); else if (f == "--device") device_name = Need(argc, argv, ++i, f); else if (f == "--encoder") encoder_path = Need(argc, argv, ++i, f); else if (f == "--prompt") prompt = Need(argc, argv, ++i, f); @@ -229,12 +260,15 @@ int main(int argc, char** argv) { // no VAEs, no conditioning, no output path. Requiring them would make the one // tool that works on a checkpoint too large to load unusable on exactly that // checkpoint. + // --prompt-image runs the vision tower ONLY (from --encoder); it needs no DiT/VAE/out. + const bool vision_probe = !prompt_image_path.empty(); const bool diag_vae_only = !decode_latent_path.empty() || !roundtrip_path.empty(); const bool need_vaes = !denoise_only && !dump_params && !diag_vae_only; const bool need_cond = !dump_params && !diag_vae_only; - // --decode-latent / --roundtrip need NO DiT and NO conditioning (their own blocks - // validate their inputs); the shared check below would otherwise reject --dit. - if (!diag_vae_only && + // --decode-latent / --roundtrip / --prompt-image need NO DiT and NO conditioning + // (their own blocks validate their inputs); the shared check below would otherwise + // reject --dit. + if (!diag_vae_only && !vision_probe && (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())))) { @@ -417,6 +451,72 @@ int main(int argc, char** argv) { return 0; } + // --- --prompt-image DIAGNOSTIC: route a real image through the shared Qwen3-VL + // image processor + the encoder's VISION tower, loaded from the encoder GGUF's + // visual.* tensors (the piece the record reconciliation found missing, spec §8.8). + // Reports the conditioning feature stats -- the REAL-WEIGHTS proof that the vision + // tower now runs. Scattering the merged/deepstack features into the DiT-conditioning + // path (merge into prompt_embeds + DeepStack inject into the device text tower) is + // the tracked residual; this probe stops after the tower. + if (!prompt_image_path.empty()) { + if (encoder_path.empty()) + throw std::runtime_error("--prompt-image needs --encoder (the vision weights live in it)"); + if (device_name != "cuda") + throw std::runtime_error( + "--prompt-image needs --device cuda (the vision tower is device-resident)"); + std::cerr << "loading encoder vision tower from " << encoder_path << "\n"; + const vllm::GgufFile ef = vllm::GgufFile::Open(encoder_path); + const vllm::multimodal::Qwen3VLVisionConfig vcfg = vllm::MiniMaxH3EncoderVisionConfig(); + const vllm::multimodal::Qwen3VLVisionWeights vw = vllm::LoadQwen3VLVisionFromGguf(ef, vcfg); + std::cerr << " vision weights: depth=" << vcfg.depth << " hidden=" << vcfg.hidden_size + << " heads=" << vcfg.num_heads << " out=" << vcfg.out_hidden_size + << " deepstack=" << vw.deepstack_mergers.size() << "\n"; + + int64_t ih = 0, iw = 0; + const std::vector rgb = ReadPpmAsHwcU8(prompt_image_path, &ih, &iw); + std::cerr << " image " << iw << "x" << ih << "\n"; + vllm::multimodal::Qwen3VLProcessorConfig pcfg; // patch16/temporal2/merge2/0.5-norm + pcfg.merge_size = static_cast(vcfg.spatial_merge_size); + const vllm::multimodal::Qwen3VLImageProcessor proc(pcfg); + const vllm::multimodal::ImageKwargs kw = proc.ProcessImage(rgb.data(), ih, iw); + const std::array grid = kw.image_grid_thw; + const int64_t tokens = grid[0] * grid[1] * grid[2]; + const int64_t merge = vcfg.spatial_merge_size * vcfg.spatial_merge_size; + std::cerr << " grid_thw=[" << grid[0] << "," << grid[1] << "," << grid[2] + << "] tokens=" << tokens << " merged=" << (tokens / merge) << "\n"; + + vt::Backend& backend = vt::GetBackend(vt::DeviceType::kCUDA); + vllm::multimodal::Qwen3VLVisionCapture cap; + const std::vector tower = vllm::multimodal::Qwen3VLVisionForward( + kw.pixel_values_bf16, grid, vw, vcfg, backend, &cap); + const int64_t nm = tokens / merge; + const int64_t width = + vcfg.out_hidden_size * + (1 + static_cast(vcfg.deepstack_visual_indexes.size())); + auto stats = [&](const std::vector& v, const char* tag) { + double s = 0, s2 = 0, mx = 0; + bool fin = true; + for (float f : v) { + if (!std::isfinite(f)) fin = false; + s += f; + s2 += double(f) * f; + mx = std::max(mx, std::fabs(static_cast(f))); + } + const double n = v.empty() ? 1.0 : static_cast(v.size()); + std::cerr << " " << tag << ": n=" << v.size() << " finite=" << (fin ? "yes" : "NO") + << " mean=" << (s / n) << " rms=" << std::sqrt(s2 / n) << " maxabs=" << mx + << "\n"; + }; + std::cerr << " tower out = [" << nm << ", " << width << "]\n"; + stats(tower, "tower_concat"); + stats(cap.merger_out, "merged"); + for (size_t d = 0; d < cap.deepstack_out.size(); ++d) + stats(cap.deepstack_out[d], ("deepstack_" + std::to_string(d)).c_str()); + std::cout << "prompt-image: vision tower RAN on real weights; merged=[" << nm << "," + << vcfg.out_hidden_size << "] + " << cap.deepstack_out.size() << " deepstack\n"; + return 0; + } + // --- 1. DiT --- std::cerr << "loading DiT " << dit_path << (keep_quant ? " (keep-quant)" : "") << "\n"; vllm::MiniMaxH3GgufDit dit; diff --git a/include/vllm/model_executor/models/minimax_h3.h b/include/vllm/model_executor/models/minimax_h3.h index cea888cb..7a82289d 100644 --- a/include/vllm/model_executor/models/minimax_h3.h +++ b/include/vllm/model_executor/models/minimax_h3.h @@ -50,6 +50,7 @@ #include #include "vllm/model_executor/models/qwen3_5_weights.h" // Nvfp4Weight (fp4 arm) +#include "vllm/model_executor/models/qwen3_vl_vision.h" // encoder vision tower reuse #include "vt/device.h" #include "vt/tensor.h" @@ -1072,6 +1073,35 @@ std::vector MiniMaxH3EncoderTextForwardDevice( MiniMaxH3EncoderQuantWeights LoadMiniMaxH3EncoderFromGguf(const GgufFile& file, int64_t max_layers = 0); +// --------------------------------------------------------------------------- +// Encoder VISION tower (image/video conditioning) — REUSE of the shared Qwen3-VL +// front end (`multimodal::Qwen3VLVisionForward` + `Qwen3VLImageProcessor`). +// +// The H3 encoder is a fine-tuned Qwen3-VL, so its ViT is the SAME architecture the +// project already ports; only the config differs. `MiniMaxH3EncoderVisionConfig` +// returns that config, measured from the real encoder GGUF + the shared +// Qwen3.6-27B vision config (state.md 2026-07-25): hidden 1152 / 16 heads / depth +// 27 / intermediate 4304 / out_hidden 5120 / patch 16 / temporal 2 / merge 2 / +// num_position_embeddings 2304. H3 differs from the 27B by carrying 3 REAL +// DeepStack mergers (`visual.deepstack_merger_list.{0,1,2}`). The one value the +// weights-only ComfyUI GGUF does NOT carry is `deepstack_visual_indexes` (WHICH +// text/vision layers the DeepStack taps sit after); the default here is inferred +// (evenly spaced) and must be confirmed against the upstream vision_config for a +// bit-correct DeepStack inject. +multimodal::Qwen3VLVisionConfig MiniMaxH3EncoderVisionConfig(); + +// Load the encoder GGUF's `visual.*` tower into the shared f32 weight struct the +// Qwen3-VL front end consumes. Mirrors the safetensors `LoadQwen3VLVisionWeights` +// (`qwen3_vl.cpp`) but (a) strips the `visual.` prefix and (b) DEQUANTIZES every +// tensor to f32 via `DequantGgufRowToF32` — the encoder GGUF stores the ViT in +// Q4_K/Q5_K (blocks) with F16 patch_embed/pos_embed. The ComfyUI export reshapes a +// non-256-aligned row (hidden 1152) to `ne0=256`; dequantizing the whole flat +// buffer preserves the row-major `[out,in]` order the tower reads (it indexes every +// weight as a flat buffer with dims taken from the config), so no reshape metadata +// is needed. `cfg.depth` / `cfg.deepstack_visual_indexes.size()` drive the loop. +multimodal::Qwen3VLVisionWeights LoadQwen3VLVisionFromGguf( + const GgufFile& file, const multimodal::Qwen3VLVisionConfig& cfg); + // Materialize the H3-Encoder (FL2VA/text_encoder, 14 shards / 1058 tensors) into // the name map both encoder forwards read. // diff --git a/src/vllm/model_executor/models/minimax_h3_vision_gguf.cpp b/src/vllm/model_executor/models/minimax_h3_vision_gguf.cpp new file mode 100644 index 00000000..d3bbd224 --- /dev/null +++ b/src/vllm/model_executor/models/minimax_h3_vision_gguf.cpp @@ -0,0 +1,127 @@ +// MiniMax-H3 encoder VISION tower — GGUF `visual.*` loader (image/video conditioning). +// +// The H3 encoder is a fine-tuned Qwen3-VL, so its ViT is the architecture the project +// already ports as `multimodal::Qwen3VLVisionForward`. This file is the ONE piece the +// record reconciliation (spec §8.8) found missing: a loader that fills the shared +// `Qwen3VLVisionWeights` from the encoder GGUF's `visual.*` tensors, so the tower can +// run on REAL weights instead of the synthetic reduced-dim fixtures it was gated on. +// +// It mirrors the safetensors `LoadQwen3VLVisionWeights` (qwen3_vl.cpp) with two GGUF +// deltas: +// * NAME: the GGUF drops the `model.` level — the ViT is `visual.*`, not +// `model.visual.*` — so the prefix is just `visual.`. +// * ENCODING: every projection is a ggml BLOCK type (Q4_K / Q5_K) and the patch/pos +// tensors are F16, so each is DEQUANTIZED to f32 via DequantGgufRowToF32. The +// ComfyUI export reshapes a non-256-aligned row (hidden 1152 -> 4.5 Q_K blocks) to +// ne0=256; dequantizing the WHOLE flat buffer is reshape-agnostic and preserves the +// row-major [out,in] order the tower reads (it indexes every weight as a flat buffer +// with dims from the config), so no `comfy.gguf.orig_shape` metadata is needed. +#include "vllm/model_executor/models/minimax_h3.h" + +#include +#include +#include +#include + +#include "vllm/model_executor/model_loader/gguf_dequant.h" +#include "vllm/model_executor/model_loader/gguf_reader.h" +#include "vt/dtype.h" + +namespace vllm { + +multimodal::Qwen3VLVisionConfig MiniMaxH3EncoderVisionConfig() { + multimodal::Qwen3VLVisionConfig cfg; + // Measured from ~/h3fp4/ckpt/qwen3vl-32B-MiniMax-H3-Q4_K_M.gguf (visual.* shapes) + + // the shared Qwen3.6-27B vision config (state.md 2026-07-25). See the header. + cfg.hidden_size = 1152; + cfg.num_heads = 16; // head_dim 72 + cfg.depth = 27; + cfg.intermediate_size = 4304; + cfg.out_hidden_size = 5120; // == encoder text dim + cfg.patch_size = 16; + cfg.temporal_patch_size = 2; + cfg.spatial_merge_size = 2; + cfg.num_position_embeddings = 2304; // 48^2 + cfg.in_channels = 3; + // H3's encoder carries 3 REAL DeepStack mergers (visual.deepstack_merger_list.{0,1,2}), + // UNLIKE the Qwen3.6-27B (empty). The WHICH-layers taps are not in the weights-only + // GGUF; these are inferred evenly-spaced over depth 27 and must be confirmed against + // the upstream vision_config for a bit-correct DeepStack inject (spec §8.8 residual). + cfg.deepstack_visual_indexes = {8, 16, 24}; + cfg.norm_eps = 1e-6f; + return cfg; +} + +namespace { + +// Dequantize one `visual.*` tensor to a flat f32 buffer, whatever its ggml encoding +// (F32 / F16 / Q4_K / Q5_K / Q6_K). numel = prod(shape); the tower reads the result as +// a flat [out, in] row-major buffer with dims taken from the config. +std::vector LoadVisionGgufF32(const GgufFile& file, const std::string& name) { + const GgufTensorInfo& info = file.Get(name); + int64_t numel = 1; + for (int64_t d : info.shape) numel *= d; + return DequantGgufRowToF32(info.ggml_type, static_cast(info.data), numel); +} + +} // namespace + +multimodal::Qwen3VLVisionWeights LoadQwen3VLVisionFromGguf( + const GgufFile& file, const multimodal::Qwen3VLVisionConfig& cfg) { + // Membership set: GgufFile has no has() query, and a missing tensor should name + // itself rather than surface as a generic "not found". + std::set present; + for (const GgufTensorInfo& info : file.Tensors()) present.insert(info.name); + auto load = [&](const std::string& name) -> std::vector { + VT_CHECK(present.count(name) != 0, + "minimax_h3 vision gguf: missing tensor " + name + + " (is this the encoder GGUF, and does it carry the visual.* tower?)"); + return LoadVisionGgufF32(file, name); + }; + + multimodal::Qwen3VLVisionWeights vw; + const std::string V = "visual."; + vw.patch_proj_w = load(V + "patch_embed.proj.weight"); + vw.patch_proj_b = load(V + "patch_embed.proj.bias"); + vw.pos_embed_w = load(V + "pos_embed.weight"); + + vw.blocks.resize(static_cast(cfg.depth)); + for (int64_t l = 0; l < cfg.depth; ++l) { + const std::string p = V + "blocks." + std::to_string(l); + multimodal::VisionBlockWeights& b = vw.blocks[static_cast(l)]; + b.norm1_w = load(p + ".norm1.weight"); + b.norm1_b = load(p + ".norm1.bias"); + b.norm2_w = load(p + ".norm2.weight"); + b.norm2_b = load(p + ".norm2.bias"); + b.qkv_w = load(p + ".attn.qkv.weight"); + b.qkv_b = load(p + ".attn.qkv.bias"); + b.proj_w = load(p + ".attn.proj.weight"); + b.proj_b = load(p + ".attn.proj.bias"); + b.fc1_w = load(p + ".mlp.linear_fc1.weight"); + b.fc1_b = load(p + ".mlp.linear_fc1.bias"); + b.fc2_w = load(p + ".mlp.linear_fc2.weight"); + b.fc2_b = load(p + ".mlp.linear_fc2.bias"); + } + + // The main merger norms the PRE-shuffle width (dim); the DeepStack mergers norm the + // POST-shuffle width (merge^2 * dim) — same split the safetensors loader encodes. + auto load_merger = [&](const std::string& prefix, bool postshuffle) { + multimodal::VisionMergerWeights m; + m.use_postshuffle_norm = postshuffle; + m.norm_w = load(prefix + ".norm.weight"); + m.norm_b = load(prefix + ".norm.bias"); + m.fc1_w = load(prefix + ".linear_fc1.weight"); + m.fc1_b = load(prefix + ".linear_fc1.bias"); + m.fc2_w = load(prefix + ".linear_fc2.weight"); + m.fc2_b = load(prefix + ".linear_fc2.bias"); + return m; + }; + vw.merger = load_merger(V + "merger", /*postshuffle=*/false); + for (size_t i = 0; i < cfg.deepstack_visual_indexes.size(); ++i) { + vw.deepstack_mergers.push_back( + load_merger(V + "deepstack_merger_list." + std::to_string(i), /*postshuffle=*/true)); + } + return vw; +} + +} // namespace vllm diff --git a/tests/vllm/models/test_minimax_h3.cpp b/tests/vllm/models/test_minimax_h3.cpp index 31eebebb..cf8f4a0c 100644 --- a/tests/vllm/models/test_minimax_h3.cpp +++ b/tests/vllm/models/test_minimax_h3.cpp @@ -4062,6 +4062,153 @@ TEST_CASE("minimax_h3: the FULL encoder vision tower matches upstream") { CHECK(grid[1] != grid[4]); } +TEST_CASE("minimax_h3: the encoder GGUF visual.* loader dequantizes the vision tower") { + // The record reconciliation (spec §8.8) found the vision tower math was gated ONLY + // at reduced dims with SYNTHETIC weights and never wired to real weights: the + // encoder GGUF loader loaded the TEXT tower only and skipped every visual.* tensor. + // This gates the new loader — that it reads a `visual.*` GGUF, DEQUANTIZES each + // ggml-block projection + F16 patch/pos to f32, and fills the shared + // multimodal::Qwen3VLVisionWeights the Qwen3-VL front end consumes — the piece that + // makes a REAL-weights vision forward possible. The loader is encoding-agnostic (it + // calls DequantGgufRowToF32 per tensor), so a Q8_0 fixture exercises the exact code + // path the shipped Q4_K/Q5_K tower takes; the real-file forward is the on-box proof. + vllm::multimodal::Qwen3VLVisionConfig cfg; // reduced dims, but the tower's exact shape math + cfg.hidden_size = 64; + cfg.num_heads = 4; + cfg.depth = 3; + cfg.intermediate_size = 128; + cfg.out_hidden_size = 96; + cfg.patch_size = 16; + cfg.temporal_patch_size = 2; + cfg.spatial_merge_size = 2; + cfg.num_position_embeddings = 2304; + cfg.in_channels = 3; + cfg.deepstack_visual_indexes = {1}; // one DeepStack merger + const int64_t dim = cfg.hidden_size; + const int64_t patch_elems = + cfg.in_channels * cfg.temporal_patch_size * cfg.patch_size * cfg.patch_size; + const int64_t merged = dim * cfg.spatial_merge_size * cfg.spatial_merge_size; + + gguf_test::GgufModelBuilder builder; + builder.AddKv(gguf_test::StrKv("general.architecture", "qwen3vl")); + std::map> orig; // logical name -> the exact f32 written + auto add_f32 = [&](const std::string& name, int64_t numel) { + const std::vector v = MakeParam("vg." + name, numel, 0.05); + orig[name] = v; + std::string bytes(reinterpret_cast(v.data()), v.size() * sizeof(float)); + builder.AddTensor(name, {static_cast(numel)}, /*ggml_type=*/0 /*F32*/, bytes); + }; + // A quantized [out, in] projection stored the way the real tower ships it (block + // encoding). Rows are whole Q8_0 blocks (in % 32 == 0 here). + auto add_q8 = [&](const std::string& name, int64_t out_dim, int64_t in_dim) { + const std::vector values = MakeParam("vg." + name, out_dim * in_dim, 0.05); + orig[name] = values; + const size_t row_bytes = vt::RowSizeBytes(vt::DType::kQ8_0, in_dim); + std::string bytes(static_cast(out_dim) * row_bytes, '\0'); + vt::cpu::FromFloatFn q = vt::cpu::BlockFromFloat(vt::DType::kQ8_0); + REQUIRE(q != nullptr); + for (int64_t r = 0; r < out_dim; ++r) + q(values.data() + r * in_dim, bytes.data() + static_cast(r) * row_bytes, in_dim); + builder.AddTensor(name, {static_cast(in_dim), static_cast(out_dim)}, + /*ggml_type=*/8 /*Q8_0*/, bytes); + }; + + const std::string V = "visual."; + add_q8(V + "patch_embed.proj.weight", dim, patch_elems); + add_f32(V + "patch_embed.proj.bias", dim); + add_f32(V + "pos_embed.weight", cfg.num_position_embeddings * dim); + for (int64_t l = 0; l < cfg.depth; ++l) { + const std::string p = V + "blocks." + std::to_string(l); + add_f32(p + ".norm1.weight", dim); + add_f32(p + ".norm1.bias", dim); + add_f32(p + ".norm2.weight", dim); + add_f32(p + ".norm2.bias", dim); + add_q8(p + ".attn.qkv.weight", 3 * dim, dim); + add_f32(p + ".attn.qkv.bias", 3 * dim); + add_q8(p + ".attn.proj.weight", dim, dim); + add_f32(p + ".attn.proj.bias", dim); + add_q8(p + ".mlp.linear_fc1.weight", cfg.intermediate_size, dim); + add_f32(p + ".mlp.linear_fc1.bias", cfg.intermediate_size); + add_q8(p + ".mlp.linear_fc2.weight", dim, cfg.intermediate_size); + add_f32(p + ".mlp.linear_fc2.bias", dim); + } + // main merger norms the pre-shuffle width (dim); deepstack norms the post-shuffle (merged) + add_f32(V + "merger.norm.weight", dim); + add_f32(V + "merger.norm.bias", dim); + add_q8(V + "merger.linear_fc1.weight", merged, merged); + add_f32(V + "merger.linear_fc1.bias", merged); + add_q8(V + "merger.linear_fc2.weight", cfg.out_hidden_size, merged); + add_f32(V + "merger.linear_fc2.bias", cfg.out_hidden_size); + add_f32(V + "deepstack_merger_list.0.norm.weight", merged); + add_f32(V + "deepstack_merger_list.0.norm.bias", merged); + add_q8(V + "deepstack_merger_list.0.linear_fc1.weight", merged, merged); + add_f32(V + "deepstack_merger_list.0.linear_fc1.bias", merged); + add_q8(V + "deepstack_merger_list.0.linear_fc2.weight", cfg.out_hidden_size, merged); + add_f32(V + "deepstack_merger_list.0.linear_fc2.bias", cfg.out_hidden_size); + + const std::string path = "/tmp/minimax_h3_visual_q.gguf"; + { + const std::string bytes = builder.Build(); + FILE* fh = std::fopen(path.c_str(), "wb"); + REQUIRE(fh != nullptr); + CHECK(std::fwrite(bytes.data(), 1, bytes.size(), fh) == bytes.size()); + std::fclose(fh); + } + + const vllm::GgufFile gguf = vllm::GgufFile::Open(path); + const vllm::multimodal::Qwen3VLVisionWeights vw = vllm::LoadQwen3VLVisionFromGguf(gguf, cfg); + + // STRUCTURE: every field is filled at the flat [out*in] size the tower reads. + CHECK(static_cast(vw.patch_proj_w.size()) == dim * patch_elems); + CHECK(static_cast(vw.patch_proj_b.size()) == dim); + CHECK(static_cast(vw.pos_embed_w.size()) == cfg.num_position_embeddings * dim); + REQUIRE(static_cast(vw.blocks.size()) == cfg.depth); + for (const vllm::multimodal::VisionBlockWeights& b : vw.blocks) { + CHECK(static_cast(b.qkv_w.size()) == 3 * dim * dim); + CHECK(static_cast(b.qkv_b.size()) == 3 * dim); + CHECK(static_cast(b.proj_w.size()) == dim * dim); + CHECK(static_cast(b.fc1_w.size()) == cfg.intermediate_size * dim); + CHECK(static_cast(b.fc2_w.size()) == dim * cfg.intermediate_size); + CHECK(static_cast(b.norm1_w.size()) == dim); + } + CHECK(vw.merger.use_postshuffle_norm == false); + CHECK(static_cast(vw.merger.norm_w.size()) == dim); // pre-shuffle width + CHECK(static_cast(vw.merger.fc2_w.size()) == cfg.out_hidden_size * merged); + REQUIRE(vw.deepstack_mergers.size() == cfg.deepstack_visual_indexes.size()); + CHECK(vw.deepstack_mergers[0].use_postshuffle_norm == true); + CHECK(static_cast(vw.deepstack_mergers[0].norm_w.size()) == merged); // post-shuffle + + // DEQUANT CORRECTNESS: an f32 tensor round-trips EXACTLY, a Q8_0 tensor within its + // block tolerance, in the SAME flat order — this is the load that the reduced-dim + // synthetic gate never exercised on real bytes. + { + const std::vector& want = orig[V + "patch_embed.proj.bias"]; + double e = MaxAbsDiff(vw.patch_proj_b, want.data(), want.size()); + INFO("patch bias f32 exact err=" << e); + CHECK(e == 0.0); + } + { + const std::vector& want = orig[V + "blocks.0.attn.qkv.weight"]; + REQUIRE(vw.blocks[0].qkv_w.size() == want.size()); + double e = MaxAbsDiff(vw.blocks[0].qkv_w, want.data(), want.size()); + INFO("qkv Q8_0 dequant err=" << e); + CHECK(e <= 5e-3); // Q8_0 block tolerance + // and it is NON-degenerate: real spread, not all-zeros. + double s2 = 0.0; + for (float f : vw.blocks[0].qkv_w) s2 += double(f) * f; + CHECK(std::sqrt(s2 / vw.blocks[0].qkv_w.size()) > 1e-3); + } + + // The production config the driver uses is the measured H3 vision geometry. + const vllm::multimodal::Qwen3VLVisionConfig prod = vllm::MiniMaxH3EncoderVisionConfig(); + CHECK(prod.hidden_size == 1152); + CHECK(prod.num_heads == 16); + CHECK(prod.depth == 27); + CHECK(prod.out_hidden_size == 5120); + CHECK(prod.num_position_embeddings == 2304); + CHECK(prod.deepstack_visual_indexes.size() == 3); +} + TEST_CASE("minimax_h3: condition-noise augmentation matches upstream") { // fl2va/ref2va pin their keyframe and reference-audio rows to a NOISED anchor. // The mix is trivial; the ROW ACCOUNTING is what this gates -- and the golden