From 37a716e18fc52d1b6925f1cfed980d72b103168f Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Fri, 7 Aug 2026 06:40:55 +0000 Subject: [PATCH] =?UTF-8?q?diag(minimax-h3):=20the=20#94=20NVFP4=20ref2va-?= =?UTF-8?q?grid=20residual=20DIAGNOSED=20=E2=80=94=20NO=20discrete=20load-?= =?UTF-8?q?path=20defect?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Ran #94's prescribed identical-weights activation diff (NVFP4-bf16 stream vs the FL2VA-GGUF-bf16 control, byte-identical inputs via an encode-once pe.f32) plus direct WEIGHT fingerprints, through an env-gated per-stage hook in MiniMaxH3DitForwardDevice (VT_H3_ACT_DUMP; byte-inert unset). The result REFUTES the "second NVFP4-render-path defect" hypothesis. Instrument (byte-inert): a reusable H3DumpFingerprint over every embed/scatter/time stage, every block's adaln+post-attn+post-mlp, the final heads, the RoPE cos/sin cache, the block-0 attention internals (qkv/split/qknorm/rope/core/out_proj), and input-independent fingerprints for every island, bias, output head, q/k-norm, and block-0/refiner-0 projection — the classes #94's oracle never sampled. Finding: EVERY weight, bias, fp32 island, output head and q/k-norm loads quant-noise-close to the coherent GGUF (qkv/out/fc/adaln 2-6% rms = Q3K-vs-NVFP4 quantizer variance, sign-correct; all biases <0.1% apart incl adaln_b), and the RoPE cache is BYTE-IDENTICAL — no scramble/transpose/mis-stride/wrong-dtype/ wrong-shape. Both arms run identical code, so the grid is 100% the per-weight NVFP4-vs-Q3K quant delta on the SAME weights. Divergence first appears at the token refiner + the block-0 attention INPUT (not RoPE, not a GEMM/norm weight) and amplifies chaotically through the 50-block stack, driven by the Qwen massive text activation (condition_proj absmax ~7.4e4); final latents are DECORRELATED (sample-rel-L2 >1, not scale). Render A/B re-confirmed in the same byte-inert build: NVFP4 t2va = pale patch grid, FL2VA-GGUF t2va = coherent orange cat. Conclusion: not a loader fix. The residual is the community NVFP4 checkpoint's quantization fidelity (Star Ultimate Model Converter Pro lineage; corr 0.85-0.94 to the coherent Q3_K) times the DiT's massive-activation sensitivity. Clean- reference disambiguation is blocked (bf16 132GiB = OOM on one GB10; a same-finetune REF2VA-GGUF control = 23G-disk-blocked); the path forward is an official modelopt-NVFP4 checkpoint. The #94 nibble fix stands. The fp4-resident Marlin arm's separate grid is untouched (wiring-gated-only residual). Records: spec §8.12 + §8.2 row, STATUS/BENCHMARKS/FEATURES H3 rows, benchmark-record, NOW, state. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] --- .agents/NOW.md | 2 +- .agents/benchmark-record.md | 14 ++ .agents/specs/minimax-h3.md | 28 ++- .agents/state.md | 20 +++ docs/BENCHMARKS.md | 2 +- docs/FEATURES.md | 4 +- docs/STATUS.md | 2 +- .../models/minimax_h3_device.cpp | 164 +++++++++++++++++- 8 files changed, 228 insertions(+), 8 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index 0940947a..b4bc37a8 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -17,7 +17,7 @@ Working head: `row/backend-rocm-w0` (#41). Prior: benchmark checkpoint | Laguna NVFP4 / DeepSeek-V4 decode | **Both CLOSED, byte-exact, default-ON**: 1.03x vLLM, 1.144x ds4 | Laguna vLLM K-run when convenient | | f32-out GEMV audit | Only laguna + ds4 bf16 tower affected; gate models unaffected | Re-verify ds4 tower same-tool | | Invocation-parity prevention | CI guard + checklist landing | Merge; build-verify `kGemvHeuristicAlgos` on dgx | -| MiniMax-H3 lane | **fl2va COHERENT; ref2va NVFP4 fp4-nibble loader bug FIXED** (byte-verified, #94) | grid PERSISTS: 2nd NVFP4-render-path defect (ckpt/params byte-match coherent GGUF) §8.11 | +| MiniMax-H3 lane | **fl2va COHERENT; ref2va NVFP4 grid DIAGNOSED (#95): NO loader bug** | weights/islands/RoPE all quant-noise-close to coherent GGUF; residual = community-NVFP4 quant fidelity §8.12 | | 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 4b01a5cc..bfd59ab2 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -14469,3 +14469,17 @@ Followed #93's residual (the ref2va NVFP4 checkpoint/loader) with an independent **RE-ATTRIBUTION (corrects #93 again).** The nibble fix CHANGES the output (low-first severe grid → high-first pale grid, so the swap IS applied) yet EVERY NVFP4 render grids — t2va, ref2va, AND fl2va-with-keyframe — while the fl2va-GGUF control (SAME weights, SAME params, verified byte-for-byte) renders a coherent cat in the same build and same task. So: the checkpoint CONTENT is sound (byte-matches the coherent GGUF), the loader dequant is now byte-correct (binary dump == oracle == GGUF-sign), the derived params are identical, and output-pinning (keyframe) does NOT rescue it (rules out free-generation divergence). The residual grid is therefore a SECOND, independent defect in the NVFP4 render PATH itself — the device stream / forward, NOT the checkpoint, NOT the fp4 nibble order, NOT free-gen. The fp4-resident Marlin arm additionally grids differently (a THIRD, Marlin-specific issue — that path was only ever wiring-gated, never correctness-gated). The nibble fix is the objectively-correct dequant (the file IS high-first) and lands default-ON (A/B via `VT_H3_NVFP4_LOWNIBBLE=1`); it is byte-verified but not yet render-validated, blocked on the second defect. Next diagnostic: layer-by-layer intermediate-activation diff of the NVFP4-bf16 stream vs the GGUF-bf16 stream (identical weights) to locate where they diverge — the streamers are structurally identical except the dequant source and the island read (bf16-disk vs f16-disk), so the divergence is a candidate. Artifacts `~/h3fp4/{t2va_fix,t2va_nofix,ctrl_t2va,ab_bf16fix,ab_fix,ab_nofix,kf_nvfp4}.mp4`. + +## MiniMax-H3 NVFP4 ref2va grid — the #94 residual DIAGNOSED: NO discrete load-path defect (identical-weights activation diff) (`row/H3-NVFP4-STREAM-DIFF` PR #95, 2026-08-07, dgx sm_121a) + +Ran #94's prescribed next diagnostic — a layer-by-layer activation diff of the NVFP4-bf16 stream vs the FL2VA-GGUF-bf16 control on BYTE-IDENTICAL inputs — plus direct WEIGHT fingerprints. The result REFUTES the "second NVFP4-render-path defect" hypothesis: there is **no discrete load-path materialization bug**. All weights load correctly; the residual grid is checkpoint-quantization fidelity × the DiT's massive-activation sensitivity. + +**Method.** Env-gated per-stage fingerprint hook inside `MiniMaxH3DitForwardDevice` (`VT_H3_ACT_DUMP`, byte-inert unset; `minimax_h3_device.cpp`): summary stats (mean/rms/absmax/finite) + a fixed positional-sample spread (catches transpose/scramble even when rms matches), at every embed/scatter/time stage, every block's adaln+post-attn+post-mlp, the final heads, the RoPE cos/sin cache, the block-0 attention internals (qkv/split/qknorm/rope/core/out_proj), and input-independent WEIGHT fingerprints for every island, bias, output head, q/k-norm and block-0/refiner-0 projection. Both arms: ONE small t2va forward (`--denoise-only --steps 2`, 256×256/22f, `--partition fl2va`) with byte-identical `--prompt-embeds pe.f32` (encoded once, reused → no encoder; both arms feed the SAME text conditioning) + identical seeded noise; final video checksum is DETERMINISTIC per arm (GGUF 9411.61, NVFP4 1160.19, reproduced across three runs). + +**Every loaded tensor is quant-noise-close; there is no scramble/transpose/mis-stride/wrong-dtype/wrong-shape.** Direct fingerprints (NVFP4 vs GGUF): `qkv_proj` rms 8.94e-2 vs 8.46e-2, `q_norm`/`k_norm` essentially identical (rms 1.086 vs 1.086), `out_proj` 8.47e-2 vs 8.14e-2, `fc1`/`fc2`/`adaln_w` within 2-6% (Q3_K-vs-NVFP4 quantizer variance, sign-correct); ALL biases identical (`adaln_b` 0.1554 vs 0.1554, `final_adaln_b`, `condition_b`, time biases all <0.1% apart); ALL fp32 islands and output heads (`video_patch_w` rms 0.22059 both, `video_out_w`, `audio_out_w`, `time_in/out_w`) quant-noise-close; and the **RoPE cos/sin cache is IDENTICAL** (same samples, rules out `rope.inv_freq`). Both arms run the IDENTICAL forward code, so the grid is 100% attributable to the per-weight NVFP4-vs-Q3_K quantization difference on the SAME weights (#94 proved sign-agreement 1.000 → same weights; dequant byte-verified). + +**Divergence profile (paired stage rms-ratio N/G, and sample-relative-L2):** embed video/audio/text-condition/adaln = 1.00 (quant noise); FIRST real divergence at the **token refiner** `text_refined` (0.92, sampRelL2 0.48); then `block.0.normed_pre_attn` (0.89 — the attention INPUT already diverges) → `qkv_out` (0.78) → `core_out` (0.78) → `out_proj` (0.60); amplifies through the 50-block stack; final latents are DECORRELATED (sampRelL2 >1, not scale-related). Driver: the Qwen massive text activation (`condition_proj` output absmax **~7.4e4**) makes the refiner/attention chaotically sensitive to the few-% weight quant delta. NVFP4 stays systematically lower-magnitude (→ "pale"). This is CHAOTIC trajectory divergence, not a fixable scramble/scale. + +**Render A/B re-confirmed in the same byte-inert build (256×256/22f/12steps, `pe.f32`, `VT_H3_ACT_DUMP` unset).** NVFP4 t2va = pale multicolour PATCH GRID every frame (0/10/21, per-patch uniform = degenerate latent); FL2VA-GGUF t2va = a COHERENT photorealistic orange cat on a windowsill. The diagnostic hooks are byte-inert (reproduce #94's outputs exactly). + +**CONCLUSION.** The mission's "2nd load-path defect" does not exist — our loader materializes the community NVFP4 file's weights faithfully (byte-verified). The residual is the **community NVFP4 checkpoint's quantization fidelity** (metadata `converted_by: "Star Ultimate Model Converter Pro"` — same dubious-converter lineage as the #94 nibble bug; corr to the coherent Q3_K only 0.85-0.94, well below the >0.99 a clean 4-bit quant gives) interacting with the DiT's massive-activation sensitivity. Definitively separating "poor community quant" from "inherent t2va-OOD sensitivity of these fl2va/ref2va finetunes" needs a clean reference — a bf16 ground truth (132 GiB host-f32 = OOM on one GB10) or a same-finetune REF2VA-GGUF control (disk-blocked, 23 GiB free) — both currently blocked; the path forward is an official modelopt-NVFP4 checkpoint, not a loader change. The #94 nibble fix stands as the objectively-correct dequant. The fp4-resident Marlin arm's separate grid stays a distinct, wiring-gated-only residual (untouched, per the mission). Artifacts `~/h3fp4/diff/{nvfp4_t2va,gguf_t2va}.mp4` + `{gguf,nvfp4}{2,3,4}.txt` fingerprints. diff --git a/.agents/specs/minimax-h3.md b/.agents/specs/minimax-h3.md index e322fdcb..9c137269 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 reference-row assembly FIXED + gated (§8.10) — ref2va still grids; the NVFP4 fp4 nibble-order loader bug is now FOUND+FIXED (byte-verified, §8.11) but the grid PERSISTS from a 2nd NVFP4-render-path defect (checkpoint content + loader dequant + params all byte-match the coherent GGUF) | **PARTIAL** (vision→conditioning scatter + ref2va assembly + NVFP4 nibble loader DONE; residual = the 2nd NVFP4-render-path defect, §8.11) | +| 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) — ref2va still grids; the NVFP4 fp4 nibble-order loader bug is now FOUND+FIXED (byte-verified, §8.11) but the grid PERSISTS — DIAGNOSED (§8.12): NO discrete load-path defect (activation diff + weight fingerprints: every weight/bias/island/head/RoPE loads quant-noise-close to the coherent GGUF, byte-identical RoPE); residual = the community NVFP4 checkpoint's quant fidelity × the DiT's Qwen massive-activation sensitivity, chaotic — not a loader fix | **PARTIAL** (vision→conditioning scatter + ref2va assembly + NVFP4 nibble loader DONE; ref2va NVFP4 render residual = checkpoint quant fidelity, no loader bug, §8.12) | | 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) | @@ -775,6 +775,32 @@ dump the ref2va target-row VAE-input latent adjacency-cosine (like #77 did for t 0.95) to confirm the target rows are white, and A/B the reference-row condition-noise vs a clean anchor. +## 8.12 THE #94 RESIDUAL DIAGNOSED — no discrete load-path defect (`row/H3-NVFP4-STREAM-DIFF` PR #95, 2026-08-07) + +Ran #94's prescribed identical-weights activation diff (NVFP4-bf16 stream vs FL2VA-GGUF-bf16 +control, byte-identical inputs) plus direct WEIGHT fingerprints via an env-gated per-stage hook +in `MiniMaxH3DitForwardDevice` (`VT_H3_ACT_DUMP`, byte-inert unset). **Result: there is NO discrete +load-path materialization bug.** Every weight, bias, fp32 island, output head, q/k-norm, and the +RoPE cos/sin cache load quant-noise-close to the coherent GGUF (no scramble/transpose/mis-stride/ +wrong-dtype/wrong-shape); the RoPE cache is byte-identical; the dequant is byte-verified (§8.11). +Both arms run the IDENTICAL forward, so the grid is 100% attributable to the per-weight +NVFP4-vs-Q3_K quantization difference on the SAME weights. + +The divergence FIRST appears (beyond quant noise) at the **token refiner** and the block-0 +attention INPUT — NOT RoPE, NOT a projection/norm weight — and amplifies chaotically through the +50-block stack, driven by the Qwen massive text activation (`condition_proj` output absmax ~7.4e4); +the two arms' final latents are DECORRELATED (sample-relative-L2 >1, not scale-related). Render A/B +re-confirmed in the same byte-inert build: NVFP4 t2va = pale patch grid, FL2VA-GGUF t2va = coherent +orange cat. **The residual is the community NVFP4 file's quantization fidelity** (same +`Star Ultimate Model Converter Pro` lineage as the #94 nibble bug; corr to the coherent Q3_K only +0.85-0.94) times the DiT's massive-activation sensitivity — a CHECKPOINT-quality issue, not a +loader fix. Definitively separating "poor community quant" from "inherent t2va-OOD sensitivity of +these fl2va/ref2va finetunes" needs a clean bf16 ground truth (132 GiB host-f32 = OOM on one GB10) +or a same-finetune REF2VA-GGUF control (disk-blocked, 23 GiB free); the path forward is an official +modelopt-NVFP4 checkpoint. The fp4-resident Marlin arm's separate grid stays a distinct, +wiring-gated-only residual (untouched). Full forensics + the divergence profile: the +`row/H3-NVFP4-STREAM-DIFF` benchmark-record entry. + ## 9. W-OAI — the `/v1/videos` OpenAI (Sora) WIRE SHAPE, 2026-08-06 Row `SERVE-VIDEOS-OAI` (engine matrix, Serving surface), claim diff --git a/.agents/state.md b/.agents/state.md index 2e6441fc..af672077 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -40716,3 +40716,23 @@ nibble, fp4, or free-gen. fp4-resident Marlin arm grids differently again (3rd, wiring-gated). NEXT: layer-by-layer activation diff of the NVFP4-bf16 stream vs the GGUF-bf16 stream (identical weights; differ only in dequant source + island read bf16-disk vs f16-disk). Records: spec §8.11 + §8.2 row, STATUS/BENCHMARKS/FEATURES H3 rows, benchmark-record, NOW. Box left clean. + +[2026-08-07] H3-NVFP4-STREAM-DIFF (#95) — the #94 residual DIAGNOSED: there is NO discrete load-path +defect. Ran #94's prescribed identical-weights activation diff (NVFP4-bf16 stream vs FL2VA-GGUF-bf16 +control, byte-identical inputs via encode-once pe.f32) + direct WEIGHT fingerprints, via an env-gated +per-stage hook in MiniMaxH3DitForwardDevice (VT_H3_ACT_DUMP, byte-inert unset). FINDING: every weight, +bias, fp32 island, output head, q/k-norm loads quant-noise-close to the coherent GGUF (qkv/out/fc/adaln +2-6% rms = Q3K-vs-NVFP4 quantizer variance, sign-correct; ALL biases <0.1% apart incl adaln_b; islands ++heads quant-noise-close), and the RoPE cos/sin cache is BYTE-IDENTICAL — no scramble/transpose/mis-stride/ +wrong-dtype/wrong-shape. Both arms run the IDENTICAL forward, so the grid is 100% from the per-weight +NVFP4-vs-Q3K quant delta on the SAME weights. Divergence FIRST appears at the token refiner (text_refined +ratio 0.92, sampRelL2 0.48) + the block-0 attention INPUT (normed_pre_attn 0.89), NOT RoPE/GEMM/norm +weights, and amplifies chaotically through the 50-block stack driven by the Qwen massive text activation +(condition_proj absmax ~7.4e4); final latents DECORRELATED (sampRelL2 >1, not scale) -> chaotic trajectory. +Render A/B re-confirmed same byte-inert build: NVFP4 t2va pale patch grid, FL2VA-GGUF t2va coherent orange +cat. CONCLUSION: not a loader fix — residual = the community NVFP4 checkpoint's quant fidelity (Star +Ultimate Model Converter Pro lineage; corr 0.85-0.94 to the coherent Q3K) x the DiT's massive-activation +sensitivity. Clean-reference disambiguation blocked (bf16 132GiB = OOM on 1 GB10; same-finetune REF2VA-GGUF += 23G-disk-blocked); path forward = an official modelopt-NVFP4 ckpt. fp4-resident Marlin arm's separate +grid untouched (wiring-gated-only). Diagnostic hooks landed byte-inert. Records: spec §8.12 + §8.2 row, +STATUS/BENCHMARKS/FEATURES H3 rows, benchmark-record, NOW. Box left clean. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 56296a43..92c27eff 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -317,7 +317,7 @@ built on it rather than keeping the flattering one. | MiniMax-H3 Thor render speed (sm_110, no FA2) | **34.6 s/step** at 864x480/124f/50 steps on Q4_K_M, **16.6x** off 574.5 (render ~28 min, was ~8 h). Landed: warp-per-query, chunked warp reduce-scatter (1.76x), bf16 `mma.sync` (9.82x) | Shared-memory K/V tiling (23% SLOWER) and register Q-blocking (-0.8%) both measured and REVERTED: memory traffic is not the bound (one head's K+V is 3.9 MB against 32 MB of L2) | | MiniMax-H3 render duration (audio halving) | **CLOSED**: a 124-frame render silently muxed as **61**. The decoded audio ran half the video's duration and the muxer passes `-shortest`. Every structural check passed: shapes were self-consistent, just halved | Gated on the duration invariant (latent steps / 40 Hz equals the video duration), not on shape self-consistency, which a halved pipeline satisfies | | MiniMax-H3 quantization floor | **Use Q4_K_M, not Q3_K_M.** H3's split-half RoPE produces channel-wise magnitude outliers 3-bit cannot hold; a controlled A/B (same prompt, seed, code) turned a murky lattice-covered silhouette into a photoreal close-up | Per-tensor mixed precision, if a smaller footprint is ever owed | -| MiniMax-H3 image conditioning (`row/H3-CONDITIONED-E2E`, `row/H3-VISION-SCATTER`, `row/H3-REF2VA-ASSEMBLY`, `row/H3-NVFP4-LOADER-DIFF`) | **fl2va COHERENT; ref2va assembly + NVFP4 fp4-nibble loader bug FIXED (byte-verified).** File packs fp4 HIGH-first, loader read low-first→scrambled (corr 0.000→1.000 sign-agree vs coherent GGUF); nibble-swap (§8.11) | grid RE-ATTRIBUTED AGAIN: with the fix the render STILL grids (t2va/ref2va/keyframe) while FL2VA-GGUF renders coherent on IDENTICAL weights+params → residual = a 2nd NVFP4-render-path defect, NOT ckpt/nibble/fp4 (§8.11) | +| MiniMax-H3 image conditioning (`row/H3-CONDITIONED-E2E`, `row/H3-VISION-SCATTER`, `row/H3-REF2VA-ASSEMBLY`, `row/H3-NVFP4-LOADER-DIFF`, `row/H3-NVFP4-STREAM-DIFF`) | **fl2va COHERENT; ref2va NVFP4 nibble loader bug FIXED (byte-verified); grid residual DIAGNOSED §8.12** | Activation diff + fingerprints (PR #95): NO load-path defect; all weights/islands/RoPE quant-noise-close to the coherent GGUF; grid = community-NVFP4 quant fidelity, not a loader fix. See benchmark-record | | 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 c4594683..afbbef08 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -133,7 +133,7 @@ they sit outside the gated list above. |---|---|---|---| | Voxtral audio (`VoxtralForConditionalGeneration`) | Voxtral-Mini-3B-2507 | near-tie-robust 16/16 vs vLLM 0.25.0 | decode 0.97x (beats vLLM); encoder TTFT ~17x, pending | | Whisper audio encoder | openai/whisper-small; whisper-large-v3 (Voxtral cfg) | encoder tower 77/77; large-v3 tower 203/203 | pending | -| MiniMax-H3 DiT (`MiniMaxH3DiTModel`, vllm-omni lane) | MiniMax-H3 (33.1B video+audio) | portable 69/69; t2va+fl2va COHERENT; ref2va NVFP4 fp4-nibble loader bug FIXED (byte-verified), grid persists from a 2nd NVFP4-render-path defect (§8.11) | FP4/Marlin landed; ref2va NVFP4 render 2nd-defect + speed pending | +| MiniMax-H3 DiT (`MiniMaxH3DiTModel`, vllm-omni lane) | MiniMax-H3 (33.1B video+audio) | portable 69/69; t2va+fl2va COHERENT; ref2va NVFP4 grid DIAGNOSED (§8.12): NO loader bug (weights/islands/RoPE all quant-noise-close to the coherent GGUF); residual = community-NVFP4 quant fidelity, not a loader fix | FP4/Marlin landed; ref2va NVFP4 render blocked on checkpoint quant (needs official modelopt NVFP4), speed pending | | MTP speculator | Qwen3.6-27B, Qwen3.6-35B-A3B | token-identical to vLLM `mtp` at c1 | ~4% faster c1; +16% output tput (MoE) | | DFlash block-diffusion | Qwen3 (DFlash draft) | near-tie e2e 27/27 vs vLLM | 2.9x over spec-off, 1.003x vs vLLM DFlash-on | | DeepSeek-V4 MTP | DeepSeek-V4-Flash (nextn head) | lossless 5/5; real-model weight-blocked | pending | @@ -161,7 +161,7 @@ model architecture is wired. | Image | ✅ correctness-gated | ✅ | ✅ | ◐ | | Video | ✅ correctness-gated | ✅ | ✅ | ☐ | | Audio | ✅ correctness-gated | ✅ | ◐ | ◐ | -| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va+fl2va COHERENT on GB10; ref2va NVFP4 fp4-nibble loader bug FOUND+FIXED (byte-verified), grid persists from a 2nd NVFP4-render-path defect (ckpt/params byte-match coherent GGUF) (§8.11) | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | +| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va+fl2va COHERENT on GB10; ref2va NVFP4 grid DIAGNOSED (§8.12): NO loader bug (weights/islands/RoPE all quant-noise-close to the coherent GGUF); residual = community-NVFP4 quant fidelity | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | | 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 5e85586d..32a11b44 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -85,7 +85,7 @@ token-for-token correctness against the pinned oracle. | OLMo-3 dense (dual rope, interleaved sliding window) | Implemented, oracle-blocked | Loads + runs in our engine (dual rope: plain sliding + YaRN full-attn, per-layer sliding window); no SACRED gate: vLLM 0.25.0 oracle cannot run OLMo-3-1025-7B (`KeyError: 'rope_theta'`; transformers 5.13.1 nests `rope_parameters` per layer-type, no flat `rope_theta`; run-verified W0 2026-07-26) | | Laguna-S-2.1 MoE (`LagunaForCausalLM`, 118B/8B) | **BINDING 2026-08-04: 87% of vLLM (37.55 vs 43.10, SAME-TOOL nsys both engines); the whole +3.1 ms/step is the bf16 M=1 GEMV bucket (2/3 o_proj, ~196-204 vs 139 us/call, identical `gemvx` kernel); attention/MoE/glue tied or ours-ahead. Invocation match (bf16-out `cublasGemmEx`) A/B'd = WASH, ruled out; ROOT CAUSE FOUND 2026-08-04 (`VT_LAGUNA_RESIDENT_BF16W`): the bf16 projections read UNIFIED/ATS host memory, not `cudaMalloc`'d device memory — staging them device-resident (byte-exact ids) gives 38.8→44.6 tok/s (o_proj 194→131, lm_head 2410→1620 us/call), parity+ vs vLLM 43.1; **default-ON** (flip smoke-verified: canonical byte-exact ids, 44.6 clean-median). Earlier ceiling/diffuse verdicts below were cross-tool artifacts.** **REAL vLLM BAR ESTABLISHED (2026-07-31, `CLAIM-LAGUNA-VLLM-NVFP4`): FIRST-EVER vLLM Laguna run** — prior numbers (incl. the correctness oracle) were all llama.cpp, never vLLM. vLLM on official `poolside/Laguna-S-2.1-NVFP4` (single GB10, greedy, eager, MARLIN backend forced via `VLLM_TEST_FORCE_FP8_MARLIN=1` because the auto-default `FLASHINFER_CUTLASS` needs an absent `nvcc`): **~18.8 tok/s** (64-tok steady) — a LOWER bound. Our GGUF-Q4_K engine = 7.7 tok/s (vLLM ~2.4×); llama.cpp GGUF = 27.8 (still fastest at batch-1). llama.cpp is now a labeled SECONDARY "beat best-in-class GGUF" note; vLLM-NVFP4 is the headline bar. TRUE apples-to-apple still owes OUR NVFP4 Laguna forward arm (same tensor-core path as 27B/35B) — bring-up W-plan SPEC'D in `.agents/specs/laguna-nvfp4-arm-2026-07-31.md` (~85% reuse of the 35B NVFP4 W4A4 MoE infra + a name-map; bf16 attn/dense + fp4 experts; N1-N5 bricks, DGX-gated). **N1-scaffold LANDED (2026-07-31):** additive `LagunaMoeWeights.experts_{gate,up,down}_fp4` + `shared_{gate,up,down}_fp4` (`Nvfp4Weight`, mirror qwen3_5), dead until the N1 loader; CPU build clean + `test_laguna_scaffold` 8/8·167 unchanged. **N1b loader IMPLEMENTED (2026-07-31, build-verified):** `LoadLagunaForCausalLMWeights` (`laguna_weights.cpp`) replaces the `VT_CHECK(false)` stub — resolver + per-layer `LoadBf16Direct` (attn/dense/norms/embed/lm_head/router/shared-expert) + F32 `e_score_correction_bias` + `LnLoadCtNvfp4Raw` W4A4 experts. Name-map + dtypes VERIFIED against the real `poolside/Laguna-S-2.1-NVFP4` index (router `mlp.gate` BF16, bias F32, experts W4A4, shared-expert BF16). **N1b RUN-VERIFIED (2026-07-31):** loader round-trips a synthetic NVFP4 checkpoint byte-identically (`test_laguna_nvfp4_loader` 2/2·29; full detail in the benchmark record). **N2 FORWARD-BRANCH LANDED + CPU-GATED (2026-07-31):** `LqGemmNvfp4Fp4` (per-expert TRUE-W4A4: `ScaledFp4Quant(input_global_scale_inv)`→`MatmulNvfp4Fp4(alpha)`, unified-memory pattern like `LqGemm`) + `LagunaFfnBlock` branches on `fp4=!experts_gate_fp4.empty()` (routed experts fp4; keep-quant grouped fast-path gated off `!fp4`; bf16 attn/dense/router/shared-expert/lm_head unchanged) + both `LagunaForwardGguf{,Cached}` guards relaxed to `has_gguf_weights||has_nvfp4_weights`. **CORRECTION:** routed experts are W4A4 ⇒ per-expert `MatmulNvfp4Fp4`, NOT the grouped W4A16 `MoeGroupedGemmNvfp4` (grouped W4A4 deferred to N5 speed). `test_laguna_nvfp4_loader` 3/3·61 (added a forward run-gate: fp4 MoE branch runs through the real `LagunaForwardGguf` → finite+deterministic logits + routed-experts-consumed); `test_laguna_scaffold` 8/8 unchanged (GGUF byte-identical). **N3 DRIVER LANDED + CPU-SMOKE-VERIFIED (2026-07-31):** `examples/laguna_gen` auto-detects a safetensors DIRECTORY (→ NVFP4: `LoadHfConfig(config.json)` + `LoadLagunaForCausalLMWeights` + `LagunaForwardGguf{,Cached}`) vs a `.gguf` FILE (→ keep-quant), sharing the greedy loop; `--token-ids` bypass the tokenizer for the id-vs-golden gate. Verified on a synthetic NVFP4 dir with a REAL config.json (exercises the `LoadHfConfig`→`ParseLagunaParams` seam the loader test bypassed) → `has_nvfp4=1`, KV-cache decode runs finite. **N4 RAN on GB10 (2026-08-01) — the arm works end-to-end; correctness coherent+near-tie, speed 120× off.** git-archived `84fab587` → clean CUDA build (`121a`) → `laguna-gen --gpu` on the real 67 GiB `ckpt` with vLLM's exact prompt ids injected (`2,785,9626,377,15360,395`, captured via the HF tokenizer). Two GB10 memory fixes landed to run: release the mmap'd shards after the loader's memcpy-copy (114→67 GiB RSS), and create the CUDA context BEFORE the load (the 67 GiB reclaimable page cache otherwise starves `cudaStreamCreate`). **Correctness:** ours `22345 83 350 71070 395 340 9626 372 1703 …` vs golden `22345 83 290 350 674 330 5541 966 340 9626 377 15360 …` — **first 2 tokens match vLLM exactly**, then near-tie divergence; coherent ("France is" = 9626/377/15360; shares golden vocab). EXPECTED: our TRUE-W4A4 (fp4 activations) vs the MARLIN golden's W4A16 (bf16 activations) — different precision, not a bug. **Speed: 6.34 s/tok (0.16 tok/s), prefill 17.3s — ~120× slower than vLLM 18.8.** ROOT CAUSE (source-confirmed): `LqGemmNvfp4Fp4` uses the generic `vt::MatmulNvfp4Fp4` = the hand-written EMULATION CUDA kernel, NOT the cutlass sm120a fp4 tensor-core path the 27B/35B W4A4 use (`MatmulNvfp4Fp4DirectD`); + per-expert loop + per-GEMM host sync + no device residency. **nsys (2026-08-01) trace-confirmed + refined:** only 2 GPU kernels — `MatmulNvfp4Fp4Naive` = 99.3% of GPU time + fp4-quant 0.7%; GPU busy only ~18% of wall. NO bf16 GEMM on the GPU ⇒ `LqGemm`'s bf16 branch runs the host `MatmulNK` reference on the CUDA queue (attention/dense/router/shared/lm_head are CPU-bound, ~4.8 s/tok) — a second lever the source scan missed. **N5 LEVER #2 LANDED (2026-08-01) — 16× decode.** Routed the bf16 tower (attention/dense/router/shared/lm_head) off the host `MatmulNK` onto the GPU (`LqGemm` bf16 branch: `vt::CastBf16` the small activation + `vt::MatmulBT` bf16×bf16→f32, weight stays bf16 — no per-token `ReadF32` of `lm_head [100352,H]`): **decode 6.34 → 0.39 s/tok (16.3×; 0.16 → 2.56 tok/s), prefill 17.3 → 2.24s**; coherence preserved (near-tie). CPU path unchanged (run-gate byte-identical). **N5 LEVER #1 LANDED (2026-08-01) — native fp4 tensor-core, another ~2×.** The engine's native sm120a fp4 tensor-core MMA (`MatmulNvfp4Fp4Native`, `mma.sync kind::mxf4nvf4`) reads the same linear scale layout `LqGemmNvfp4Fp4` produces — it was gated OFF behind `VT_NVFP4_FP4_NATIVE`; the Laguna driver now defaults it ON (scoped; 27B/35B untouched). **decode 0.39 → ~0.20-0.24 s/tok (~2×; ~4.2-5.0 tok/s)**; coherent (byte-identical ids to the emulation path — numerically equivalent), first token matches the golden. **Cumulative N5: 0.16 → ~4.5 tok/s (~28×), now ~4× from vLLM 18.8.** **Device-resident MoE block LANDED + MEASURED (2026-08-01, `LagunaMoeResidentFp4`, `VT_LAGUNA_RESIDENT_MOE` default-ON):** the whole token's routed experts as ONE async device chain (fp4-quant→GEMM gate/up, `MoeSiluMul`, →down stacked, ONE `MoeCombine`), draining once vs ~Pk×3 syncs. **Speed EAGER-NEUTRAL (0.20 s/tok)** — empirically confirms the ds4 precedent (per-op syncs overlap GPU compute; wall is GPU-serial-bound; the graph is the payoff). **CORRECTNESS WIN: golden-token match 2 → 13** (the device `MoeSiluMul`/`MoeCombine` mirror vLLM's fused MoE faithfully). Lands default-ON (better correctness, no speed cost, graph prerequisite). **CORRECTED CEILING (from the measured state): a perfect decode graph caps at ~5.9 tok/s** (GPU already ~87% busy at 0.20 s/tok), still 3.3× short of vLLM 18.8 — the graph is necessary but NOT sufficient; the remaining 3.3× is KERNEL EFFICIENCY (native fp4 MMA ~302µs/M=1 expert GEMM vs vLLM's tuned cutlass sm120a fp4 + fused norm/quant/silu). Parity = TWO campaigns: (A) device-resident+graph → ~5.9; (B) cutlass DirectD experts + fused ops + M=1-tuned GEMV → the rest. **CAMPAIGN-B FIRST BRICK LANDED (2026-08-01): coalesced M=1 fp4 GEMV** (`MatmulNvfp4Fp4Gemv`, one warp/column, coalesced weight-row reads, `VT_NVFP4_FP4_GEMV` default-ON) — same-binary A/B: **decode 0.20 → 0.15 s/tok (1.33×; → ~6.7 tok/s), prefill 1.14 → 0.86s**, coherent+near-tie. **Cumulative this session: 0.16 → ~6.7 tok/s (~42×), now ~2.8× from vLLM 18.8.** (ILP variant `kCpw=4` measured SLOWER — 0.21 s/tok, occupancy loss > activation-reuse gain — reverted to `kCpw=1`; kernel kept templated as a re-measurable knob.) **ncu of the GEMV (sudo): sm__throughput 35-71%, DRAM n/a — COMPUTE/LATENCY-bound, not BW-bound.** Corrects the earlier "~6× BW → ~16-17 tok/s" estimate: the next GEMV lever is HARDWARE fp4 dequant (`cvt.e2m1x2`), not more bandwidth. Parity (18.8) is a multi-brick campaign (decode graph + fused norm/quant + hardware-dequant GEMV), not one more kernel. **B0 hw-fp8 SCALE-decode: MEASURED NEGATIVE, reverted (2026-08-01, `ab7a1c1e`).** Replacing the GEMV's per-byte software fp8-e4m3 group-scale decode (`F8E4M3ToF32Dev`/`ldexpf`) with hardware `cvt.rn.f16.e4m3` (`__nv_fp8_e4m3`→float) is bit-exact (ids byte-identical on the real ckpt) but paging-immune ncu shows it NEUTRAL-to-slightly-WORSE (grid768 41.2 vs 41.9µs tie; mean 53.6 vs 49.4µs) — GPU `ldexpf` is a cheap exponent-bit add, not a libcall. NOTE this is the fp8 SCALE decode, NOT the fp4-e2m1 WEIGHT dequant (the `kE2M1` `__constant__` LUT); the LUT→arithmetic/`cvt.e2m1x2` weight-dequant is a SEPARATE still-open lever (spec brick B1). Also: end-to-end wall-clock is unusable for kernel A/B here (67 GiB unified reload swings TPOT 0.16↔1.08 s/tok run-to-run) — kernel-duration ncu is the only honest anchor. **★ B2 SCOPED + DE-RISKED (2026-08-01, zero-DGX) — the real 18.8 lever:** vLLM's 18.8 bar is MARLIN W4A16 (`VLLM_TEST_FORCE_FP8_MARLIN=1`), which is LOW-M-optimized (decode-correct, unlike a tensor-core W4A4 GEMM that wastes M=1 tile rows). The engine already ships the EXACT kernel `vt::MoeGroupedGemmNvfp4Marlin` (1:1 lift of vLLM `moe_wna16_marlin_gemm`) + shared `MarlinRepackExpertWeight`, and qwen3_5 (27B/35B) already routes its NVFP4 experts through it (default-ON `VT_NVFP4_MARLIN`, 16/16-vs-oracle, +22% gate/+80% decode) via `BuildMoeMarlinResident`. So B2 = mirror that for `LagunaMoeWeights.experts_*_fp4` (a `BuildLagunaMoeMarlinResident` reusing the shared repack + route `LagunaFfnBlock`'s fp4 branch to the Marlin grouped GEMM, GEMV kept as the `=0` escape hatch) — pure reuse, no new kernel, matches vLLM's exact W4A16 numerics. **B2 IMPLEMENTED (2026-08-01, `3c49ef37`) — COMPILES CLEAN on GB10 sm_121a, runtime bug pending.** `LagunaMoeResidentMarlin` + `BuildLagunaMoeMarlinResident` (laguna.cpp, `#ifdef VT_MARLIN_NVFP4`) reconstruct the MoE Marlin path over the SHARED `dense_nvfp4::Dev`/`DBuf`/`ResidentNvfp4` + shared `vt::cuda` Marlin repack/align ops + `vt::MoeGroupedGemmNvfp4Marlin`; SACRED 27B/35B path BYTE-UNTOUCHED; gated `VT_LAGUNA_MARLIN_MOE=1` **default-OFF** (zero regression to the default GEMV path). Compiles clean on the full CUDA build. RUN: loads OK (48 layers, 256 experts) but the FIRST FORWARD device-faults silently on the Marlin path — a layout/param bug (suspects: `MoeCombine` bf16-in/f32-out dtype, the down-GEMM reusing the gate/up align, or the fp4-original free omitted → mem ~doubles). NEXT: `compute-sanitizer` localize → fix → near-tie vs the vLLM-Marlin golden + kernel-duration ncu → flip default-ON. Default path unaffected. **UPDATE (`22d6e146`): added the qwen3_5-style fp4-original free after repack** (device transients + host bytes; peak was ~3× the expert tower → past the 119 GiB pool → null-alloc → silent fault the likely cause); compiles clean. The runtime gate stayed INCONCLUSIVE this session (contended/orphaned processes on the shared box, no captured ids) — rerun on a clean uncontended session, compute-sanitizer if it still faults. **★★ B2 VALIDATED on GB10 (2026-08-01, with the mem-free fix): RUN_EXIT=0, coherent, first 13 generated tokens MATCH the vLLM-Marlin golden EXACTLY** (`22345 83 290 350 674 330 5541 966 340 9626 377 15360 81` — the best Laguna-NVFP4 correctness yet, W4A16 matching vLLM's config). **Steady-state decode 0.10 s/tok = ~10 tok/s** (steps 10-17 all 0.10; the TPOT-0.56 average is warmup-polluted — the DevicePool warms over ~9 decode steps then reuses). vs the GEMV path's 6.7 tok/s = **~1.5× faster; the gap to vLLM 18.8 closes from ~3× to ~1.9×.** Memory flat (7.9 GiB host RSS — the fp4-original free worked; it also fixed the first-forward fault). Still `VT_LAGUNA_MARLIN_MOE=1` default-OFF. TO DONE: move the lazy Marlin-resident build (216s first-forward, 48L×256E repack) to model-LOAD time → clean warm A/B + ncu → flip default-ON → matrix/roadmap. Remaining ~1.9×: vLLM graphs its decode (ours still eager) — decode CUDA-graph is the next lever. **REPRODUCED 3× (reproduction gate MET): GB10 runs deterministic — first 18-20 tokens byte-identical, steady-state 0.10 s/tok confirmed each — so the ~10 tok/s + golden-match is gated, not a single sample.** **#234 item (1) — load-time resident-build LANDED (`LagunaBuildMarlinResidents`, called from the example after load; mirrors vLLM process_weights_after_loading): builds all 48L×256E Marlin residents at LOAD so the repack is not a first-token TTFT spike. Fixed an anon-namespace linkage bug (public fn was defined with internal linkage → moved outside the anon namespace); BUILD CLEAN + links on GB10 sm_121a, default-OFF. Runtime prewarm-fires-at-load timing UNVERIFIED this session (repeated ssh-drops ate the run capture); the forward's lazy build is the validated fallback so it cannot regress. Owed: one clean run to confirm the build moved to load + then flip default-ON.** **★★ DONE (2026-08-01): Marlin is now the UNCONDITIONAL DEFAULT (`LagunaMarlinMoeEnabled` default-ON; `=0` is a code-level A/B opt-out no user needs) — "it just works" with NO env. Confirmed in a no-env GB10 run captured via tmux: `MARLIN residents built at load in 238.4s`, prefill 14.78s (build moved OUT of first-forward), golden-matching ids, steady-state 0.10 s/tok = ~10 tok/s (4th reproduction), RSS ~5-8 GiB. So a default Laguna-NVFP4 load on GB10 gets vLLM's own W4A16 Marlin decode (~10 tok/s, ~1.9× from vLLM 18.8) with zero flags. The 238s load-time repack is a one-time cost (mirrors vLLM process_weights_after_loading); optimizing its 48×256 per-expert sync count is a follow-up. Residual to 18.8 = decode CUDA-graph (deferred; user refocusing on DeepSeek next).** Post-lever-1 nsys: the remaining ~4× is HOST-SYNC-bound — 22,115 `cudaStreamSynchronize` (78.6% of API time, ~2,760/token, the per-GEMM `DrainQueue`), GPU kernels fast. Remaining levers: grouped W4A4 MoE (design input: `vt::MoeGroupedGemmNvfp4` is W4A16, so true-W4A4 grouped needs a new fp4×fp4 op or the `use_a16` mode + expert-stacking — needs a spike), device-resident decode (RECOMMENDED — the current forward is host-style so every GEMM drains; keep activations on-device, drain once/step; reuse qwen3_5's `Dev`/`Nvfp4Dev`/`ResidentNvfp4`/device-SwiGLU machinery; kills the 22k syncs; converges with the pending GGUF #228 and lifts both quant paths), decode CUDA-graph. Binding number needs a clean 2-3× re-run. See `docs/BENCHMARKS.md` + the spec N5 plan. See `docs/BENCHMARKS.md` `CLAIM-LAGUNA-VLLM-NVFP4`. Prior W7 nsys attribution: host-orchestration-bound, levers ranked (spec `laguna-s21-w7-speed-2026-07-31.md`, ledger `CLAIM-LAGUNA-W7-SPEED`). Prior RUNNABLE + FAST DECODE (W6, 2026-07-31): a per-layer K/V cache + single-token incremental decode replaces W5's O(n²) STATELESS full-recompute — TOKEN-IDENTICAL (byte-equal ids, md5 match, == the W5 golden) and 5.05× faster per token: decode 3.33 → 0.66 s/tok on the real 3-shard UD-Q4_K_XL GGUF (GB10, `--gpu`, keep-quant), same "The capital of France is" → " Paris.\n\nThe user is seeking a detailed explanation of the concept of \"cultural capital\"…". `LagunaKvCache` (mirrors `DeepseekV4KvCache`, MLA-latent → GQA multi-head K/V) caches post-QK-RMSNorm/post-RoPE K + raw V at f32 (bit-exact by construction: RoPE/QK-norm are position-only and attention is causal). MIXED attention handled per-layer: 12 GLOBAL layers grow the cache unbounded (full causal); 36 SLIDING-WINDOW-512 layers EVICT the oldest rows beyond the 512 window (gemma2/3 `is_sliding`), capping their K/V. `LagunaForwardGgufCached` + shared `LagunaAttention`/`LagunaFfnBlock` helpers used by BOTH forwards (identical float ops — the recompute path's ids are unchanged after the refactor); `examples/laguna_gen --stateless` forces the W5 recompute for the A/B gate. No cache bug: bit-exact on the first run. Next speed: grouped-expert GEMM + device-resident decode (both in-tree from ds4). See `.agents/specs/laguna-s21-w6-2026-07-31.md`. Prior RUNNABLE (W5, 2026-07-31): our engine greedy-generates COHERENT text on the REAL 3-shard UD-Q4_K_XL GGUF (GB10, keep-quant). `laguna-gen` "The capital of France is" → " Paris.\n\nThe user is seeking a detailed explanation of the concept of \"cultural capital\" as developed by French soci…" — the FIRST token is "Paris.", matching the llama.cpp-Poolside reference on the identical bytes. Multi-shard GGUF reader (LagunaGgufCtx routes each of 814 tensors to its shard; shard-1 = header only) + keep-quant tower (attn/dense/shared/experts/lm_head stay Q8_0/Q4_K/Q5_K COMPRESSED, consumed via `vt::MatmulBT`; norms/router/bias/embed → f32) + `LagunaForwardGguf` (the f32 composition with the ~9 GEMM sites swapped to keep-quant Gemm/GemmRowSlice, ds4 precedent) + `examples/laguna_gen`. Real GGUF metadata verified: dual-RoPE freq_base 500000/10000, dims 64/128, YaRN factor 32, sigmoid ungrouped-noaux router (scale 2.5), per-layer Q-head [48 global/72 sliding], per-head softplus out-gate, QK-RMSNorm. Load 20.6s, peak 71 GiB (fits 119 pool). Prior W4 IN PROGRESS (2026-07-31): 73.4 GiB UD-Q4_K_XL GGUF FETCHED + read authoritatively (814 tensors); 3 CPU-verified fidelity corrections grounded in the real GGUF + llama.cpp — per-head QK-RMSNorm (`attn_q/k_norm`, the scope MISSED it), GGUF-authoritative dual-RoPE mscale (llama.cpp `yarn_attn_factor·(1+0.1·ln(factor))`, factor 32 not HF 128), separate `ffn_gate/up_exps`. Keep-quant tower materialization + `ForwardGguf` + the real-model greedy run vs llama.cpp-laguna same-quant oracle = W5 close. Prior: W3 REAL host-reference forward + 3 new ops (`laguna_ops.cpp`, CPU `-Werror` clean, `test_laguna_scaffold` unit-gated)** | Poolside Laguna: 48 layers (12 global + 36 sliding-window-512), 256 routed top-10 + 1 shared expert, per-head **softplus attention output gate**, sigmoid `noaux_tc` router, dual per-layer RoPE (YaRN full-attn / plain sliding), GQA 8 KV / 128 head-dim, 1M ctx. **W3 (2026-07-31):** the 3 genuinely-NEW small host ops landed in `laguna_ops.cpp` — per-head softplus attn out-gate (`LagunaSoftplusHeadGate`), ungrouped sigmoid-noaux router (`LagunaUngroupedRouterTopK`, ds3 noaux_tc MINUS the group step + tie-break razor), dual per-layer RoPE cos/sin builders (`BuildLaguna{FullYarn,Sliding}CosSin`, reusing the pinned YaRN inv_freq over the partial-64 dims); `LagunaModel::Forward` is now a REAL runnable host-reference composition (variable-Q-head GQA + dual RoPE + sliding-window mask + softplus gate + dense L0 / ungrouped-MoE L1..47 + untied lm_head) replacing the `VT_CHECK(false)` stub; `test_laguna_scaffold` **8/8·166** (softplus math, router selection+tie-break RED-first, dual-RoPE bit-match, variable-Q-head shapes, forward composition on synthetic weights), `test_model_registry` 24/24. **W2 (2026-07-30):** registered, `ParseLagunaParams`, GGUF `blk.N.*` name-map + UD-Q4_K_XL quant-mix (Q4_K/Q5_K/Q6_K/Q8_0 ALL already decoded → ZERO new kernel). **W1 oracle DECISION:** vLLM NATIVE `laguna.py` (in pin → config constructs); dual-oracle = vLLM-NVFP4/-FP8 (fits GB10 119 GiB; BF16 235 GiB does NOT) + llama.cpp-Q4_K token-exact. ~85–90% reuse (ds4-MoE + Gemma-sliding + OLMo-3-dual-rope + Q4_K keep-quant, ALREADY landed). DEFERRED (W4): GGUF keep-quant tower materialization + device/paged production forward (loaders still LOUDLY throw) + strict dual-oracle greedy gate on a fetched checkpoint + `poolside_v1` parser. See `.agents/specs/laguna-s21-w3-2026-07-31.md` (+ W1/W2 `laguna-s21-w1w2-2026-07-30.md`, W0 `laguna-s21-scope-2026-07-30.md`). **Decode attention-glue fusion LANDED (2026-08-02, `CLAIM-LAGUNA-GLUE-FUSED`, default-ON `VT_LAGUNA_GLUE_FUSED`, `=0` A/B):** BYTE-EXACT L1 (softplus out-gate → `DecodeAttnCombineKernel` store) + L4 (residual-Add+RMSNorm pairs → the shared `vt::FusedChain(kFusedAddRmsNormStd)` seam) on the resident decode-graph — same-binary A/B ids byte-identical (159/159 @160), paging-immune nsys steady decode **−4.2% GPU-busy (28.90→27.69 ms/step), −120 graph nodes/step (−10%)**, wall drop_caches-tied (no regression). C shared-into-MoeCombine SKIPPED (Laguna's bf16 `MoeCombine` → not byte-exact); L2 qk-norm+RoPE preamble DEFERRED (needs a device-position kernel variant). See BENCHMARKS.md `CLAIM-LAGUNA-GLUE-FUSED`. **On-device greedy sample LANDED (2026-08-02, `CLAIM-LAGUNA-ONDEV-SAMPLE`, default-ON `VT_LAGUNA_ONDEV_SAMPLE`, `=0` A/B):** the resident decode graph used to Synchronize, return the whole `[100352]` logits, and argmax on the HOST between replays (+ host embed-gather of the next token) — the off-framework "born-on-host" seam the decode-framework-routing audit flagged. Now BOTH run ON-DEVICE inside the captured graph: `vt::GreedyArgmax` (lowest-index tie = the exact host winner) → 1-elem device token buffer, + a new capture-safe `embed_gather` kernel gathers the next input embedding from it (the stock `vt::Embedding` is NOT capture-safe: per-call event-sync + D2H ring). BYTE-EXACT (160-id stream identical `=0`/`=1` on `~/laguna-xs-nvfp4`) + faster: paired drop_caches decode wall **+0.28% median** (8/8 reps ≥0; removes ~150 us/step host argmax) at GPU-busy parity (nsys 2-length 27.44→27.42 ms/step). Aligns Laguna decode with vLLM on-device sampling. **Lever 2 (lm_head GEMV DRAM eff) MEASURED, NOT landed:** `[M=1,100352,2048]` bf16 = **170 GB/s (2.41 ms)** = ~91% of the cuBLAS M=1×large-N reference (~187 GB/s / 2.2 ms) — at the M=1 practical floor (the 273 GB/s ceiling is streaming-only, unreachable for a once-read GEMV); ≤0.7%-of-step headroom needs a reduction reorder (near-tie re-gate) ⇒ not chased, per prior "lm_head optimal". See BENCHMARKS.md `CLAIM-LAGUNA-ONDEV-SAMPLE`. **MoE add_rms_norm fold LANDED (2026-08-02, `CLAIM-LAGUNA-MOE-ADDNORM`, default-ON `VT_LAGUNA_MOE_ADDNORM_FUSED`, `=0` A/B):** the glue-fused MoE tail ran its residual update as TWO graph nodes — `vt::Add(hidden,routed)` [`AddKernel`] + `FusedChain(kFusedAddRmsNormStd)` [shared-add+RMSNorm, `RmsNormRowKernel`] — now ONE `fused_add2_rmsnorm` device node/MoE-layer (`hidden=(hidden+routed)+shared; hn=rms_norm(hidden)*w`). BYTE-EXACT (IEEE add commutes + the identical 256-thread shared-tree norm reduction; 160-id stream byte-identical `=0`/`=1` on `~/laguna-xs-nvfp4`) + faster: **−39 `AddKernel` graph nodes/step** (2.63ms→0 over 69 steps), paging-immune nsys 2-length **~−46 us/tok GPU (27339→27293)**, nsys wall **+0.4% (34.00→34.14 tok/s @70-tok)**. Small (byte-exact node-count trim on the graph-captured, GPU-bound decode; the dominant ~72% cost is the bf16 projection GEMVs — see the Lever-B negative in BENCHMARKS.md). See BENCHMARKS.md `CLAIM-LAGUNA-MOE-ADDNORM`. **Shared expert kept fp4 LANDED (2026-08-03, `CLAIM-LAGUNA-SHARED-FP4`, default-ON `VT_LAGUNA_SHARED_FP4`, `=0` A/B):** the XS-NVFP4 shared expert was DEQUANTIZED to bf16 at load (`LnLoadSharedExpertBf16`) → the M=1 decode GEMV read 4× the DRAM bytes of vLLM (which keeps it fp4). Now kept fp4-resident and routed through the SAME Marlin W4A16 single-expert (num_experts=1) grouped GEMM the routed experts win on (`dense_nvfp4::GateUpFusedMarlinD`+`MatmulNvfp4MarlinD`); the decode GEMV drops to router-ONLY (`moe.router`), shared gate/up/down go fp4. ADDITIVE new `laguna_shared_fp4.cpp` re-reads the on-disk fp4 from the gen driver before shard release (does NOT touch SACRED `laguna_weights.cpp`); bf16 shared KEPT for the T>1 prefill. NEAR-TIE (fp4≠bf16): coherent, first-20 ids == documented golden, byte-identical to bf16 for ~85 tokens then diverges; **DISTRIBUTIONAL GATE PASS 40/40** (ours' first-40 ids ∈ vLLM's 8-run greedy candidate set; vLLM XS-greedy is bf16-non-det, 8 unique of 8). FASTER: paging-immune nsys 2-length **GPU 27.24→26.53 ms/step (−2.6%)**, wall drop_caches **35.8→36.3 tok/s (+1.4%, fp4 wins all 3 reps)**; shared-expert kernel bucket ~1.68→~0.90 ms/step (halved); vs vLLM ~43 tok/s 83.3%→84.4%; RSS 22.2→22.1 GiB (freed the decode-only fused router-shared projection). Modest by design — XS's shared expert is small (`shared_expert_intermediate_size==moe_intermediate_size==512`). Default-ON per parity (matches vLLM's fp4 shared). See BENCHMARKS.md `CLAIM-LAGUNA-SHARED-FP4`. **qk-norm+RoPE preamble fusion LANDED (2026-08-03, `CLAIM-LAGUNA-PREAMBLE-FUSED`, default-ON `VT_LAGUNA_PREAMBLE_FUSED`, `=0` A/B):** closes the `CLAIM-LAGUNA-GLUE-FUSED` L2 deferral — the decode graph ran the per-layer attention preamble as FOUR under-occupied M=1 nodes (`rms_norm_seq(q)`+`rms_norm_seq(k)`+`rope_from_cache_g(q)`+`rope_from_cache_g(k)`); now ONE capture-safe `fused_qk_norm_rope_g` node/layer (`FusedQkNormRopeGKernel`, one block/head, reads the decode position from DEVICE `*pos_buf`, handles the per-layer dual-RoPE 64/128 + `Hq` 48/64). BYTE-EXACT BY CONSTRUCTION: it replicates the composed path's f32 MEMORY round-trip (Phase A 256-thread Σx² == `RmsNormSeqKernel`; Phase B the same `(x*inv)*w` store; `__syncthreads`; Phase C the `RopeFromCacheGKernel` rope read back) — an earlier register-only recompute was numerically-equivalent but diverged at a token-110 near-tie via compiler fma-contraction; the memory boundary forces bit-identity. 160-id stream byte-identical `=0`/`=1` on `~/laguna-xs-nvfp4` (determinism verified `=0`×3/`=1`×3 each run-to-run identical). FASTER: preamble norm+rope kernels **160→40 launches/tok, 326→154 us/tok (−0.17 ms/step)**; all decode-scaling kernels 26.53→26.37 ms/step; wall drop_caches **36.42→36.64 tok/s (+0.6%, fused wins all 3 paired reps)**; vs vLLM ~43 84.7%→85.2%. Modest (preamble ~1.2% of the 26.5 ms/step decode; the dominant cost stays the bf16 projection GEMVs at cuBLAS parity) — a byte-exact graph-node/launch trim (the glue-fusion residual mechanism). Default-ON per parity. See BENCHMARKS.md `CLAIM-LAGUNA-PREAMBLE-FUSED`. **W7 two-front pass LANDED (2026-08-03, `CLAIM-LAGUNA-W7-DECODE`):** FRONT 1 — the example driver logged `[gen] step N …(RSS)` EVERY decode step, and the RSS arg calls `CurResidentGiB()` (a `/proc/self/status` read) + an unbuffered stderr write in the GPU-idle gap between replays; guarded behind `VT_LAGUNA_STEP_LOG` (default OFF) + added a `decode_wall` line (TRUE end-to-end throughput incl. per-step gaps) next to the gap-free `decode_hp`. Since the fprintf sat OUTSIDE the `s0→s1` timer, `decode_hp` was ALREADY honest; with the log off `decode_wall == decode_hp` (within 0.001 tok/s, every LOG_OFF rep) and the recovered host tax is only ~0.1% (drop_caches noise floor). CONCLUSION: the ~86% gap to vLLM 43 is genuine device compute, NOT a harness artifact. FRONT 2 — `VT_LAGUNA_MOE_ONECAST` (default ON): a MoE layer cast the same `hn[1,H]` f32→bf16 THREE times (router GEMV + routed Marlin + shared Marlin); now cast ONCE into a persistent buffer and reuse (`CastHnBf16`/`GemmBf16Pre` + optional pre-cast param on both `…Into` helpers). BYTE-EXACT (deterministic truncation; `=1` vs `=0` byte-identical 300-tok ids); `CastBf16` **200→122 nodes/step (−78 = 2×39 MoE layers)**, GPU-busy parity within nsys noise, decode_hp +0.29%. Combined (onecast on + log off) **36.97 tok/s = 86.0% of vLLM-NVFP4 43** (from 36.64/85.2%). See BENCHMARKS.md `CLAIM-LAGUNA-W7-DECODE`. **Tail-fold follow-up LANDED (2026-08-03, `CLAIM-LAGUNA-TAIL-FUSED`, default-ON `VT_LAGUNA_TAIL_FUSED`, `=0` A/B):** a fresh node-ranking of the baseline decode graph found the routed-MoE `CastF32` as the one clean byte-exact fold left; it folds into the trailing `fused_add2_rmsnorm` via a new bf16-x1 sibling kernel (`AddAdd2RmsNormStdBf16Kernel` — `MoeCombine` writes bf16 straight to a persistent buffer, widened in-kernel by `__bfloat162float`). BYTE-EXACT (`=1` vs `=0` byte-identical 160-tok ids), `CastF32` **78→39 nodes/step**, total graph nodes **919→880**, GPU-busy parity; decode_hp a WASH (median +0.14% / mean −0.04%, at the drop_caches noise floor). Lands on the deterministic node-count basis (like onecast/preamble/addnorm), NOT a wall win; combined headline UNCHANGED **36.97 tok/s = 86.0%**. The ranking confirms the byte-exact decode-tail fold tier is now essentially EXHAUSTED (residual tail = already-folded norms + attention compute + cuBLAS-adjacent router/topk + ported-Marlin `MoeAlign`/`SiluAndMul`/`MoeCombine`); the gap to vLLM 43 is genuine device compute at the practical ceiling. See BENCHMARKS.md `CLAIM-LAGUNA-TAIL-FUSED`. **KERNEL-EFFICIENCY tier (2026-08-03, `VT_LAGUNA_FAST_NORM` default ON + f32 ext of `VT_RMSNORM_DECODE_FAST`):** the fold tier was exhausted but the residual-stream norm KERNELS were still under-occupied — `ncu` on the shipped `<<<1,256>>>` `AddAdd2RmsNormStdBf16`/`RmsNormRow` decode norms: `launch__waves_per_multiprocessor≈0.00`, `sm__throughput≈0.06%` (one 256-thread block on 1 SM of ~100+, latency-bound). Porting the PROVEN bit-identical `RmsNormRowFastKernel` structure (1024-thread float4 memory passes; 256-strided-partial + tree reduction reproduced byte-for-byte) to the f32 kernels cut each **286→~155 µs/tok (1.85×)**, **byte-exact** (160-tok ids identical `=1`vs`=0`; the f32 fix vs the bf16 sibling: store `v` not `v²` and square in the reduction so nvcc emits shipped's `acc += v*v` **fma** — a pre-squared f32 `v²` is not exact and flipped an XS near-tie at tok 108). **−0.81% decode-step GPU time** (paging-immune 70-vs-20 2-length diff, 26192→25980 µs/step); wall-clock ON/OFF overlap (noise floor). Residual: the byte-exact 256-strided reduction can't reach vLLM's per-kernel norm floor (~2.4× vLLM) without breaking byte-exactness → that remainder is byte-exactness-BLOCKED. See BENCHMARKS.md `CLAIM-LAGUNA-FAST-NORM`. **Router top-k warp-shuffle LANDED (2026-08-03, `CLAIM-LAGUNA-TOPK-SHFL`, default-ON `VT_LAGUNA_TOPK_SHFL`, `=0` A/B): BYTE-EXACT** — an nsys 2-length rank of the remaining small kernels (past the at-parity `gemvx` projection GEMVs ~69% of step + Marlin MoE) put the router `SigmoidTopKKernel` top (415 µs/step); `ncu` showed it `<<<1,256>>>` at `waves≈0.000`/`sm≈0.2%` — pure latency (8 serially-dependent rounds × a ~10-sync `sh[256]` argmax tree). New `SigmoidTopKShflKernel` reduces each round by warp-shuffle argmax (2 syncs/round; argmax over the total order is associative ⇒ SAME winner) → **`SigmoidTopK` 414.6→248.8 µs/step (1.67×)**, decode-step GPU **−0.57%** (26.018→25.869 ms/step), 37.39→37.49 tok/s decode_hp (**87.2% of vLLM-NVFP4 43**); 160-id stream byte-identical `=1`vs`=0`. **NOT landed — norm warp-shuffle (`VT_LAGUNA_NORM_SHFL`):** a near-tie register-accumulate+shuffle reduce for the Laguna `AddAdd2RmsNormStd{,Bf16}Fast` norms PASSED the distributional gate (coherent, in-set 38/40 = baseline, one near-tie fork at pos 37) and was −19.3% per-kernel (`AddAdd2RmsNormStdBf16` 150.3→121.3 µs/step) BUT washed at whole-step (0.6% of step; +0.02% within noise) — a near-tie fork isn't justified by a below-noise gain, so it was dropped. The small-kernel norm tail is at its occupancy floor; the decode step is dominated by the at-parity projection GEMVs. See BENCHMARKS.md `CLAIM-LAGUNA-TOPK-SHFL`. **Shared-expert 2-stream overlap LANDED (2026-08-03, `CLAIM-LAGUNA-SHARED-AUX`, default-ON `VT_LAGUNA_SHARED_AUX`, `=0` A/B):** mirror of vLLM's `MULTI_STREAM_OVERLAPPED` — in `LagunaGraph::RunChain` the fp4-shared arm's shared expert is EARLY-forked onto a second CUDA stream from the post-attn hidden `hn` BEFORE the router GEMV (aux reads `hn` f32 + does its own byte-identical cast; scratch from `AuxPool`), overlapping router+`sigmoid_topk`+routed grouped GEMM, joined before the combine — the SAME machinery the 35B ships default-ON (ENG-MOE-SHARED-AUX, runs inside the captured graph). This is the EARLY fork the prior fused-`router_shared_gu` attempt (`89e0d074`, −0.35% wash) could not reach. Capture-safe (aux stream+2 events in the ctor; gstate-0 warm-run builds residents + warms `AuxPool`). **BYTE-EXACT** (`=1`vs`=0` byte-identical 63-tok ids). REAL concurrency: nsys `--cuda-graph-trace=node` 20↔70 sum-vs-union → OVERLAP **2.34 ms/step** (SUM/UNION 1.092) vs `=0`'s 0.0004 ms; net GPU-busy wall **26.213→25.467 ms/step (−2.9%, 38.15→39.27 tok/s)**, wall @200 37.08→37.93 (+2.3%). Net #include #include +#include #include #include "vllm/model_executor/models/dense_device_glue.h" @@ -93,6 +94,54 @@ const minimax_h3::MiniMaxH3DeviceKernels* Glue(const Dev& d) { return minimax_h3::MiniMaxH3Device(d.q.device.type); } +// DIAGNOSTIC helper: write a per-stage activation/weight fingerprint of a CONTIGUOUS +// device tensor -- summary stats plus a fixed spread of positional sample values that +// catches a transpose/scramble even when the summary matches. Byte-inert unless a +// caller passes a non-null file. Shared by the forward's stage hooks and AttentionDev. +void H3DumpFingerprint(std::FILE* f, vt::Backend& backend, vt::Queue& q, const char* stage, + const Tensor& t) { + if (f == nullptr || t.data == nullptr) return; + int64_t n = 1; + for (int r = 0; r < t.rank; ++r) n *= t.shape[r]; + if (n <= 0) return; + std::vector host(static_cast(n)); + if (t.dtype == DType::kF32) { + backend.Copy(q, host.data(), t.data, static_cast(n) * sizeof(float)); + backend.Synchronize(q); + } else { // bf16 stream tensor -> widen on the host + std::vector raw(static_cast(n)); + backend.Copy(q, raw.data(), t.data, static_cast(n) * sizeof(uint16_t)); + backend.Synchronize(q); + for (int64_t i = 0; i < n; ++i) { + const uint32_t widened = static_cast(raw[static_cast(i)]) << 16; + std::memcpy(&host[static_cast(i)], &widened, sizeof(float)); + } + } + double s = 0.0, s2 = 0.0, amax = 0.0; + bool finite = true; + for (int64_t i = 0; i < n; ++i) { + const double v = static_cast(host[static_cast(i)]); + if (!std::isfinite(v)) finite = false; + s += v; + s2 += v * v; + if (std::fabs(v) > amax) amax = std::fabs(v); + } + const double denom = static_cast(n); + std::fprintf(f, "%-26s n=%-9lld dt=%s finite=%d mean=%+.6e rms=%.6e absmax=%.6e |", stage, + static_cast(n), t.dtype == DType::kF32 ? "f32" : "bf16", finite ? 1 : 0, + s / denom, std::sqrt(s2 / denom), amax); + const int64_t samp[] = {0, 1, 2, 3, n / 8, n / 4, + 3 * n / 8, n / 2, 5 * n / 8, 3 * n / 4, 7 * n / 8, n - 1}; + for (int64_t idx0 : samp) { + int64_t idx = idx0; + if (idx < 0) idx = 0; + if (idx >= n) idx = n - 1; + std::fprintf(f, " %+.5e", host[static_cast(idx)]); + } + std::fprintf(f, "\n"); + std::fflush(f); +} + // vt::MatmulBT + optional rank-1 bias, the device twin of the reference `Linear`. // Weight is [out_features, in_features] — every {Column,Row,QKV,MergedColumn} // ParallelLinear at TP=1. @@ -173,18 +222,28 @@ struct AttnWeightsDev { void AttentionDev(Dev d, const MiniMaxH3DitParams& params, const AttnWeightsDev& w, const Tensor& in, int64_t rows, const Tensor* rope_cache, const Tensor* rope_positions, const int32_t* cu_seqlens, int num_reqs, - const DeviceStreamDtype& dt, Tensor& out) { + const DeviceStreamDtype& dt, Tensor& out, std::FILE* dbg = nullptr, + const char* dbg_pfx = nullptr) { const int64_t heads = params.num_attention_heads; const int64_t head_dim = params.attention_head_dim; const int64_t inner = heads * head_dim; + auto dbgd = [&](const char* leaf, const Tensor& t) { + if (dbg == nullptr) return; + char nm[96]; + std::snprintf(nm, sizeof(nm), "%s.%s", dbg_pfx ? dbg_pfx : "attn", leaf); + H3DumpFingerprint(dbg, vt::GetBackend(d.q.device.type), d.q, nm, t); + }; DBuf qkv(d, dt.S(), {rows, 3 * inner}); LinearDev(d, in, rows, params.hidden_size, *w.qkv, nullptr, qkv.t(), w.qkv_fp4); + dbgd("qkv_out", qkv.t()); DBuf qb(d, dt.S(), {rows, inner}); DBuf kb(d, dt.S(), {rows, inner}); DBuf vb(d, dt.S(), {rows, inner}); vt::QkvSplit(d.q, qb.t(), kb.t(), vb.t(), qkv.t()); + dbgd("q_split", qb.t()); + dbgd("v_split", vb.t()); // Per-head RMSNorm over head_dim: [rows, heads, head_dim] -> [rows*heads, head_dim]. Tensor qn = dense_attn::Reshape(qb.t(), {rows * heads, head_dim}); @@ -193,6 +252,7 @@ void AttentionDev(Dev d, const MiniMaxH3DitParams& params, const AttnWeightsDev& norm_args.eps = static_cast(params.qk_norm_eps); vt::RmsNorm(d.q, qn, qn, *w.q_norm, norm_args); vt::RmsNorm(d.q, kn, kn, *w.k_norm, norm_args); + dbgd("q_qknorm", qb.t()); if (rope_cache != nullptr) { Tensor q3 = dense_attn::Reshape(qb.t(), {rows, heads, head_dim}); @@ -202,6 +262,7 @@ void AttentionDev(Dev d, const MiniMaxH3DitParams& params, const AttnWeightsDev& rope.is_neox_style = true; vt::RopeFromCache(d.q, q3, &k3, *rope_positions, *rope_cache, rope); } + dbgd("q_roped", qb.t()); Tensor tq = dense_attn::Reshape(qb.t(), {rows, heads, head_dim}); Tensor tk = dense_attn::Reshape(kb.t(), {rows, heads, head_dim}); @@ -214,9 +275,11 @@ void AttentionDev(Dev d, const MiniMaxH3DitParams& params, const AttnWeightsDev& args.cu_seqlens = cu_seqlens; args.num_reqs = num_reqs; vt::DFlashBlockAttention(d.q, attn.t(), tq, tk, tv, args); + dbgd("core_out", attn.t()); Tensor flat = dense_attn::Reshape(attn.t(), {rows, inner}); LinearDev(d, flat, rows, inner, *w.out_proj, nullptr, out, w.out_fp4); + dbgd("out_proj", out); } // MiniMaxH3MLP.forward (minimax_h3_transformer.py:512-517): silu(gate) * up. @@ -498,6 +561,52 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, vt::Backend& backend = vt::GetBackend(queue.device.type); Dev d{backend, queue}; const auto* glue = Glue(d); + + // DIAGNOSTIC (env-gated, byte-identical when unset): VT_H3_ACT_DUMP= writes a + // per-STAGE activation fingerprint of THIS forward -- stats (mean/rms/absmax/finite) + // plus a fixed set of positional sample values -- so two weight arms that run the + // SAME graph on the SAME inputs (e.g. the NVFP4-bf16 stream vs the FL2VA-GGUF-bf16 + // control, whose weights differ only by quant noise) can be diffed layer-by-layer to + // localise WHERE they diverge. A JUMP at a stage names the guilty tensor class; a + // scramble/transpose is caught by the positional samples even when rms matches. Only + // the forward whose 0-based call index equals VT_H3_ACT_CALL (default 0) dumps, so a + // single small render (--denoise-only --steps 1) captures exactly one clean forward. + std::FILE* act_f = nullptr; + { + const char* act_path = std::getenv("VT_H3_ACT_DUMP"); + if (act_path != nullptr) { + static int act_call_no = -1; // process-global; resets per process (per arm) + const int this_call = ++act_call_no; + const char* want = std::getenv("VT_H3_ACT_CALL"); + const int want_call = (want != nullptr) ? std::atoi(want) : 0; + if (this_call == want_call) act_f = std::fopen(act_path, "w"); + } + } + auto dump_act = [&](const char* stage, const Tensor& t) { + H3DumpFingerprint(act_f, backend, queue, stage, t); + }; + + // Input-independent WEIGHT fingerprints for the islands + every bias + the output + // heads -- the classes #94's oracle never sampled (it sampled adaln/qkv/out/fc + // WEIGHTS only, never the biases or islands). Comparing these between the two arms + // isolates a bias/island materialization difference (bf16-disk vs f16-disk). + auto W = [&](const char* n, const Tensor& t) { if (t.data != nullptr) dump_act(n, t); }; + W("Wg.video_patch_w", weights.video_patch_proj_w); + W("Wg.video_patch_b", weights.video_patch_proj_b); + W("Wg.audio_patch_w", weights.audio_patch_proj_w); + W("Wg.audio_patch_b", weights.audio_patch_proj_b); + W("Wg.condition_b", weights.condition_proj_b); + W("Wg.time_in_w", weights.time_proj_in_w); + W("Wg.time_in_b", weights.time_proj_in_b); + W("Wg.time_out_w", weights.time_proj_out_w); + W("Wg.time_out_b", weights.time_proj_out_b); + W("Wg.final_norm", weights.final_norm); + W("Wg.final_adaln_b", weights.final_adaln_b); + W("Wg.video_out_w", weights.video_out_w); + W("Wg.video_out_b", weights.video_out_b); + W("Wg.audio_out_w", weights.audio_out_w); + W("Wg.audio_out_b", weights.audio_out_b); + // kBF16 reproduces upstream's PRODUCTION dtype policy: the block stream is bf16 // while the patch projections, the time embedder and both output heads stay fp32 // islands (minimax_h3_transformer.py:85-101). Both dtypes run the SAME code; only @@ -550,6 +659,7 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, DBuf d_rope_cache_f32(d, DType::kF32, {seq_len, rot_dim}, rope_cache_host.data()); DBuf d_rope_cache(d, dt.S(), {seq_len, rot_dim}); CastTo(d, d_rope_cache.t(), d_rope_cache_f32.t()); + dump_act("rope.cache_f32", d_rope_cache_f32.t()); std::vector arange(static_cast(seq_len)); for (int64_t i = 0; i < seq_len; ++i) arange[static_cast(i)] = static_cast(i); DBuf d_rope_pos(d, DType::kI32, {seq_len}, arange.data()); @@ -561,6 +671,8 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, DBuf video_embed(d, DType::kF32, {inputs.num_img_pos, hidden}); LinearDev(d, video_rows.t(), inputs.num_img_pos, video_width, weights.video_patch_proj_w, &weights.video_patch_proj_b, video_embed.t()); + dump_act("embed.video_rows_in", video_rows.t()); + dump_act("embed.video_patch_out", video_embed.t()); DBuf d_audio_x(d, DType::kF32, {seq_len, audio_width}, inputs.audio_x); DBuf audio_rows(d, DType::kF32, {inputs.num_audio_pos, audio_width}); @@ -568,6 +680,7 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, DBuf audio_embed(d, DType::kF32, {inputs.num_audio_pos, hidden}); LinearDev(d, audio_rows.t(), inputs.num_audio_pos, audio_width, weights.audio_patch_proj_w, &weights.audio_patch_proj_b, audio_embed.t()); + dump_act("embed.audio_patch_out", audio_embed.t()); // text rows enter as the stream dtype before the BF16 condition projection. // text rows enter as the stream dtype before the BF16 condition projection. @@ -577,6 +690,7 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, DBuf text_embed(d, dt.S(), {inputs.num_text_pos, hidden}); LinearDev(d, text_rows.t(), inputs.num_text_pos, params.text_dim, weights.condition_proj_w, &weights.condition_proj_b, text_embed.t(), &weights.condition_fp4); + dump_act("embed.text_condition_out", text_embed.t()); // Token refiner: a plain pre-norm stack, no AdaLN and no RoPE (:564-623), on the // REPLICATED text rows, so it uses the refiner's own cu_seqlens. @@ -586,7 +700,19 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, DBuf tmp(d, dt.S(), {rows, hidden}); vt::RmsNormArgs args; args.eps = static_cast(params.norm_eps); + int64_t ref_idx = -1; for (const MiniMaxH3DitBlockWeights& block : weights.refiner) { + ++ref_idx; + if (ref_idx == 0) { + dump_act("W.ref0.norm1", block.norm1); + if (block.qkv_proj.data != nullptr) dump_act("W.ref0.qkv_proj", block.qkv_proj); + dump_act("W.ref0.q_norm", block.q_norm); + dump_act("W.ref0.k_norm", block.k_norm); + if (block.out_proj.data != nullptr) dump_act("W.ref0.out_proj", block.out_proj); + if (block.fc1.data != nullptr) dump_act("W.ref0.fc1", block.fc1); + if (block.fc2.data != nullptr) dump_act("W.ref0.fc2", block.fc2); + dump_act("W.ref0.norm2", block.norm2); + } vt::RmsNorm(d.q, normed.t(), text_embed.t(), block.norm1, args); AttentionDev(d, params, AttnWeightsDev{&block.qkv_proj, &block.q_norm, &block.k_norm, &block.out_proj, @@ -611,6 +737,7 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, vt::RmsNorm(d.q, normed.t(), text_embed.t(), weights.refiner_final_norm, final_args); backend.Copy(d.q, text_embed.ptr(), normed.ptr(), normed.bytes()); } + dump_act("embed.text_refined_out", text_embed.t()); // index_add_ scatter of the three modality embeddings into the packed stream. // The three position sets are DISJOINT by construction (each packed row carries @@ -643,6 +770,7 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, vt::IndexCopy(d.q, stream.t(), text_embed.t(), d_text_pos.t()); vt::IndexCopy(d.q, stream.t(), video_s.t(), d_img_pos.t()); vt::IndexCopy(d.q, stream.t(), audio_s.t(), d_audio_pos.t()); + dump_act("embed.stream_post_scatter", stream.t()); // --- time embedding (minimax_h3_transformer.py:272-285) --- DBuf t_emb(d, DType::kF32, {m, params.time_embed_dim}); @@ -669,6 +797,7 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, LinearDev(d, mid.t(), m, params.time_embed_hidden_size, weights.time_proj_out_w, &weights.time_proj_out_b, t_emb.t()); } + dump_act("time.t_emb", t_emb.t()); // silu(t_emb), hoisted: t_emb is loop-invariant, so the reference's per-block // silu inside AdalnProject is redundant work here. @@ -699,9 +828,29 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, return t; }; + int64_t blk_idx = -1; for (const MiniMaxH3DitBlockWeights& block : weights.blocks) { + ++blk_idx; + // Direct WEIGHT fingerprints for block 0 -- comparing these between the two arms + // isolates a loader/materialization bug (a weight class that differs) from a + // compute divergence. The bf16 arm has real bf16 tensors here (fp4-resident leaves + // the projection Empty, so those print n=0). + if (blk_idx == 0) { + dump_act("W.blk0.norm1", block.norm1); + if (block.qkv_proj.data != nullptr) dump_act("W.blk0.qkv_proj", block.qkv_proj); + dump_act("W.blk0.q_norm", block.q_norm); + dump_act("W.blk0.k_norm", block.k_norm); + if (block.out_proj.data != nullptr) dump_act("W.blk0.out_proj", block.out_proj); + if (block.fc1.data != nullptr) dump_act("W.blk0.fc1", block.fc1); + if (block.fc2.data != nullptr) dump_act("W.blk0.fc2", block.fc2); + dump_act("W.blk0.norm2", block.norm2); + if (block.adaln_w.data != nullptr) dump_act("W.blk0.adaln_w", block.adaln_w); + if (block.adaln_b.data != nullptr) dump_act("W.blk0.adaln_b", block.adaln_b); + } AdalnProjectDev(d, t_emb_s.t(), m, params.time_embed_dim, block.adaln_w, block.adaln_b, projected.t(), &block.adaln_fp4); + if (act_f != nullptr) dump_act(("block." + std::to_string(blk_idx) + ".adaln_proj").c_str(), + projected.t()); const Tensor shift_msa = chunk_view(projected, adaln_rows, 6, 0); const Tensor scale_msa = chunk_view(projected, adaln_rows, 6, 1); const Tensor gate_msa = chunk_view(projected, adaln_rows, 6, 2); @@ -713,13 +862,17 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, glue->modulate_scale_shift(d.q, normed.t().data, shift_msa.data, scale_msa.data, d_combined.t().Ptr(), seq_len, hidden, 6 * hidden, dt.S()); + if (blk_idx == 0) dump_act("block.0.normed_pre_attn", normed.t()); AttentionDev(d, params, AttnWeightsDev{&block.qkv_proj, &block.q_norm, &block.k_norm, &block.out_proj, &block.qkv_fp4, &block.out_fp4}, normed.t(), seq_len, &d_rope_cache.t(), &d_rope_pos.t(), inputs.cu_seqlens, - num_reqs, dt, tmp.t()); + num_reqs, dt, tmp.t(), blk_idx == 0 ? act_f : nullptr, "block.0.attn"); + if (blk_idx == 0) dump_act("block.0.attn_contrib", tmp.t()); glue->modulate_gate(d.q, stream.t().data, gate_msa.data, tmp.t().data, d_combined.t().Ptr(), seq_len, hidden, 6 * hidden, dt.S()); + if (act_f != nullptr) + dump_act(("block." + std::to_string(blk_idx) + ".post_attn").c_str(), stream.t()); vt::RmsNorm(d.q, normed.t(), stream.t(), block.norm2, block_args); glue->modulate_scale_shift(d.q, normed.t().data, shift_mlp.data, scale_mlp.data, @@ -729,6 +882,8 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, &block.fc2_fp4); glue->modulate_gate(d.q, stream.t().data, gate_mlp.data, tmp.t().data, d_combined.t().Ptr(), seq_len, hidden, 6 * hidden, dt.S()); + if (act_f != nullptr) + dump_act(("block." + std::to_string(blk_idx) + ".post_mlp").c_str(), stream.t()); } // --- final layer (minimax_h3_transformer.py:724-743) --- @@ -743,6 +898,8 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, // The final layer is single-modality, so it indexes by inverse_indices directly. glue->modulate_scale_shift(d.q, normed.t().data, final_shift.data, final_scale.data, d_inverse.t().Ptr(), seq_len, hidden, 2 * hidden, dt.S()); + dump_act("final.adaln_proj", final_projected.t()); + dump_act("final.normed", normed.t()); // Cast UP before both output heads: they are fp32 ISLANDS. DBuf head_in(d, DType::kF32, {seq_len, hidden}); CastTo(d, head_in.t(), normed.t()); @@ -753,6 +910,9 @@ MiniMaxH3DitOutputs MiniMaxH3DitForwardDevice(vt::Queue& queue, DBuf audio_all(d, DType::kF32, {seq_len, audio_width}); LinearDev(d, head_in.t(), seq_len, hidden, weights.audio_out_w, &weights.audio_out_b, audio_all.t()); + dump_act("final.video_all", video_all.t()); + dump_act("final.audio_all", audio_all.t()); + if (act_f != nullptr) std::fclose(act_f); // Select the inference-output rows ON DEVICE, so only the selected rows cross // the bus (:1087-1101).