diff --git a/.agents/NOW.md b/.agents/NOW.md index b5ed4e4f3..3a6a84f52 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 | **vision→cond scatter WIRED+gated; fl2va COHERENT via encoder path** (`H3-VISION-SCATTER` PR#90; indices {8,16,24} confirmed) | ref2va grids; residual = ref2va ref-row assembly (§8.9) | +| MiniMax-H3 lane | **fl2va COHERENT; ref2va assembly bug FIXED+gated** (`H3-REF2VA-ASSEMBLY`) | grid RE-ATTRIBUTED to the ref2va NVFP4 checkpoint/loader, NOT assembly/fp4 (§8.10) | | 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 22a705f54..93a60b6a0 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -14362,3 +14362,41 @@ rows. So the residual is the **ref2va reference-row conditioning ASSEMBLY** noised-anchor layout + how the denoise loop conditions the un-pinned target rows), NOT the prompt and NOT the DiT forward. Next diagnostic: dump the ref2va target-row VAE-input adjacency-cosine (like #77's 0.95 for coherent fl2va) to confirm the target rows are white, and A/B the reference-row condition-noise. + +## MiniMax-H3 ref2va ASSEMBLY fix + grid RE-ATTRIBUTED to the NVFP4 checkpoint (`row/H3-REF2VA-ASSEMBLY` PR #93, 2026-08-07, dgx sm_121a) + +Followed §8.9's residual (ref2va reference-row assembly), diffed it vs upstream, fixed a real bug, +gated it, re-rendered — and the render DISPROVED the §8.9 attribution. + +**Real bug fixed.** `MiniMaxH3EncodeReferenceImages`/`Video` emitted PATCHED ref-block dims +(`ls.h/patch_size_h`); `BuildMiniMaxH3PackedSequenceRef2va` divides `block.latent_h/kPatchH` AGAIN +(mirroring upstream `packed_sequence.py:328-330`, which takes the UNPATCHED latent; upstream feeds the +raw latent `visual_shape=(1,height//16,width//16)`, `pipeline_minimax_h3.py:1141-1145`). Double-division +under-allocated the reference span by patch_h*patch_w (=4); the denoise pin-loop `>=` check silently +TRUNCATED the oversized encoded reference to its first quarter. Fix = emit raw `ls.{t,h,w}`. + +**Gates.** goldens section 5c (ref2va-shaped DiT-forward rung: image + video+audio ref blocks, 8×8 +geometry, ref2va timestep partition + audio update mask, RefDiT) → C++ case "DiT-forward REF2VA rung +matches upstream (reference rows, mixing)"; RED-first encoded-vs-layout row-count invariant +(reintroducing the bug fails 128==512 = 16 encoded vs 4 allocated; restored green). Suite 69/69, 52377 +assertions. Goldens regen purely additive (635 inserts, 0 deletes). + +**GB10 render A/B + ISOLATION (256×256/22f/12steps, fixed incremental binary).** + +| Render | Checkpoint | Quant | Assembly | Result | +|---|---|---|---|---| +| ref2va (`--ref-image`+`--cond-image`) | ref2va NVFP4 | fp4-resident | full | GRID | +| ref2va (`--ref-image`+`--cond-image`) | ref2va NVFP4 | bf16 dequant | full | GRID (fp4 eliminated) | +| t2va (no refs, no cond-image, `--partition fl2va`) | ref2va NVFP4 | bf16 dequant | NONE | GRID (assembly eliminated) | +| fl2va (keyframe + `--cond-image`) | FL2VA GGUF | bf16 | keyframe | COHERENT (cat→windowsill; no regression) | + +**Re-attribution (corrects §8.9 + §8.6).** The grid is NOT the ref2va assembly and NOT fp4 — it appears +with the ENTIRE reference assembly removed, in both fp4 and bf16. It correlates 1:1 with the ref2va +NVFP4 CHECKPOINT: every render loading `minimax_h3_ref2va_nvfp4_full` grids; every FL2VA-GGUF render is +coherent. §8.9's A/B varied only the prompt, §8.6's only the task — neither varied the checkpoint/quant, +so both misattributed a checkpoint/loader defect. True residual = the NVFP4 DiT loader for this file +(`StreamMiniMaxH3Nvfp4ToDeviceBf16/Fp4`): suspect fp32-island preservation (patch/time/output layers, +`minimax_h3_transformer.py:898-904`), `weight_scale_2` double-dequant, or the 1051-tensor name mapping. +Synthetic-NVFP4 gates proved the dequant MATH byte-exact but never loaded THIS file vs a coherent oracle. +Next: a REF2VA GGUF (bf16, known-good loader) as checkpoint oracle — dgx-disk-blocked (23 GiB free). +Artifacts `~/h3fp4/out_{vs_ref2va,bf16_ref2va,t2va_nvfp4,vs_fl2va}.mp4`. diff --git a/.agents/specs/minimax-h3.md b/.agents/specs/minimax-h3.md index ee734bb99..ef5d3d4e0 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; vision tower LOADS real `visual.*` + runs; merged→prompt_embeds scatter + DeepStack→device text tower WIRED 1:1 + gated (§8.9); fl2va COHERENT via BOTH the VAE-keyframe AND the encoder vision path; ref2va still grids (residual RE-ATTRIBUTED to the ref2va reference-row assembly, NOT the conditioning) | **PARTIAL** (vision→conditioning scatter DONE §8.9; residual = ref2va reference-row assembly) | +| Modalities in: text/image/video/audio | `pipeline:1036-1104` | t2va (text) done; vision tower LOADS real `visual.*` + runs; merged→prompt_embeds scatter + DeepStack→device text tower WIRED 1:1 + gated (§8.9); fl2va COHERENT via BOTH the VAE-keyframe AND the encoder vision path; ref2va reference-row assembly FIXED + gated (§8.10, the block-dim double-division) — but ref2va still grids, now RE-ATTRIBUTED to the ref2va NVFP4 CHECKPOINT/loader (t2va-zero-assembly grids too), NOT the assembly | **PARTIAL** (vision→conditioning scatter + ref2va assembly DONE; residual = the NVFP4 DiT loader for the ref2va checkpoint, §8.10) | | 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) | @@ -774,3 +774,73 @@ the un-pinned target rows on them), NOT the prompt_embeds and NOT the DiT forwar dump the ref2va target-row VAE-input latent adjacency-cosine (like #77 did for the coherent fl2va, 0.95) to confirm the target rows are white, and A/B the reference-row condition-noise vs a clean anchor. + +## 8.10 REF2VA ASSEMBLY — a real assembly bug FIXED + gated, and the grid RE-ATTRIBUTED to the CHECKPOINT (2026-08-07, `row/H3-REF2VA-ASSEMBLY` PR #93) + +Took §8.9's residual (the ref2va reference-row assembly) at face value, diffed the whole assembly +against upstream on CPU, found + fixed a real bug, extended the permanent gate, then re-rendered on +real weights — and the render **DISPROVED §8.9's attribution**: the grid is NOT the assembly. + +**The real assembly bug FOUND + FIXED (suspect #1, `MiniMaxH3EncodeReferenceImages`).** +`MiniMaxH3EncodeReferenceImages` (`minimax_h3_pipeline.cpp:162`) and `MiniMaxH3EncodeReferenceVideo` +(`:83`) populated the `MiniMaxH3RefBlock` with **PATCHED** dims (`ls.h / patch_size_h`), but +`BuildMiniMaxH3PackedSequenceRef2va` (`minimax_h3_packing.cpp:453`) divides `block.latent_h / kPatchH` +**again** — it mirrors upstream `minimax_h3_packed_sequence_ref2va_blocks` (`packed_sequence.py:328-330`), +which takes the **UNPATCHED** latent and divides by `_PATCH_H` once (upstream feeds it the raw latent: +`pipeline_minimax_h3.py:1141-1145` sets `visual_shape = (1, height//16, width//16)`, and the target +`latent_h = height//16` is fed to the same builder). The double-division under-allocated the reference +span by `patch_h*patch_w` (=4): the layout claimed **1/4** the reference rows the VAE encode produced, +and the denoise loop's `keyframe_cond_rows.size() >= …` check is a `>=`, so the oversized encoded +reference was **silently truncated to its first quarter** — no throw, a coherence-mangling layout bug. +Fix: emit the RAW `ls.{t,h,w}` in both encode functions (byte-for-byte upstream's `visual_shape`). The +other three suspects were CLEARED by the same read: the per-token timesteps mirror `denoise_loop.py:109-118` +exactly (pinned refs at `max(t_v,0.999)` / `max(t_a,1.0)` — `minimax_h3.cpp:843-854`); the packed layout, +position grid and token tags are byte-exact vs upstream (gated below); the #77 output bookkeeping is +correct (upstream `video_rows[update_mask]` == our trailing-target slice, refs front-loaded). + +**Permanent gates (the recurring blind-spot class closed).** +1. `scripts/gen-minimax-h3-goldens.py` grew **section 5c**: a REF2VA-shaped DiT-forward rung built from + the upstream `minimax_h3_packed_sequence_ref2va_blocks` (an image reference + a video+audio reference + prepended to the target) at 8×8 spatial geometry, forwarded through `RefDiT` with the ref2va per-token + timestep partition (target `t_v`/`t_a`, pinned visual refs at 0.999, pinned audio refs at 1.0) and its + audio update mask. The C++ case `test_minimax_h3.cpp :: "DiT-forward REF2VA rung matches upstream + (reference rows, mixing)"` rebuilds it with `BuildMiniMaxH3PackedSequenceRef2va` and gates the layout, + host+device video/audio logits (≤2e-5), the reference-row output masking, and the target-row + spatial-mixing probe (mix_fraction 1.0). No prior DiT-forward gate ever forwarded a ref2va layout — + the §8.9 blind spot. Needed a 1-line RefDiT extension (honor `audio_update_mask`, mirroring + `minimax_h3_transformer.py:1099-1101`; our port already did). Goldens regen is purely additive (635 + inserts, 0 deletes — no drift). +2. A RED-first **encoded-vs-layout row-count invariant** in the ref2va image + video subcases: after + `MiniMaxH3EncodeReferenceImages`, `ref_visual_rows * video_row_width == ref_rows.size()`. With the bug + reintroduced it fails `128 == 512` (16 encoded rows vs 4 allocated — the exact 4× patch double-count); + restored → green. Nothing else in the suite couples encoded rows to layout rows (the section-2 layout + gate hand-builds unpatched blocks; the denoise round-trip only checks finiteness/motion). Suite + **69/69, 52377 assertions** (CPU), the fix RED→GREEN proven. + +**GB10 render A/B + isolation (2026-08-07, dgx sm_121a, 256×256/22f/12steps, fixed binary).** +- **ref2va, fp4-resident** (`--ref-image` + `--cond-image` + `--partition ref2va`): still a multicolour + patch grid. `~/h3fp4/out_vs_ref2va.mp4`. +- **ref2va, bf16** (same, NVFP4→bf16 dequant, no `--fp4-resident`): **also grids** → fp4 is NOT the cause. + `~/h3fp4/out_bf16_ref2va.mp4`. +- **t2va with ZERO reference assembly** (ref2va NVFP4, bf16, `--partition fl2va` to bypass the guard, no + refs, no cond-image, plain text prompt): **still grids** → the reference-row assembly is NOT the cause; + the grid appears with the entire assembly removed. `~/h3fp4/out_t2va_nvfp4.mp4`. +- **fl2va control** (FL2VA GGUF `--dequant-bf16` + keyframe + `--cond-image`): **COHERENT** — frame 0 a + photorealistic orange cat matching the keyframe, frame 21 the cat on a windowsill in warm sunlight + (evolved toward the prompt). No regression from the fix. `~/h3fp4/out_vs_fl2va.mp4`. + +**RE-ATTRIBUTION (corrects §8.9 and §8.6).** The grid is **NOT the ref2va reference-row assembly** and +**NOT fp4**. It correlates 1:1 with the **ref2va NVFP4 checkpoint** (`minimax_h3_ref2va_nvfp4_full`): +EVERY render that loads it grids (t2va #74, ref2va-fp4, ref2va-bf16, t2va-zero-assembly), while EVERY +render that loads the FL2VA GGUF is coherent (t2va #77, fl2va #90, fl2va here). §8.9's A/B varied only the +prompt and §8.6's only the task — neither ever varied the checkpoint or the quant, so both misattributed +a checkpoint/loader defect (to "assembly" and to "wrong partition"). The DiT MATH, the ref2va assembly +math, and the vision scatter are now all gate-proven correct vs upstream; the real residual is the +**NVFP4 DiT LOADER for this checkpoint** (`StreamMiniMaxH3Nvfp4ToDeviceBf16/Fp4`, `minimax_h3_nvfp4*`) — +suspect the fp32-island preservation (patch/time/output layers, `minimax_h3_transformer.py:898-904`), +the `weight_scale_2` double-dequant, or the tensor-name mapping for this specific 1051-tensor file. The +synthetic-NVFP4 gates (§8.4/§8.6) proved the dequant MATH byte-exact but never loaded THIS file end to +end against a coherent oracle. Next: a REF2VA GGUF (bf16, known-good loader) as the checkpoint oracle — +if ref2va-on-REF2VA-GGUF is coherent, the NVFP4 loader for this file is the bug; else the checkpoint +itself. The REF2VA GGUF download is dgx-disk-blocked (23 GiB free, 100% full; large dirs belong to other +campaigns — not prunable). diff --git a/.agents/state.md b/.agents/state.md index e8d5f80c8..238e97b4b 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -40304,3 +40304,50 @@ cosine (vs #77's 0.95 for coherent fl2va) + A/B the reference-row condition-nois Box left clean (renders exited, GPU idle, gpu.lock released, worker stays parked, temp PNGs pruned, ckpts kept). Records: spec §8.9 + §8.2 row, STATUS/BENCHMARKS/FEATURES H3 rows, benchmark-record, NOW. + +## 2026-08-07T05:30 - MiniMax-H3 ref2va assembly bug FIXED + gated; render grid RE-ATTRIBUTED to the ref2va NVFP4 CHECKPOINT (NOT the assembly) (row/H3-REF2VA-ASSEMBLY, helper, PR #93) + + +Took §8.9's residual (the ref2va reference-row assembly) and ran the §70 latent-bisection discipline at +the assembly boundary: diffed the whole assembly vs upstream on CPU, found + fixed a real bug, extended +the permanent gate to a ref2va-shaped rung, re-rendered on real weights — and the render DISPROVED +§8.9's attribution. + +**REAL BUG FIXED (suspect #1).** `MiniMaxH3EncodeReferenceImages` (minimax_h3_pipeline.cpp:162) and +`MiniMaxH3EncodeReferenceVideo` (:83) emitted the ref block with PATCHED dims (`ls.h/patch_size_h`), but +`BuildMiniMaxH3PackedSequenceRef2va` (minimax_h3_packing.cpp:453) divides `block.latent_h/kPatchH` AGAIN, +mirroring upstream packed_sequence.py:328-330 which takes the UNPATCHED latent (upstream feeds the raw +latent: pipeline_minimax_h3.py:1141-1145 `visual_shape=(1,height//16,width//16)`). Double-division +under-allocated the reference span by patch_h*patch_w (=4); the denoise pin-loop's `>=` check let the +oversized encoded reference be SILENTLY TRUNCATED to its first quarter. Fix: emit raw `ls.{t,h,w}` in +both encode fns. The other 3 suspects CLEARED by the same read: per-token timesteps mirror +denoise_loop.py:109-118 exactly (minimax_h3.cpp:843-854); layout/grid/tags byte-exact vs upstream; #77 +output bookkeeping correct (upstream `video_rows[update_mask]` == our trailing-target slice). + +**GATES.** (1) goldens section 5c: a ref2va-shaped DiT-forward rung (image + video+audio reference blocks, +8×8 geometry) forwarded through RefDiT with the ref2va timestep partition + audio update mask; C++ case +"DiT-forward REF2VA rung matches upstream (reference rows, mixing)" gates layout + host/device logits +(≤2e-5) + ref-row masking + target mixing (frac 1.0). Needed a 1-line RefDiT extension (honor +audio_update_mask, upstream minimax_h3_transformer.py:1099-1101; our port already did). Goldens regen +purely additive (635 inserts, 0 deletes). (2) RED-first encoded-vs-layout row-count invariant in the +ref2va image+video subcases: reintroducing the bug fails 128==512 (16 encoded vs 4 allocated), restored → +green. Suite 69/69, 52377 assertions (CPU Release, standalone — no line-3503 flake). + +**GB10 render A/B + ISOLATION (256×256/22f/12steps, fixed binary, incremental build in ~/h3fp4/src).** +ref2va fp4-resident → grids; ref2va bf16 (no --fp4-resident) → ALSO grids (fp4 eliminated); t2va with +ZERO reference assembly on the ref2va NVFP4 (--partition fl2va bypass, no refs/cond-image) → STILL grids +(assembly eliminated); fl2va control (FL2VA GGUF + keyframe + --cond-image) → COHERENT (frame 0 orange +cat matching keyframe, frame 21 cat on a windowsill in warm sunlight — no regression from the fix). +Artifacts ~/h3fp4/out_{vs_ref2va,bf16_ref2va,t2va_nvfp4,vs_fl2va}.mp4. + +**RE-ATTRIBUTION (corrects §8.9 + §8.6).** The grid is NOT the assembly and NOT fp4. It correlates 1:1 +with the ref2va NVFP4 CHECKPOINT (minimax_h3_ref2va_nvfp4_full): every render loading it grids, every +FL2VA-GGUF render is coherent. §8.9 varied only the prompt; §8.6 only the task — neither varied the +checkpoint/quant, so both misattributed a checkpoint/loader defect. True residual = the NVFP4 DiT LOADER +for this file (StreamMiniMaxH3Nvfp4ToDeviceBf16/Fp4): suspect fp32-island preservation (patch/time/output +layers, minimax_h3_transformer.py:898-904), weight_scale_2 double-dequant, or tensor-name mapping for the +1051-tensor file. Synthetic-NVFP4 gates proved the dequant MATH byte-exact but never loaded THIS file vs +a coherent oracle. NEXT: a REF2VA GGUF (bf16, known-good loader) as checkpoint oracle — blocked on dgx +disk (23 GiB free, 100% full; large dirs belong to other campaigns, not prunable). Fix + gates LAND +regardless; box left clean (renders exited, gpu.lock released, worker stays parked, ckpts kept). Records: +spec §8.10 + §8.2 row, STATUS/BENCHMARKS/FEATURES H3 rows, benchmark-record, NOW. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index d09ba9679..a8a8608b5 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -311,7 +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`, `row/H3-VISION-SCATTER`) | **Vision→cond scatter WIRED+gated; fl2va COHERENT via encoder path.** merged→prompt_embeds + DeepStack→device text tower (device==host 3.8e-4); fl2va+`--cond-image` matched cat → prompt scene; ref2va STILL grids | grid RE-ATTRIBUTED to the ref2va reference-row assembly, NOT the conditioning (DiT math byte-exact; vision sound per the coherent fl2va A/B); indices {8,16,24} confirmed. §8.9 | +| MiniMax-H3 image conditioning (`row/H3-CONDITIONED-E2E`, `row/H3-VISION-SCATTER`, `row/H3-REF2VA-ASSEMBLY`) | **fl2va COHERENT; ref2va assembly bug FIXED+gated.** vision→cond scatter gated; ref2va block-dim double-division fixed + RED-first gated (128 vs 512) + a permanent ref2va DiT-forward rung (§8.10) | grid RE-ATTRIBUTED: with the fix ref2va grids in fp4 AND bf16, and t2va with no refs on the ref2va NVFP4 also grids while FL2VA-GGUF renders, so it is the **NVFP4 checkpoint/loader**, NOT assembly/fp4 (§8.10) | | 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 a29a37651..7c3469160 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 66/66 (ladder + CUDA-vs-host seq 1920); t2va+fl2va COHERENT on GB10 (fl2va via VAE-keyframe + encoder path); merged→prompt_embeds + DeepStack→text tower wired+gated; ref2va grids (residual = ref2va ref rows) | FP4/Marlin landed; ref2va reference-row conditioning + speed pending | +| MiniMax-H3 DiT (`MiniMaxH3DiTModel`, vllm-omni lane) | MiniMax-H3 (33.1B video+audio) | portable 69/69 (ladder + ref2va DiT rung + seq 1920); t2va+fl2va COHERENT; vision→cond scatter gated; ref2va assembly FIXED+gated but grids, RE-ATTRIBUTED to the NVFP4 checkpoint/loader, NOT assembly (§8.10) | FP4/Marlin landed; ref2va NVFP4-checkpoint load path + 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+fl2va COHERENT on GB10 (fl2va via VAE-keyframe AND encoder vision path); merged→prompt_embeds + DeepStack→text tower wired+gated (§8.9); guard mirrors `_resolve_task`; ref2va grids (residual = ref2va ref rows) | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | +| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va+fl2va COHERENT on GB10 (fl2va via keyframe+encoder path); vision→cond scatter gated; ref2va assembly FIXED+gated but grids, RE-ATTRIBUTED to the NVFP4 checkpoint/loader, NOT assembly (§8.10) | ✅ (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 c625b3546..e9f10a918 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 None: emit_f64(out, f"{p}MixFraction", [r["mix_fraction"]]) +# --------------------------------------------------------------------------- +# Section 5c: the DiT-forward REF2VA rung (reference rows). +# +# Every existing DiT-forward gate (section 5 + the 5b geometry ladder) runs the +# FL2VA packed layout: a keyframe-cond prefix that shares the TARGET frame grid, +# no audio reference rows, one update mask. The REF2VA assembly is structurally +# different and was NEVER exercised through the DiT forward -- exactly the +# recurring blind-spot class (spec section 8.9): a bug in the reference-row +# assembly (VAE-reference row count, the ref-block position grid, the ref token +# tags, the audio reference update mask) is invisible until a REF2VA layout is +# actually forwarded. This rung builds a real ref2va layout via the upstream +# minimax_h3_packed_sequence_ref2va_blocks (an image reference block + a +# video+audio reference block prepended to the target video+audio) at real 8x8 +# spatial geometry, runs the RefDiT forward with the ref2va per-token timestep +# partition (target video t_video, pinned video refs at the imgvid cond anchor +# 0.999, target audio t_audio, pinned audio refs at 1.0), and gates the layout, +# the video/audio logits and the target-row spatial-mixing probe. The C++ side +# rebuilds it with BuildMiniMaxH3PackedSequenceRef2va, so the port's ref2va +# packing AND its DiT forward over reference rows are both pinned to upstream. +# --------------------------------------------------------------------------- + +# text_len, latent_t, latent_h, latent_w, audio_t, audio_channel +REF2VA_DIT = dict(text_len=5, latent_t=2, latent_h=8, latent_w=8, audio_t=3, audio_channel=2) +# One image reference + one video+audio reference, prepended in request order. +REF2VA_DIT_BLOCKS = [ + {"kind": "image", "latent_h": 8, "latent_w": 8}, + {"kind": "video_audio", "ref_audio_t": 2, "latent_t": 2, "latent_h": 8, "latent_w": 8}, +] + + +def _run_ladder_dit_ref2va(model, arch, packed, name): + """Mirror _run_ladder_dit but for a ref2va layout with reference rows. + + The one structural difference is the AUDIO update mask: ref2va pins audio + reference rows (audio_update_mask False) at the audio-ref cond timestep and + masks their DiT output, so the timestep partition and the audio masking both + have to carry it. Everything else (scatter, refiner cu, the video mixing + probe) is identical to the fl2va helper. + """ + seq_len = int(packed["seq_len"]) + img_pos = packed["img_pos"].to(torch.long) + audio_pos = packed["audio_pos"].to(torch.long) + text_pos = packed["text_pos"].to(torch.long) + token_tags = packed["token_tags"].to(torch.long) + update_mask = packed["update_mask"].to(torch.bool) + audio_update_mask = packed["audio_update_mask"].to(torch.bool) + cu_seqlens = packed["cu_seqlens"].to(torch.int64) + img_position_ids = packed["img_position_ids"].to(torch.float32)[None] + + video_width = arch["latents_dim"] * int(np.prod(arch["patch_size"])) + x = torch.zeros(1, seq_len, video_width, dtype=torch.float32) + x[0].index_copy_( + 0, + img_pos, + torch.from_numpy( + h3_rand(f"ref2va_dit.{name}.video_rows", img_pos.shape[0] * video_width).astype(np.float32) + ).reshape(img_pos.shape[0], video_width), + ) + audio_x = torch.zeros(1, seq_len, arch["audio_latents_dim"], dtype=torch.float32) + audio_x[0].index_copy_( + 0, + audio_pos, + torch.from_numpy( + h3_rand(f"ref2va_dit.{name}.audio_rows", + audio_pos.shape[0] * arch["audio_latents_dim"]).astype(np.float32) + ).reshape(audio_pos.shape[0], arch["audio_latents_dim"]), + ) + prompt_embeds = torch.from_numpy( + h3_rand(f"ref2va_dit.{name}.prompt_embeds", text_pos.shape[0] * arch["text_dim"]).astype(np.float32) + ).reshape(text_pos.shape[0], arch["text_dim"]) + + # ref2va per-token timestep partition (denoise_loop.py:109-118): target rows at + # t_video/t_audio, pinned visual refs at the imgvid cond anchor, pinned audio + # refs at the audio-ref cond anchor, everything else at t_video. + t_video, t_audio, imgvid_cond_t, audio_ref_cond_t = 0.4, 0.35, 0.999, 1.0 + timesteps = torch.full((seq_len,), t_video, dtype=torch.float32) + timesteps[img_pos[update_mask]] = t_video + timesteps[img_pos[~update_mask]] = imgvid_cond_t + timesteps[audio_pos[audio_update_mask]] = t_audio + timesteps[audio_pos[~audio_update_mask]] = audio_ref_cond_t + unique_timesteps, inverse_indices = torch.unique(timesteps, sorted=True, return_inverse=True) + refiner_cu = torch.tensor([0, text_pos.shape[0], text_pos.shape[0]], dtype=torch.int64) + + def forward(x_in): + return model( + x=x_in, + audio_x=audio_x, + img_position_ids=img_position_ids, + unique_timesteps=unique_timesteps, + inverse_indices=inverse_indices, + update_mask=update_mask, + token_tags=token_tags, + prompt_embeds=prompt_embeds, + img_pos=img_pos, + audio_pos=audio_pos, + text_pos=text_pos, + infer_out_pos=img_pos, + cu_seqlens=cu_seqlens, + refiner_cu_seqlens=refiner_cu, + audio_update_mask=audio_update_mask, + ) + + video_logits, audio_logits = forward(x) + + # spatial-mixing probe over the video TARGET rows (identical to the fl2va helper). + target_local = [r for r in range(img_pos.shape[0]) if bool(update_mask[r])] + first_target = target_local[0] + x_pert = x.clone() + x_pert[0, int(img_pos[first_target]), :] += 1.0 + video_pert, _ = forward(x_pert) + delta = (video_pert - video_logits).abs().amax(dim=-1) + other_targets = [r for r in target_local if r != first_target] + responded = sum(1 for r in other_targets if float(delta[r]) > _LADDER_MIX_EPS) + mix_fraction = responded / len(other_targets) if other_targets else -1.0 + + return dict( + seq_len=seq_len, img_pos=img_pos, audio_pos=audio_pos, text_pos=text_pos, + update_mask=update_mask, audio_update_mask=audio_update_mask, token_tags=token_tags, + cu_seqlens=cu_seqlens, img_position_ids=packed["img_position_ids"], + unique_timesteps=unique_timesteps, inverse_indices=inverse_indices, + video_logits=video_logits, audio_logits=audio_logits, video_width=video_width, + first_target=first_target, mix_delta=delta, mix_fraction=mix_fraction, + ) + + +def emit_dit_ladder_ref2va(out, packed_sequence) -> None: + out.write("// --- section 5c: DiT-forward REF2VA rung (reference rows, section 8.9) ---\n") + torch.manual_seed(0) + arch = ARCH + model = RefDiT(arch) + + cfg = REF2VA_DIT + packed = packed_sequence.minimax_h3_packed_sequence_ref2va_blocks( + text_len=cfg["text_len"], + latent_t=cfg["latent_t"], + latent_h=cfg["latent_h"], + latent_w=cfg["latent_w"], + audio_t=cfg["audio_t"], + ref_blocks=REF2VA_DIT_BLOCKS, + audio_channel=cfg["audio_channel"], + ) + r = _run_ladder_dit_ref2va(model, arch, packed, "r2v") + p = "kH3Ref2vaDit" + + for key, value in cfg.items(): + emit_scalar(out, f"{p}_{key}", value) + emit_scalar(out, f"{p}_seq_len", r["seq_len"]) + emit_scalar(out, f"{p}_num_img", int(r["img_pos"].shape[0])) + emit_scalar(out, f"{p}_num_audio", int(r["audio_pos"].shape[0])) + emit_scalar(out, f"{p}_num_text", int(r["text_pos"].shape[0])) + emit_scalar(out, f"{p}_num_unique", int(r["unique_timesteps"].shape[0])) + emit_scalar(out, f"{p}_video_width", int(r["video_width"])) + emit_scalar(out, f"{p}_first_target", int(r["first_target"])) + # The reference blocks, so the C++ side reconstructs the identical layout: + # kind is 0=image, 1=audio, 2=video_audio (MiniMaxH3RefBlock::Kind order). + kind_map = {"image": 0, "audio": 1, "video": 2, "video_audio": 2} + block_kinds, block_ra, block_lt, block_lh, block_lw = [], [], [], [], [] + for b in REF2VA_DIT_BLOCKS: + block_kinds.append(kind_map[b["kind"]]) + block_ra.append(int(b.get("ref_audio_t", 0))) + block_lt.append(int(b.get("latent_t", 0))) + block_lh.append(int(b.get("latent_h", 0))) + block_lw.append(int(b.get("latent_w", 0))) + emit_scalar(out, f"{p}_num_blocks", len(REF2VA_DIT_BLOCKS)) + out.write("\n") + emit_i64(out, f"{p}BlockKinds", block_kinds) + emit_i64(out, f"{p}BlockRefAudioT", block_ra) + emit_i64(out, f"{p}BlockLatentT", block_lt) + emit_i64(out, f"{p}BlockLatentH", block_lh) + emit_i64(out, f"{p}BlockLatentW", block_lw) + # Layout tensors (gates BuildMiniMaxH3PackedSequenceRef2va through the DiT path). + emit_i64(out, f"{p}CuSeqlens", r["cu_seqlens"]) + emit_i64(out, f"{p}ImgPos", r["img_pos"]) + emit_i64(out, f"{p}AudioPos", r["audio_pos"]) + emit_i64(out, f"{p}TextPos", r["text_pos"]) + emit_i64(out, f"{p}UpdateMask", r["update_mask"].to(torch.int64)) + emit_i64(out, f"{p}AudioUpdateMask", r["audio_update_mask"].to(torch.int64)) + emit_i64(out, f"{p}TokenTags", r["token_tags"]) + emit_f64(out, f"{p}ImgPositionIds", r["img_position_ids"]) + emit_f32(out, f"{p}UniqueTimesteps", r["unique_timesteps"]) + emit_i64(out, f"{p}InverseIndices", r["inverse_indices"]) + # RefDiT forward + the spatial-mixing probe. + emit_f32(out, f"{p}VideoLogits", r["video_logits"]) + emit_f32(out, f"{p}AudioLogits", r["audio_logits"]) + emit_f32(out, f"{p}VideoMixDelta", r["mix_delta"]) + emit_f64(out, f"{p}MixFraction", [r["mix_fraction"]]) + + def emit_condition_noise(out, condition_noise, packed_tokens) -> None: """Section 7: condition-noise augmentation (condition_noise.py, VERBATIM). @@ -1349,6 +1543,7 @@ def main() -> int: emit_scheduler(out, scheduling) emit_dit(out, packed) emit_dit_ladder(out, packed_sequence) + emit_dit_ladder_ref2va(out, packed_sequence) emit_planner(out, time_request) emit_condition_noise(out, condition_noise, packed_tokens) emit_reference_video(out) diff --git a/src/vllm/model_executor/models/minimax_h3_pipeline.cpp b/src/vllm/model_executor/models/minimax_h3_pipeline.cpp index e60b9d829..ccd03e989 100644 --- a/src/vllm/model_executor/models/minimax_h3_pipeline.cpp +++ b/src/vllm/model_executor/models/minimax_h3_pipeline.cpp @@ -110,12 +110,16 @@ std::vector MiniMaxH3EncodeReferenceVideo( if (out_block != nullptr) { // kVideoAudio is the only kind that carries a temporal extent -- kImage counts // exactly one frame regardless of latent_t. ref_audio_t stays 0: silent. + // Dims are the RAW VAE latent (t, h, w). BuildMiniMaxH3PackedSequenceRef2va + // applies the DiT [1,2,2] patch division ITSELF (it mirrors upstream + // minimax_h3_packed_sequence_ref2va_blocks, which takes the unpatched latent and + // divides by _PATCH_{H,W} once) -- pre-dividing here would double-count the patch. MiniMaxH3RefBlock b; b.kind = MiniMaxH3RefBlock::Kind::kVideoAudio; b.ref_audio_t = 0; - b.latent_t = ls.t / dit_params.patch_size_t; - b.latent_h = ls.h / dit_params.patch_size_h; - b.latent_w = ls.w / dit_params.patch_size_w; + b.latent_t = ls.t; + b.latent_h = ls.h; + b.latent_w = ls.w; *out_block = b; } return rows; @@ -185,12 +189,22 @@ std::vector MiniMaxH3EncodeReferenceImages( dit_params.patch_size_h, dit_params.patch_size_w); rows.insert(rows.end(), patched.begin(), patched.end()); if (out_blocks != nullptr) { - // The block declares the PATCHED grid, which is what occupies packed rows. + // The block declares the RAW VAE-latent grid (t, h, w). + // BuildMiniMaxH3PackedSequenceRef2va applies the DiT [1,2,2] patch division + // ITSELF, mirroring upstream minimax_h3_packed_sequence_ref2va_blocks, which + // takes the unpatched latent and divides by _PATCH_{H,W} once + // (pipeline_minimax_h3.py:1141-1145 sets visual_shape = (1, height//16, + // width//16), the RAW latent, NOT the patched grid). Pre-dividing here + // double-counted the patch: it under-allocated the reference rows by + // patch_h*patch_w (=4), silently truncating the pinned reference to its first + // quarter and positioning it on a shrunken grid -- the ref2va patch grid + // (spec section 8.9). The encoded row COUNT is invariant-gated in + // test_minimax_h3.cpp so this convention cannot silently drift again. MiniMaxH3RefBlock b; b.kind = MiniMaxH3RefBlock::Kind::kImage; - b.latent_t = ls.t / dit_params.patch_size_t; - b.latent_h = ls.h / dit_params.patch_size_h; - b.latent_w = ls.w / dit_params.patch_size_w; + b.latent_t = ls.t; + b.latent_h = ls.h; + b.latent_w = ls.w; out_blocks->push_back(b); } } diff --git a/tests/vllm/models/minimax_h3_goldens.inc b/tests/vllm/models/minimax_h3_goldens.inc index d5f9bb788..f35819dbf 100644 --- a/tests/vllm/models/minimax_h3_goldens.inc +++ b/tests/vllm/models/minimax_h3_goldens.inc @@ -7257,6 +7257,641 @@ inline constexpr double kH3LadderR6MixFraction[] = { 1.0, }; +// --- section 5c: DiT-forward REF2VA rung (reference rows, section 8.9) --- +inline constexpr int64_t kH3Ref2vaDit_text_len = 5; +inline constexpr int64_t kH3Ref2vaDit_latent_t = 2; +inline constexpr int64_t kH3Ref2vaDit_latent_h = 8; +inline constexpr int64_t kH3Ref2vaDit_latent_w = 8; +inline constexpr int64_t kH3Ref2vaDit_audio_t = 3; +inline constexpr int64_t kH3Ref2vaDit_audio_channel = 2; +inline constexpr int64_t kH3Ref2vaDit_seq_len = 128; +inline constexpr int64_t kH3Ref2vaDit_num_img = 80; +inline constexpr int64_t kH3Ref2vaDit_num_audio = 10; +inline constexpr int64_t kH3Ref2vaDit_num_text = 5; +inline constexpr int64_t kH3Ref2vaDit_num_unique = 4; +inline constexpr int64_t kH3Ref2vaDit_video_width = 32; +inline constexpr int64_t kH3Ref2vaDit_first_target = 48; +inline constexpr int64_t kH3Ref2vaDit_num_blocks = 2; + +inline constexpr int64_t kH3Ref2vaDitBlockKinds[] = { + 0, 2, +}; + +inline constexpr int64_t kH3Ref2vaDitBlockRefAudioT[] = { + 0, 2, +}; + +inline constexpr int64_t kH3Ref2vaDitBlockLatentT[] = { + 0, 2, +}; + +inline constexpr int64_t kH3Ref2vaDitBlockLatentH[] = { + 8, 8, +}; + +inline constexpr int64_t kH3Ref2vaDitBlockLatentW[] = { + 8, 8, +}; + +inline constexpr int64_t kH3Ref2vaDitCuSeqlens[] = { + 0, 95, 128, +}; + +inline constexpr int64_t kH3Ref2vaDitImgPos[] = { + 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, + 17, 18, 19, 20, 25, 26, 27, 28, 29, 30, 31, 32, + 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, + 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, + 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, + 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, + 87, 88, 89, 90, 91, 92, 93, 94, +}; + +inline constexpr int64_t kH3Ref2vaDitAudioPos[] = { + 21, 22, 23, 24, 57, 58, 59, 60, 61, 62, +}; + +inline constexpr int64_t kH3Ref2vaDitTextPos[] = { + 0, 1, 2, 3, 4, +}; + +inline constexpr int64_t kH3Ref2vaDitUpdateMask[] = { + 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, +}; + +inline constexpr int64_t kH3Ref2vaDitAudioUpdateMask[] = { + 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, +}; + +inline constexpr int64_t kH3Ref2vaDitTokenTags[] = { + 1, 1, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0, 0, 2, 2, 2, + 2, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0, 0, 2, 2, 2, + 2, 2, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, -1, + -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, + -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, + -1, -1, -1, -1, -1, -1, -1, -1, +}; + +inline constexpr double kH3Ref2vaDitImgPositionIds[] = { + 0.0, 0.0, 0.0, 1.0, 0.0, 0.0, + 2.0, 0.0, 0.0, 3.0, 0.0, 0.0, + 4.0, 0.0, 0.0, 5.0, 0.0, 0.0, + 5.0, 0.0, 8.0, 5.0, 0.0, 16.0, + 5.0, 0.0, 24.0, 5.0, 8.0, 0.0, + 5.0, 8.0, 8.0, 5.0, 8.0, 16.0, + 5.0, 8.0, 24.0, 5.0, 16.0, 0.0, + 5.0, 16.0, 8.0, 5.0, 16.0, 16.0, + 5.0, 16.0, 24.0, 5.0, 24.0, 0.0, + 5.0, 24.0, 8.0, 5.0, 24.0, 16.0, + 5.0, 24.0, 24.0, 6.0, 0.0, 0.0, + 7.0, 0.0, 0.0, 6.0, 0.0, 24.0, + 7.0, 0.0, 24.0, 6.0, 0.0, 0.0, + 6.0, 0.0, 8.0, 6.0, 0.0, 16.0, + 6.0, 0.0, 24.0, 6.0, 8.0, 0.0, + 6.0, 8.0, 8.0, 6.0, 8.0, 16.0, + 6.0, 8.0, 24.0, 6.0, 16.0, 0.0, + 6.0, 16.0, 8.0, 6.0, 16.0, 16.0, + 6.0, 16.0, 24.0, 6.0, 24.0, 0.0, + 6.0, 24.0, 8.0, 6.0, 24.0, 16.0, + 6.0, 24.0, 24.0, 7.666666666666667, 0.0, 0.0, + 7.666666666666667, 0.0, 8.0, 7.666666666666667, 0.0, 16.0, + 7.666666666666667, 0.0, 24.0, 7.666666666666667, 8.0, 0.0, + 7.666666666666667, 8.0, 8.0, 7.666666666666667, 8.0, 16.0, + 7.666666666666667, 8.0, 24.0, 7.666666666666667, 16.0, 0.0, + 7.666666666666667, 16.0, 8.0, 7.666666666666667, 16.0, 16.0, + 7.666666666666667, 16.0, 24.0, 7.666666666666667, 24.0, 0.0, + 7.666666666666667, 24.0, 8.0, 7.666666666666667, 24.0, 16.0, + 7.666666666666667, 24.0, 24.0, 14.333333333333334, 0.0, 0.0, + 15.333333333333334, 0.0, 0.0, 16.333333333333336, 0.0, 0.0, + 14.333333333333334, 0.0, 24.0, 15.333333333333334, 0.0, 24.0, + 16.333333333333336, 0.0, 24.0, 14.333333333333334, 0.0, 0.0, + 14.333333333333334, 0.0, 8.0, 14.333333333333334, 0.0, 16.0, + 14.333333333333334, 0.0, 24.0, 14.333333333333334, 8.0, 0.0, + 14.333333333333334, 8.0, 8.0, 14.333333333333334, 8.0, 16.0, + 14.333333333333334, 8.0, 24.0, 14.333333333333334, 16.0, 0.0, + 14.333333333333334, 16.0, 8.0, 14.333333333333334, 16.0, 16.0, + 14.333333333333334, 16.0, 24.0, 14.333333333333334, 24.0, 0.0, + 14.333333333333334, 24.0, 8.0, 14.333333333333334, 24.0, 16.0, + 14.333333333333334, 24.0, 24.0, 16.0, 0.0, 0.0, + 16.0, 0.0, 8.0, 16.0, 0.0, 16.0, + 16.0, 0.0, 24.0, 16.0, 8.0, 0.0, + 16.0, 8.0, 8.0, 16.0, 8.0, 16.0, + 16.0, 8.0, 24.0, 16.0, 16.0, 0.0, + 16.0, 16.0, 8.0, 16.0, 16.0, 16.0, + 16.0, 16.0, 24.0, 16.0, 24.0, 0.0, + 16.0, 24.0, 8.0, 16.0, 24.0, 16.0, + 16.0, 24.0, 24.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, + 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, +}; + +inline constexpr float kH3Ref2vaDitUniqueTimesteps[] = { + 0.349999994f, 0.400000006f, 0.999000013f, 1.0f, +}; + +inline constexpr int64_t kH3Ref2vaDitInverseIndices[] = { + 1, 1, 1, 1, 1, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 3, 3, 3, + 3, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 0, 0, 0, + 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, +}; + +inline constexpr float kH3Ref2vaDitVideoLogits[] = { + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, -0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, -0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, 0.0f, -0.0f, -0.0f, + -0.0f, 0.0f, 0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, -0.0f, + 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, 0.0f, + -0.0f, -0.0f, 0.0f, 0.0f, 0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + -0.0f, -0.0f, 0.0f, -0.0f, -0.0f, -0.0f, + 0.0683853701f, -0.302417517f, -0.038751632f, 0.164269686f, 0.144495726f, 0.144483447f, + 0.100064225f, -0.120993569f, 0.24778907f, -0.0682222322f, -0.418082118f, 0.0550465211f, + 0.117635213f, -0.0350158215f, 0.201566875f, -0.269323349f, 0.143784687f, 0.439153582f, + 0.0689799413f, -0.126906976f, -0.262566924f, -0.208282739f, 0.162980795f, -0.216301113f, + -0.0509345904f, -0.288753808f, -0.160164252f, -0.35787797f, -0.164826363f, 0.346313477f, + 0.0152547453f, -0.162321746f, 0.147275582f, 0.0136015005f, -0.265664726f, -0.0501339063f, + -0.122233264f, -0.161970884f, -0.320988685f, -0.229727238f, 0.0238773078f, -0.373492241f, + 0.219281703f, 0.353419513f, 0.0820959583f, -0.290405542f, 0.236092895f, -0.232293069f, + -0.00797994528f, -0.195480913f, -0.404459476f, -0.16594553f, -0.0188375488f, 0.352597982f, + -0.0845848918f, 0.310537666f, 0.712878048f, 0.121931776f, 0.185330972f, -0.0474758409f, + -0.197808892f, -0.0836850256f, 0.0308345184f, -0.0812817514f, -0.211990595f, -0.0598963127f, + -0.663606226f, -0.0941413268f, -0.112184413f, 0.0805223137f, 0.353173345f, 0.236871123f, + 0.297253847f, -0.195954576f, 0.171666846f, -0.00247376133f, -0.217715636f, -0.0656633079f, + 0.197745726f, -0.241577417f, 0.341899782f, 0.133379176f, -0.183493152f, 0.274353564f, + 0.315879643f, -0.0837267637f, 0.141871229f, -0.230157375f, -0.16525878f, 0.123054147f, + 0.224795386f, -0.00915654935f, 0.388626635f, 0.36176306f, 0.0155411419f, 0.119069643f, + -0.29774037f, -0.336172312f, 0.189236134f, -0.131834626f, -0.0352207199f, 0.122107059f, + -0.038571503f, -0.0323736966f, -0.0521909632f, -0.0498498566f, -0.0529495627f, -0.235594153f, + 0.401605099f, -0.439865947f, 0.205289096f, 0.0929486826f, 0.380461097f, -0.175701171f, + -0.0382301584f, -0.0852285475f, -0.371930003f, -0.453418165f, 0.0812322572f, -0.0572505742f, + -0.211210132f, 0.284584612f, -0.31840679f, -0.0300634932f, -0.16785565f, -0.169625521f, + -0.0845148414f, 0.083503902f, -0.344625205f, -0.259694904f, -0.31962502f, -0.0723061189f, + 0.0434791856f, 0.225984335f, 0.0831942856f, 0.194909558f, 0.116919443f, -0.0229526255f, + -0.130863279f, 0.159771845f, -0.314576417f, 0.579768538f, -0.369040817f, 0.152375638f, + 0.167203724f, 0.00234315358f, 0.222248435f, -0.210842103f, 0.0525912531f, -0.374099731f, + 0.205424964f, 0.191809133f, -0.538637102f, -0.191430435f, 0.375532955f, 0.00743314251f, + -0.214858234f, 0.575971544f, 0.571781158f, 0.644073546f, -0.0727062598f, 0.329578906f, + 0.189360857f, 0.0490010493f, 0.18988727f, 0.424117237f, -0.0970942304f, 0.211925164f, + -0.457020879f, 0.0823278129f, 0.0228636097f, 0.0318842232f, -0.178942636f, 0.0309770517f, + -0.231418177f, 0.0424621627f, -0.0783122107f, -0.10549514f, 0.0736531541f, -0.0615496337f, + -0.43353498f, 0.00853781682f, 0.200606927f, 0.251038671f, 0.149479374f, -0.0423985012f, + 0.209346578f, 0.109508693f, -0.0976776928f, -0.143117517f, -0.0228238236f, -0.0958668217f, + -0.15824908f, 0.324258715f, 0.486981481f, 0.0123231411f, 0.292581886f, 0.0486408956f, + 0.119241826f, -0.135720536f, -0.0225959029f, 0.384353161f, -0.0677870587f, -0.275358915f, + 0.072921209f, -0.134437501f, -0.180474356f, -0.237697124f, -0.22815153f, 0.114845335f, + -0.0209006462f, 0.232880563f, 0.0124225738f, -0.0186875407f, 0.0798427686f, -0.395429343f, + 0.224594921f, 0.122757837f, -0.391109943f, 0.359779686f, -0.205828473f, -0.337036282f, + -0.237919688f, -0.284536421f, 0.282937706f, -0.122111931f, -0.463010877f, 0.117776662f, + -0.437004238f, -0.0837378055f, -0.0620765835f, -0.16449751f, 0.197769552f, -0.0143989306f, + 0.0837777853f, 0.633967578f, -0.424199462f, -0.272275627f, 0.294273823f, 0.102676652f, + 0.0154044796f, -0.199395254f, -0.358154416f, 0.052111648f, 0.291472942f, 0.186706215f, + -0.0807699263f, -0.0150378542f, 0.490030438f, -0.185432598f, 0.301605284f, 0.0358163081f, + -0.183107764f, 0.051285319f, -0.00640778802f, 0.134327263f, -0.251707643f, -0.307934135f, + -0.232302234f, 0.339599192f, 0.228932679f, -0.0516497865f, 0.0973690152f, 0.203088775f, + -0.0729266554f, 0.0511626899f, 0.000144550577f, 0.0338785164f, -0.387063324f, -0.216417551f, + -0.0992284715f, -0.0988189653f, 0.355625302f, 0.176144943f, -0.0491423979f, -0.127766266f, + 0.174693793f, -0.128540531f, -0.0853508264f, -0.328206599f, 0.00799026527f, -0.343757004f, + 0.0113086775f, -0.429229915f, 0.183775663f, 0.416798532f, 0.219785362f, -0.00861916319f, + 0.0766684189f, -0.0289828833f, 0.0124675659f, -0.144761413f, 0.208851889f, -0.042010285f, + -0.10018646f, 0.475548536f, -0.177227288f, -0.138401985f, 0.0543179587f, 0.0337854475f, + -0.140232295f, 0.266066313f, -0.461909264f, 0.12423408f, 0.0517351255f, -0.104465306f, + -0.186212137f, 0.149358675f, -0.477459401f, 0.22264035f, 0.121317066f, 0.263963252f, + -0.239272878f, -0.249439046f, 0.329441458f, 0.0869507939f, -0.279319435f, -0.17153962f, + 0.222491369f, 0.247594535f, -0.253046632f, 0.114690177f, -0.229009047f, 0.105326943f, + 0.24800387f, -0.0421342328f, -0.341411918f, -0.0309396945f, 0.0786876753f, 0.22674723f, + 0.178922445f, 0.0254796669f, 0.0279762521f, -0.0293078031f, 0.191333547f, -0.103343137f, + -0.210484281f, 0.22041665f, -0.316559494f, -0.0701857731f, 0.316479951f, -0.0534916967f, + -0.34901908f, -0.0668046474f, 0.136767641f, -0.368777752f, -0.208156958f, -0.133148685f, + -0.220756993f, 0.353303581f, 0.214356169f, 0.0446658768f, -0.135181546f, 0.0288653895f, + -0.574238241f, -0.0529859215f, -0.113185771f, 0.083465375f, -0.0400385782f, 0.394520044f, + 0.0563473366f, -0.211410508f, 0.227124974f, 0.195309356f, -0.315367937f, 0.150496691f, + -0.159350276f, 0.263272732f, 0.0848181248f, -0.315585256f, -0.0558170751f, -0.162136257f, + 0.272469342f, -0.0959581807f, -0.164749771f, 0.111178055f, -0.381809533f, -0.0964191109f, + 0.412722886f, 0.0383799709f, 0.131018683f, 0.144461587f, 0.354236603f, 0.486262858f, + -0.141378924f, -0.0846856982f, -0.431161821f, 0.240915135f, 0.144220352f, 0.355637401f, + 0.0150045194f, 0.149889365f, -0.0680524409f, 0.458886802f, -0.121422686f, -0.293058723f, + 0.0740835294f, -0.350761682f, -0.000894539058f, 0.384265631f, 0.501691937f, -0.568690419f, + -0.197880268f, 0.383607358f, -0.133265138f, -0.246917233f, 0.142275542f, -0.359980673f, + -0.20341593f, 0.0267644711f, 0.00455209054f, 0.191284716f, 0.0763618648f, 0.191542014f, + -0.0750670657f, 0.10258618f, 0.0322781354f, 0.0786032453f, 0.0372859389f, 0.315133661f, + 0.0527720973f, -0.0613773391f, 0.0218562018f, -0.242040932f, -0.138092279f, 0.0125654712f, + -0.0323075242f, 0.0119286245f, 0.179391429f, -0.376417994f, 0.123768032f, 0.0479975641f, + 0.0794122145f, -0.0943062007f, 0.293318123f, -0.289999366f, -0.00604027323f, 0.0492574275f, + -0.159546897f, 0.259360373f, 0.321536422f, 0.269150138f, 0.0425514504f, -0.112295106f, + 0.210114896f, -0.324028224f, -0.276506633f, -0.502986908f, 0.125669211f, 0.486499399f, + 0.215647027f, 0.0588063002f, 0.517121494f, 0.499570757f, -0.208031461f, -0.422520757f, + 0.195213497f, -0.0162700489f, -0.271749109f, -0.362408638f, -0.367944896f, 0.365849257f, + -0.224571928f, -0.0726166591f, -0.127671674f, -0.130294561f, 0.0724604204f, 0.123845138f, + -0.0142857963f, -0.351632327f, -0.0872258395f, -0.433053374f, 0.069206439f, -0.0635657087f, + -0.45191592f, -0.153549835f, -0.437477082f, -0.124400243f, -0.053455241f, -0.122584388f, + -0.225324452f, 0.254033536f, 0.0675216839f, 0.24424836f, 0.36300981f, 0.473406583f, + 0.00394695438f, -0.125183031f, 0.0919116214f, 0.25931567f, -0.373253047f, -0.295568824f, + -0.0877878293f, -0.093072176f, 0.0155781908f, 0.0134879816f, 0.120657958f, -0.308228433f, + 0.45113036f, 0.11469648f, 0.0884344801f, -0.505948842f, -0.0522445031f, -0.132650673f, + 0.293378562f, 0.189735577f, 0.0570731014f, -0.0358962305f, -0.0815464184f, -0.0155379828f, + 0.033543013f, -0.108646438f, 0.39084217f, 0.140258029f, -0.322957665f, 0.241837144f, + -0.079322435f, -0.190860569f, -0.0950520039f, -0.214459315f, 0.11007873f, -0.0272033475f, + -0.233143091f, 0.083770752f, -0.224430591f, 0.0297636762f, -0.0858487412f, 0.148154438f, + 0.274925381f, -0.516174912f, -0.0558241084f, 0.0203452967f, 0.0610451698f, 0.12354885f, + 0.135156929f, -0.369889885f, -0.0544916615f, 0.195365489f, -0.262851804f, 0.163814068f, + -0.00773257948f, -0.360242367f, -0.211752772f, -0.111801088f, -0.0420685299f, -0.114264913f, + -0.230256557f, 0.176710069f, 0.0181461405f, 0.355189413f, -0.0895749852f, 0.202274293f, + -0.333303422f, 0.208028451f, 0.0367314219f, -0.280425996f, 0.234683663f, -0.18459259f, + 0.0016088523f, 0.183981031f, 0.343330503f, -0.307200402f, -0.482079893f, 0.0588888898f, + -0.367795885f, -0.46325466f, 0.217538252f, -0.229203641f, -0.02572513f, -0.166667357f, + -0.0954447836f, -0.0504280254f, 0.234223723f, 0.351125568f, 0.419561714f, 0.189305723f, + 0.433872432f, 0.1925807f, 0.230249897f, 0.208006889f, -0.234283373f, 0.398693293f, + -0.0266196653f, -0.735273957f, 0.0375122651f, 0.531214535f, -0.152077883f, 0.0566421375f, + -0.156760558f, 0.144407481f, 0.19984585f, 0.203780904f, -0.311629772f, -0.19840008f, + -0.0637202561f, -0.00841059629f, -0.085103251f, -0.255225062f, -0.0157641433f, -0.455667108f, + 0.215079248f, -0.214075074f, -0.324859917f, 0.0850959197f, -0.100063421f, 0.0245425124f, + -0.0617397651f, -0.0284083616f, 0.117511086f, -0.430113375f, -0.245704815f, 0.266748369f, + -0.426474184f, 0.191904724f, -0.0902626738f, -0.179249242f, -0.14185667f, 0.269174814f, + 0.0731702298f, 0.0244926065f, 0.304561794f, -0.178739473f, 0.404347539f, 0.104836911f, + 0.0794076174f, 0.182192028f, -0.244228855f, -0.0985781401f, -0.0711366609f, -0.154120564f, + 0.00148715172f, -0.322248906f, -0.171978265f, -0.35612008f, -0.173487872f, -0.0974601135f, + -0.0349931195f, 0.495550543f, -0.0840475708f, -0.0137973446f, 0.210874677f, 0.530388415f, + 0.327586949f, 0.087290816f, 0.251690567f, 0.12922065f, -0.265099019f, -0.0897473767f, + 0.168217152f, 0.294207513f, -0.230185747f, -0.10079743f, -0.24121961f, 0.621004403f, + -0.277811617f, 0.304204434f, -0.480657458f, -0.073165141f, 0.227232575f, -0.353338093f, + 0.108307078f, -0.267820895f, -0.0586439036f, -0.140542626f, -0.506769955f, -0.207242727f, + -0.293678731f, 0.140986368f, -0.321905941f, -0.178248376f, -0.0189245921f, 0.233141184f, + 0.0144207329f, 0.0374977924f, -0.0298009776f, -0.0187456645f, -0.178520069f, 0.0900144652f, + -0.136663333f, -0.22344844f, -0.161729604f, -0.178276613f, 0.180460215f, -0.143943295f, + 0.264885068f, -0.131917179f, 0.0760837197f, 0.00199598446f, 0.169522703f, -0.182114229f, + -0.240528062f, -0.292006582f, -0.257208854f, -0.184911445f, -0.180202171f, -0.0955756456f, + 0.0585438125f, 0.299170345f, -0.275772959f, -0.0881355256f, 0.167331561f, 0.118079729f, + 0.123744898f, 0.0943271369f, -0.15085952f, -0.0310158413f, 0.0696849748f, -0.369367093f, + 0.345203668f, 0.0206907596f, -0.0575700179f, -0.0594177842f, 0.20023118f, -0.247949317f, + -0.113475315f, 0.0662569404f, -0.0854699835f, 0.0255322456f, 0.038116768f, 0.054551892f, + 0.0530406013f, -0.061619021f, 0.275688112f, 0.31479308f, -0.431890756f, 0.149548814f, + 0.0106597329f, 0.51549691f, -0.117748447f, 0.381832868f, 0.129150569f, -0.014040567f, + -0.367831856f, -0.405496061f, 0.206991881f, 0.186802089f, -0.079846926f, 0.0672726631f, + -0.417520225f, 0.0177549105f, -0.187768459f, 0.0967462659f, 0.0470340177f, -0.217641011f, + 0.109605163f, 0.262091219f, -0.217447892f, 0.358378381f, -0.220027924f, -0.0725047067f, + 0.486974627f, -0.0757337958f, -0.0932847038f, -0.188518077f, 0.412554622f, 0.123037584f, + 0.00904804375f, 0.0266182069f, -0.0333613716f, 0.283378035f, 0.306822538f, 0.0844236538f, + 0.325418681f, 0.24778606f, 0.06994766f, -0.159247726f, -0.05939392f, -0.0376499742f, + -0.275105f, -0.302713692f, 0.177701116f, 0.019407386f, 0.0974135473f, -0.132893205f, + -0.144380867f, -0.187692612f, 0.287970036f, 0.142574504f, -0.0875463486f, 0.124142841f, + 0.44675824f, -0.106950924f, 0.209746674f, -0.0239889566f, -0.0525392964f, 0.105237357f, + 0.517884016f, -0.252345651f, -0.159279168f, -0.104454607f, -0.235742241f, 0.049866911f, + -0.0346819498f, 0.136104912f, -0.21668154f, 0.0271730237f, -0.097074382f, -0.473531455f, + -0.263651907f, -0.196691155f, -0.231367067f, -0.0498870723f, 0.190022916f, 0.216294527f, + 0.0497019142f, 0.136194825f, 0.127996832f, 0.000966975698f, -0.346956402f, 0.391804755f, + -0.159567311f, -0.356071234f, 0.308894366f, -0.431172103f, 0.0154399378f, -0.0902391896f, + 0.2527439f, -0.384141892f, 0.0580948219f, 0.234232366f, -0.157120019f, 0.00522045884f, + 0.185017392f, -0.379306257f, 0.332881451f, 0.303770751f, -0.14273864f, 0.391424984f, + 0.0640040487f, -0.367125124f, -0.263305843f, -0.171271145f, 0.0866907313f, 0.0189590063f, + -0.0693977028f, -0.395112038f, -0.0552371219f, 0.472774297f, -0.0595461875f, -0.595905721f, + 0.0499059148f, -0.0943939984f, -0.148128241f, -0.0782312006f, 0.127581924f, 0.349330634f, + 0.188588336f, 0.191385821f, -0.109989338f, -0.204765335f, 0.0953036919f, 0.292454451f, + -0.140175134f, -0.259315729f, 0.0543871857f, -0.0723760575f, -0.107576102f, 0.184027806f, + -0.130524263f, 0.132918626f, -0.565447927f, -0.277017653f, -0.262732267f, 0.280904531f, + -0.0255430304f, -0.283022583f, -0.0675111711f, 0.088514559f, -0.0600302517f, -0.371696413f, + 0.35115099f, -0.118149601f, 0.0927317888f, -0.0773862973f, 0.302174479f, -0.181334913f, + 0.120502181f, -0.510554135f, 0.383172333f, -0.125465766f, 0.400373548f, 0.0924518108f, + 0.0103693027f, 0.256517053f, 0.0399101041f, -0.0209869202f, -0.0347967334f, -0.0394569449f, + 0.0497096479f, 0.0862422287f, -0.0743332654f, -0.343686134f, 0.434015214f, -0.0332979411f, + -0.212395355f, -0.618067265f, -0.274069846f, -0.151017576f, -0.769001305f, -0.0707277432f, + -0.0773043558f, 0.224555731f, 0.183413833f, 0.357738137f, 0.185269415f, -0.3884269f, + 0.328703761f, 0.118487f, -0.277752787f, -0.309168428f, -0.000840255991f, -0.0585820451f, + 0.318006754f, -0.0996575952f, -0.183232978f, 0.135472789f, 0.409136266f, -0.123814963f, + -0.0102677746f, 0.00986894406f, -0.0138626127f, 0.188263178f, 0.482776135f, -0.070366092f, + 0.309928894f, 0.147729412f, 0.204476506f, 0.13174209f, -0.0776322335f, -0.0236048661f, + -0.090790607f, -0.23091656f, 0.225184396f, -0.250982314f, -0.0154102109f, 0.35865429f, + 0.272865057f, -0.135762826f, -0.124283724f, 0.101410985f, -0.0650683418f, 0.395460278f, + -0.270016313f, 0.16310963f, -0.164707735f, 0.115821786f, 0.281719029f, -0.0899642855f, + 0.342835158f, 0.20235087f, 0.290058732f, 0.148220435f, -0.549945116f, -0.0722222626f, + 0.0630169138f, 0.333838522f, -0.272376925f, -0.30472216f, 0.167422429f, 0.352758825f, + -0.0453360751f, -0.0451698154f, -0.13505964f, 0.319697738f, -0.0716901794f, -0.216980904f, + 0.170809716f, 0.246790707f, -0.440044165f, 0.301941216f, -0.164368525f, 0.251053154f, + -0.0263977163f, -0.0241019037f, -0.332880497f, -0.194535911f, 0.422218442f, -0.191390306f, + -0.198879272f, 0.0121599156f, -0.151236311f, 0.254273176f, 0.261875808f, 0.0575069524f, + 0.448803067f, 0.133254528f, 0.224921077f, 0.320049226f, 0.0270879604f, -0.123557568f, + 0.0152334347f, -0.155950576f, -0.244281605f, -0.389123112f, -0.118373983f, 0.030262487f, + -0.098713249f, 0.0529759303f, 0.0852231681f, -0.197771206f, 0.12738581f, -0.158155128f, + -0.22960636f, 0.121003941f, -0.00396165531f, 0.259542286f, -0.389531016f, 0.114301428f, + 0.543102324f, -0.430179417f, 0.261698157f, -0.168545619f, -0.271545202f, -0.172932088f, + 0.196557343f, 0.0356157683f, -0.114286683f, 0.307875961f, 0.0889413729f, 0.0934487954f, + -0.173833415f, 0.0831822753f, 0.153423592f, 0.326139539f, +}; + +inline constexpr float kH3Ref2vaDitAudioLogits[] = { + 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, 0.0f, + 0.0f, -0.0f, -0.0f, -0.0f, -0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, -0.0f, 0.0f, -0.0f, + 0.176047862f, -0.190033063f, -0.0444384366f, -0.155746758f, 0.049266465f, -0.188856632f, + 0.286717683f, 0.213227376f, 0.0361472815f, -0.0795138106f, 0.182128355f, 0.197248593f, + 0.239325434f, -0.363577902f, -0.232493371f, -0.130726367f, -0.138408616f, -0.115149826f, + 0.0115625989f, -0.0590781905f, 0.174234748f, -0.355059713f, 0.38107577f, -0.399034411f, + 0.026262613f, -0.482597589f, -0.187431455f, -0.121354371f, -0.338573933f, -0.241146058f, + -0.0321198553f, 0.145789146f, 0.186963856f, 0.122371733f, -0.0511230156f, 0.220132664f, +}; + +inline constexpr float kH3Ref2vaDitVideoMixDelta[] = { + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, + 0.316139281f, 4.83095646e-05f, 0.000112984329f, 4.98592854e-05f, 8.40723515e-05f, 0.000100553036f, + 4.6774745e-05f, 5.96940517e-05f, 2.23815441e-05f, 0.000121861696f, 4.44948673e-05f, 0.000111013651f, + 0.000124149024f, 4.42266464e-05f, 8.45491886e-05f, 0.000102758408f, 0.000112757087f, 8.92579556e-05f, + 0.000140696764f, 2.83420086e-05f, 5.90085983e-05f, 6.74165785e-05f, 7.19353557e-05f, 0.000113494694f, + 7.69495964e-05f, 6.07967377e-05f, 8.75890255e-05f, 5.80847263e-05f, 0.000112004578f, 8.77827406e-05f, + 0.000183090568f, 2.85357237e-05f, +}; + +inline constexpr double kH3Ref2vaDitMixFraction[] = { + 1.0, +}; + // --- section 6: request planning (time_request.py + shape resolution) --- inline constexpr int64_t kH3PlanFrameCases = 14; diff --git a/tests/vllm/models/test_minimax_h3.cpp b/tests/vllm/models/test_minimax_h3.cpp index 708e8e8f1..3d28ac7d3 100644 --- a/tests/vllm/models/test_minimax_h3.cpp +++ b/tests/vllm/models/test_minimax_h3.cpp @@ -934,6 +934,188 @@ TEST_CASE("minimax_h3: DiT device-vs-host forward holds at the REAL head_dim=128 } } +// The DiT-forward REF2VA rung (goldens section 5c). Every other DiT-forward gate +// runs the FL2VA layout (keyframe-cond prefix sharing the target frame grid, no +// audio reference rows). The ref2va reference-row assembly -- image/video/audio +// reference blocks PREPENDED with their OWN position grid, a separate audio update +// mask, the pinned reference rows carrying the CONDITION timestep -- was never +// forwarded through the DiT, which is exactly how the section 8.9 ref2va grid hid: +// the pure-math layout was gated (goldens section 2) but nothing forwarded a +// ref2va layout end-to-end. This case builds the SAME upstream ref2va layout with +// BuildMiniMaxH3PackedSequenceRef2va, runs the host AND device DiT forward against +// the RefDiT oracle, and asserts the reference rows are masked out and the target +// rows spatially couple. It pins the port's ref2va packing and its DiT forward over +// reference rows to upstream in one gate. +TEST_CASE("minimax_h3: DiT-forward REF2VA rung matches upstream (reference rows, mixing)") { + const std::unique_ptr weights = BuildGoldenWeights(); + const MiniMaxH3DitParams& p = weights->params; + vt::Queue q{Cpu(), nullptr}; + const MiniMaxH3DitDeviceWeights staged = StageMiniMaxH3DitWeights(q, p, weights->views); + + // (1) Reconstruct the reference blocks the generator emitted (kind: 0=image, + // 1=audio, 2=video_audio) and rebuild the identical ref2va layout. + std::vector blocks(static_cast(vllm_test::kH3Ref2vaDit_num_blocks)); + for (size_t i = 0; i < blocks.size(); ++i) { + blocks[i].kind = static_cast(vllm_test::kH3Ref2vaDitBlockKinds[i]); + blocks[i].ref_audio_t = vllm_test::kH3Ref2vaDitBlockRefAudioT[i]; + blocks[i].latent_t = vllm_test::kH3Ref2vaDitBlockLatentT[i]; + blocks[i].latent_h = vllm_test::kH3Ref2vaDitBlockLatentH[i]; + blocks[i].latent_w = vllm_test::kH3Ref2vaDitBlockLatentW[i]; + } + const MiniMaxH3PackedSequence packed = vllm::BuildMiniMaxH3PackedSequenceRef2va( + vllm_test::kH3Ref2vaDit_text_len, vllm_test::kH3Ref2vaDit_latent_t, + vllm_test::kH3Ref2vaDit_latent_h, vllm_test::kH3Ref2vaDit_latent_w, + vllm_test::kH3Ref2vaDit_audio_t, blocks, vllm_test::kH3Ref2vaDit_audio_channel); + + const int64_t seq_len = packed.seq_len; + const int64_t num_img = static_cast(packed.img_pos.size()); + const int64_t num_audio = static_cast(packed.audio_pos.size()); + const int64_t num_text = static_cast(packed.text_pos.size()); + const int64_t video_width = p.video_row_width(); + REQUIRE(seq_len == vllm_test::kH3Ref2vaDit_seq_len); + REQUIRE(num_img == vllm_test::kH3Ref2vaDit_num_img); + REQUIRE(num_audio == vllm_test::kH3Ref2vaDit_num_audio); + REQUIRE(num_text == vllm_test::kH3Ref2vaDit_num_text); + + CheckI64(packed.cu_seqlens, vllm_test::kH3Ref2vaDitCuSeqlens, + std::size(vllm_test::kH3Ref2vaDitCuSeqlens)); + CheckI64(packed.img_pos, vllm_test::kH3Ref2vaDitImgPos, std::size(vllm_test::kH3Ref2vaDitImgPos)); + CheckI64(packed.audio_pos, vllm_test::kH3Ref2vaDitAudioPos, + std::size(vllm_test::kH3Ref2vaDitAudioPos)); + CheckI64(packed.text_pos, vllm_test::kH3Ref2vaDitTextPos, + std::size(vllm_test::kH3Ref2vaDitTextPos)); + CheckI64(packed.update_mask, vllm_test::kH3Ref2vaDitUpdateMask, + std::size(vllm_test::kH3Ref2vaDitUpdateMask)); + CheckI64(packed.audio_update_mask, vllm_test::kH3Ref2vaDitAudioUpdateMask, + std::size(vllm_test::kH3Ref2vaDitAudioUpdateMask)); + CheckI64(packed.token_tags, vllm_test::kH3Ref2vaDitTokenTags, + std::size(vllm_test::kH3Ref2vaDitTokenTags)); + // The FP64 position grid over reference + target rows feeds RoPE; gate BIT-EXACT. + REQUIRE(packed.img_position_ids.size() == std::size(vllm_test::kH3Ref2vaDitImgPositionIds)); + for (size_t i = 0; i < packed.img_position_ids.size(); ++i) { + CHECK(packed.img_position_ids[i] == vllm_test::kH3Ref2vaDitImgPositionIds[i]); + } + + // (2) Build the DiT inputs: scatter the same rung-seeded random rows, and carry + // the ref2va audio_update_mask + the emitted timestep partition. + auto build_inputs = [&](std::vector& x, std::vector& audio_x, + std::vector& prompt, std::vector& unique, + std::vector& inverse, std::vector& refiner_cu, + MiniMaxH3DitInputs& in) { + x.assign(static_cast(seq_len * video_width), 0.0f); + const std::vector vrows = MakeParam("ref2va_dit.r2v.video_rows", num_img * video_width, 1.0); + for (int64_t r = 0; r < num_img; ++r) { + std::memcpy(x.data() + packed.img_pos[static_cast(r)] * video_width, + vrows.data() + r * video_width, static_cast(video_width) * sizeof(float)); + } + audio_x.assign(static_cast(seq_len * p.audio_latents_dim), 0.0f); + const std::vector arows = + MakeParam("ref2va_dit.r2v.audio_rows", num_audio * p.audio_latents_dim, 1.0); + for (int64_t r = 0; r < num_audio; ++r) { + std::memcpy(audio_x.data() + packed.audio_pos[static_cast(r)] * p.audio_latents_dim, + arows.data() + r * p.audio_latents_dim, + static_cast(p.audio_latents_dim) * sizeof(float)); + } + prompt = MakeParam("ref2va_dit.r2v.prompt_embeds", num_text * p.text_dim, 1.0); + unique.assign(vllm_test::kH3Ref2vaDitUniqueTimesteps, + vllm_test::kH3Ref2vaDitUniqueTimesteps + std::size(vllm_test::kH3Ref2vaDitUniqueTimesteps)); + inverse.assign(vllm_test::kH3Ref2vaDitInverseIndices, + vllm_test::kH3Ref2vaDitInverseIndices + std::size(vllm_test::kH3Ref2vaDitInverseIndices)); + refiner_cu = {0, static_cast(num_text), static_cast(num_text)}; + in = MiniMaxH3DitInputs{}; + in.seq_len = seq_len; + in.x = x.data(); + in.audio_x = audio_x.data(); + in.img_position_ids = packed.img_position_ids.data(); + in.unique_timesteps = unique.data(); + in.num_unique_timesteps = static_cast(unique.size()); + in.inverse_indices = inverse.data(); + in.token_tags = packed.token_tags.data(); + in.prompt_embeds = prompt.data(); + in.img_pos = packed.img_pos.data(); + in.num_img_pos = num_img; + in.audio_pos = packed.audio_pos.data(); + in.num_audio_pos = num_audio; + in.text_pos = packed.text_pos.data(); + in.num_text_pos = num_text; + in.infer_out_pos = packed.img_pos.data(); + in.num_infer_out_pos = num_img; + in.update_mask = packed.update_mask.data(); + in.audio_update_mask = packed.audio_update_mask.data(); + in.cu_seqlens = packed.cu_seqlens.data(); + in.num_cu_seqlens = static_cast(packed.cu_seqlens.size()); + in.refiner_cu_seqlens = refiner_cu.data(); + in.num_refiner_cu_seqlens = static_cast(refiner_cu.size()); + }; + + std::vector x, audio_x, prompt, unique; + std::vector inverse; + std::vector refiner_cu; + MiniMaxH3DitInputs in; + build_inputs(x, audio_x, prompt, unique, inverse, refiner_cu, in); + REQUIRE(static_cast(unique.size()) == vllm_test::kH3Ref2vaDit_num_unique); + + // (3) HOST + DEVICE forward vs the RefDiT oracle. + const MiniMaxH3DitOutputs host = MiniMaxH3DitForward(Cpu(), p, weights->views, in, vt::DType::kF32); + const double hv = MaxAbsDiff(host.video_logits, vllm_test::kH3Ref2vaDitVideoLogits, host.video_logits.size()); + const double ha = MaxAbsDiff(host.audio_logits, vllm_test::kH3Ref2vaDitAudioLogits, host.audio_logits.size()); + INFO("host: video max|diff| = " << hv << ", audio max|diff| = " << ha); + CHECK(hv <= 2e-5); + CHECK(ha <= 2e-5); + + const MiniMaxH3DitOutputs dev = MiniMaxH3DitForwardDevice(q, p, staged.weights, in, vt::DType::kF32); + const double dv = MaxAbsDiff(dev.video_logits, vllm_test::kH3Ref2vaDitVideoLogits, dev.video_logits.size()); + const double da = MaxAbsDiff(dev.audio_logits, vllm_test::kH3Ref2vaDitAudioLogits, dev.audio_logits.size()); + INFO("device: video max|diff| = " << dv << ", audio max|diff| = " << da); + CHECK(dv <= 2e-5); + CHECK(da <= 2e-5); + + // The pinned reference rows (video AND audio) must be masked to zero in the output. + for (int64_t r = 0; r < num_img; ++r) { + if (packed.update_mask[static_cast(r)]) continue; + for (int64_t i = 0; i < video_width; ++i) + CHECK(host.video_logits[static_cast(r * video_width + i)] == 0.0f); + } + for (int64_t r = 0; r < num_audio; ++r) { + if (packed.audio_update_mask[static_cast(r)]) continue; + for (int64_t i = 0; i < p.audio_latents_dim; ++i) + CHECK(host.audio_logits[static_cast(r * p.audio_latents_dim + i)] == 0.0f); + } + + // (4) SPATIAL-MIXING probe: perturb the first video-TARGET token and require the + // response to match the oracle and to couple every other target token through the + // packed bidirectional attention (the property the section 8.9 grid lacked). + std::vector xp, ap, pp, up; + std::vector ip; + std::vector rc; + MiniMaxH3DitInputs inp; + build_inputs(xp, ap, pp, up, ip, rc, inp); + const int64_t pert_pos = packed.img_pos[static_cast(vllm_test::kH3Ref2vaDit_first_target)]; + for (int64_t i = 0; i < video_width; ++i) xp[static_cast(pert_pos * video_width + i)] += 1.0f; + const MiniMaxH3DitOutputs pert = MiniMaxH3DitForward(Cpu(), p, weights->views, inp, vt::DType::kF32); + std::vector delta(static_cast(num_img), 0.0f); + for (int64_t r = 0; r < num_img; ++r) { + float worst = 0.0f; + for (int64_t i = 0; i < video_width; ++i) { + const size_t idx = static_cast(r * video_width + i); + worst = std::max(worst, std::abs(pert.video_logits[idx] - host.video_logits[idx])); + } + delta[static_cast(r)] = worst; + } + CHECK(MaxAbsDiff(delta, vllm_test::kH3Ref2vaDitVideoMixDelta, delta.size()) <= 2e-5); + int64_t responded = 0, others = 0; + for (int64_t r = 0; r < num_img; ++r) { + if (!packed.update_mask[static_cast(r)]) continue; + if (r == vllm_test::kH3Ref2vaDit_first_target) continue; + ++others; + if (delta[static_cast(r)] > kLadderMixEps) ++responded; + } + const double frac = others > 0 ? static_cast(responded) / static_cast(others) : -1.0; + INFO("ref2va mix fraction port=" << frac << " oracle=" << vllm_test::kH3Ref2vaDitMixFraction[0]); + CHECK(frac == doctest::Approx(vllm_test::kH3Ref2vaDitMixFraction[0]).epsilon(1e-9)); + CHECK(vllm_test::kH3Ref2vaDitMixFraction[0] >= 0.999); +} + // Brick H3-2b. The DEVICE-RESIDENT forward runs the same graph with every // activation in device memory. It is NOT bit-identical to the CPU reference and // does not claim to be -- it reuses the tuned SHARED vt:: ops (vt::RmsNorm reduces @@ -2740,6 +2922,29 @@ TEST_CASE("minimax_h3: fl2va keyframe conditioning is wired and load-bearing") { CHECK(blocks[0].latent_t >= 1); REQUIRE(!ref_rows.empty()); + // INVARIANT (RED-first, the spec section 8.9 ref2va-grid root cause): the + // encoded reference row COUNT must equal exactly what the packed layout + // allocates for the reference span. MiniMaxH3EncodeReferenceImages returns the + // RAW VAE-latent grid in its block; BuildMiniMaxH3PackedSequenceRef2va (mirroring + // upstream) applies the DiT [1,2,2] patch division itself. If the encode + // PRE-divides (the bug that shipped the grid), the layout under-allocates the + // reference by patch_h*patch_w and the denoise loop silently truncates the + // pinned reference to its first quarter -- coherence-destroying but non-throwing, + // because keyframe_cond_rows is then LONGER than the layout and passes the + // pin-loop's `>=` check. Nothing else in the suite couples encoded-rows to + // layout-rows: the section-2 layout gate hand-builds unpatched blocks, and the + // denoise round-trip below only checks finiteness + motion. This is the coupling. + const vllm::MiniMaxH3PackedSequence ref_layout = + vllm::BuildMiniMaxH3PackedSequenceRef2va(req.text_len, req.latent_t, req.latent_h, + req.latent_w, req.audio_t, blocks, + req.audio_channel); + int64_t layout_ref_rows = 0; + for (const uint8_t upd : ref_layout.update_mask) layout_ref_rows += (upd == 0 ? 1 : 0); + const int64_t block_grid_rows = + (blocks[0].latent_h / p.patch_size_h) * (blocks[0].latent_w / p.patch_size_w); + CHECK(layout_ref_rows == block_grid_rows); + CHECK(layout_ref_rows * p.video_row_width() == static_cast(ref_rows.size())); + const int64_t frame_rows = (req.latent_h / p.patch_size_h) * (req.latent_w / p.patch_size_w); const std::vector prompt = MakeParam("ref2va.prompt", req.text_len * p.text_dim, 0.2); const std::vector nv = @@ -2814,6 +3019,18 @@ TEST_CASE("minimax_h3: fl2va keyframe conditioning is wired and load-bearing") { REQUIRE(iblocks.size() == 1); CHECK(vrows.size() > irows.size()); + // Same encoded-vs-layout row-count invariant as the image path (section 8.9): + // the RAW-latent block dims the video encode emits must produce a layout whose + // reference span holds exactly the encoded rows. A patch-pre-division here would + // silently truncate the pinned video reference, so pin it in the suite too. + { + const vllm::MiniMaxH3PackedSequence vlayout = vllm::BuildMiniMaxH3PackedSequenceRef2va( + 4, 2, 4, 4, 4, {vb}, 2); + int64_t vref_rows = 0; + for (const uint8_t upd : vlayout.update_mask) vref_rows += (upd == 0 ? 1 : 0); + CHECK(vref_rows * p.video_row_width() == static_cast(vrows.size())); + } + vllm::MiniMaxH3T2vaRequest req; req.text_len = 4; req.latent_t = 2;