Skip to content
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion .agents/NOW.md
Original file line number Diff line number Diff line change
Expand Up @@ -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/` |
Expand Down
14 changes: 14 additions & 0 deletions .agents/benchmark-record.md
Original file line number Diff line number Diff line change
Expand Up @@ -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.
28 changes: 27 additions & 1 deletion .agents/specs/minimax-h3.md
Original file line number Diff line number Diff line change
Expand Up @@ -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) |
Expand Down Expand Up @@ -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
Expand Down
20 changes: 20 additions & 0 deletions .agents/state.md
Original file line number Diff line number Diff line change
Expand Up @@ -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.
2 changes: 1 addition & 1 deletion docs/BENCHMARKS.md
Original file line number Diff line number Diff line change
Expand Up @@ -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 |
Expand Down
Loading