fix(minimax-h3): half of every video was discarded, plus the missing user docs - #68
Merged
Merged
Conversation
mudler
force-pushed
the
fix/h3-audio-duration-and-readme
branch
from
August 6, 2026 16:41
458c813 to
a07403d
Compare
mudler
added a commit
that referenced
this pull request
Aug 6, 2026
…lash-attn build (#68 W2/W3) QUANT-CT-MXFP4-FLASH-AUDIT names the flash +807us/step (+22.4us/call) MXFP4-8B decode term: vLLM's flash-attn (_vllm_fa2_C) is compiled --use_fast_math; our vendored FA2 TUs (same 2c839c33 source) were not. Without it nvcc emits the precise div/rsqrt sequences: 5448 SASS instrs / REG246 vs 4832 / REG255 with it — matching vLLM's 4880 / REG255, HMMA(768)/LDSM(408)/LDGSTS(120) byte-identical. At the tiny latency-bound decode tile that ~11.6% scalar bloat ≈ the +14% per-call time. Non-byte-exact (fast-math div/rsqrt) => SACRED battery owed before any default claim; measured recovery + gates in follow-up. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
mudler
added a commit
that referenced
this pull request
Aug 6, 2026
…4 flash +807us decode term (#68) QUANT-CT-MXFP4-FLASH-AUDIT closes the #67-OWED flash audit: why the IDENTICAL-grid flash_fwd_splitkv MXFP4-8B decode kernel runs ~+22us/call slower than the vLLM oracle. W1 (fresh same-tool reconciliation, analyze_decode.py on the #57 c8 traces): flash MAIN ours 178.8 vs vLLM 156.3 us/call (+807us/step), combine tied, totals reconcile — REAL, not a reconciliation artifact. W2 (cuobjdump/ptxas arbiter, nvcc/ptxas 13.0.88 both sides): kernel-version REFUTED — vLLM v0.25.0 pins vllm-flash-attn @ 2c839c33, the exact commit we vendored, so flash SOURCE is byte-identical. The real difference is a BUILD FLAG: vLLM's flash-attn is compiled --use_fast_math; our vendored FA2 TUs were not. Decode kernel <128,64,128,4>: ours 5448 instrs/REG246 -> 4832/REG255 with fast-math (= vLLM's 4880/REG255), with HMMA(768)/LDSM(408)/LDGSTS(120) byte-identical. The +616 scalar instrs (precise div/rsqrt) = +11.6% ~ the +14.3% per-call time at the latency-bound decode tile. hd256 (27B/35B) also leans 4928->4560. Ruled out: -maxrregcount=255 (5424/246); our own nvcc-13 compute_80 PTX (5440). FIX: CMakeLists flash TU COMPILE_OPTIONS += -use_fast_math (mirrors flash-attn's own build). Real sm_121a build's flash object VERIFIED = 4832 instrs/REG255. Non-byte-exact (fast-math div/rsqrt) => runtime nsys A/B + SACRED battery OWED (box-contended by the active H3 campaign); NO default flipped. Full forensics in .agents/benchmark-record.md (QUANT-CT-MXFP4-FLASH-AUDIT) + spec; STATUS/BENCHMARKS/NOW/state updated. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
mudler
added a commit
that referenced
this pull request
Aug 6, 2026
…ash regression); gap is occupancy/L2, not the SASS QUANT-CT-MXFP4-FLASH-AUDIT runs the #67-OWED flash audit to a MEASURED verdict. W1: fresh same-tool c8 decode-window — the flash gap on CURRENT main is +12.5us/call (+450us/step) ours 168.8 vs vLLM 156.3, SMALLER than #57's +807 (leaner marlin/glue). W2 lens 1 (cuobjdump HYPOTHESIS): vLLM's flash-attn is --use_fast_math, ours was not (same 2c839c33 src; +fast-math -> 4832/REG255 = vLLM's 4880/255). W2 lens 2 (MEASURED, the arbiter): a controlled same-build nsys A/B REFUTES it — -use_fast_math makes flash +21us/call SLOWER (168.8->189.8); the kernel is memory-latency-bound (ncu occ 8.3%, register-limited), so the higher reg count (246->255) lowers occupancy, which dominates. cuobjdump was necessary but NOT sufficient. So the flag is REVERTED (a CMakeLists NOTE records why); no functional code ships. Real residual = occupancy/L2 (ncu: 38% smem-scoreboard + 37% barrier stalls); vLLM-side ncu owed (box OOM-rebooted twice under 3-way contention). No default flip (no throughput win). Records + spec updated. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
mudler
added a commit
that referenced
this pull request
Aug 6, 2026
…IDENTICAL, gap is an irreducible-for-us ptxas SASS-quality difference; no lever (#75) Runs the #69-owed ours-vs-vLLM flash decode ncu diff to a MEASURED verdict on an idle box, and CORRECTS #69's "occupancy/L2-bound" framing. W1 (matched-c8 ncu, both engines, grid 1x3x64, full section set): - vLLM runs at the SAME 8.33% occupancy as ours — both smem-limited to 1 CTA/SM by the byte-identical 81.92 KB kSmemSize (the 216-vs-241 reg delta is moot). - L2 hit ~1% on BOTH (KV streamed from DRAM, no adjacency/warmth to win), identical short-scoreboard stall structure. Only measured diff: +13% instr. - The prior #69 "8.3% occupancy" was the batch-1 short-context num_splits=1 kernel; the prior ncu pair was context-mismatched. Now measured on the real c8 kernel with everything else held identical. W2 (mirror-first levers, all refuted DECISIVELY by building vLLM's exact recipe): - arch-mirror (compute_80 PTX driver-JIT): neutral (no change from native). - compute_80 + -use_fast_math (vLLM's EXACT recipe): reproduces vLLM's SASS profile EXACTLY (241 reg, 17,008 instr vs vLLM 241/17,020) and is STILL ~167us = the same as native ~165us and ~10us slower than vLLM's ~157us. So matching vLLM's arch, fast-math, register AND instruction count does NOT close the gap: the kernel is not instruction-bound; the residual is vLLM's wheel-ptxas SASS-scheduling quality, un-reachable from nvcc 13.0. (native+fast-math REGRESSES to 189.8us per #68; compute_80+fast-math is neutral.) W3 VERDICT: MXFP4 stays BELOW-FLOOR at c2-c8 (binding unchanged 1.020/0.962/0.966/0.969). NO lever exists on our stack — occupancy, L2, num_splits, reg count, instruction count, __launch_bounds__, arch-mirror and fast-math ALL refuted by measurement. The flash term (+450us/step, ~40% of the c8 residual) is an irreducible-for-us ptxas quality gap. NO default flip owed; no functional code ships. CMakeLists NOTE + benchmark-record #75 record the closed levers so they are not re-tried. No byte-exact razor / SACRED battery owed (nothing shipped to gate). check-fusion-consistency (minimax_h3_video_vae_device) is a pre-existing red, untouched by this change. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
mudler
force-pushed
the
fix/h3-audio-duration-and-readme
branch
from
August 6, 2026 20:47
f032e0b to
ae47f25
Compare
mudler
added a commit
that referenced
this pull request
Aug 7, 2026
…user docs PR #68. A real bug found by RENDERING, not by the suite, plus the user-facing documentation the H3 lane never had. THE BUG. A 124-frame render silently muxed as 61 frames. `audio_t` is the PER-CHANNEL latent length (the planner sets it from 40 Hz * duration) and the packed layout carries `audio_t * audio_channel` ROWS, one per (channel, step). The denormalize step divided by the channel count, so the decoded audio ran half the video's duration; because the muxer passes `-shortest`, that silently truncated the VIDEO to half its frames too. The fix is one line, and it now agrees with the same function's own `target_audio_rows = request.audio_t * request.audio_channel` a few lines above. WHY THE SUITE WAS BLIND, which is the part worth keeping. Every existing gate asserts shape SELF-CONSISTENCY, and a uniformly halved pipeline is perfectly self-consistent: the shapes agreed with each other, they were just half as long as the request asked for. ffprobe on the artifact exposed it. The gate added is the one that is NOT self-consistency: the structural t2va e2e case now decodes `audio_t` latent steps INDEPENDENTLY and requires the pipeline to have produced exactly that many samples. RED-first: restoring the divide turns that case RED (1 failed assertion), removing it goes GREEN. Suite 68/68, 46586 assertions. A SECOND BUG FROM THE SAME SESSION, DROPPED RATHER THAN LANDED TWICE. Conditioning PREPENDS rows to img_pos, so the denoise loop returns condition rows followed by the targets, and unpatchify was handed the whole set. This branch fixed it by carrying the layout's own `update_mask` out of the loop; main fixed the SAME bug independently in row/H3-RENDER-CLOSE by taking the TRAILING target_video_rows with a VT_CHECK on the row accounting. On rebase the update_mask mechanism and its test were DROPPED: landing a second solution to a fixed bug is churn, and the positional form already ships gated. The alternative is recorded in .agents/state.md rather than silently lost. DOCS. The video+audio generation path had no user-facing documentation at all, and the multimodal INPUT interface was documented WRONG: `vllm-cli --image/--video/--audio` does not exist, the CLI is text-only. Multimodal input is served over the OpenAI API as `image_url`, `video_url` and `input_audio`/`audio_url` content parts on /v1/chat/completions (src/vllm/entrypoints/openai/chat_mm.cpp), which was documented NOWHERE, so docs/USAGE.md now states it rather than the wrong claim merely being removed. README gains a News section led by video+audio generation, paid for INSIDE the 30,000-char landing-page budget (29,997) rather than by raising it. PUBLIC DOCS. The 19-line narrative this branch originally appended to docs/BENCHMARKS.md was a non-canonical H2 section and pushed the page over its prose budget: converted into three keyed ROWS (Thor render speed, render duration, quantization floor), with the forensics moved verbatim into .agents/benchmark-record.md, which is what that page's own checker instructs. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
mudler
force-pushed
the
fix/h3-audio-duration-and-readme
branch
from
August 7, 2026 03:28
ae47f25 to
8f4dd83
Compare
localai-bot
changed the title
mudler
added a commit
that referenced
this pull request
Aug 7, 2026
…user docs PR #68. A real bug found by RENDERING, not by the suite, plus the user-facing documentation the H3 lane never had. THE BUG. A 124-frame render silently muxed as 61 frames. `audio_t` is the PER-CHANNEL latent length (the planner sets it from 40 Hz * duration) and the packed layout carries `audio_t * audio_channel` ROWS, one per (channel, step). The denormalize step divided by the channel count, so the decoded audio ran half the video's duration; because the muxer passes `-shortest`, that silently truncated the VIDEO to half its frames too. The fix is one line, and it now agrees with the same function's own `target_audio_rows = request.audio_t * request.audio_channel` a few lines above. WHY THE SUITE WAS BLIND, which is the part worth keeping. Every existing gate asserts shape SELF-CONSISTENCY, and a uniformly halved pipeline is perfectly self-consistent: the shapes agreed with each other, they were just half as long as the request asked for. ffprobe on the artifact exposed it. The gate added is the one that is NOT self-consistency: the structural t2va e2e case now decodes `audio_t` latent steps INDEPENDENTLY and requires the pipeline to have produced exactly that many samples. RED-first: restoring the divide turns that case RED (1 failed assertion), removing it goes GREEN. Suite 68/68, 46586 assertions. A SECOND BUG FROM THE SAME SESSION, DROPPED RATHER THAN LANDED TWICE. Conditioning PREPENDS rows to img_pos, so the denoise loop returns condition rows followed by the targets, and unpatchify was handed the whole set. This branch fixed it by carrying the layout's own `update_mask` out of the loop; main fixed the SAME bug independently in row/H3-RENDER-CLOSE by taking the TRAILING target_video_rows with a VT_CHECK on the row accounting. On rebase the update_mask mechanism and its test were DROPPED: landing a second solution to a fixed bug is churn, and the positional form already ships gated. The alternative is recorded in .agents/state.md rather than silently lost. DOCS. The video+audio generation path had no user-facing documentation at all, and the multimodal INPUT interface was documented WRONG: `vllm-cli --image/--video/--audio` does not exist, the CLI is text-only. Multimodal input is served over the OpenAI API as `image_url`, `video_url` and `input_audio`/`audio_url` content parts on /v1/chat/completions (src/vllm/entrypoints/openai/chat_mm.cpp), which was documented NOWHERE, so docs/USAGE.md now states it rather than the wrong claim merely being removed. README gains a News section led by video+audio generation, paid for INSIDE the 30,000-char landing-page budget (29,997) rather than by raising it. PUBLIC DOCS. The 19-line narrative this branch originally appended to docs/BENCHMARKS.md was a non-canonical H2 section and pushed the page over its prose budget: converted into three keyed ROWS (Thor render speed, render duration, quantization floor), with the forensics moved verbatim into .agents/benchmark-record.md, which is what that page's own checker instructs. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
mudler
force-pushed
the
fix/h3-audio-duration-and-readme
branch
from
August 7, 2026 03:31
8f4dd83 to
7f44da1
Compare
…user docs PR #68. A real bug found by RENDERING, not by the suite, plus the user-facing documentation the H3 lane never had. THE BUG. A 124-frame render silently muxed as 61 frames. `audio_t` is the PER-CHANNEL latent length (the planner sets it from 40 Hz * duration) and the packed layout carries `audio_t * audio_channel` ROWS, one per (channel, step). The denormalize step divided by the channel count, so the decoded audio ran half the video's duration; because the muxer passes `-shortest`, that silently truncated the VIDEO to half its frames too. The fix is one line, and it now agrees with the same function's own `target_audio_rows = request.audio_t * request.audio_channel` a few lines above. WHY THE SUITE WAS BLIND, which is the part worth keeping. Every existing gate asserts shape SELF-CONSISTENCY, and a uniformly halved pipeline is perfectly self-consistent: the shapes agreed with each other, they were just half as long as the request asked for. ffprobe on the artifact exposed it. The gate added is the one that is NOT self-consistency: the structural t2va e2e case now decodes `audio_t` latent steps INDEPENDENTLY and requires the pipeline to have produced exactly that many samples. RED-first: restoring the divide turns that case RED (1 failed assertion), removing it goes GREEN. Suite 69/69, 47416 assertions. A SECOND BUG FROM THE SAME SESSION, DROPPED RATHER THAN LANDED TWICE. Conditioning PREPENDS rows to img_pos, so the denoise loop returns condition rows followed by the targets, and unpatchify was handed the whole set. This branch fixed it by carrying the layout's own `update_mask` out of the loop; main fixed the SAME bug independently in row/H3-RENDER-CLOSE by taking the TRAILING target_video_rows with a VT_CHECK on the row accounting. On rebase the update_mask mechanism and its test were DROPPED: landing a second solution to a fixed bug is churn, and the positional form already ships gated. The alternative is recorded in .agents/state.md rather than silently lost. DOCS. The video+audio generation path had no user-facing documentation at all, and the multimodal INPUT interface was documented WRONG: `vllm-cli --image/--video/--audio` does not exist, the CLI is text-only. Multimodal input is served over the OpenAI API as `image_url`, `video_url` and `input_audio`/`audio_url` content parts on /v1/chat/completions (src/vllm/entrypoints/openai/chat_mm.cpp), which was documented NOWHERE, so docs/USAGE.md now states it rather than the wrong claim merely being removed. README gains a News section led by video+audio generation, paid for INSIDE the 30,000-char landing-page budget rather than by raising it. PUBLIC DOCS. The 19-line narrative this branch originally appended to docs/BENCHMARKS.md was a non-canonical H2 section and pushed the page over its prose budget: converted into three keyed ROWS (Thor render speed, render duration, quantization floor), with the forensics moved verbatim into .agents/benchmark-record.md, which is what that page's own checker instructs. ALSO REPAIRS RECORD DAMAGE THAT LEFT check-agent-record RED ON MAIN. #93's squash-merge removed 130 lines from specs/minimax-h3.md: its branch predated #71/#92, so landing it DELETED sections 9 and 10 (the /v1/videos rows) and left two orphaned table fragments with no heading and no separator. The checker reports that as "table has 4 pipes; expected 3" plus "no linked spec names exact stable token SERVE-VIDEOS-REFS", the engine-matrix row pointing at a section that no longer existed. Restored sections 9 and 10 VERBATIM from 548b000 and cut the orphans; #93's own content is preserved byte-for-byte. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context)
mudler
force-pushed
the
fix/h3-audio-duration-and-readme
branch
from
August 7, 2026 04:20
7f44da1 to
10b4b62
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
A real bug found by rendering, not by the suite, plus the user-facing documentation the H3 lane never had.
The bug
A 124-frame render silently muxed as 61 frames.
audio_tis the per-channel latent length (the planner sets it from 40 Hz * duration) and the packed layout carriesaudio_t * audio_channelrows, one per (channel, step). The denormalize step divided by the channel count, so the decoded audio ran half the video's duration; because the muxer passes-shortest, that silently truncated the video to half its frames too.The fix is one line, and it now agrees with the same function's own
target_audio_rows = request.audio_t * request.audio_channela few lines above.Why the suite was blind, which is the part worth keeping
Every existing gate asserts shape self-consistency, and a uniformly halved pipeline is perfectly self-consistent: the shapes agreed with each other, they were just half as long as the request asked for.
ffprobeon the artifact exposed it.The gate added is the one that is not self-consistency: the structural t2va e2e case now decodes
audio_tlatent steps independently and requires the pipeline to have produced exactly that many samples. RED-first: restoring the divide turns that case red (1 failed assertion), removing it goes green. Suite 69/69, 47416 assertions.A second bug from the same session, dropped rather than landed twice
Conditioning prepends rows to
img_pos, so the denoise loop returns condition rows followed by the targets, and unpatchify was handed the whole set. This branch fixed it by carrying the layout's ownupdate_maskout of the loop.mainfixed the same bug independently inrow/H3-RENDER-CLOSE, by taking the trailingtarget_video_rowswith aVT_CHECKon the row accounting.On rebase the
update_maskmechanism and its test were dropped: landing a second solution to a fixed bug is churn, and the positional form already ships gated. The alternative is recorded in.agents/state.mdrather than silently lost.Docs
The video+audio generation path had no user-facing documentation at all, and the multimodal input interface was documented wrong:
vllm-cli --image/--video/--audiodoes not exist, the CLI is text-only. Multimodal input is served over the OpenAI API asimage_url,video_urlandinput_audio/audio_urlcontent parts on/v1/chat/completions(src/vllm/entrypoints/openai/chat_mm.cpp), which was documented nowhere, sodocs/USAGE.mdnow states it rather than the wrong claim merely being removed.README gains a News section led by video+audio generation, paid for inside the 30,000-char landing-page budget (29,987) rather than by raising it.
Public docs
The 19-line narrative this branch originally appended to
docs/BENCHMARKS.mdwas a non-canonical H2 section and pushed the page over its prose budget. Converted into three keyed rows (Thor render speed, render duration, quantization floor), with the forensics moved verbatim into.agents/benchmark-record.md, which is what that page's own checker instructs.FOLLOWING_AGENTS_PROTOCOL