feat(minimax-h3): port the MiniMax-H3 omni-modal video+audio DiT (portable path complete; speed GPU-gated) - #26
feat(minimax-h3): port the MiniMax-H3 omni-modal video+audio DiT (portable path complete; speed GPU-gated)#26localai-bot wants to merge 85 commits into
Conversation
MiniMax-H3 (`MiniMaxAI/MiniMax-H3`) is the project's FIRST diffusion
architecture: a 33.1B CFG-distilled joint video+audio transformer served by
vLLM-Omni over `/v1/videos`. One request runs a 50-step flow-matching denoise
loop, forwarding the DiT ONCE PER STEP over the whole packed sequence, then
decodes latents to 24 FPS frames plus 32 kHz stereo through two VAEs. It is NOT
autoregressive: no KV cache, no sampler, no logits. It is therefore deliberately
NOT registered in the causal-LM registry, and the born-on-the-runner seam does
not apply to it by construction.
Ported from vllm-project/vllm-omni, vllm_omni/diffusion/models/minimax_h3/:
minimax_h3_transformer.py -> minimax_h3.{h,cpp} (arch config, 3D MM-RoPE,
time embedder, AdaLN proj, DiT block, token
refiner, final layer, packed forward, weight
contract, grouped-qkv reorder)
denoise_loop.py -> minimax_h3.cpp (CFG-distilled driver)
packed_sequence.py -> minimax_h3_packing.cpp (fl2va + ref2va layouts)
packed_tokens.py -> minimax_h3_packing.cpp (latent <-> token packing)
scheduling_..._euler_...py -> minimax_h3_packing.cpp (euler eta0 scheduler)
No new kernel was added. The packed varlen NON-CAUSAL attention routes through
the shared `vt::DFlashBlockAttention(causal=false)` -- its per-document
bidirectional contract is exactly upstream's cu_seqlens varlen FA call -- and
every projection through `vt::MatmulBT`.
HARDWARE VERDICT (recorded, not worked around): the checkpoint is ~354 GB and
upstream validates on 4x NVIDIA B300 at ~133 GB peak per rank. One GB10 has 119
GiB UNIFIED memory, so CPU offload cannot help. End-to-end H3 is impossible on
this project's hardware; no e2e result and no speed number is claimed.
What IS gated, and exactly: upstream's H3 modules are pure Python, so they are
executed at REDUCED DIMENSIONS as the oracle. gen-minimax-h3-goldens.py imports
packed_sequence/packed_tokens/scheduling by file path (bypassing the package
__init__, so neither vllm nor aenum is needed) and restates the DiT at TP=1;
both sides rebuild weights and inputs from an identical FNV-1a + splitmix64
stream, so not one weight byte is checked in.
test_minimax_h3: 10/10 cases, 2539 assertions, clean CPU build (0 warnings)
- fl2va + ref2va packed layouts EXACT, fp64 position grid BIT-EXACT
- patchify / unpatchify / audio pack EXACT + round-trip identity
- euler-ancestral eta0 scheduler EXACT
- DiT forward max abs diff 1.6e-7 (video) / 1.5e-7 (audio)
- denoise-loop invariants (pinned rows reset per step, targets advance)
The fp64 grid is gated bit-exact because it feeds RoPE, which required matching
upstream's arithmetic ORDER: numpy linspace(endpoint=False) evaluates
i*step + start, and upstream keeps a numpy-PAIRWISE and a Python-SEQUENTIAL span
sum deliberately separate (packed_sequence.py:101-113).
Not yet built, recorded honestly in the spec: device-resident/bf16 forward and
the fusion folds (W2b, where speed work begins -- upstream reports the DiT at
88% of request latency), the H3-Encoder on our existing Qwen3-VL tower (W3), the
two VAEs -- which are checkpoint REMOTE CODE under trust_remote_code and must be
REIMPLEMENTED in C++, not adapted (W4/W5) -- pipeline/tasks (W6), /v1/videos plus
MP4 muxing, which needs a new dependency decision (W7), and USP multi-GPU (W8).
OPEN: there is no vllm-omni parity PIN; the upstream-sync protocol covers only
the vLLM repo, and H3 is both beyond the pin and outside that repository.
Record: .agents/specs/minimax-h3.md, model-matrix row
MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit (PARTIAL) + checklist/rollup,
roadmap ROAD-V1-H3, docs/STATUS.md, docs/BENCHMARKS.md (PENDING, hardware
blocked), parity-ledger and state entries.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
Three additions, and one correction to the previous commit's hardware verdict.
CORRECTION FIRST. The previous commit concluded MiniMax-H3 end-to-end was
"impossible on this project's hardware". That was wrong. It reasoned from ONE
artifact -- the BF16 release (~354 GB, validated on 4x B300). Quantized H3
checkpoints exist and they fit one GB10 (119 GiB unified):
realrebelai/MiniMax-H3_GGUFs DiT Q3_K_M 15.6 GB + Qwen3-VL encoder Q4_K_M
14.6 GB + VAEs ~11 GB => ~41 GB working set
lilcheaty/MiniMax-H3-NVFP4 NVFP4 DiT + AWQ encoder + both VAEs
So e2e AND a speed comparison are REACHABLE; they are gated on the remaining
bricks (encoder, VAEs, pipeline), not on hardware. NVFP4 is the likely speed
path: sm_121 has native FP4 tensor cores and our NVFP4 stack is the most tuned
one we own. Still no e2e or speed NUMBER is claimed here.
1. BF16 PRODUCTION STREAM (minimax_h3.cpp). The DiT forward now runs upstream's
production dtype policy as well as the f32 parity path: the block stream is
bf16 while the fp32 islands (both patch projections, the time embedder, both
output heads -- minimax_h3_transformer.py:85-101) stay fp32, with the explicit
casts of _modulate_scale_shift / _modulate_gate reproduced at their sites.
Gated against a bf16 upstream golden: max abs diff 2.4e-3 (bf16 scale; same
cast points, different GEMM accumulation order).
2. REQUEST PLANNING (minimax_h3_planner.cpp) <- time_request.py:5-61 and
pipeline_minimax_h3.py:121-122, 207-222, 374-434. Frame snapping to 17n+5,
video/audio latent shapes, the rectified-flow time-shift sigma schedule,
canvas resolution, reference-image rescale, and t2va/fl2va/ref2va dispatch.
EXACT vs upstream. Note Python's round() is half-to-EVEN and both
_align_multiple and _audio_latent_t depend on it.
3. GGUF ARM (minimax_h3_gguf.cpp). The ComfyUI-format H3 GGUF keeps the
checkpoint's own parameter names, so THE NAME MAP IS THE IDENTITY: all 535
tensors of MiniMax-H3-FL2VA-Q3_K_M.gguf match the contract
EnumerateMiniMaxH3DitTensors derived from upstream SOURCE -- the weight
contract is now validated against a real checkpoint. Two shape rules:
- GGUF `ne` is reversed relative to torch;
- `comfy.gguf.orig_shape.<name>` overrides it where ComfyUI reshaped a
tensor for quant-block alignment (the 50 AdaLN projections are logical
[96768, 2688], and 2688 is not a multiple of the 256-element Q3_K block).
Geometry is derived from SHAPES ALONE, because a ComfyUI GGUF ships no
transformer config -- that is what lets a GGUF load without the original repo.
The gate needs no download: a GGUF header is self-delimiting and at the front
of the file, so a 4 MiB HTTP range request yields the whole manifest, which
scripts/gen-minimax-h3-gguf-manifest.py freezes into a fixture (names, dims,
types, orig_shape -- no weight bytes).
test_minimax_h3: 13/13 cases, 3907 assertions, clean CPU build (0 warnings).
Record updated with the corrected verdict throughout: .agents/specs/minimax-h3.md
(section 0 rewritten, W9/W10 bricks added), the model-matrix row + checklist,
roadmap ROAD-V1-H3, docs/STATUS.md, docs/BENCHMARKS.md (the "not reproducible"
disposition is WITHDRAWN), plus ledger and state entries.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
Update: quantized arms land, and the hardware verdict is correctedPushed
So e2e and a speed comparison are reachable, gated on the remaining bricks (encoder → VAEs → pipeline), not on hardware. NVFP4 is the likely speed path: sm_121 has native FP4 tensor cores and our NVFP4 stack is the most tuned one we own. What this push adds1. BF16 production stream. The DiT forward now runs upstream's production dtype policy alongside the f32 parity path — bf16 block stream with the fp32 islands (both patch projections, time embedder, both output heads) preserved. Gated against a bf16 upstream golden at max abs diff 2.4e-3. 2. Request planning ( 3. GGUF arm. The best result here: the name map is the identity. All 535 tensors of
Geometry is derived from shapes alone, since a ComfyUI GGUF ships no transformer config — that's what lets a GGUF load without the original repo. The gate needed no download: a GGUF header is self-delimiting and at the front of the file, so a 4 MiB HTTP range request yields the whole manifest, which Status
Next: download a quantized checkpoint and close the e2e loop (encoder on our existing Qwen3-VL tower → the two VAEs → pipeline), then the NVFP4 arm for speed. |
The biggest unknown in this lane was that H3's two VAEs are checkpoint REMOTE
CODE: they ship inside the HF repo (FL2VA/{audio,video}_vae/*.py) and are loaded
through get_class_from_dynamic_module under trust_remote_code. vLLM-Omni only
ADAPTS them (vae.py:41-53), so a no-Python engine has to reimplement them.
The remote code is now in hand (~130 KB of Python, NOT vendored here -- it ships
under the MiniMax H3 Community License with the checkpoint), and the AUDIO side
is done.
WHAT IT IS. A DAC-lineage BigVGAN vocoder, per the checkpoint's config.yaml +
metadata.json: dec_in_proj (32 -> 2048, k=1) then conv_pre -> 7 upsample stages
(ConvTranspose1d; rates 5,5,2,2,2,2,2 / kernels 9,9,4,4,4,4,4), each followed by
3 AMPBlock1 residual stacks (kernels 3,7,11, dilations 1,3,5) whose outputs are
AVERAGED -> anti-aliased SnakeBeta -> conv_post (1 ch, k=7, no bias) -> clamp to
[-1, 1] (H3 sets use_tanh_at_final=false). 32 kHz, 2 channels.
Two details that are easy to get wrong, both gated:
* Every conv is WEIGHT-NORMALIZED, so the checkpoint stores (g, v) pairs as
parametrizations.weight.original0/original1 and the loader materializes
g * v / norm(v) with the norm over every dim except dim 0. ConvTranspose1d
weight is [in, out, k], so its weight-norm dim 0 is the INPUT channel.
* The anti-aliased activation is up 2x -> SnakeBeta -> down 2x through a
KAISER-SINC filter COMPUTED at load time and never loaded -- needing a
Bessel I0, torch's periodic=false kaiser window, and REPLICATE padding. The
filter is gated separately so a filter bug cannot masquerade as a decoder bug.
GATE. scripts/gen-minimax-h3-audio-vae-goldens.py imports the CHECKPOINT'S OWN
modules and runs them at reduced dimensions as the oracle, with weights from the
shared H3Rand stream so no weight byte is checked in:
waveform max abs diff 4.2e-9 (f32 round-off)
kaiser filter max abs diff 3.0e-8
The first golden had 18 of 32 samples pinned at the final clamp, which would have
HIDDEN errors; the generator's weight scale is tuned so the output is fully
unsaturated, and the test asserts non-saturation explicitly.
test_minimax_h3: 14/14 cases, 3983 assertions, clean CPU build (0 warnings).
NOT done: the encode-side determinism-context semantics, and the VIDEO VAE (W4),
which is the largest remaining brick -- klvae.py alone is ~48 KB, plus a CNN/ViT
hybrid, tiling, and a parallel path.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
Update: the audio VAE is reimplemented (
|
| max abs diff | |
|---|---|
| waveform | 4.2e-9 (f32 round-off) |
| kaiser-sinc filter | 3.0e-8 |
One methodology note worth flagging: the first golden had 18 of 32 samples pinned at the final clamp, which would have hidden errors. The weight scale is tuned so the output is fully unsaturated, and the test now asserts non-saturation explicitly.
test_minimax_h3: 14/14 cases, 3983 assertions, clean build, all checkers green.
Remaining
Video VAE (W4) is the largest brick left — klvae.py alone is ~48 KB, plus a CNN/ViT hybrid, tiling, and a parallel path. Then the encoder (W3, mostly reuse of our Qwen3-VL tower) and the pipeline (W6). After those, an e2e run on a quantized checkpoint is reachable — and only then is a speed number meaningful.
…m real manifests
Generalizes the GGUF header trick: a safetensors header is ALSO front-loaded (an
8-byte length plus JSON) and only tens of KB even for a 10 GB file, so one HTTP
range request captures the entire tensor manifest. scripts/gen-minimax-h3-
safetensors-manifest.py turns that into a C++ fixture, and two real MiniMax-H3
checkpoints are now gated without downloading a byte of payload.
1. THE NVFP4 CHECKPOINT IS OUR LAYOUT, EXACTLY.
lilcheaty/MiniMax-H3-NVFP4 (1051 tensors) is the textbook compressed-tensors
triple:
weight U8 FP4 packed 2-per-byte; [21504, 2688] for a logical
[21504, 5376] fused qkv
weight_scale F8_E4M3 one per group of 16 along K; [21504, 336]
weight_scale_2 F32 one global scalar
258 quantized projections, each carrying all three. The fp32/bf16 ISLANDS
(both patch projections, the time embedder, both output heads, the norms and
rope.inv_freq) are left unquantized, so the DiT's dtype policy survives. And
the names are IDENTICAL to the contract we derived from upstream source.
So W10 is loader WIRING onto the NVFP4 stack this project already tuned, not a
new quantization scheme. That matters because W10 is the speed path: sm_121
has native FP4 tensor cores.
2. THE VIDEO VAE DECODER IS A ViT, NOT A CNN.
W4 was scoped as "port a 48 KB klvae.py" and treated as the scary brick. The
real 560-tensor manifest says otherwise: the ENCODER is the 3D CNN (116
tensors, rank-5 Conv3d down blocks), but the DECODER -- the only half
generation needs -- is a 36-block TRANSFORMER (440 tensors: attn.to_qkv /
attn.to_out, ff.w1 / ff.w2, two norms and two learned residual scales per
block, plus x_embedder, mask_token, register_tokens, norm_out, proj_out). fp32
throughout. We have every primitive for that, so W4 is materially smaller than
previously recorded.
test_minimax_h3: 16/16 cases, 5908 assertions, clean CPU build (0 warnings).
Still NO end-to-end result and NO speed number. What changed is that the two
remaining unknowns on the critical path are now measured rather than guessed.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
Update: NVFP4 arm grounded, video VAE re-scoped (
|
| tensor | dtype | shape (fused qkv, block 0) |
|---|---|---|
weight |
U8 | [21504, 2688] — FP4 packed 2-per-byte, logical [21504, 5376] |
weight_scale |
F8_E4M3 | [21504, 336] — one per group of 16 along K |
weight_scale_2 |
F32 | scalar global |
258 quantized projections, each carrying all three. The fp32/bf16 islands (patch projections, time embedder, output heads, norms, rope) are left unquantized, so the DiT's dtype policy survives. Names identical to the contract derived from source.
So W10 is loader wiring onto the NVFP4 stack we already tuned for Laguna — not a new quant scheme. That matters because this is the speed path: sm_121 has native FP4 tensor cores.
2. The video VAE decoder is a ViT, not a CNN
I had scoped W4 as "port a 48 KB klvae.py" and called it the scary brick. The real 560-tensor manifest says otherwise:
- Encoder = the 3D CNN (116 tensors, rank-5 Conv3d down blocks)
- Decoder = a 36-block transformer (440 tensors:
attn.to_qkv/to_out,ff.w1/w2, two norms + two learned residual scales per block, plusx_embedder,mask_token,register_tokens,norm_out,proj_out), fp32 throughout
Generation only needs the decoder. We have every primitive for a ViT, so W4 is materially smaller than previously recorded.
test_minimax_h3: 16/16 cases, 5908 assertions, clean build, all checkers green.
Honest status
Still no e2e result and no speed number. What changed is that the two biggest unknowns left on the critical path are now measured instead of guessed — and both came back smaller than feared. Remaining: W4 video-VAE decoder (ViT port), W3 encoder (Qwen3-VL reuse), W6 pipeline, then W10 loader wiring. Only after those does a speed measurement mean anything.
…rtial)
The video VAE is checkpoint REMOTE CODE like the audio one, so it must be
reimplemented rather than adapted. Its real 560-tensor manifest showed the
decoder -- the only half generation needs -- is a 36-block transformer, and this
ports that repeated unit.
Per block, all fp32 (base_module.py:200-281):
h += scale1 * Attention(RMSNorm(h))
h += scale2 * GatedSiLU_FeedForward(RMSNorm(h))
with scale1/scale2 LEARNED PER-CHANNEL vectors (not scalars) and per-head RMS
q/k normalization carrying NO affine weight.
THE TRAP THIS CATCHES. This ViT's to_qkv output is PER-HEAD INTERLEAVED: upstream
does qkv.view(B, S, -1, 3*dim_head) then chunk(3, dim=-1), so the layout is
[head0_q | head0_k | head0_v | head1_q | ...] -- NOT the [q_all | k_all | v_all]
that the H3 DiT's fused qkv uses. Both layouts are the same SIZE, so reading it
the DiT way raises no error and no shape mismatch; it just produces a
plausible-but-wrong image. Only a numeric gate against the real module finds it.
GATE. scripts/gen-minimax-h3-video-vae-goldens.py executes the CHECKPOINT'S OWN
base_module.TransformerBlock at reduced dimensions as the oracle, with weights
from the shared H3Rand stream (no weight bytes checked in):
video VAE decoder block max abs diff 6.0e-8
The bundle imports a handful of diffusers symbols (a logger, two mixin bases, two
no-op decorators); the generator STUBS them rather than taking the whole diffusers
dependency just to run an oracle.
test_minimax_h3: 17/17 cases, 5911 assertions, clean CPU build (0 warnings).
W4 REMAINS PARTIAL. Still to do: the 36-block stack surround (x_embedder,
mask_token / register_tokens, 3D RoPE, norm_out / proj_out, unpatchify, tiling)
and the 3D-CNN encoder, which is only needed for image/video CONDITIONING rather
than for generation output. Still no e2e run and no speed number.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
Update: video-VAE decoder block ported (
|
Both VAE DECODERS are now done. This adds the stack surround around the already gated TransformerBlock, so the whole generation-critical half of the video VAE reproduces the checkpoint's own ViT3DDecoder: _pack_tensors_3d (channels-last flatten) x_embedder suffix = register tokens + a ZERO cls token 3D RoPE 36-block stack norm_out <- LAYER norm (the blocks use RMS) proj_out _unpack_tensors_3d -> [C, T*pt, H*ps, W*ps] Run at the REAL hyperparameters, read from the checkpoint's source/config.json::vit_decoder_kwargs: 36 layers, 32 heads x 64, rms_norm affine, qk rms_norm WITHOUT affine, gated SiLU, rope_theta 100.0, rope_dim_ratio 0.75. 3D ROPE DETAIL. RotaryEmbeddingND builds angles from LENGTH-NORMALIZED token ids ((i + 0.5)/n mapped into [-1, 1)), scales them by 2*pi (use_angle=True), and concatenates the three per-axis frequency blocks before TILING the result twice to fill rot_dim. The suffix tokens carry id 0, so their cos/sin are 1/0 -- an identity rotation -- which falls out of initializing the tables that way. GATE (vs the checkpoint's own module at reduced dimensions): video VAE full ViT3D decoder max abs diff 8.9e-8 video VAE decoder block max abs diff 6.0e-8 audio VAE decoder max abs diff 4.2e-9 test_minimax_h3: 18/18 cases, 5918 assertions, clean CPU build (0 warnings). REMAINING on the VAE side: video tiling (vae_tile_size 256, overlap 64) and the 3D-CNN ENCODER -- and the encoder is only needed for image/video CONDITIONING (fl2va/ref2va), not for producing output frames, so a t2va path does not need it. Still no e2e run and no speed number. W3 (H3-Encoder on our Qwen3-VL tower) and W6 (pipeline) are what stand between here and an end-to-end t2va run. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
Update: video-VAE ViT3D decoder complete (
|
| max abs diff | |
|---|---|
| video VAE full ViT3D decoder | 8.9e-8 |
| video VAE decoder block | 6.0e-8 |
| audio VAE decoder | 4.2e-9 |
test_minimax_h3: 18/18 cases, 5918 assertions, clean build, all checkers green.
What's left
On the VAE side: video tiling (vae_tile_size 256 / overlap 64) and the 3D-CNN encoder — and the encoder is only needed for image/video conditioning (fl2va/ref2va), not for producing output frames, so a t2va path doesn't need it.
On the critical path to e2e: W3 (H3-Encoder on our existing Qwen3-VL tower) and W6 (pipeline). Then W10 NVFP4 loader wiring.
Still no e2e run and no speed number — and neither is reachable from this machine (no GPU). That's a run on the DGX once W3+W6 close.
The H3-Encoder produces the [seq, 5120] prompt_embeds the DiT consumes. Its
architecture is a Qwen3-VL, which this project already ports, so the value of
this change is pinning down and GATING the three H3-specific deltas:
1. LAYER TRUNCATION. num_layers = min(config.num_hidden_layers, 50). The gate's
config deliberately declares MORE layers than are selected so truncation is
actually exercised, plus an assertion that it never EXTENDS a shallower model.
2. UNNORMALIZED OUTPUT -- the load-bearing one. H3 consumes the hidden state
straight out of layer 49 with NO final RMSNorm, unlike a stock Qwen3-VL text
model. Applying one raises no error and changes no shape; it just silently
shifts every conditioning vector.
3. DEEPSTACK. Visual features are added at the visual token positions after each
of the first len(deepstack_visual_embeds) layers. The test asserts DeepStack
actually CHANGES the result, so a no-op injection cannot pass.
The layer itself is the familiar pre-norm block: RMSNorm -> fused-QKV attention
with per-head q/k RMSNorm and interleaved M-RoPE -> causal GQA -> o_proj ->
residual; RMSNorm -> gated-SiLU MLP -> residual.
GATE. scripts/gen-minimax-h3-encoder-goldens.py runs the UPSTREAM
MiniMaxH3Qwen3VLTextModel at reduced dimensions. encoder.py imports exactly one
vllm symbol (vllm.logger), so a one-line stub lets the oracle run without vllm or
any of its dependencies -- the same pattern as the diffusers stub for the video
VAE.
encoder text tower, plain path max abs diff 1.2e-7
encoder text tower, DeepStack path max abs diff 1.2e-7
test_minimax_h3: 19/19 cases, 5930 assertions, clean CPU build (0 warnings).
REMAINS for W3: the encoder's VISION tower (a reuse of our qwen3_vl_vision.cpp
rather than a new port) and the MM processor.
Still no e2e run and no speed number. W6 (pipeline assembly) is now the main
thing standing between here and an end-to-end t2va run.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
Update: encoder text tower ported (
|
| Component | Gate |
|---|---|
| packed layout (fl2va + ref2va) | exact, fp64 grid bit-exact |
| latent packing | exact + round-trip |
| flow-matching scheduler | exact |
| DiT forward (f32) | 1.6e-7 |
| DiT forward (bf16 production stream) | 2.4e-3 |
| request planning | exact |
| GGUF manifest (535 real tensors) | exact |
| NVFP4 manifest (1051 real tensors) | exact, layout = ours |
| audio VAE decoder | 4.2e-9 |
| video VAE ViT3D decoder | 8.9e-8 |
| encoder text tower | 1.2e-7 |
Remaining: encoder vision tower + MM processor, W6 pipeline assembly, W7 serving/MP4, W10 NVFP4 loader wiring, W2b device-resident forward, video tiling.
Still no e2e run and no speed number — and neither is reachable from this machine (no GPU, 90 GB free vs a ~41 GB checkpoint). W6 is now the main thing between here and a t2va run on the DGX.
…es (W6)
MiniMaxH3GenerateT2va wires the separately-gated stages into one path:
prompt_embeds -> packed layout -> sigma schedules (video shift 12, audio 3)
-> denoise loop (one DiT forward per step, euler-eta0 update)
-> unpatchify / audio unpack -> per-channel denormalize
-> video ViT3D decoder + audio BigVGAN decoder
-> frames [3, T*pt, H*ps, W*ps] + stereo waveform at 32 kHz
A structural end-to-end test runs the whole thing at reduced dimensions with
random weights. That is explicitly NOT a quality result -- it is proof the stages
COMPOSE: shapes are right, every value is finite, the waveform lands inside
[-1, 1], and the denoise loop demonstrably moves the latents rather than passing
noise straight through.
ASSEMBLING IT CAUGHT A REAL GAP, which is why doing this before the checkpoint
arrives was worth it: the audio decode was missing the checkpoint's dec_in_proj
(Conv1d k=1, vae_latent_channels -> num_mels) ahead of BigVGAN. The DiT emits a
32-wide audio latent while BigVGAN expects 2048 mels; without that projection the
two never meet. It is now applied when the weight is present, leaving the
standalone BigVGAN gate untouched.
DESIGN NOTE: noise is an INPUT, not generated internally. Upstream seeds a torch
CPU generator (pipeline_minimax_h3.py:813-843); reproducing torch's RNG
bit-exactly decides WHICH sample you get, not whether the pipeline is correct, so
it is recorded as an open item rather than guessed at.
test_minimax_h3: 20/20 cases, 6370 assertions, clean CPU build (0 warnings).
REMAINING before a real generation: the encoder's VISION tower (reuse of
qwen3_vl_vision.cpp) and MM processor, fl2va/ref2va conditioning, the quantized
loader wiring (GGUF dequant / NVFP4), and a GPU with the checkpoint on it. The
device-resident forward (W2b) and NVFP4 wiring (W10) are where speed work begins.
Still no run on a real checkpoint and no speed number.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
Milestone: the whole t2va path composes end to end (
|
| Component | Gate |
|---|---|
| packed layout (fl2va + ref2va) | exact, fp64 grid bit-exact |
| latent packing / scheduler / request planning | exact |
| DiT forward (f32 / bf16 production stream) | 1.6e-7 / 2.4e-3 |
| GGUF manifest (535 real tensors) | exact, identity name map |
| NVFP4 manifest (1051 real tensors) | exact, layout = our stack |
| audio VAE decoder | 4.2e-9 |
| video VAE ViT3D decoder | 8.9e-8 |
| encoder text tower | 1.2e-7 |
| whole t2va path | composes; correct shapes, finite, in range |
test_minimax_h3: 20/20 cases, 6370 assertions, clean build, all checkers + mutation suites green. No weight bytes checked in anywhere.
Remaining
Encoder vision tower (reuse of qwen3_vl_vision.cpp) + MM processor · fl2va/ref2va conditioning · quantized loader wiring (GGUF dequant / NVFP4) · /v1/videos + MP4 muxing (needs a dependency decision) · device-resident forward (W2b) · video tiling.
No run on a real checkpoint and no speed number. Those need a GPU with the ~41 GB checkpoint on it — a DGX operation, not something reachable from this box.
LoadMiniMaxH3DitFromGguf turns a ComfyUI-format GGUF into a runnable DiT:
* resolve the manifest -- identity name map, `ne` reversal, and the
`comfy.gguf.orig_shape` reshape rule (all three already gated against the
real 535-tensor MiniMax-H3-FL2VA-Q3_K_M.gguf manifest);
* derive the geometry from SHAPES ALONE, because a ComfyUI GGUF ships no
transformer config;
* dequantize every tensor to f32 through the SHARED DequantGgufRowToF32, so
the Q2_K / Q3_K / Q4_K families the H3 GGUFs use are handled by the same code
path every other GGUF model in this tree uses -- no new quant code;
* bind the forward's non-owning views, with a missing tensor throwing BY NAME
rather than yielding a null view the forward would silently read as zeros.
Gated by a synthetic-file LOAD-AND-RUN test rather than a shape check: the
geometry comes back exactly, a loaded weight carries the LOGICAL (torch) shape
rather than the reversed ne, and a REAL DiT forward executes off the loaded
weights with finite, correctly-sized outputs.
test_minimax_h3: 21/21 cases, 6668 assertions, clean CPU build (0 warnings).
Full scripts suite: 128 passed.
NOT covered: a load of the ACTUAL 15.6 GB file (needs the download) and NVFP4
loader wiring (W10 -- the layout is already proven identical to ours).
ENVIRONMENT NOTE: dgx.casa is currently unreachable ("no route to host") and this
workstation has no GPU, so no run on a real checkpoint -- and therefore no speed
number -- is possible from here regardless of how much of the port is done.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
GGUF arm complete — load into a runnable DiT (
|
| Component | Gate |
|---|---|
| packed layout (fl2va + ref2va) | exact, fp64 grid bit-exact |
| latent packing / scheduler / request planning | exact |
| DiT forward (f32 / bf16 production stream) | 1.6e-7 / 2.4e-3 |
| audio VAE decoder | 4.2e-9 |
| video VAE ViT3D decoder | 8.9e-8 |
| encoder text tower | 1.2e-7 |
| GGUF manifest (535 real tensors) | exact, identity map |
| NVFP4 manifest (1051 real tensors) | exact, layout = our stack |
| GGUF load → runnable DiT | geometry from shapes; forward runs |
| whole t2va path | composes; correct, finite, in range |
No weight bytes checked in anywhere — every gate runs against upstream code, the checkpoint's own remote code, or real checkpoint headers fetched by range request.
What remains
Encoder vision tower (reuse of qwen3_vl_vision.cpp) + MM processor · fl2va/ref2va conditioning · NVFP4 loader wiring (W10 — layout already proven identical to ours) · /v1/videos + MP4 muxing (needs a dependency decision) · device-resident forward (W2b) · video tiling.
Environment
dgx.casa is currently unreachable ("no route to host") and this workstation has no GPU. So a run on a real checkpoint — and therefore any speed number — is not possible from here regardless of how much of the port is finished. That's the gating constraint on the remaining goal, not the code.
LoadMiniMaxH3DitFromNvfp4 completes the loader half of the NVFP4 arm. The
compressed-tensors triple
<name>.weight U8 FP4 packed 2-per-byte, [out, in/2]
<name>.weight_scale F8_E4M3 one per group of 16 along K, [out, in/16]
<name>.weight_scale_2 F32 one global scalar
goes through this project's EXISTING DequantNvfp4ToBf16 -- no new quant code,
because the manifest gate had already proven the real checkpoint's layout IS
ours. The fp32/bf16 islands are read as-is, quant sidecars are excluded from the
model tensor set, and the geometry is recovered from the dequantized shapes.
The view binding is now SHARED between the GGUF and NVFP4 arms
(BindMiniMaxH3DitViews), since both land on the same weight contract.
Gated by a synthetic-file LOAD-AND-RUN test: a packed [out, in/2] weight comes
back as the logical [out, in], weight_scale/weight_scale_2 never appear as model
tensors, and a REAL DiT forward executes off the loaded weights.
A MISTAKE WORTH RECORDING. Extracting the shared binder, I applied a blind
`out.` -> `out->` rewrite that corrupted six tensor NAME STRINGS --
"time_embedder.proj_out.bias" became "...proj_out->bias". Two lessons: never
mechanically rewrite across code and string literals in one pass; and the
loader's throw-BY-NAME-on-missing-tensor design caught it instantly, where a
null-view default would have silently bound zeros and produced a
plausible-but-wrong forward. That design choice paid for itself the first time it
was exercised.
test_minimax_h3: 22/22 cases, 6971 assertions, clean CPU build (0 warnings).
NOT done: the DEVICE path that keeps FP4 packed and routes projections through
the cutlass FP4 GEMM -- that is where the speed actually is -- and a run on the
real file. dgx.casa is unreachable and this workstation has no GPU, so no speed
number is obtainable here.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
…one (W3) The repeated unit of the H3-Encoder's Qwen3-VL vision tower now matches upstream at 6.0e-8, so both encoder towers' cores are ported (text tower was 1.2e-7). THE TWO ViTs IN THIS MODEL DISAGREE ON QKV LAYOUT. The video VAE decoder's ViT is PER-HEAD INTERLEAVED ([head][q,k,v]); this vision tower is [q_all, k_all, v_all]. Same tensor size, no error either way -- just a wrong result. Having now ported both, it is worth stating plainly: never assume a qkv layout carries across ViTs inside one checkpoint; read the reshape/permute in the source each time. Other deltas from the TEXT tower, all exercised: * LayerNorm WITH BIAS rather than RMSNorm (eps 1e-6); * rotary applied in fp32, cos/sin shared across heads; * the TANH-approximate GELU (gelu_pytorch_tanh), not exact erf; * NON-CAUSAL attention segmented by cu_seqlens. The test PROVES the segmentation rather than assuming it: perturbing a token in the second packed image leaves every output of the first BIT-IDENTICAL, while the second's outputs do change. A segmentation bug would otherwise slip through a plain tolerance check. test_minimax_h3: 23/23 cases, 7040 assertions, clean CPU build (0 warnings). Full scripts suite: 128 passed. REMAINING for the encoder: the vision surround (Conv3d patch embed, learned pos-embed interpolation, the 2D rotary table, patch + DeepStack mergers) and the MM processor. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
Encoder vision block ported — both tower cores done (
|
| ViT | qkv layout |
|---|---|
| video VAE decoder | per-head interleaved — [head][q,k,v] |
| encoder vision tower | [q_all, k_all, v_all] |
Same tensor size, no error either way — just a wrong result. Never assume a qkv layout carries across ViTs inside one checkpoint; read the reshape/permute in the source each time.
Other deltas from the text tower, all exercised: LayerNorm with bias (not RMSNorm), fp32 rotary, tanh-approximate GELU (not exact erf), and cu_seqlens-segmented non-causal attention.
The test proves the segmentation rather than assuming it: perturbing a token in the second packed image leaves every output of the first bit-identical, while the second's do change. A segmentation bug would slip past a plain tolerance check.
test_minimax_h3: 23/23 cases, 7040 assertions; scripts suite 128 passed; clean build, all checkers green.
Branch: 11 commits
| Component | Gate |
|---|---|
| packed layout (fl2va + ref2va) | exact, fp64 grid bit-exact |
| latent packing / scheduler / request planning | exact |
| DiT forward (f32 / bf16) | 1.6e-7 / 2.4e-3 |
| audio VAE decoder | 4.2e-9 |
| video VAE ViT3D decoder | 8.9e-8 |
| encoder text tower | 1.2e-7 |
| encoder vision block | 6.0e-8 |
| GGUF + NVFP4 manifests (real checkpoints) | exact, identity map |
| GGUF load → runnable DiT | geometry from shapes; forward runs |
| NVFP4 load → runnable DiT | triple dequantized; forward runs |
| whole t2va path | composes; correct, finite, in range |
Remaining
Vision surround (Conv3d patch embed, pos-embed interpolation, 2D rotary table, patch + DeepStack mergers) · MM processor · fl2va/ref2va conditioning · video tiling · /v1/videos + MP4 muxing (needs a dependency decision) · device-resident FP4 forward (where speed actually comes from).
No speed number. dgx.casa is unreachable and this workstation has no GPU.
The full Qwen3-VL vision tower lands, so both encoder towers are ported end to
end: Conv3d patch embed -> bilinear-interpolated position embedding -> 2D rotary
-> block stack over per-frame cu_seqlens -> DeepStack mergers + final patch
merger.
Gated over a RAGGED two-image batch (different h/w), which is what actually
exercises the position-embedding interpolation and the per-frame cu_seqlens.
Merged output and DeepStack features are both within 1e-4 of upstream.
DETAILS THAT WERE LOAD-BEARING, now pinned by the gate:
* the patch embed's Conv3d has kernel == stride, so it is a plain linear over
the flattened patch -- no sliding window;
* torch.linspace(0, n-1, 1) returns the START, not the end (single-row grids);
* .int() TRUNCATES when picking the bilinear corners, it does not round;
* the two merger flavours differ -- the FINAL merger norms the PRE-shuffle
width while the DEEPSTACK mergers norm the POST-shuffle width, and both use
exact-erf GELU rather than the block MLP's tanh approximation. Three
different GELU/norm conventions inside one encoder.
test_minimax_h3: 24/24 cases, 7050 assertions, clean CPU build (0 warnings).
ENCODER STATUS: complete apart from the MM processor (image/video preprocessing
into patches), which is input plumbing rather than model math.
Remaining in the lane: MM processor, fl2va/ref2va conditioning, video tiling,
/v1/videos + MP4 muxing (blocked on a dependency decision), and the
device-resident forward -- which is where speed work begins and needs a GPU to be
worth doing.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
Encoder complete — full vision tower (
|
| Component | Gate |
|---|---|
| packed layout (fl2va + ref2va) | exact, fp64 grid bit-exact |
| latent packing / scheduler / request planning | exact |
| DiT forward (f32 / bf16) | 1.6e-7 / 2.4e-3 |
| audio VAE decoder | 4.2e-9 |
| video VAE ViT3D decoder | 8.9e-8 |
| encoder text tower | 1.2e-7 |
| encoder vision tower (full) | 6.0e-8 / ≤1e-4 |
| GGUF + NVFP4 manifests (real checkpoints) | exact, identity map |
| GGUF → runnable DiT · NVFP4 → runnable DiT | geometry from shapes; forward runs |
| whole t2va path | composes; correct, finite, in range |
No weight bytes checked in anywhere — every gate runs against upstream code, the checkpoint's own remote code, or real checkpoint headers fetched by range request.
Remaining
MM processor (input plumbing) · fl2va/ref2va conditioning · video tiling · /v1/videos + MP4 muxing (blocked on a dependency decision — the tree has no muxer or A/V encoder) · device-resident FP4 forward.
No speed number. The speed work is the device-resident forward; the reference forward is deliberately CPU and unoptimized. dgx.casa is unreachable and this workstation has no GPU, so there is nothing to measure on.
fl2va and ref2va pin their keyframe and reference-audio rows to a NOISED anchor
rather than the clean latent: out = noise_aug*clean + (1-noise_aug)*noise.
The mix is one line. The reason this needed a gate is the ROW ACCOUNTING around
it -- three parts, each easy to get subtly wrong and none of which would raise an
error:
* each visual condition draws noise of length
target_latent_t + imgvid_cond_num_frames and slices the PREFIX matching its
own latent_t -- a shorter condition does NOT get a shorter draw;
* every condition RESTARTS the same seed, so concatenating all conditions and
drawing once would be numerically different for multi-reference requests;
* rows advance by that condition's own patchified row count.
Gated EXACT (<= 1e-6) against upstream with the noise SUPPLIED, so the comparison
isolates the accounting from torch's RNG. That keeps the RNG a single tracked
open item rather than smearing it across every conditioning path -- the t2va
pipeline takes noise as an input for the same reason.
noise_aug == 1.0 is asserted as the documented identity, and shape/row-count
disagreements throw rather than silently mis-slicing.
TOOLING: condition_noise.py uses a RELATIVE import, so the golden generator's
by-path loader now registers a synthetic package whose __path__ is the upstream
directory -- still bypassing the real vllm_omni __init__ (which would drag in
vllm and aenum) while letting relative imports resolve.
test_minimax_h3: 25/25 cases, 8787 assertions, clean CPU build (0 warnings).
Full scripts suite: 128 passed.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
Ports the PURE-MATH half of reference_video.py:
* the canvas pipeline -- aspect clamp to [1:4, 4:1], 768 short edge, max-pixel
rescale, nearest multiple of 32 (Python round-half-to-even);
* the frame schedule -- 24 FPS resampled to the 2 FPS Qwen video rate with
duplicate indices dropped, then timestamps averaged per temporal patch with
the tail padded by REPEATING the last.
Both gated EXACT against upstream.
A TEST-AUTHORING LESSON. I added an invariant of my own -- "the snapped canvas
respects the max-pixel budget" -- and it failed on 3840x1080 (1920x544 =
1,044,480 > 1,032,192). The port was right; MY invariant was wrong: upstream
applies the budget BEFORE snapping to 32 and never re-checks. I corrected the
test rather than the implementation. When a self-invented invariant fails, check
whether the REFERENCE actually holds it before touching the port.
DEPENDENCY BOUNDARY, now explicit. The rest of reference_video.py -- probe,
transcode, frame extraction, audio decode -- shells out to ffmpeg/soundfile.
That is the SAME blocker as /v1/videos MP4 muxing: one dependency decision
unlocks reference-video INPUT decode and generated-video OUTPUT encode together.
It is the single item in this lane that needs a project decision rather than more
porting.
test_minimax_h3: 26/26 cases, 8908 assertions, clean CPU build (0 warnings).
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
Ports klvae.py's split_tiles and blend, so large canvases can be decoded in overlapping tiles. With this the video VAE side is complete for GENERATION -- only the 3D-CNN encoder remains, and that is needed for image/video CONDITIONING, not for producing output frames. The plan is NOT a simple stride, which is the whole reason it needed a gate: it takes the SMALLEST tile count whose MINIMUM overlaps still cover the axis, then distributes the leftover slack in whole vae_ratio units ROUND-ROBIN across the seams. Get that distribution wrong and every tile after the first shifts -- surfacing as seam artifacts in the output, not as an error anywhere. Shipped config: tile_size 256, tile_overlap_min 64, vae_ratio 16 (= prod(space_down) [2,2,2,2,1,1] -- the "f16" in f16t4; vae_ratio_t = 4 is the "t4"). Gated EXACT over six cases (tiled, exactly-one-tile, smaller-than-tile, and a non-default tile/overlap pair), plus structural invariants: tiles cover the axis, every seam meets the minimum overlap, overlaps stay congruent mod vae_ratio, and the cross-fade starts fully on the previous tile and ends fully on the next. test_minimax_h3: 27/27 cases, 9036 assertions, clean CPU build (0 warnings). Full scripts suite: 128 passed. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
…ride) Closes a gap flagged when the denoise loop was ported: its contract says "token_tags must already carry any fl2va vision-span overrides", and nothing in the port produced them. Now it does. THE LOAD-BEARING DETAIL. A vision block is vision_start + pad*count + vision_end, and the WHOLE block -- markers included -- is tagged VIDEO(0). Tagging only the pads leaves two markers as TEXT(1) and shifts every AdaLN modulation index after them: no error, no shape change, just wrong modulation for the rest of the sequence. The test proves each VIDEO run in the output equals a whole emitted vision span, so an off-by-two cannot pass. Tokenization deliberately stays with the CALLER (it owns the tokenizer); this owns the span -> tag mapping, which is the part that must agree with the packed layout. The gate drives upstream with a stub tokenizer, since only span LENGTHS affect tags. test_minimax_h3: 28/28 cases, 9129 assertions, clean CPU build (0 warnings). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
… GroupNorm3D, ResnetBlock3D)
Ports the repeated unit of the video VAE's ENCODER stack. The encoder serves
image/video CONDITIONING (fl2va keyframes, ref2va references); a t2va generation
path does not need it.
Two details pinned by the gate:
* the convolution is CAUSAL in time -- padding[0]*2 frames on the LEFT, none on
the right, with CONSTANT (zero) temporal padding and `reflect` SPATIAL
padding. A symmetric temporal pad would let a frame see the future.
* GroupNorm's statistics span TIME as well as space (32 groups, eps 1e-6), so a
per-frame normalization would silently differ.
CAUSALITY IS PROVEN, NOT ASSUMED. On the bare convolution, changing the last
frame leaves earlier frames BIT-IDENTICAL while the last frame's own output
moves. At block level GroupNorm legitimately mixes across time, so exact equality
would be the wrong assertion there; the weaker claim (the perturbed frame moves
strictly more than the first) is asserted instead, with the strict proof done at
the layer where it actually holds.
test_minimax_h3: 29/29 cases, 9137 assertions, clean CPU build (0 warnings).
REMAINS on the encoder: the Downsample3D + EncoderFCN3D assembly.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
The strided convolution between VAE encoder levels, plus stride support in the shared causal Conv3d it builds on. THE SUBTLETY: when the spatial stride is 2, the input is padded by ONE pixel on the RIGHT of W and the BOTTOM of H -- F.pad(x, (0,1,0,1,0,0)) -- BEFORE a stride-2 conv with padding (1, 0, 0). That asymmetric pre-pad is what keeps the sampling lattice aligned; padding symmetrically instead shifts everything by half a pixel, which is a silent wrong latent rather than an error anywhere. Gated EXACT against the checkpoint's own module, with output extents asserted (H/W halve; T halves under the causal pad). test_minimax_h3: 30/30 cases, 9143 assertions, clean CPU build (0 warnings). REMAINS on the VAE encoder: only the EncoderFCN3D level-loop assembly. The encoder is conditioning-only -- a t2va generation path does not use it. HOUSEKEEPING: a clean rebuild hit ENOSPC partway through; ~15 throwaway build trees at 4.7 GB each had accumulated across this session. Removed, 66 GB free, re-verified from scratch. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
The encoder assembly lands, so BOTH halves of the video VAE are now ported and gated: the ViT3D decoder (8.9e-8) and this 3D-CNN encoder. conv_in -> per level [ResnetBlock3D x N, then a Downsample3D or a 1x1x1 channel match] -> GroupNorm -> SiLU -> conv_out The channel plan is the fiddly part and is now pinned: block_mid[i] = ch * ch_mult[i] block_in[0] = block_mid[0]; block_in[i>0] = block_mid[i-1] A level gets a Downsample3D when space_down[i]*time_down[i] > 1; otherwise a 1x1x1 conv ONLY if its channel count changes, and nothing at all when it does not. THE FAILURE WAS MINE, IN THE TEST. The first run mismatched at 2.25 with correct shapes. Cause: the generator's scale rule tested `".norm" in name`, which silently misses `norm_out.weight` (no leading dot), so generator and test seeded that group-norm gain differently. Rule corrected to `"norm" in name`; the port needed no change. Same lesson as the reference-video invariant earlier -- when a gate fails, establish WHICH side is wrong before touching the implementation. Two for two this session, the fixture was at fault. test_minimax_h3: 31/31 cases, 9150 assertions, clean CPU build (0 warnings). Full scripts suite: 128 passed. The encoder serves image/video CONDITIONING; a t2va generation path does not call it. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-5 [ClaudeCode]
…o VAE
Both measured on the Thor sm_110 board:
* DiT step loop 509.78 -> 18.68 s/step at seq ~3.2k (~27x), and the scaling is
now near-linear in sequence rather than quadratic.
* video VAE decoder device-resident at 1.19e-07 vs the checkpoint's own remote
code, retiring the last CPU-only stage in the video path.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Opus 5 (1M context)
…windows" A real prompted 512x512 run came out globally correct (recognisable subject, right scene layout) but covered in a grid of small squares. That is not quantization noise, and it was not argued away — it was A/B'd: MiniMaxH3SplitTiles plans 512 px as 3 tiles per axis, but plans 256 px as exactly ONE tile. So 256x256 is the one size where tiled and untiled are the SAME computation. 256x256 came out clean; 512x512 did not. Root cause: MiniMaxH3SplitTiles/MiniMaxH3BlendTiles were implemented and gated, and nothing ever called them. Tiling here is NOT a memory strategy. The ViT3D's RoPE coordinates are LENGTH-NORMALIZED — 2*((i+0.5)/n)-1 over whatever grid it is handed — so the grid EXTENT is part of the input. Handing the decoder a 32x32 latent when it was trained on 16x16 tiles gives every token a position the model has never seen, and the patches stop cohering with their neighbours. Every reduced-dimension gate in the suite is smaller than one tile, where tiled and untiled coincide. That is exactly why the whole suite missed this, and why the new gate covers BOTH ends: a single-tile canvas stays bit-identical to the untiled decode, and on a 2x2-tile canvas each tile's un-blended interior equals a standalone decode of that tile's own latent slice exactly. Both failure modes otherwise produce plausible finite frames. The blend extent converts through latent units rather than reusing the plan's canvas pixels directly: vae_ratio and patch_size are both 16 on the real checkpoint but differ in the reduced-dimension configs the gates run at. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
Comparing two DiT checkpoints on output quality only means something if the text conditioning is identical. Two separate encoder runs are not identical, and on a unified-memory box the second run cannot always afford the 13 GB tower alongside a DiT that dequantizes to ~33 GB host plus ~33 GB device. Writing the conditioning once and replaying it via --prompt-embeds fixes both: the comparison isolates the checkpoint, and the second run drops the encoder entirely. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
The NVFP4 arm was OOM-killed during load on Thor: kernel oom-kill, anon-rss 125 GB, box survived. Cause is not subtle once measured -- LoadMiniMaxH3DitFromNvfp4 materializes EVERY weight as host f32, which for this 18.75 GB checkpoint is ~132 GB, and then staging would double it. The GGUF arm never hit this because it already had StreamMiniMaxH3DitToDeviceBf16. StreamMiniMaxH3Nvfp4ToDeviceBf16 is that function's twin: dequantize ONE tensor from its NVFP4 triple, upload it, and let the host buffer die before the next, so peak is the device copy plus one tensor. Same fp32 ISLAND split (vt::MatmulBT rejects an f32-activation/bf16-weight pair, so this is load-bearing), and the same host-resident rope.inv_freq (the forward reads it before any kernel; a device pointer there segfaults). The name-by-name view binding is now factored into BindStreamedDitViews and shared by both streamers, split on the one seam that genuinely differs -- rope.inv_freq, which each loader reads from its own container. Two copies of a ~30-entry binding would drift the moment a tensor is added to one path only. Gated, because "it loads without OOM" is not the same claim as "it loads the right numbers": the streamed weights match the reference loader EXACTLY (max|diff| == 0, NVFP4 values land on bf16 without rounding), the geometry agrees, and rope.inv_freq is asserted to point at host memory. The reference loader stays as-is, and stays the thing this is gated against. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
The NVFP4 arm produces PURE NOISE while reporting the same layers/hidden/heads as the working GGUF, but a different packed seq_len for an identical request (3264 vs 3224). That gap has to come from geometry the run never prints -- patch sizes, latent dims, adaln widths. Reads names and shapes ONLY, no payload, so it is safe on a checkpoint whose weights do not fit: the NVFP4 reference loader is ~132 GB of host f32 and is an OOM kill, which is exactly why "just load it and print" was not an option. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
It was still failing the usage gate for missing VAEs and an output path, which made the one tool that works on a checkpoint too large to LOAD unusable on exactly that checkpoint. Caught because the dump files came out EMPTY and the diff of two empty files reported IDENTICAL -- a pass that meant nothing, because stderr had been sent to /dev/null. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
…mporal) Found by READING upstream's klvae.py instead of inferring from output, after two wrong diagnoses. `decode_base` routes video through `decode_temporal`, never a single pass: the ViT is handed `tokens_chunk_size + token_overlap` temporal tokens at a time. Shipped config (clip_length 17, token_drop 3, vae_ratio_t 4) gives chunk 5 / overlap 2 / pre-pad 3 / frame-overlap 5, so a 12-token latent decodes as 2 chunks of 7 tokens -- never 12. It matters for the same reason the spatial extent does: the decoder's RoPE is LENGTH-NORMALIZED over the grid it is handed, so the TEMPORAL extent is part of the input, and a whole-latent pass gives every token a position the model never saw. This also CORRECTS an earlier commit of mine. Upstream does NOT spatially tile the ViT decode: `decode()` calls `self.decoder(z2)` directly, `_adaptive_decode` tiles only when `decoder_tiling` is set, that defaults FALSE, and the shipped video_vae_config.json carries no tiling key at all. The pipeline now routes to the temporal path; the spatial tiling entry point stays but is no longer on the default path. Head/tail frame isolation is deliberately NOT implemented: isolated_first_frame / isolated_last_frame default false and the shipped config sets neither, so those branches would be untested code. Gated on the chunk ARITHMETIC, which is the part that was wrong -- including that the modulos are Python's, since (-clip_length) % ratio is 3 in Python and -1 in C and a negative pre-padding slices from the wrong end. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
2c70809 to
788ad91
Compare
Temporal chunking alone made the artifact worse, not better (period-16 4.20, the highest measured). Re-reading upstream's chunk loop explains why: it calls `self._adaptive_decode(clip_z)`, so with decoder tiling on, each TEMPORAL chunk is ALSO spatially tiled. The two compose; they are not alternatives, and running either one alone is a third thing that upstream never does. Spatial tiling is a RUNTIME choice rather than a checkpoint property -- the wrapper reads `vae_decoder_tiling` from caller-supplied config, and neither the VAE config.json nor video_vae/source/config.json carries any tiling key -- so both modes are legitimate upstream. Defaulted ON because the tiled path falls through to the untiled one BIT-IDENTICALLY when the canvas fits one tile, making it a no-op below 256 px, and it measures better above. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
The warp-per-query kernel re-reads all of K and V from GLOBAL memory for every query, so traffic is O(seq^2 * d) per head. Invisible at short sequences, dominant at long ones: MiniMax-H3 at 864x480/124f measured 557.74 s/step at seq_len 15424, where 4.7x the sequence of the 512x512 run cost 29x the time. So ~85% of the step was attention, and memory-bound rather than FLOP-bound. The warps of a block now cooperatively stage a K/V tile into shared memory and each evaluates its own query against it, cutting global K/V traffic by the warps per block (8x here). The online-softmax recurrence and the KEY ORDER are unchanged, so this is identical to the untiled kernel rather than merely close. GUARDED to non-causal, single-document, no-window, seq >= 2048 -- only then do all warps in a block provably share one key range, which is what makes a shared tile correct at all. Everything else keeps the per-warp path. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
… the tiled kernel Two holes, both found by running rather than reasoning. 1. Temporal decode aborted on a latent SHORTER than one chunk. latent_t 2 gives pseudo 5 = exactly one chunk, minus the token_drop chunk => num_chunks 0, and the VT_CHECK fired. That is the regime every reduced-dimension gate runs in, so it broke the t2va composition test -- on CUDA only, since CPU does not take the device path. Now decodes directly when there is no chunking to do. 2. The shared-memory tiled attention kernel was UNEXERCISED by the suite that appeared to gate it. It is guarded to seq >= 2048 and every existing case in test_ops_dflash_block_attn is far shorter, so 198,412 assertions passed without the new kernel running once. Added a 2560-key non-causal single-document case comparing the CUDA tiled path against the CPU reference over identical inputs. Same failure family as the NVFP4 gate: green that proves something other than what it appears to. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
The shared-memory tiled kernel measured EXACTLY the same step time as before
(557.82 s vs 557.74 s), which is the signature of a path not being taken. It was
guarded to num_reqs == 1, but H3's packed sequence is {0, used, seq_len} -- content
plus a padding tail, TWO documents -- so the kernel never ran on the workload it
was written for, while the suite reported green.
Generalized instead of widening blindly: the tile is only correct if every warp in
a block shares a key range, so the kernel now derives its request range from the
block's first query and checks the block's last query falls in the same request.
Uniform blocks take the shared tile over [rs, re); a block straddling a boundary
falls back to per-warp global reads with the identical recurrence. Requests are
contiguous, so at most one block per boundary takes the slow path.
Gated on H3's REAL shape (T 3000, cu {0, 2317, 3000}, `used` deliberately not a
multiple of the 8-warp block so the straddling branch is exercised), against the
CPU reference.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Opus 5 (1M context)
Kept honest rather than kept. At 864x480/124f (seq_len 15424) a step is 557.74 s and ~85% of that is attention. Shared-memory K/V tiling measured 686.91 s/step, so it went backwards and is removed. Why it did not help: tiling changed WHERE K comes from, but the inner loop still does a full warp-shuffle reduction for EVERY key (5 shuffles + a broadcast), which is the actual cost. Meanwhile 32 KB of shared memory per block crushed occupancy, and K/V were likely already served from L2. The direction that should work is one key PER LANE -- each lane computing an entire d-length dot product independently with Q in shared memory -- which removes the cross-lane reduction from the inner loop and needs one reduction per 32 keys instead of one per key. Not attempted here. The long-sequence GATES are kept: nothing in the suite previously exercised attention above seq 2048, which is how the first version of the tiled kernel came to be guarded to num_reqs == 1 -- never running on H3's two-document packed shape -- while the suite stayed green. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
Every fl2va primitive was already ported and gated -- the packed keyframe layout,
the VAE 3D-CNN encoder, condition-noise augmentation, the denoise loop's pinned
condition rows -- and none of it was reachable, because the only pipeline passed
{} for both conditioning arguments. This connects them.
Two real gaps had to be closed:
* the decoder loader deliberately SKIPS `encoder.*`, so there was no way to turn
an image into a latent at all. LoadMiniMaxH3VideoVaeEncoderWeights loads that
half (strips `encoder.`, and KEEPS `quant_conv`, which belongs to the encoder).
* MiniMaxH3VideoVaeEncodeToLatent takes the distribution MEAN, not a sample: a
sample would make one reference image condition differently on every run.
fl2va is a PARAMETERIZATION of the t2va pipeline, not a parallel path, which is how
upstream models it too -- t2va is just the empty-keyframe case. The pipeline also
absorbs the row-count difference: keyframes add condition rows to the layout, and
callers can only reasonably supply noise for the TARGET rows.
Gated on the property that matters: with identical prompt, noise and schedule,
conditioning MOVES the video rows (measured 1.20). Wiring that accepted the rows
and ignored them would produce finite, correctly shaped output and pass every
structural check.
Found while doing it that `z_channels` is the MOMENTS width (48 = mean|logvar), not
the latent width (24).
Driver takes binary PPM -- the format this example already writes -- so a frame
from one run can condition the next.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Opus 5 (1M context)
Reference blocks now build the ref2va packed layout, and their visual rows ride the
same pinned-condition mechanism keyframes use -- so this reuses the image->latent
path added for fl2va instead of duplicating it. MiniMaxH3EncodeReferenceImages also
reports each image's latent geometry, because a ref2va block must declare the shape
it occupies in the packed layout.
Gated on the same load-bearing property as fl2va: with identical prompt, noise and
schedule an image reference MOVES the video rows (measured 0.98).
Two refusals are gated as BEHAVIOUR rather than left to chance:
* an AUDIO reference block throws. The audio VAE's ENCODER is not ported (only
its decoder), so honouring one is impossible, and silently conditioning on
nothing is the worse failure -- it looks like it worked.
* fl2va keyframes and ref2va references together throw as mutually exclusive:
one pins frames OF THE OUTPUT, the other prepends whole reference blocks.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Opus 5 (1M context)
No new porting was needed: the 3D CNN is CAUSAL in time, so a reference clip is the image path called with t > 1. What did need care was the block KIND. A video reference must be kVideoAudio, because kImage counts exactly ONE frame however large its latent_t (packing: image blocks add (h/ph)*(w/pw) rows and ignore latent_t entirely). A clip returned as kImage would silently lose its temporal extent and still produce plausible output, so the gate asserts a video reference occupies MORE rows than an image one. ref_audio_t stays 0 -- a SILENT video reference. That is precisely the part that can be honoured without the audio-VAE encoder, which is not ported; a video reference WITH audio throws, like a bare audio reference does. Moves the video rows by 0.83 under identical prompt, noise and schedule. `--ref-video DIR` reads DIR/frame_%06d.ppm, which is exactly what this example writes, so a previous run's workdir feeds straight back in -- clip chaining with no new format. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
DFlashAttnQBlockKernel carries kQ queries per warp so each K/V row is loaded ONCE and reused kQ times, cutting global traffic by kQ. At H3's default canvas the one-query-per-warp kernel moves ~341 TB per step against Thor's 273 GB/s, and attention is ~85% of a 557 s step. Uses NO SHARED MEMORY by design: the previous attempt staged K/V tiles in shared and measured 23% SLOWER, because 32 KB per block against a 48 KB/block limit left roughly one block per SM. Registers are the right store here. Per-query key ranges are precomputed, so causal, sliding-window and multi-request layouts all work and the per-query skip is warp-UNIFORM (no divergence). Uninstantiated so far -- no launcher dispatches it yet, so this commit cannot change behaviour. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
…eo+audio
The last two unwired ref2va conditioning modes were blocked on the same thing:
only the audio VAE's DECODER (the DAC/BigVGAN vocoder, gated at 4.2e-9) was
ported, so MiniMaxH3DenoiseT2va THREW on any reference block carrying audio.
This ports the analysis half and connects it.
THE ENCODE PATH IS NOT A METHOD ON THE SHIPPED MODULE. DacAudioVAE exposes only
decode (dac_audio_vae.py:211-225); vLLM-Omni composes the encode by hand
(vae.py:317-325), and this mirrors that composition exactly:
preprocess right-pad to a whole hop_length dac_audio_vae.py:201-209
-> Encoder strided conv stack dac_audio_vae.py:90-117
-> pre_block AttnProjection dac_attn_proj.py:69-88
-> mean_proj Conv1d(32 -> 32, k=1) dac_audio_vae.py:157
Ported, each citing the upstream line it came from: Snake1d
(dac_audio_vae.py:25-40 -- one parameter, never log-scaled, and it both scales
the sine and forms the reciprocal, so it is NOT the decoder's SnakeBeta),
ResidualUnit (:50-66), EncoderBlock (:69-87), Encoder (:90-117), and the
AttnProjection block: causal scaled-dot-product attention with an ASSEMBLED
[q_bias | zero_k_bias | v_bias] qkv bias (dac_attn_proj.py:52-66), the MEAN over
heads, the adaptive-average-pool narrowing 2048 -> 32, and the GeGLU MLP
(:8-25). Only the narrowing branch (in_dim > out_dim) is implemented -- it is
the one the checkpoint ships; the widening branch refuses loudly rather than
running untested.
mean_proj and never logs_proj: conditioning takes the distribution MEAN, the
same rule MiniMaxH3VideoVaeEncodeToLatent follows, so a reference conditions
identically on every run. logs_proj is not even loaded.
GATED AGAINST UPSTREAM, STAGE BY STAGE.
scripts/gen-minimax-h3-audio-vae-encoder-goldens.py imports the CHECKPOINT'S OWN
remote-code modules and runs them at reduced dimensions, rebuilding every
parameter from the shared FNV-1a + splitmix64 stream keyed by state_dict NAME,
exactly as the DiT and decoder generators do -- so the golden is reproducible
from source alone and no weight byte is checked in. Three stages are compared,
not one end-to-end number:
conv stack max|diff| 2.98e-08
AttnProjection max|diff| 1.64e-07
whole encode max|diff| 1.86e-08
The reduced input is 25 samples against a hop of 4, deliberately NOT a multiple:
a port that skipped preprocess returns 6 frames instead of 7 and fails.
LOADER. LoadMiniMaxH3AudioVaeEncoderWeights takes the half the decoder loader
skips, with the same rules because it is the same file: strip encoder., keep
top-level pre_block.*/mean_proj.*, and accept all three weight-norm spellings
(legacy weight_g/weight_v, parametrizations.weight.originalN, and a plain
materialized weight, the last reconstructed exactly -- round-trip <= 1e-6).
Gated against the REAL 1087-tensor manifest, which also confirms the SHIPPED
geometry from shapes alone: encoder_rates [2,4,4,5,5], latent_dim 2048,
attn_proj_dim 32, qkv 3x the INPUT width. A plain Linear's .weight must not be
mistaken for a materialized weight-norm, and that is asserted.
WIRING, GATED ON THE PROPERTY THAT MATTERS. The two VT_CHECK refusals are gone;
what is refused now is an audio-bearing block with no encoded rows behind it,
because the layout would grow around rows nothing ever wrote. Reference rows
reach the loop's audio_ref_rows (previously always {}), and the initial-audio
scatter mirrors the video side (pipeline_minimax_h3.py:937-955). Under identical
prompt, noise and schedule:
audio reference moves the AUDIO rows by 0.51, video rows 5.5e-3
video+audio reference moves the AUDIO rows by 0.71, video rows 5.5e-3
(against the SILENT same-clip control)
a DIFFERENT waveform moves the AUDIO rows by 7.1e-4
a DIFFERENT clip, same audio moves the VIDEO rows by 3.7e-2
The last two are the ones a wiring that reserved the rows and then pinned a
constant would fail; every structural check would still pass.
DRIVER. --ref-audio f.wav, attaching to a --ref-video block when there is one
(one kVideoAudio block carrying both) and standing alone as a kAudio block
otherwise. The WAV reader lives in the library next to the writer it inverts
(MiniMaxH3ReadWav) so it is unit-gated: 16-bit PCM, chunk-list walk rather than
a fixed 44-byte header, mono REPEATED to stereo, and a non-32 kHz file REFUSED
rather than mis-encoded -- there is no audio dependency here to resample with.
Suite: 61/61 cases, 29756 assertions, CPU build. The encoder gate is proven to
BITE, not merely pass: dropping the causal mask in the AttnProjection moves it
from 1.64e-7 to 1.05e-3 and fails.
Record: spec, model-matrix row, parity ledger and state log carry the same
numbers, and record what is NOT gated -- no real-checkpoint render conditioned
on a real waveform has been run.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Opus 5 (1M context)
…iagnosis corrected
Carrying kQ queries per warp reuses each K/V row kQ times, in REGISTERS with no
shared memory (ptxas: 64 regs at kQ=4, 108 at kQ=8, zero stack, zero local),
deliberately avoiding what sank the earlier shared-memory attempt.
Paired reps with the arm order reversed between them, box noise <=0.07%:
864x480 seq 15424 baseline 557.81/557.78 s
kQ=4 553.13/553.15 s -0.84%
kQ=8 765.57 s +37.2%
512x512 seq 3264 baseline 28.05/28.03 s
kQ=4 28.99 s +3.32%
Output checksums identical across every arm. Reverted: a wash at the production
canvas does not buy a reproducible regression at the short one.
THE PREMISE WAS WRONG, and it is the same premise that sank the tiling attempt.
One head's K+V at seq 15424 is 3.9 MB against Thor's 32 MB L2, and every warp on a
given blockIdx.y streams the same rows -- so the ~341 TB/step traffic figure was
never the bound; that traffic was already L2-resident. Q-blocking removes re-reads
that cost nothing, does NOT remove the per-key warp-shuffle reduction (kQ of them
per K row, the same total), and adds a serial expf chain of kQ softmax updates,
which is why kQ=8 is far worse than kQ=4 rather than twice as good.
The open lever is the per-key REDUCTION -- one key per lane -- not the memory
source.
Kernel kept in-tree, uninstantiated and annotated with the numbers: the verdict is
hardware-specific and a part with small L2 relative to one head's K/V could flip it.
Tests KEPT, and they are the durable result. A green suite does not prove a kernel
ran: a RED proof (perturbing the Q-block store) turned 6/11 cases red, and three
new LONG cases cover causal, causal sliding-window and ragged multi-request. The
two pre-existing long cases are both non-causal single-document -- the one mask
where every query sees the same key range -- so they could not have caught a
per-query mask bug.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Opus 5 (1M context)
…oth conversions This is why renders were dark and washed out. Found by reading ComfyUI's implementation rather than by more staring at output. Upstream's VAE wrapper converts on BOTH sides, and both steps sit OUTSIDE ViT3DDecoder: decode: dec*std + mean, clamp [0,1], *2 - 1 (comfy/ldm/minimax/vae.py:693) encode: ((x + 1)/2 - mean) / std (vae.py:659) We did neither. Raw decoder output went straight to a writer expecting [-1, 1], which compounds two errors: the per-channel means differ (0.485/0.456/0.406) so it CASTS COLOUR, and a std of ~0.22 compresses the true dynamic range ~4.4x. The 864x480 render measured a mean pixel of 18/255. Exactly the same shape of omission as post_quant_conv: a wrapper stage outside the decoder, which is why the decoder's own 1.19e-07 gate against the checkpoint's remote code could never have caught it. Gating the component proved the component. Fixed on the ENCODE side too -- keyframe and reference conditioning had been handing the CNN [-1,1] pixels where it wanted ImageNet space, so every reference mode was feeding it the wrong distribution. Gated PER CHANNEL rather than on a round trip: a shared mean/std would round-trip perfectly and still cast colour. Also gates the clamp ORDER, since clamping after the [-1,1] map would let out-of-gamut values survive rescaled instead of saturating. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
…p (sm_110)
The portable attention kernel was 47x slower than this tree's own GEMMs on the
same chip: MEASURED 0.63 TFLOP/s of attention against 30.0 TFLOP/s of cuBLASLt
GEMM on Thor (sm_110, 20 SMs, no FlashAttention-2), which put attention at 97%
of a MiniMax-H3 DiT forward at the production canvas and made a 50-step render
~7.8 h.
The warp-per-query kernel reduced ONE KEY AT A TIME: ~4 useful FMAs per lane,
then a 5-stage __shfl_down_sync tree plus a broadcast, and every key's
(m, l, acc) update read the previous key's — so the warp was LATENCY-bound at
~8% of the CUDA-core ceiling rather than throughput-bound at the ~32% its
instruction mix allowed.
DFlashAttnChunkKernel processes keys 32 AT A TIME:
- each lane keeps its 32 per-key partials in registers (indices are
compile-time, so p[32] never spills to local memory);
- ONE butterfly REDUCE-SCATTER — 5 rounds, 31 shuffles against 192 — turns
the 32x32 partial matrix into "lane L holds the complete score for key L",
each round exchanging only the half of the live range the partner owns;
- ONE online-softmax update per 32 keys instead of 32, i.e. 2 expf per chunk
against 64, which is the 32x shortening of the dependency chain;
- loads stay lane==element and COALESCED, with the per-lane element partition
chosen per head_dim (contiguous+vector at 64/128, strided at 96).
Same-binary A/B, --denoise-only --steps 3, per FORWARD on Thor:
seq_len 3264: 28.08 -> 16.33 s (1.72x)
seq_len 6080: 91.61 -> 51.97 s (1.76x)
seq_len 15424: 558.39 -> 316.95 s (1.76x)
The warp arm reproduces the standing 28.16/91.81/559.46 s baselines to within
0.3%. A joint t = G*seq + A*seq^2 fit with a SHARED linear term (the GEMMs are
untouched) attributes 542.81 -> 301.01 s to attention at the production canvas:
1.80x, and 0.63 -> 1.13 TFLOP/s over the 341.1 TFLOP that canvas costs. A
50-step 864x480/124f render goes ~7.8 h -> ~4.4 h. Chunked max/sum is the same
softmax but not bitwise identical to the per-key recurrence, so this is gated on
tolerance against the CPU reference, not bit-equality.
RED PROOF, because a green suite does not prove a kernel ran: perturbing the new
kernel's store by 1.001x turns 7 of 12 cases red (193,127 assertions) — exactly
the CUDA-touching ones, including LONG causal, LONG sliding-window, LONG ragged
multi-request and H3's real {0, used, seq_len} TWO-document packing — and
reverting returns 12/12. head_dim 96 had NO coverage at all (every case was
16/32/64/128), so the fast path's head_dim/32 == 3 instantiation was shipping
unexercised; two short and two LONG d=96 cases come with this change.
Also recorded, not deleted: one-key-per-lane (a whole d-length dot product per
lane, Q broadcast from shared) removes the reduction completely and measured
43.35 s/forward against 28.08 at seq_len 3264 — 54% SLOWER — because a whole K
row per lane makes every load instruction 32-way scattered; the work moved from
the shuffle pipe to the memory pipe. It stays in-tree behind
VT_DFLASH_ATTN_KEYLANE=1 (with VT_DFLASH_ATTN_WARP=1 for the incumbent) so the
three-way verdict reproduces on ONE binary. That negative is what produced the
winning design: kill the per-key reduction WITHOUT surrendering coalescing.
examples/minimax-h3-gen --denoise-only had a real reporting bug: an N-sigma
schedule runs N-1 forwards but elapsed time was divided by N, under-reporting
per-forward cost by 1.5x at --steps 3 — enough to make an A/B look like a
regression. It now prints forwards= and per_forward=.
Gates (Thor sm_110, all green): test_ops_dflash_block_attn 12/12 (3,260,884
assertions), test_ops_dflash_paged_block_attn 1/1 (795,648), test_ops_attention
9/9, test_qwen27_dflash_spec_decode 3/3, test_minimax_h3 62/62.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Opus 5 (1M context)
Rounds 1 and 2 of the sm_110 attention work rescheduled the kernel; both left it
on the CUDA CORES. After the chunked reduce-scatter it still ran at 1.09 TFLOP/s
against 30.0 TFLOP/s for our own cuBLASLt GEMMs on the same chip, which is not a
scheduling problem — it is the wrong MATH UNIT.
DFlashAttnMmaKernel puts both GEMMs of attention on the bf16 tensor cores with
mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32. That instruction was PROBED
to assemble, launch and accumulate on sm_110 before anything was written: the
kind::mxf4nvf4 restriction recorded for this box is specific to the block-scaled
FP4 MMA and does not reach standard bf16 MMA.
One warp owns 16 query rows and the full head_dim of output; 4 warps per block
share a staged 32-key K/V tile. S, P and the ENTIRE O accumulator stay in
registers, because the m16n8k16 C fragment (rows gid/gid+8, cols tig*2+{0,1}) is
exactly the next MMA's A fragment — so P costs zero shuffles and zero shared
round-trip, and the online-softmax (m, l) state reduces across a quad with two
__shfl_xor instead of a barrier. Only K and V are staged: 17,408 B/block at
head_dim 128, half of what the round-1 shared-memory attempt spent when it
measured 23% slower. 168 registers/thread, no spill, 3 blocks x 4 warps = 25%
occupancy — the trade a tensor-core kernel makes.
Same-binary A/B (VT_DFLASH_ATTN_MMA=0 selects the round-2 arm), --denoise-only
--steps 3, per FORWARD:
seq_len 3264 16.34 -> 3.25 s (5.03x)
seq_len 6074 51.96 -> 7.11 s (7.31x)
seq_len 15421 321.13 -> 32.71 s (9.82x)
The disabled arm reproduces the standing 16.33/51.97/316.95 baselines to within
1.3%. A joint t = G*seq + A*seq^2 fit over all six points with a shared linear
term puts attention at 311.6 -> 23.1 s at the production canvas: 13.5x, and
1.09 -> 14.75 TFLOP/s over that canvas's 341.1 TFLOP. A 50-step 864x480/124f
render drops from ~4.4 h to ~27 min.
NUMERICS. Q/K/V are already bf16 in the production stream so QK^T loses nothing,
but the P*V GEMM must round the probabilities to bf16 exactly as FlashAttention
does. The tensor-core path is therefore gated at a bf16 tolerance (5e-3 max,
1e-3 RMS — the bar test_minimax_h3 already uses for its bf16 stream) and f32
inputs deliberately stay on DFlashAttnChunkKernel at 2e-5. Nothing was loosened.
Every pre-existing CUDA case in test_ops_dflash_block_attn uploads f32, so this
kernel would have shipped NEVER-EXECUTED. Added with it: 7 short bf16 semantic
corners, 6 LONG bf16 cases (H3's two-document {0, used, seq_len} packing at
d=128, d=96, causal, SWA, ragged causal, ragged non-causal) and a bf16 RED case
pinning causal-vs-non-causal separation at >0.5. RED PROOF: perturbing the store
by +0.25 turns exactly those 2 cases red (26 assertions) plus 1 of 62
test_minimax_h3 cases — the CUDA bf16 DEVICE FORWARD, i.e. the real production
path — while the two suites that do not route here stay green; reverted, 15/15
and 62/62.
Gates on the sm_110 board: test_ops_dflash_block_attn 15/15 (6,352,687
assertions), test_ops_dflash_paged_block_attn 1/1 (795,648), test_ops_attention
9/9, test_qwen27_dflash_spec_decode 3/3, test_minimax_h3 62/62.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Opus 5 (1M context)
…error build repaired (#61) row/BUILD-CPU-WERROR-MXFP4-TESTFIX squash. The #54 MXFP4 test helper is used only inside the VT_MARLIN_NVFP4 region but was defined unguarded, so the plain CPU-only Release build (-Werror, no marlin) fails with unused-function at test_linear_method.cpp:77 — reproduced on a clean checkout of main 8d55550 before the fix; the TU compiles clean after. 2-line guard, no behavior change on any configuration that compiled before. Found while build-gating the PR #26 merge. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
…s model
Controlled A/B at 864x480/124f/50 steps: identical prompt, seed, code and VAE, only
the DiT quantisation changed. Q3_K_M produces a murky silhouette under a visible
lattice; Q4_K_M produces a photoreal facial close-up -- correct anatomy, hair
detail, layered robes, bamboo in mist, snow as real bokeh.
Measured on the same frame, against a ~1.00 control period:
VAE-patch lattice p16 1.18 -> 1.07
DiT-token lattice p32 1.19 -> 1.13
mean exposure 50.7 -> 60.7
Upstream predicted this. ComfyUI PR 15298 reports H3's partial split-half RoPE
creates CHANNEL-WISE MAGNITUDE OUTLIERS that corrupt INT8 quantization, so 3-bit
K-quant is well past the edge. The artifact sitting at the DIT-TOKEN period (32 px
= one token) is the signature of per-token weight error, not a layout bug -- which
is consistent with the scan that cleared patchify order, the area-normalized
spatial grid, the {1,4,4,4,4} temporal grid, the qkv split and the qk-norm against
ComfyUI's implementation.
This A/B cost ~28 min per arm only because of the 16.6x attention work. At the
original 574 s/forward it would have been 8 h per arm, which is why the quality
question stayed open as long as it did.
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Opus 5 (1M context)
…ION lane lands (#26) feat/minimax-h3 squash (84 commits): the vLLM-Omni MiniMax-H3 33.1B CFG-distilled joint video+audio diffusion transformer, ported DERIVE-AND-SHIP. Packed fl2va/ref2va layout (fp64 position grid BIT-EXACT), euler-ancestral scheduler, 50-block AdaLN DiT forward (1.6e-7 vs upstream at reduced dims), BOTH VAEs reimplemented from the checkpoint's remote code (audio 4.2e-9, video ViT3D 8.9e-8), truncated Qwen3-VL encoder (1.2e-7), t2va/fl2va/ref2va pipelines, ComfyUI-GGUF + NVFP4 loaders, /v1/videos API logic, WAV/PPM/MP4-argv output, device- resident f32 forward GPU-VERIFIED on Thor sm_110. Attention routes through shared vt::DFlashBlockAttention(causal=false); projections through vt::MatmulBT; kMiniMaxH3 appended after main's landed enum ids. Merge resolutions: model-matrix keeps main's newer Laguna/kimi-k3 rows + both H3 rows (summary 359, engaged 45); pre-cutover BENCHMARKS sections rolled into the benchmark record; H3 STATUS row compacted to one in-budget cell with the full narrative preserved in state.md; ratchet paid by collapsing the superseded Laguna W7 block to a ledger pointer. Gate: clean full CPU -Werror build + full ctest on the merged tree (test_async_llm parallel-starve flake passes serially, the documented class). The two pre-existing plain-CPU -Werror breaks found by this gate landed separately as #61/#62. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
|
Landed on main as 574399c (mudler-authored squash of the 84-commit lane, rebased across today's six landings). Merge gate: all doc/record gates + a clean full CPU -Werror build + full ctest on the merged tree (the two pre-existing plain-CPU -Werror breaks it surfaced landed separately as #61/#62; test_async_llm is the documented parallel-starve flake, passes serially). Next per the lane: the FP4 speed path on GB10 (sm_121). 🤖 Generated with Claude Code |
Ports MiniMax-H3 — a 33.1B CFG-distilled joint video+audio diffusion transformer — into vllm.cpp. It is the project's first DIFFUSION architecture: no KV cache, no sampler, no logits, so it is not a causal-LM registry model. Upstream is
vllm-project/vllm-omni, not the vLLM repo, and it sits beyond the parity pin (555967922).Every portable piece is done. What remains needs a GPU, not more porting.
How this is gated (no weight bytes checked in)
Two techniques carry the whole PR:
trust_remote_codeand therefore had to be reimplemented in C++) and execute them at reduced dimensions. Both sides rebuild weights from a shared FNV-1a + splitmix64 PRNG, so a golden is reproducible from source alone.Correctness gates
rf_v_to_x0/v1/videoscontract, job store, route dispatchThe fp64 grid is gated bit-exact because it feeds RoPE. Matching it required reproducing upstream's arithmetic order, not just its formulas:
numpy.linspace(endpoint=False)evaluatesi*step + start, and upstream keeps a numpy-pairwise and a Python-sequential span sum deliberately separate because they diverge in the last ulp from n=16 (packed_sequence.py:101-113).Attention reuses the shared
vt::DFlashBlockAttention(causal=false)for packed varlen — no new kernel.The ffmpeg boundary
/v1/videosreturns an MP4, and upstream shells out to ffmpeg. This library has no subprocess precedent, so the split was put to the project owner and ratified: "re: ffmpeg invocation, correct - let's keep in the examples only".src/vllm/builds the artifacts (PPM frames, WAV) and the argv — and spawns nothing.examples/minimax_h3_mux/performs the invocation./v1/videostherefore takes a caller-suppliedVideoRunnercallback.Serving
POST /v1/videos(async),POST /v1/videos/sync,GET /v1/videos/{id}, registered onApiServervia the server's existing additive/opt-in pattern: they appear only whenset_video_runnerhas been called, so a server built without video support is byte-identical to before — no new constructor parameter, no existing caller touched.Async uses a joinable worker drained in
~ApiServer; a detached thread would outlivethisand write into a destroyed job store. A throwing runner fails the job, not the process.Hardware verdict (corrected mid-PR)
An early reading called this hardware-blocked. That was wrong and the developer corrected it: the bf16 release (~354 GB) does not fit one box, but the quantized arms do — GGUF (DiT 15.6 GB + encoder 14.6 GB + VAEs ~11 GB ≈ 41 GB) and NVFP4 both load. Both quantized loaders are implemented and gated against the real manifests.
The device-resident forward (W2b) — landed and GPU-verified
MiniMaxH3DitForwardDeviceruns the whole DiT graph with every activation in device memory, so the block stack — and above it the 50-step denoise loop — never round-trips through the host. Verified on a real GPU, not just compiled: the CUDA case is proven to have run (220 assertions execute; 9673 total on GPU vs 9453 on CPU, where it skips).Three H3 kernels were needed, and only three, because the port reuses the tuned shared ops (
MatmulBT,RmsNorm,QkvSplit,SiluAndMul,Add,IndexSelect/IndexCopy,DFlashBlockAttention). H3's RoPE looks exotic but is plain NeoX rotate_half — only the angles are unusual (three axes off the fp64 position grid), so a per-row cos/sin cache feedsvt::RopeFromCachewith no bespoke kernel. That left two indexed AdaLN modulates and an ungated SiLU, in akMiniMaxH3glue table mirroring thekLagunaprecedent — but registered on both kCPU and kCUDA (Laguna's is CUDA-only), so the whole device path is gated in CPU CI too.Not bit-identical to the CPU reference, deliberately:
vt::RmsNormreduces in f32 where the reference accumulates in double, and f32 is what upstream torch does. Held to the same goldens instead.Not yet built — honestly recorded
vt::FusedChainglue folds on the device forward (the bf16 fold also clears its merged-GEMM allowlist entry, added with a specific reason rather than silently).Open: there is no vllm-omni parity pin — the upstream-sync protocol covers only the vLLM repo, and H3 is both beyond the pin and outside that repository.
Record
.agents/specs/minimax-h3.md, model-matrix rowMODEL-DIFFUSION-minimax-h3-mini-max-h3-dit+ checklist/rollup, roadmapROAD-V1-H3,docs/STATUS.md,docs/BENCHMARKS.md, parity-ledger and state entries.Checkers green on committed HEAD:
check-agent-record,check-model-checklist,check-doc-checkpoint,check-readme-structure,check-fusion-consistency,check-runner-routing-consistency,check-device-leakage,check-env-doc, plus the checker mutation suites.🤖 Generated with Claude Code