From 10b4b620f94398d0fc206fa1a83e1de3d701a390 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Fri, 7 Aug 2026 00:00:05 +0000 Subject: [PATCH] fix(minimax-h3): half of every video was discarded, plus the missing 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 548b0000 and cut the orphans; #93's own content is preserved byte-for-byte. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context) --- .agents/NOW.md | 2 +- .agents/benchmark-record.md | 45 ++++++ .agents/specs/minimax-h3.md | 128 +++++++++++++++++- .agents/state.md | 72 ++++++++++ README.md | 35 ++++- docs/BENCHMARKS.md | 3 + docs/FEATURES.md | 2 +- docs/STATUS.md | 2 +- docs/USAGE.md | 68 ++++++++++ .../models/minimax_h3_pipeline.cpp | 7 +- tests/vllm/models/test_minimax_h3.cpp | 60 ++++++++ 11 files changed, 415 insertions(+), 9 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index f4d7a462..c687296d 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -17,7 +17,7 @@ Working head: `row/backend-rocm-w0` (#41). Prior: benchmark checkpoint | Laguna NVFP4 / DeepSeek-V4 decode | **Both CLOSED, byte-exact, default-ON**: 1.03x vLLM, 1.144x ds4 | Laguna vLLM K-run when convenient | | f32-out GEMV audit | Only laguna + ds4 bf16 tower affected; gate models unaffected | Re-verify ds4 tower same-tool | | Invocation-parity prevention | CI guard + checklist landing | Merge; build-verify `kGemvHeuristicAlgos` on dgx | -| MiniMax-H3 lane | **fl2va COHERENT; ref2va assembly bug FIXED+gated** (`H3-REF2VA-ASSEMBLY`) | grid RE-ATTRIBUTED to the ref2va NVFP4 checkpoint/loader, NOT assembly/fp4 (§8.10) | +| MiniMax-H3 lane | **fl2va COHERENT; ref2va assembly bug FIXED+gated** (`H3-REF2VA-ASSEMBLY`) | +audio-duration fix (#68); grid RE-ATTRIBUTED to the ref2va NVFP4 ckpt/loader (§8.10) | | Kimi-Linear-48B (KDA+NoPE-MLA+MoE) | **e2e RUNS** (bf16-resident §13): 13/13·656. Token gate **NEAR-TIE 106/128** | device GDN/MLA islands; 1.59 tok/s; default OFF | | 35B fresh grid | **BOUND** @`1ea26427`: 0.93-1.03x, c16 0.93x. INTAKE + Option A both NEGATIVE | Lever left: prefill glue (#61) | | Qwen3.5-4B revalidation | 0.9971x @`59674cf1` (#35); TTFT/PSS pass, TPOT/ITL open | `docs/bench-evidence/` | diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 93a60b6a..b21472f5 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -14400,3 +14400,48 @@ so both misattributed a checkpoint/loader defect. True residual = the NVFP4 DiT Synthetic-NVFP4 gates proved the dequant MATH byte-exact but never loaded THIS file vs a coherent oracle. Next: a REF2VA GGUF (bf16, known-good loader) as checkpoint oracle — dgx-disk-blocked (23 GiB free). Artifacts `~/h3fp4/out_{vs_ref2va,bf16_ref2va,t2va_nvfp4,vs_fl2va}.mp4`. + +## MiniMax-H3 — the Thor render-speed leg, the audio-duration bug, and the Q3_K_M quantization floor (2026-08-06) + +Moved out of `docs/BENCHMARKS.md` on landing: the scoreboard is a keyed table, so +the measured results live there as three rows and the forensics live here. + +**Thor render speed (sm_110, no FlashAttention-2).** One Jetson Thor, 864x480 / +124 frames / 50 steps, Q4_K_M DiT: + +| stage | measured | +|---|---| +| DiT forward | **34.6 s/step** (was 574.5 s before the attention work: **16.6x**) | +| attention share of a step | 96% before, and the whole of that 16.6x | +| full 50-step render | **~28 min** (was ~8 h) | + +The 16.6x is three landed changes and two MEASURED NEGATIVES: warp-per-query, +then a chunked warp reduce-scatter (1.76x), then bf16 tensor cores via +`mma.sync` (9.82x). Shared-memory K/V tiling (**23% SLOWER**) and register +Q-blocking (**-0.8%**) were both measured and REVERTED. Both chased memory +traffic, which is not the bound: one head's K+V is 3.9 MB against 32 MB of L2. +Do not re-run either without a new reason. + +**Quantization floor: use Q4_K_M, not Q3_K_M.** H3's split-half RoPE produces +channel-wise magnitude outliers that 3 bits cannot hold. A controlled A/B (same +prompt, same seed, same code, only the DiT encoding changed) turned a murky +lattice-covered silhouette into a photoreal close-up. + +**The audio-duration 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. + +Worth recording is WHY the suite was blind to it. 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 invariant that is NOT self-consistency: latent steps / 40 Hz must equal +the video duration, to within one latent step. + +(The sibling reference-leak bug found in the same render session, where +conditioning rows reached unpatchify, was independently fixed on `main` by +`row/H3-RENDER-CLOSE` taking the TRAILING target rows; this row's alternative +`update_mask` mechanism was dropped rather than landed twice.) diff --git a/.agents/specs/minimax-h3.md b/.agents/specs/minimax-h3.md index 4604770e..0bbe9df5 100644 --- a/.agents/specs/minimax-h3.md +++ b/.agents/specs/minimax-h3.md @@ -775,6 +775,78 @@ dump the ref2va target-row VAE-input latent adjacency-cosine (like #77 did for t 0.95) to confirm the target rows are white, and A/B the reference-row condition-noise vs a clean anchor. +## 9. W-OAI — the `/v1/videos` OpenAI (Sora) WIRE SHAPE, 2026-08-06 + +Row `SERVE-VIDEOS-OAI` (engine matrix, Serving surface), claim +`CLAIM-SERVE-VIDEOS-OAI`, branch `row/SERVE-VIDEOS-OAI`. + +Developer-directed: an unmodified OpenAI client must work against `/v1/videos`. +ADDITIVE — the vLLM-Omni-derived fields keep working, and every body that parsed +before means exactly what it meant before. + +SPLIT, deliberately: this row is the REQUEST/RESPONSE SHAPE only (`model`, +`size`, `seconds`, and the MP4 download route). It touches no generation code and +loads no VAE. The REFERENCE CONDITIONING half (`input_reference` -> fl2va, plus +the two `metadata` reference modalities -> ref2va) is §10, row +`SERVE-VIDEOS-REFS`, because it is a separate capability that pulls in the VAE +encoder halves and the runner. Each half is independently reviewable and gated. + +### 9.0 Spike contract (`SERVE-VIDEOS-OAI`) + +| Section | Content | +|---|---| +| Scope | IN: the OpenAI (Sora) REQUEST SPELLINGS `model`, `size`, `seconds` on `/v1/videos`, their precedence against the native fields, the `model`-mismatch warning on the job, and `GET /v1/videos/{id}/content`. OUT: reference conditioning of any modality (§10); OpenAI's status vocabulary / id shape / `progress` / multipart upload; any change to generation, the DiT, the VAEs or the muxer. | +| Upstream chain | OpenAI's published video API (`POST /v1/videos`, `GET /v1/videos/{video_id}/content`, `size` "WxH", `seconds` string enum) is the request CONTRACT; vLLM-Omni's `/v1/videos` async+sync job pair is the endpoint shape we already mirror. | +| Our baseline | `ParseVideoRequest` took only the native spellings (`duration`, `height`/`width`, `num_frames`, `num_inference_steps`, `flow_shift`, `audio_flow_shift`, `seed`, plus `extra_params`); an OpenAI client's body parsed to DEFAULT geometry and duration. `VideoJobStore` had no `model`/`warning`. The routes stopped at status: the produced .mp4 was reachable only through the filesystem. | +| Port map | Request contract -> `include/vllm/entrypoints/openai/video_api.h` (`VideoRequest::model` + `ParseVideoSize`) and `src/vllm/entrypoints/openai/video_api.cpp` (`ParseVideoRequest`, `ReadDuration`, `ParseWholeNumber`). Job record -> `VideoJobStore::Create(model, warning)` + `VideoJobStatusJson`. Download route -> `ApiServer::handle_video_content` + `video_model_warning` + their registration in `src/vllm/entrypoints/openai/api_server.cpp`. | +| Tests to port | No upstream test module exists for this surface (OpenAI publishes an API, not tests; vLLM-Omni's video endpoint has no ported test). The contract is gated in-tree instead, extending the existing files: `tests/vllm/entrypoints/openai/test_video_api.cpp` (parsing, precedence, the job record) and `tests/vllm/entrypoints/openai/test_api_server.cpp` (routes, content behaviour, additivity over a real socket). Every assertion uses values that DIFFER from the field default. | +| Gates | CPU, foreground: `test_video_api` 11/11 (125 assertions), `test_openai_api_server` 40/40 (509), `server` builds clean. Content route: 404 unknown / 409 unfinished (no bytes leaked) / 500 failed / 500 vanished / 200 byte-exact `video/mp4`. Additivity: with no `VideoRunner`, `POST /v1/videos` is 404 over a real socket with no `ErrorResponse` envelope; with one, it is 200 and the unknown-id 404 IS ours. Commands: `cmake --build build --target test_video_api test_openai_api_server server -j12`. Real-weights e2e rides §8's GB10/disk window. | +| Dependencies | Row IDs: the MiniMax-H3 model rows and `row/H3-FP4-SPEED` (UNTOUCHED - no generation code changed); `SERVE-VIDEOS-REFS` (§10) stacks on this row. No new download, no GPU, no toolchain change for the CPU gate. | +| Work breakdown | (1) alias parsing + precedence + `ParseVideoSize`; (2) `model` recording + the job `warning`; (3) `handle_video_content` + its route; (4) both test files; (5) docs + record. | +| Risks/decisions | NATIVE-wins precedence: the only direction that leaves every previously-parsing body meaning what it meant. `model` mismatch WARNS rather than 404s: a Sora client cannot know the local model name, so a rejection would defeat the compatibility; silence would hide it. A 409 (never bytes) on an unfinished job: a partially muxed file would reach the client as a valid-looking, truncated MP4. No vLLM-defined behaviour is reopened. | + +### 9.1 The aliases and their precedence + +| OpenAI | Lands on | Notes | +|---|---|---| +| `model` | `VideoRequest::model` | Recorded + echoed; an unserved name is a job `warning`, never a rejection (a Sora client cannot know the local model's name) | +| `size` | `width`, `height` | `"x"`, whole positive pixels, one `x`/`X` | +| `seconds` | `duration_seconds` | Number OR numeric string — OpenAI types it as a string enum ("4"/"8"/"12") | + +PRECEDENCE: the NATIVE field WINS (`width`/`height` over `size`, `duration` over +`seconds`). Both spellings are VALIDATED whichever wins, so a malformed `size` is +a 400 even when explicit `width`/`height` override it. Precedence is PER-AXIS: an +explicit `width` alone still lets `size` supply the height it did not specify. + +### 9.2 `GET /v1/videos/{id}/content` + +Returns the finished MP4 as `video/mp4`. Without it a caller could start and poll +a job but never fetch the result over HTTP. Unknown id -> 404; queued/running -> +409 naming the status (never a truncated file); failed -> 500 carrying the +failure; a vanished output -> 500, not a 200 with zero bytes. + +### 9.3 Status + +- **CPU-LANDED + gated.** `test_video_api` 11/11 (125 assertions), + `test_openai_api_server` 40/40 (509), `server` builds clean. Additivity is + gated over a REAL socket: without a `VideoRunner` all four routes are absent + (a 404 with no `ErrorResponse` envelope), with one they serve. +- **Residuals, named.** OpenAI's status vocabulary is not mirrored (ours stays + queued/running/succeeded/failed, ids `vid_N`, no `object`/`progress`/ + `created_at`); reference conditioning is §10 (`SERVE-VIDEOS-REFS`), not this row. +- **Real-weights leg** rides the same GB10/disk window as §8. + +## 10. W-REFS — reference conditioning over `/v1/videos`, 2026-08-06 + +Row `SERVE-VIDEOS-REFS` (engine matrix, Serving surface), claim +`CLAIM-SERVE-VIDEOS-REFS`, branch `row/SERVE-VIDEOS-REFS`, stacked on §9. + +§9 made an OpenAI client's request PARSE. This row makes its REFERENCES do +something: an image the video starts from, a clip it continues, a voice it +carries. Before it, no reference modality was reachable over HTTP at all. + +### 10.0 Spike contract (`SERVE-VIDEOS-REFS`) + | Section | Content | |---|---| | Scope | IN: OpenAI's `input_reference` mapped to fl2va first-frame conditioning; the two reference modalities OpenAI has no slot for carried in `metadata` (`input_reference_video`, `input_reference_audio`) mapped to ref2va blocks; the fl2va/ref2va combination rule enforced at the request boundary; the `examples/server` runner wiring (PPM decode, frame-directory clip, WAV, lazily-loaded VAE encoder halves). OUT: the ref2va IMAGE modality (reachable via the native `task` + the CLI, deliberately not bound to `input_reference`); any change to generation, the DiT, the VAEs or the muxer; OpenAI's multipart upload. | @@ -786,7 +858,55 @@ anchor. | Dependencies | Row `SERVE-VIDEOS-OAI` (§9), stacked. Code: `MiniMaxH3Encode{KeyframeCondRows,ReferenceVideo,ReferenceAudio}`, `MiniMaxH3ReadWav`, `DecodeDataUri`. Runtime: `--video-vae` for an image or video reference, `--audio-vae` for an audio reference (both encoder halves, loaded lazily and once). No new download, no GPU. | | Work breakdown | (1) `input_reference` parsing (path or `data:` URL) -> fl2va, with the geometry refusal; (2) the `metadata` map + the video/audio reference keys; (3) the combination rule in the parser; (4) the `examples/server` runner branches; (5) both test files; (6) docs + record. | | Risks/decisions | `input_reference` -> fl2va, NOT ref2va: OpenAI documents it as the frame the video starts from; ref2va would silently change what the API promises. The two extra modalities go in `metadata` rather than new top-level fields, so a strict client's schema validation still passes. Combination legality is enforced in the PARSER, not left to the pipeline, so a supplied reference is never silently dropped. | -| OpenAI | Lands on | Notes | -| `model` | `VideoRequest::model` | Recorded + echoed; an unserved name is a job `warning`, never a rejection (a Sora client cannot know the local model's name) | -| `size` | `width`, `height` | `"x"`, whole positive pixels, one `x`/`X` | -| `seconds` | `duration_seconds` | Number OR numeric string — OpenAI types it as a string enum ("4"/"8"/"12") | + +### 10.1 `input_reference` is fl2va, not ref2va + +OpenAI documents it as the image the video STARTS FROM, which is what +`MiniMaxH3EncodeKeyframeCondRows` expresses (frame 0 of the output pinned to the +image, `imgvid_noise_aug = 1.0`). `MiniMaxH3EncodeReferenceImages` prepends whole +reference images as their own blocks — guidance that never becomes a frame — so +mapping it there would have changed what the API promises. With a reference image +and no explicit `task`, the task IS `fl2va`, and the image's aspect drives the +default resolution through `MiniMaxH3ResolveShape`. + +Two limits, both refused up front rather than deep in the denoise: the image must +be a binary PPM (P6), because no PNG/JPEG codec is vendored (the same NAMED +residual the chat multimodal path carries), and it must already be at the +resolved output geometry, because no image resampler is vendored. The refusal +names both geometries. + +### 10.2 The two reference modalities OpenAI has no slot for + +H3 has three (image, silent video, audio); the Sora schema carries one. The other +two enter through `metadata`, the standard OpenAI free-form string map that +strict clients tolerate, rather than invented top-level fields that would fail a +client's schema validation. The whole map is kept verbatim. + +- `metadata.input_reference_video` — a DIRECTORY of `frame_%06d.ppm`, the exact + layout `minimax-h3-gen` and the server WRITE, so clips chain. No demuxer is + vendored, hence a frame directory rather than a container; a `data:` URL cannot + name a directory and is refused by name. + `MiniMaxH3EncodeReferenceVideo` emits `ref_audio_t == 0`: the clip is SILENT. +- `metadata.input_reference_audio` — a 16-bit PCM WAV path or `data:` URL, read + by the existing `MiniMaxH3ReadWav` and encoded by + `MiniMaxH3EncodeReferenceAudio`. Supplied with a video reference it ATTACHES to + that block (one `kVideoAudio` block carrying both, the layout + `packed_sequence.py` builds); alone it is its own block. + +LEGALITY is the pipeline's own rule (`minimax_h3_pipeline.cpp:251`: fl2va +keyframes and ref2va blocks are exclusive), enforced in the PARSER so it is a 400 +naming the pair rather than a failed job — and never a silently dropped +reference, which is the failure that looks like it worked. Legal: none / image / +video / audio / video+audio. Illegal: `input_reference` with either metadata +reference. + +### 10.3 Status + +- **CPU-LANDED + gated.** `test_video_api` 14/14 (167 assertions), + `test_openai_api_server` 41/41 (525), `server` builds clean. Each modality is + gated as ARRIVING at the runner, and an illegal pair is a 400 that generates + nothing (`calls == 0`). +- **Residuals, named.** Reference images are binary PPM at the output resolution; + a video reference is a frame DIRECTORY; OpenAI's real `input_reference` upload + is multipart, ours is the JSON spelling. +- **Real-weights leg** rides the same GB10/disk window as §8. diff --git a/.agents/state.md b/.agents/state.md index 3458ef1f..c5bb21c4 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -40623,3 +40623,75 @@ a coherent oracle. NEXT: a REF2VA GGUF (bf16, known-good loader) as checkpoint o disk (23 GiB free, 100% full; large dirs belong to other campaigns, not prunable). Fix + gates LAND regardless; box left clean (renders exited, gpu.lock released, worker stays parked, ckpts kept). Records: spec §8.10 + §8.2 row, STATUS/BENCHMARKS/FEATURES H3 rows, benchmark-record, NOW. + +## 2026-08-07T05:35 - MiniMax-H3: half of every video was being discarded (audio duration), plus the user-facing docs the lane never had + + + +Branch `fix/h3-audio-duration-and-readme` (PR #68). Found by RENDERING, not by +the suite. + +**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 +(`audio_steps = request.audio_t`), and it now agrees with the same function's own +`target_audio_rows = request.audio_t * request.audio_channel` two lines up. + +**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 invariant that is NOT self-consistency: latent steps / 40 Hz must equal the +video duration, to within one latent step, and the halved value is asserted to be +off by more than a second. + +**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. Recorded +here so the alternative is not 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. + +**STATUS, paid for in place.** The page sits within 20 chars of its 283470 +ratchet, so the note was made to FIT rather than granted headroom: the H3 row's +"(t2va-no-refs grids too)" forensic aside was dropped (the detail is spec §8.10, +where #93 recorded it) and the duration result stated as "audio full-duration". +Net -4 chars; the ratchet is untouched because it did not need to move. + +**THIRD-PARTY RECORD DAMAGE REPAIRED (out of row, carried here because it is the +same file and it 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 `## 9` and `## 10` (the `/v1/videos` rows) and left two +orphaned table fragments with no heading and no separator, which the record +checker reports as `table has 4 pipes; expected 3` plus `no linked spec names +exact stable token SERVE-VIDEOS-REFS` (the engine-matrix row pointed at a section +that no longer existed). Repaired by restoring `## 9` and `## 10` VERBATIM from +`548b0000` and cutting the orphans; `#93`'s own content is preserved byte-for-byte. +This is the hazard AGENTS.md names: a keyed/structured record file merged from a +stale base silently loses another row's section. + +**GATE (CPU, foreground).** `test_minimax_h3` full suite, plus the new duration +case. No numbers changed by the doc work; the Thor speed figures are prior +measurements now recorded in the scoreboard rather than new ones. diff --git a/README.md b/README.md index 436247c4..0ff65f36 100644 --- a/README.md +++ b/README.md @@ -33,6 +33,14 @@ > appending fields whose zero value keeps existing behavior byte-identical, and only bumps on an > incompatible change. If you embed us, embed through that header. +## News + +- **2026-08** **MiniMax-H3 video + audio generation lands.** A 33 B joint video+audio diffusion + transformer, prompt to MP4 with a stereo track, on one Jetson Thor. Attention got **16.6x** + faster on the way (bf16 tensor cores): a 50-step 864x480 render went ~8 h to ~28 min. +- **2026-08** Audio input (Whisper, Voxtral) and image+video input (Qwen3-VL, Qwen3.6-27B) + correctness-complete, single-sequence path. + vllm.cpp is a from-scratch C++20 inference engine chasing three things at once: be the **smallest** thing you can deploy, be the **fastest** on the hardware you already own, and still carry **every feature people actually want**. No Python and no PyTorch at inference time. @@ -270,6 +278,15 @@ and Voxtral (audio). | Qwen3-VL (image + video) | Qwen3-VL-4B-Instruct | - | Strict token-exact 32/32 (image) | Speed-pending | | Qwen3.6-27B vision (image + video) | Qwen3.6-27B | - | Strict token-exact 32/32 | Speed-pending | | Voxtral (audio) | Voxtral-Mini-3B-2507 | - | Near-tie-robust (decoder 48/48 exact) | Speed-pending | +| **MiniMax-H3 (video + audio GENERATION)** | MiniMaxAI/MiniMax-H3 | Q4_K_M / NVFP4 | Renders 864x480 / 124f with audio | **34.6 s/step, one Jetson Thor** | + +**Video + audio GENERATION is supported**, not just video *input*. MiniMax-H3 renders end to +end: prompt -> Qwen3-VL-32B encoder -> DiT denoise -> ViT3D video VAE + DAC/BigVGAN audio VAE +-> MP4 with a stereo track. The project's first DIFFUSION architecture (no KV cache, no +sampler, no logits); upstream is `vllm-project/vllm-omni`, beyond the parity pin. Five +conditioning modes, each gated on the conditioning CHANGING the output rather than merely +being accepted, and `POST /v1/videos`. Use **Q4_K_M**: 3 bits cannot hold the channel-wise +outliers H3's split-half RoPE produces. Detail: [docs/STATUS.md](docs/STATUS.md). Compressed-tensors NVFP4A16 (W4A16) dense weights also load and compute natively (RedHatAI/Qwen3-32B-NVFP4A16). Long-context RoPE (YaRN, Llama-3, LongRoPE, dynamic-NTK) and @@ -286,7 +303,7 @@ hardware-blocked and why, is in [docs/STATUS.md](docs/STATUS.md). | Backend | Hardware | State | |---|---|---| | **CUDA** | GB10 / DGX Spark (sm_121a) | Runtime-gated. 27B at/above vLLM throughput, 35B prefill-pending | -| **CUDA** | Blackwell, Hopper, Ampere, Ada (sm_80 through sm_121a) | Build-supported, compiles to real machine code, fast GDN path build-verified per-arch. Not runtime-proven here (no such boards) | +| **CUDA** | Blackwell, Hopper, Ampere, Ada (sm_80 through sm_121a) | Build-supported, fast GDN path build-verified per-arch. Not runtime-proven here (no such boards) | | **CPU** | x86-64, arm64 | Correctness / CI reference. At or ahead of llama.cpp on every GGUF axis, Arm i8mm quant-GEMM tier | | **Metal** | Apple Silicon | Two models end to end, 18 of 75 ops native. Prefill ahead of MLX-LM, warm total 97.6% with the MLX provider | | **Vulkan** | Portable GPU | Skeleton: 8 ops plus the fusion catalogue cross-check against CPU and CUDA. No model runs yet | @@ -325,6 +342,22 @@ build/examples/vllm-cli --model /path/to/Qwen3.6-27B --prompt "The capital of Fr `vllm-bench` (throughput/latency harness) and `tokenize` (tokenizer smoke tool) ship alongside it. All flags, including `--speculative-config`: [docs/USAGE.md](docs/USAGE.md). +### Multimodal INPUT and video GENERATION + +Multimodal INPUT goes through `/v1/chat/completions` content parts (`image_url`, +`video_url`, `input_audio`). Video GENERATION: + +```sh +build/examples/minimax-h3-gen --dit MiniMax-H3-FL2VA-Q4_K_M.gguf --dequant-bf16 \ + --encoder qwen3vl-32B-MiniMax-H3-Q4_K_M.gguf --tokenizer tokenizer.json \ + --video-vae video_vae.safetensors --audio-vae audio_vae.safetensors \ + --prompt "A golden retriever runs across a sunlit beach" \ + --frames 124 --height 480 --width 864 --steps 50 --device cuda --out out.mp4 +``` + +Conditioning flags, the PPM convention that chains clips, and serving: +[docs/USAGE.md](docs/USAGE.md). + ## OpenAI-compatible server ```sh diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 12ebaa2a..c3cdebf6 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -314,6 +314,9 @@ built on it rather than keeping the flattering one. | MiniMax-H3 FP4 speed (W-FP4a) | **Measured GB10 (`row/H3-FP4-GPU-E2E`).** Marlin W4A16 byte-exact vs bf16; fp4 a memory win, 0.8x bf16/forward. Real-ckpt fp4-resident e2e RUNS (mp4/wav) | fp4 speed CLOSED. Detail: benchmark-record + spec §8 | | MiniMax-H3 render coherence (`row/H3-RENDER-CLOSE` #77) | **CLOSED: a COHERENT scene on GB10.** #70/#74 white was wrong-PARTITION usage (t2va on the ref2va ckpt); t2va on the FL2VA GGUF renders a prompt-matched orange cat (adj-cos 0.95 vs 0.06, no patch-grid) | Verified first: t2va inputs byte-exact vs upstream; CUDA device==host at seq 1920. Follow-up `H3-TASK-PARTITION-GUARD`: the task/partition mismatch now RAISES 1:1 with `_resolve_task` (spec §8.6-8.7) | | MiniMax-H3 image conditioning (`row/H3-CONDITIONED-E2E`, `row/H3-VISION-SCATTER`, `row/H3-REF2VA-ASSEMBLY`) | **fl2va COHERENT; ref2va assembly bug FIXED+gated.** vision→cond scatter gated; ref2va block-dim double-division fixed + RED-first gated (128 vs 512) + a permanent ref2va DiT-forward rung (§8.10) | grid RE-ATTRIBUTED: with the fix ref2va grids in fp4 AND bf16, and t2va with no refs on the ref2va NVFP4 also grids while FL2VA-GGUF renders, so it is the **NVFP4 checkpoint/loader**, NOT assembly/fp4 (§8.10) | +| MiniMax-H3 Thor render speed (sm_110, no FA2) | **34.6 s/step** at 864x480/124f/50 steps on Q4_K_M, **16.6x** off 574.5 (render ~28 min, was ~8 h). Landed: warp-per-query, chunked warp reduce-scatter (1.76x), bf16 `mma.sync` (9.82x) | Shared-memory K/V tiling (23% SLOWER) and register Q-blocking (-0.8%) both measured and REVERTED: memory traffic is not the bound (one head's K+V is 3.9 MB against 32 MB of L2) | +| MiniMax-H3 render duration (audio halving) | **CLOSED**: a 124-frame render silently muxed as **61**. The decoded audio ran half the video's duration and the muxer passes `-shortest`. Every structural check passed: shapes were self-consistent, just halved | Gated on the duration invariant (latent steps / 40 Hz equals the video duration), not on shape self-consistency, which a halved pipeline satisfies | +| MiniMax-H3 quantization floor | **Use Q4_K_M, not Q3_K_M.** H3's split-half RoPE produces channel-wise magnitude outliers 3-bit cannot hold; a controlled A/B (same prompt, seed, code) turned a murky lattice-covered silhouette into a photoreal close-up | Per-tensor mixed precision, if a smaller footprint is ever owed | | MXFP4 Qwen3-8B (W4A16 Marlin) | **`KERNEL-MARLIN-DENSE-EXEC` x3 (dense-ON default): c1 1.020, c2/c4/c8 0.962/0.966/0.969, GPU mem 2.63x less** (beats #51 1.005/0.925/0.939/0.953 EVERY axis); #44 3/3, 32B-NVFP4A16 6/6; -Werror test-guard fixes x2 | **VT_MARLIN_DENSE default-ON** (+951us). `FLASH-PTXAS` #82: cuModule A/B ties our+vLLM PTX across ptxas 13.0/13.2/driver-JIT (~144us); +10us is engine CONTEXT not codegen, no ptxas lever/flip (retires #75) | | Vulkan vs llama.cpp Vulkan (`BENCH-VK-LLAMA`) | **NOT APPLICABLE: no number measured, claimed or owed.** Vulkan registers 8 of the CPU backend's 83 ops and runs no model. The 2026-08-06 campaign spec is PLAN ONLY. [Detail](../.agents/specs/vulkan-full-support.md) | `VK-E`: build llama.cpp `-DGGML_VULKAN=ON` at pin `237ad9b96` on dgx, `llama-bench` on the same GGUF, record three columns (ours-Vulkan, llama.cpp-Vulkan, ours-CUDA). Blocked by `VK-B` and by `glslc` on the gate box | | Memory footprint vs declared workload (`ROAD-V1-MEM`, #83) | **Never measured, and not measurable today**: there is no auto-sizing to compare against, because the KV pool is a hand-typed `--num-blocks`, so "what the run actually needed" has no number | Once M1's `MemoryBudget` lands: predicted-vs-actual bytes per allocation class, then peak footprint ours-auto vs vLLM at its 0.9 default on the same model and config | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 7c5a24b3..fc9ef5b2 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -161,7 +161,7 @@ model architecture is wired. | Image | ✅ correctness-gated | ✅ | ✅ | ◐ | | Video | ✅ correctness-gated | ✅ | ✅ | ☐ | | Audio | ✅ correctness-gated | ✅ | ◐ | ◐ | -| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va+fl2va COHERENT on GB10 (fl2va via keyframe+encoder path); vision→cond scatter gated; ref2va assembly FIXED+gated but grids, RE-ATTRIBUTED to the NVFP4 checkpoint/loader, NOT assembly (§8.10) | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | +| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va+fl2va COHERENT on GB10 at FULL requested duration (fl2va via keyframe+encoder); vision→cond scatter gated; ref2va assembly FIXED+gated but grids, RE-ATTRIBUTED to the NVFP4 checkpoint/loader, NOT assembly (§8.10) | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | | Multimodal over the OpenAI server | ☐ | ✅ | ✅ | ◐ | Image, video and audio are correct through the CLI and library. Serving them diff --git a/docs/STATUS.md b/docs/STATUS.md index 06f5cb81..0cddb9c5 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -85,7 +85,7 @@ token-for-token correctness against the pinned oracle. | OLMo-3 dense (dual rope, interleaved sliding window) | Implemented, oracle-blocked | Loads + runs in our engine (dual rope: plain sliding + YaRN full-attn, per-layer sliding window); no SACRED gate: vLLM 0.25.0 oracle cannot run OLMo-3-1025-7B (`KeyError: 'rope_theta'`; transformers 5.13.1 nests `rope_parameters` per layer-type, no flat `rope_theta`; run-verified W0 2026-07-26) | | Laguna-S-2.1 MoE (`LagunaForCausalLM`, 118B/8B) | **BINDING 2026-08-04: 87% of vLLM (37.55 vs 43.10, SAME-TOOL nsys both engines); the whole +3.1 ms/step is the bf16 M=1 GEMV bucket (2/3 o_proj, ~196-204 vs 139 us/call, identical `gemvx` kernel); attention/MoE/glue tied or ours-ahead. Invocation match (bf16-out `cublasGemmEx`) A/B'd = WASH, ruled out; ROOT CAUSE FOUND 2026-08-04 (`VT_LAGUNA_RESIDENT_BF16W`): the bf16 projections read UNIFIED/ATS host memory, not `cudaMalloc`'d device memory — staging them device-resident (byte-exact ids) gives 38.8→44.6 tok/s (o_proj 194→131, lm_head 2410→1620 us/call), parity+ vs vLLM 43.1; **default-ON** (flip smoke-verified: canonical byte-exact ids, 44.6 clean-median). Earlier ceiling/diffuse verdicts below were cross-tool artifacts.** **REAL vLLM BAR ESTABLISHED (2026-07-31, `CLAIM-LAGUNA-VLLM-NVFP4`): FIRST-EVER vLLM Laguna run** — prior numbers (incl. the correctness oracle) were all llama.cpp, never vLLM. vLLM on official `poolside/Laguna-S-2.1-NVFP4` (single GB10, greedy, eager, MARLIN backend forced via `VLLM_TEST_FORCE_FP8_MARLIN=1` because the auto-default `FLASHINFER_CUTLASS` needs an absent `nvcc`): **~18.8 tok/s** (64-tok steady) — a LOWER bound. Our GGUF-Q4_K engine = 7.7 tok/s (vLLM ~2.4×); llama.cpp GGUF = 27.8 (still fastest at batch-1). llama.cpp is now a labeled SECONDARY "beat best-in-class GGUF" note; vLLM-NVFP4 is the headline bar. TRUE apples-to-apple still owes OUR NVFP4 Laguna forward arm (same tensor-core path as 27B/35B) — bring-up W-plan SPEC'D in `.agents/specs/laguna-nvfp4-arm-2026-07-31.md` (~85% reuse of the 35B NVFP4 W4A4 MoE infra + a name-map; bf16 attn/dense + fp4 experts; N1-N5 bricks, DGX-gated). **N1-scaffold LANDED (2026-07-31):** additive `LagunaMoeWeights.experts_{gate,up,down}_fp4` + `shared_{gate,up,down}_fp4` (`Nvfp4Weight`, mirror qwen3_5), dead until the N1 loader; CPU build clean + `test_laguna_scaffold` 8/8·167 unchanged. **N1b loader IMPLEMENTED (2026-07-31, build-verified):** `LoadLagunaForCausalLMWeights` (`laguna_weights.cpp`) replaces the `VT_CHECK(false)` stub — resolver + per-layer `LoadBf16Direct` (attn/dense/norms/embed/lm_head/router/shared-expert) + F32 `e_score_correction_bias` + `LnLoadCtNvfp4Raw` W4A4 experts. Name-map + dtypes VERIFIED against the real `poolside/Laguna-S-2.1-NVFP4` index (router `mlp.gate` BF16, bias F32, experts W4A4, shared-expert BF16). **N1b RUN-VERIFIED (2026-07-31):** loader round-trips a synthetic NVFP4 checkpoint byte-identically (`test_laguna_nvfp4_loader` 2/2·29; full detail in the benchmark record). **N2 FORWARD-BRANCH LANDED + CPU-GATED (2026-07-31):** `LqGemmNvfp4Fp4` (per-expert TRUE-W4A4: `ScaledFp4Quant(input_global_scale_inv)`→`MatmulNvfp4Fp4(alpha)`, unified-memory pattern like `LqGemm`) + `LagunaFfnBlock` branches on `fp4=!experts_gate_fp4.empty()` (routed experts fp4; keep-quant grouped fast-path gated off `!fp4`; bf16 attn/dense/router/shared-expert/lm_head unchanged) + both `LagunaForwardGguf{,Cached}` guards relaxed to `has_gguf_weights||has_nvfp4_weights`. **CORRECTION:** routed experts are W4A4 ⇒ per-expert `MatmulNvfp4Fp4`, NOT the grouped W4A16 `MoeGroupedGemmNvfp4` (grouped W4A4 deferred to N5 speed). `test_laguna_nvfp4_loader` 3/3·61 (added a forward run-gate: fp4 MoE branch runs through the real `LagunaForwardGguf` → finite+deterministic logits + routed-experts-consumed); `test_laguna_scaffold` 8/8 unchanged (GGUF byte-identical). **N3 DRIVER LANDED + CPU-SMOKE-VERIFIED (2026-07-31):** `examples/laguna_gen` auto-detects a safetensors DIRECTORY (→ NVFP4: `LoadHfConfig(config.json)` + `LoadLagunaForCausalLMWeights` + `LagunaForwardGguf{,Cached}`) vs a `.gguf` FILE (→ keep-quant), sharing the greedy loop; `--token-ids` bypass the tokenizer for the id-vs-golden gate. Verified on a synthetic NVFP4 dir with a REAL config.json (exercises the `LoadHfConfig`→`ParseLagunaParams` seam the loader test bypassed) → `has_nvfp4=1`, KV-cache decode runs finite. **N4 RAN on GB10 (2026-08-01) — the arm works end-to-end; correctness coherent+near-tie, speed 120× off.** git-archived `84fab587` → clean CUDA build (`121a`) → `laguna-gen --gpu` on the real 67 GiB `ckpt` with vLLM's exact prompt ids injected (`2,785,9626,377,15360,395`, captured via the HF tokenizer). Two GB10 memory fixes landed to run: release the mmap'd shards after the loader's memcpy-copy (114→67 GiB RSS), and create the CUDA context BEFORE the load (the 67 GiB reclaimable page cache otherwise starves `cudaStreamCreate`). **Correctness:** ours `22345 83 350 71070 395 340 9626 372 1703 …` vs golden `22345 83 290 350 674 330 5541 966 340 9626 377 15360 …` — **first 2 tokens match vLLM exactly**, then near-tie divergence; coherent ("France is" = 9626/377/15360; shares golden vocab). EXPECTED: our TRUE-W4A4 (fp4 activations) vs the MARLIN golden's W4A16 (bf16 activations) — different precision, not a bug. **Speed: 6.34 s/tok (0.16 tok/s), prefill 17.3s — ~120× slower than vLLM 18.8.** ROOT CAUSE (source-confirmed): `LqGemmNvfp4Fp4` uses the generic `vt::MatmulNvfp4Fp4` = the hand-written EMULATION CUDA kernel, NOT the cutlass sm120a fp4 tensor-core path the 27B/35B W4A4 use (`MatmulNvfp4Fp4DirectD`); + per-expert loop + per-GEMM host sync + no device residency. **nsys (2026-08-01) trace-confirmed + refined:** only 2 GPU kernels — `MatmulNvfp4Fp4Naive` = 99.3% of GPU time + fp4-quant 0.7%; GPU busy only ~18% of wall. NO bf16 GEMM on the GPU ⇒ `LqGemm`'s bf16 branch runs the host `MatmulNK` reference on the CUDA queue (attention/dense/router/shared/lm_head are CPU-bound, ~4.8 s/tok) — a second lever the source scan missed. **N5 LEVER #2 LANDED (2026-08-01) — 16× decode.** Routed the bf16 tower (attention/dense/router/shared/lm_head) off the host `MatmulNK` onto the GPU (`LqGemm` bf16 branch: `vt::CastBf16` the small activation + `vt::MatmulBT` bf16×bf16→f32, weight stays bf16 — no per-token `ReadF32` of `lm_head [100352,H]`): **decode 6.34 → 0.39 s/tok (16.3×; 0.16 → 2.56 tok/s), prefill 17.3 → 2.24s**; coherence preserved (near-tie). CPU path unchanged (run-gate byte-identical). **N5 LEVER #1 LANDED (2026-08-01) — native fp4 tensor-core, another ~2×.** The engine's native sm120a fp4 tensor-core MMA (`MatmulNvfp4Fp4Native`, `mma.sync kind::mxf4nvf4`) reads the same linear scale layout `LqGemmNvfp4Fp4` produces — it was gated OFF behind `VT_NVFP4_FP4_NATIVE`; the Laguna driver now defaults it ON (scoped; 27B/35B untouched). **decode 0.39 → ~0.20-0.24 s/tok (~2×; ~4.2-5.0 tok/s)**; coherent (byte-identical ids to the emulation path — numerically equivalent), first token matches the golden. **Cumulative N5: 0.16 → ~4.5 tok/s (~28×), now ~4× from vLLM 18.8.** **Device-resident MoE block LANDED + MEASURED (2026-08-01, `LagunaMoeResidentFp4`, `VT_LAGUNA_RESIDENT_MOE` default-ON):** the whole token's routed experts as ONE async device chain (fp4-quant→GEMM gate/up, `MoeSiluMul`, →down stacked, ONE `MoeCombine`), draining once vs ~Pk×3 syncs. **Speed EAGER-NEUTRAL (0.20 s/tok)** — empirically confirms the ds4 precedent (per-op syncs overlap GPU compute; wall is GPU-serial-bound; the graph is the payoff). **CORRECTNESS WIN: golden-token match 2 → 13** (the device `MoeSiluMul`/`MoeCombine` mirror vLLM's fused MoE faithfully). Lands default-ON (better correctness, no speed cost, graph prerequisite). **CORRECTED CEILING (from the measured state): a perfect decode graph caps at ~5.9 tok/s** (GPU already ~87% busy at 0.20 s/tok), still 3.3× short of vLLM 18.8 — the graph is necessary but NOT sufficient; the remaining 3.3× is KERNEL EFFICIENCY (native fp4 MMA ~302µs/M=1 expert GEMM vs vLLM's tuned cutlass sm120a fp4 + fused norm/quant/silu). Parity = TWO campaigns: (A) device-resident+graph → ~5.9; (B) cutlass DirectD experts + fused ops + M=1-tuned GEMV → the rest. **CAMPAIGN-B FIRST BRICK LANDED (2026-08-01): coalesced M=1 fp4 GEMV** (`MatmulNvfp4Fp4Gemv`, one warp/column, coalesced weight-row reads, `VT_NVFP4_FP4_GEMV` default-ON) — same-binary A/B: **decode 0.20 → 0.15 s/tok (1.33×; → ~6.7 tok/s), prefill 1.14 → 0.86s**, coherent+near-tie. **Cumulative this session: 0.16 → ~6.7 tok/s (~42×), now ~2.8× from vLLM 18.8.** (ILP variant `kCpw=4` measured SLOWER — 0.21 s/tok, occupancy loss > activation-reuse gain — reverted to `kCpw=1`; kernel kept templated as a re-measurable knob.) **ncu of the GEMV (sudo): sm__throughput 35-71%, DRAM n/a — COMPUTE/LATENCY-bound, not BW-bound.** Corrects the earlier "~6× BW → ~16-17 tok/s" estimate: the next GEMV lever is HARDWARE fp4 dequant (`cvt.e2m1x2`), not more bandwidth. Parity (18.8) is a multi-brick campaign (decode graph + fused norm/quant + hardware-dequant GEMV), not one more kernel. **B0 hw-fp8 SCALE-decode: MEASURED NEGATIVE, reverted (2026-08-01, `ab7a1c1e`).** Replacing the GEMV's per-byte software fp8-e4m3 group-scale decode (`F8E4M3ToF32Dev`/`ldexpf`) with hardware `cvt.rn.f16.e4m3` (`__nv_fp8_e4m3`→float) is bit-exact (ids byte-identical on the real ckpt) but paging-immune ncu shows it NEUTRAL-to-slightly-WORSE (grid768 41.2 vs 41.9µs tie; mean 53.6 vs 49.4µs) — GPU `ldexpf` is a cheap exponent-bit add, not a libcall. NOTE this is the fp8 SCALE decode, NOT the fp4-e2m1 WEIGHT dequant (the `kE2M1` `__constant__` LUT); the LUT→arithmetic/`cvt.e2m1x2` weight-dequant is a SEPARATE still-open lever (spec brick B1). Also: end-to-end wall-clock is unusable for kernel A/B here (67 GiB unified reload swings TPOT 0.16↔1.08 s/tok run-to-run) — kernel-duration ncu is the only honest anchor. **★ B2 SCOPED + DE-RISKED (2026-08-01, zero-DGX) — the real 18.8 lever:** vLLM's 18.8 bar is MARLIN W4A16 (`VLLM_TEST_FORCE_FP8_MARLIN=1`), which is LOW-M-optimized (decode-correct, unlike a tensor-core W4A4 GEMM that wastes M=1 tile rows). The engine already ships the EXACT kernel `vt::MoeGroupedGemmNvfp4Marlin` (1:1 lift of vLLM `moe_wna16_marlin_gemm`) + shared `MarlinRepackExpertWeight`, and qwen3_5 (27B/35B) already routes its NVFP4 experts through it (default-ON `VT_NVFP4_MARLIN`, 16/16-vs-oracle, +22% gate/+80% decode) via `BuildMoeMarlinResident`. So B2 = mirror that for `LagunaMoeWeights.experts_*_fp4` (a `BuildLagunaMoeMarlinResident` reusing the shared repack + route `LagunaFfnBlock`'s fp4 branch to the Marlin grouped GEMM, GEMV kept as the `=0` escape hatch) — pure reuse, no new kernel, matches vLLM's exact W4A16 numerics. **B2 IMPLEMENTED (2026-08-01, `3c49ef37`) — COMPILES CLEAN on GB10 sm_121a, runtime bug pending.** `LagunaMoeResidentMarlin` + `BuildLagunaMoeMarlinResident` (laguna.cpp, `#ifdef VT_MARLIN_NVFP4`) reconstruct the MoE Marlin path over the SHARED `dense_nvfp4::Dev`/`DBuf`/`ResidentNvfp4` + shared `vt::cuda` Marlin repack/align ops + `vt::MoeGroupedGemmNvfp4Marlin`; SACRED 27B/35B path BYTE-UNTOUCHED; gated `VT_LAGUNA_MARLIN_MOE=1` **default-OFF** (zero regression to the default GEMV path). Compiles clean on the full CUDA build. RUN: loads OK (48 layers, 256 experts) but the FIRST FORWARD device-faults silently on the Marlin path — a layout/param bug (suspects: `MoeCombine` bf16-in/f32-out dtype, the down-GEMM reusing the gate/up align, or the fp4-original free omitted → mem ~doubles). NEXT: `compute-sanitizer` localize → fix → near-tie vs the vLLM-Marlin golden + kernel-duration ncu → flip default-ON. Default path unaffected. **UPDATE (`22d6e146`): added the qwen3_5-style fp4-original free after repack** (device transients + host bytes; peak was ~3× the expert tower → past the 119 GiB pool → null-alloc → silent fault the likely cause); compiles clean. The runtime gate stayed INCONCLUSIVE this session (contended/orphaned processes on the shared box, no captured ids) — rerun on a clean uncontended session, compute-sanitizer if it still faults. **★★ B2 VALIDATED on GB10 (2026-08-01, with the mem-free fix): RUN_EXIT=0, coherent, first 13 generated tokens MATCH the vLLM-Marlin golden EXACTLY** (`22345 83 290 350 674 330 5541 966 340 9626 377 15360 81` — the best Laguna-NVFP4 correctness yet, W4A16 matching vLLM's config). **Steady-state decode 0.10 s/tok = ~10 tok/s** (steps 10-17 all 0.10; the TPOT-0.56 average is warmup-polluted — the DevicePool warms over ~9 decode steps then reuses). vs the GEMV path's 6.7 tok/s = **~1.5× faster; the gap to vLLM 18.8 closes from ~3× to ~1.9×.** Memory flat (7.9 GiB host RSS — the fp4-original free worked; it also fixed the first-forward fault). Still `VT_LAGUNA_MARLIN_MOE=1` default-OFF. TO DONE: move the lazy Marlin-resident build (216s first-forward, 48L×256E repack) to model-LOAD time → clean warm A/B + ncu → flip default-ON → matrix/roadmap. Remaining ~1.9×: vLLM graphs its decode (ours still eager) — decode CUDA-graph is the next lever. **REPRODUCED 3× (reproduction gate MET): GB10 runs deterministic — first 18-20 tokens byte-identical, steady-state 0.10 s/tok confirmed each — so the ~10 tok/s + golden-match is gated, not a single sample.** **#234 item (1) — load-time resident-build LANDED (`LagunaBuildMarlinResidents`, called from the example after load; mirrors vLLM process_weights_after_loading): builds all 48L×256E Marlin residents at LOAD so the repack is not a first-token TTFT spike. Fixed an anon-namespace linkage bug (public fn was defined with internal linkage → moved outside the anon namespace); BUILD CLEAN + links on GB10 sm_121a, default-OFF. Runtime prewarm-fires-at-load timing UNVERIFIED this session (repeated ssh-drops ate the run capture); the forward's lazy build is the validated fallback so it cannot regress. Owed: one clean run to confirm the build moved to load + then flip default-ON.** **★★ DONE (2026-08-01): Marlin is now the UNCONDITIONAL DEFAULT (`LagunaMarlinMoeEnabled` default-ON; `=0` is a code-level A/B opt-out no user needs) — "it just works" with NO env. Confirmed in a no-env GB10 run captured via tmux: `MARLIN residents built at load in 238.4s`, prefill 14.78s (build moved OUT of first-forward), golden-matching ids, steady-state 0.10 s/tok = ~10 tok/s (4th reproduction), RSS ~5-8 GiB. So a default Laguna-NVFP4 load on GB10 gets vLLM's own W4A16 Marlin decode (~10 tok/s, ~1.9× from vLLM 18.8) with zero flags. The 238s load-time repack is a one-time cost (mirrors vLLM process_weights_after_loading); optimizing its 48×256 per-expert sync count is a follow-up. Residual to 18.8 = decode CUDA-graph (deferred; user refocusing on DeepSeek next).** Post-lever-1 nsys: the remaining ~4× is HOST-SYNC-bound — 22,115 `cudaStreamSynchronize` (78.6% of API time, ~2,760/token, the per-GEMM `DrainQueue`), GPU kernels fast. Remaining levers: grouped W4A4 MoE (design input: `vt::MoeGroupedGemmNvfp4` is W4A16, so true-W4A4 grouped needs a new fp4×fp4 op or the `use_a16` mode + expert-stacking — needs a spike), device-resident decode (RECOMMENDED — the current forward is host-style so every GEMM drains; keep activations on-device, drain once/step; reuse qwen3_5's `Dev`/`Nvfp4Dev`/`ResidentNvfp4`/device-SwiGLU machinery; kills the 22k syncs; converges with the pending GGUF #228 and lifts both quant paths), decode CUDA-graph. Binding number needs a clean 2-3× re-run. See `docs/BENCHMARKS.md` + the spec N5 plan. See `docs/BENCHMARKS.md` `CLAIM-LAGUNA-VLLM-NVFP4`. Prior W7 nsys attribution: host-orchestration-bound, levers ranked (spec `laguna-s21-w7-speed-2026-07-31.md`, ledger `CLAIM-LAGUNA-W7-SPEED`). Prior RUNNABLE + FAST DECODE (W6, 2026-07-31): a per-layer K/V cache + single-token incremental decode replaces W5's O(n²) STATELESS full-recompute — TOKEN-IDENTICAL (byte-equal ids, md5 match, == the W5 golden) and 5.05× faster per token: decode 3.33 → 0.66 s/tok on the real 3-shard UD-Q4_K_XL GGUF (GB10, `--gpu`, keep-quant), same "The capital of France is" → " Paris.\n\nThe user is seeking a detailed explanation of the concept of \"cultural capital\"…". `LagunaKvCache` (mirrors `DeepseekV4KvCache`, MLA-latent → GQA multi-head K/V) caches post-QK-RMSNorm/post-RoPE K + raw V at f32 (bit-exact by construction: RoPE/QK-norm are position-only and attention is causal). MIXED attention handled per-layer: 12 GLOBAL layers grow the cache unbounded (full causal); 36 SLIDING-WINDOW-512 layers EVICT the oldest rows beyond the 512 window (gemma2/3 `is_sliding`), capping their K/V. `LagunaForwardGgufCached` + shared `LagunaAttention`/`LagunaFfnBlock` helpers used by BOTH forwards (identical float ops — the recompute path's ids are unchanged after the refactor); `examples/laguna_gen --stateless` forces the W5 recompute for the A/B gate. No cache bug: bit-exact on the first run. Next speed: grouped-expert GEMM + device-resident decode (both in-tree from ds4). See `.agents/specs/laguna-s21-w6-2026-07-31.md`. Prior RUNNABLE (W5, 2026-07-31): our engine greedy-generates COHERENT text on the REAL 3-shard UD-Q4_K_XL GGUF (GB10, keep-quant). `laguna-gen` "The capital of France is" → " Paris.\n\nThe user is seeking a detailed explanation of the concept of \"cultural capital\" as developed by French soci…" — the FIRST token is "Paris.", matching the llama.cpp-Poolside reference on the identical bytes. Multi-shard GGUF reader (LagunaGgufCtx routes each of 814 tensors to its shard; shard-1 = header only) + keep-quant tower (attn/dense/shared/experts/lm_head stay Q8_0/Q4_K/Q5_K COMPRESSED, consumed via `vt::MatmulBT`; norms/router/bias/embed → f32) + `LagunaForwardGguf` (the f32 composition with the ~9 GEMM sites swapped to keep-quant Gemm/GemmRowSlice, ds4 precedent) + `examples/laguna_gen`. Real GGUF metadata verified: dual-RoPE freq_base 500000/10000, dims 64/128, YaRN factor 32, sigmoid ungrouped-noaux router (scale 2.5), per-layer Q-head [48 global/72 sliding], per-head softplus out-gate, QK-RMSNorm. Load 20.6s, peak 71 GiB (fits 119 pool). Prior W4 IN PROGRESS (2026-07-31): 73.4 GiB UD-Q4_K_XL GGUF FETCHED + read authoritatively (814 tensors); 3 CPU-verified fidelity corrections grounded in the real GGUF + llama.cpp — per-head QK-RMSNorm (`attn_q/k_norm`, the scope MISSED it), GGUF-authoritative dual-RoPE mscale (llama.cpp `yarn_attn_factor·(1+0.1·ln(factor))`, factor 32 not HF 128), separate `ffn_gate/up_exps`. Keep-quant tower materialization + `ForwardGguf` + the real-model greedy run vs llama.cpp-laguna same-quant oracle = W5 close. Prior: W3 REAL host-reference forward + 3 new ops (`laguna_ops.cpp`, CPU `-Werror` clean, `test_laguna_scaffold` unit-gated)** | Poolside Laguna: 48 layers (12 global + 36 sliding-window-512), 256 routed top-10 + 1 shared expert, per-head **softplus attention output gate**, sigmoid `noaux_tc` router, dual per-layer RoPE (YaRN full-attn / plain sliding), GQA 8 KV / 128 head-dim, 1M ctx. **W3 (2026-07-31):** the 3 genuinely-NEW small host ops landed in `laguna_ops.cpp` — per-head softplus attn out-gate (`LagunaSoftplusHeadGate`), ungrouped sigmoid-noaux router (`LagunaUngroupedRouterTopK`, ds3 noaux_tc MINUS the group step + tie-break razor), dual per-layer RoPE cos/sin builders (`BuildLaguna{FullYarn,Sliding}CosSin`, reusing the pinned YaRN inv_freq over the partial-64 dims); `LagunaModel::Forward` is now a REAL runnable host-reference composition (variable-Q-head GQA + dual RoPE + sliding-window mask + softplus gate + dense L0 / ungrouped-MoE L1..47 + untied lm_head) replacing the `VT_CHECK(false)` stub; `test_laguna_scaffold` **8/8·166** (softplus math, router selection+tie-break RED-first, dual-RoPE bit-match, variable-Q-head shapes, forward composition on synthetic weights), `test_model_registry` 24/24. **W2 (2026-07-30):** registered, `ParseLagunaParams`, GGUF `blk.N.*` name-map + UD-Q4_K_XL quant-mix (Q4_K/Q5_K/Q6_K/Q8_0 ALL already decoded → ZERO new kernel). **W1 oracle DECISION:** vLLM NATIVE `laguna.py` (in pin → config constructs); dual-oracle = vLLM-NVFP4/-FP8 (fits GB10 119 GiB; BF16 235 GiB does NOT) + llama.cpp-Q4_K token-exact. ~85–90% reuse (ds4-MoE + Gemma-sliding + OLMo-3-dual-rope + Q4_K keep-quant, ALREADY landed). DEFERRED (W4): GGUF keep-quant tower materialization + device/paged production forward (loaders still LOUDLY throw) + strict dual-oracle greedy gate on a fetched checkpoint + `poolside_v1` parser. See `.agents/specs/laguna-s21-w3-2026-07-31.md` (+ W1/W2 `laguna-s21-w1w2-2026-07-30.md`, W0 `laguna-s21-scope-2026-07-30.md`). **Decode attention-glue fusion LANDED (2026-08-02, `CLAIM-LAGUNA-GLUE-FUSED`, default-ON `VT_LAGUNA_GLUE_FUSED`, `=0` A/B):** BYTE-EXACT L1 (softplus out-gate → `DecodeAttnCombineKernel` store) + L4 (residual-Add+RMSNorm pairs → the shared `vt::FusedChain(kFusedAddRmsNormStd)` seam) on the resident decode-graph — same-binary A/B ids byte-identical (159/159 @160), paging-immune nsys steady decode **−4.2% GPU-busy (28.90→27.69 ms/step), −120 graph nodes/step (−10%)**, wall drop_caches-tied (no regression). C shared-into-MoeCombine SKIPPED (Laguna's bf16 `MoeCombine` → not byte-exact); L2 qk-norm+RoPE preamble DEFERRED (needs a device-position kernel variant). See BENCHMARKS.md `CLAIM-LAGUNA-GLUE-FUSED`. **On-device greedy sample LANDED (2026-08-02, `CLAIM-LAGUNA-ONDEV-SAMPLE`, default-ON `VT_LAGUNA_ONDEV_SAMPLE`, `=0` A/B):** the resident decode graph used to Synchronize, return the whole `[100352]` logits, and argmax on the HOST between replays (+ host embed-gather of the next token) — the off-framework "born-on-host" seam the decode-framework-routing audit flagged. Now BOTH run ON-DEVICE inside the captured graph: `vt::GreedyArgmax` (lowest-index tie = the exact host winner) → 1-elem device token buffer, + a new capture-safe `embed_gather` kernel gathers the next input embedding from it (the stock `vt::Embedding` is NOT capture-safe: per-call event-sync + D2H ring). BYTE-EXACT (160-id stream identical `=0`/`=1` on `~/laguna-xs-nvfp4`) + faster: paired drop_caches decode wall **+0.28% median** (8/8 reps ≥0; removes ~150 us/step host argmax) at GPU-busy parity (nsys 2-length 27.44→27.42 ms/step). Aligns Laguna decode with vLLM on-device sampling. **Lever 2 (lm_head GEMV DRAM eff) MEASURED, NOT landed:** `[M=1,100352,2048]` bf16 = **170 GB/s (2.41 ms)** = ~91% of the cuBLAS M=1×large-N reference (~187 GB/s / 2.2 ms) — at the M=1 practical floor (the 273 GB/s ceiling is streaming-only, unreachable for a once-read GEMV); ≤0.7%-of-step headroom needs a reduction reorder (near-tie re-gate) ⇒ not chased, per prior "lm_head optimal". See BENCHMARKS.md `CLAIM-LAGUNA-ONDEV-SAMPLE`. **MoE add_rms_norm fold LANDED (2026-08-02, `CLAIM-LAGUNA-MOE-ADDNORM`, default-ON `VT_LAGUNA_MOE_ADDNORM_FUSED`, `=0` A/B):** the glue-fused MoE tail ran its residual update as TWO graph nodes — `vt::Add(hidden,routed)` [`AddKernel`] + `FusedChain(kFusedAddRmsNormStd)` [shared-add+RMSNorm, `RmsNormRowKernel`] — now ONE `fused_add2_rmsnorm` device node/MoE-layer (`hidden=(hidden+routed)+shared; hn=rms_norm(hidden)*w`). BYTE-EXACT (IEEE add commutes + the identical 256-thread shared-tree norm reduction; 160-id stream byte-identical `=0`/`=1` on `~/laguna-xs-nvfp4`) + faster: **−39 `AddKernel` graph nodes/step** (2.63ms→0 over 69 steps), paging-immune nsys 2-length **~−46 us/tok GPU (27339→27293)**, nsys wall **+0.4% (34.00→34.14 tok/s @70-tok)**. Small (byte-exact node-count trim on the graph-captured, GPU-bound decode; the dominant ~72% cost is the bf16 projection GEMVs — see the Lever-B negative in BENCHMARKS.md). See BENCHMARKS.md `CLAIM-LAGUNA-MOE-ADDNORM`. **Shared expert kept fp4 LANDED (2026-08-03, `CLAIM-LAGUNA-SHARED-FP4`, default-ON `VT_LAGUNA_SHARED_FP4`, `=0` A/B):** the XS-NVFP4 shared expert was DEQUANTIZED to bf16 at load (`LnLoadSharedExpertBf16`) → the M=1 decode GEMV read 4× the DRAM bytes of vLLM (which keeps it fp4). Now kept fp4-resident and routed through the SAME Marlin W4A16 single-expert (num_experts=1) grouped GEMM the routed experts win on (`dense_nvfp4::GateUpFusedMarlinD`+`MatmulNvfp4MarlinD`); the decode GEMV drops to router-ONLY (`moe.router`), shared gate/up/down go fp4. ADDITIVE new `laguna_shared_fp4.cpp` re-reads the on-disk fp4 from the gen driver before shard release (does NOT touch SACRED `laguna_weights.cpp`); bf16 shared KEPT for the T>1 prefill. NEAR-TIE (fp4≠bf16): coherent, first-20 ids == documented golden, byte-identical to bf16 for ~85 tokens then diverges; **DISTRIBUTIONAL GATE PASS 40/40** (ours' first-40 ids ∈ vLLM's 8-run greedy candidate set; vLLM XS-greedy is bf16-non-det, 8 unique of 8). FASTER: paging-immune nsys 2-length **GPU 27.24→26.53 ms/step (−2.6%)**, wall drop_caches **35.8→36.3 tok/s (+1.4%, fp4 wins all 3 reps)**; shared-expert kernel bucket ~1.68→~0.90 ms/step (halved); vs vLLM ~43 tok/s 83.3%→84.4%; RSS 22.2→22.1 GiB (freed the decode-only fused router-shared projection). Modest by design — XS's shared expert is small (`shared_expert_intermediate_size==moe_intermediate_size==512`). Default-ON per parity (matches vLLM's fp4 shared). See BENCHMARKS.md `CLAIM-LAGUNA-SHARED-FP4`. **qk-norm+RoPE preamble fusion LANDED (2026-08-03, `CLAIM-LAGUNA-PREAMBLE-FUSED`, default-ON `VT_LAGUNA_PREAMBLE_FUSED`, `=0` A/B):** closes the `CLAIM-LAGUNA-GLUE-FUSED` L2 deferral — the decode graph ran the per-layer attention preamble as FOUR under-occupied M=1 nodes (`rms_norm_seq(q)`+`rms_norm_seq(k)`+`rope_from_cache_g(q)`+`rope_from_cache_g(k)`); now ONE capture-safe `fused_qk_norm_rope_g` node/layer (`FusedQkNormRopeGKernel`, one block/head, reads the decode position from DEVICE `*pos_buf`, handles the per-layer dual-RoPE 64/128 + `Hq` 48/64). BYTE-EXACT BY CONSTRUCTION: it replicates the composed path's f32 MEMORY round-trip (Phase A 256-thread Σx² == `RmsNormSeqKernel`; Phase B the same `(x*inv)*w` store; `__syncthreads`; Phase C the `RopeFromCacheGKernel` rope read back) — an earlier register-only recompute was numerically-equivalent but diverged at a token-110 near-tie via compiler fma-contraction; the memory boundary forces bit-identity. 160-id stream byte-identical `=0`/`=1` on `~/laguna-xs-nvfp4` (determinism verified `=0`×3/`=1`×3 each run-to-run identical). FASTER: preamble norm+rope kernels **160→40 launches/tok, 326→154 us/tok (−0.17 ms/step)**; all decode-scaling kernels 26.53→26.37 ms/step; wall drop_caches **36.42→36.64 tok/s (+0.6%, fused wins all 3 paired reps)**; vs vLLM ~43 84.7%→85.2%. Modest (preamble ~1.2% of the 26.5 ms/step decode; the dominant cost stays the bf16 projection GEMVs at cuBLAS parity) — a byte-exact graph-node/launch trim (the glue-fusion residual mechanism). Default-ON per parity. See BENCHMARKS.md `CLAIM-LAGUNA-PREAMBLE-FUSED`. **W7 two-front pass LANDED (2026-08-03, `CLAIM-LAGUNA-W7-DECODE`):** FRONT 1 — the example driver logged `[gen] step N …(RSS)` EVERY decode step, and the RSS arg calls `CurResidentGiB()` (a `/proc/self/status` read) + an unbuffered stderr write in the GPU-idle gap between replays; guarded behind `VT_LAGUNA_STEP_LOG` (default OFF) + added a `decode_wall` line (TRUE end-to-end throughput incl. per-step gaps) next to the gap-free `decode_hp`. Since the fprintf sat OUTSIDE the `s0→s1` timer, `decode_hp` was ALREADY honest; with the log off `decode_wall == decode_hp` (within 0.001 tok/s, every LOG_OFF rep) and the recovered host tax is only ~0.1% (drop_caches noise floor). CONCLUSION: the ~86% gap to vLLM 43 is genuine device compute, NOT a harness artifact. FRONT 2 — `VT_LAGUNA_MOE_ONECAST` (default ON): a MoE layer cast the same `hn[1,H]` f32→bf16 THREE times (router GEMV + routed Marlin + shared Marlin); now cast ONCE into a persistent buffer and reuse (`CastHnBf16`/`GemmBf16Pre` + optional pre-cast param on both `…Into` helpers). BYTE-EXACT (deterministic truncation; `=1` vs `=0` byte-identical 300-tok ids); `CastBf16` **200→122 nodes/step (−78 = 2×39 MoE layers)**, GPU-busy parity within nsys noise, decode_hp +0.29%. Combined (onecast on + log off) **36.97 tok/s = 86.0% of vLLM-NVFP4 43** (from 36.64/85.2%). See BENCHMARKS.md `CLAIM-LAGUNA-W7-DECODE`. **Tail-fold follow-up LANDED (2026-08-03, `CLAIM-LAGUNA-TAIL-FUSED`, default-ON `VT_LAGUNA_TAIL_FUSED`, `=0` A/B):** a fresh node-ranking of the baseline decode graph found the routed-MoE `CastF32` as the one clean byte-exact fold left; it folds into the trailing `fused_add2_rmsnorm` via a new bf16-x1 sibling kernel (`AddAdd2RmsNormStdBf16Kernel` — `MoeCombine` writes bf16 straight to a persistent buffer, widened in-kernel by `__bfloat162float`). BYTE-EXACT (`=1` vs `=0` byte-identical 160-tok ids), `CastF32` **78→39 nodes/step**, total graph nodes **919→880**, GPU-busy parity; decode_hp a WASH (median +0.14% / mean −0.04%, at the drop_caches noise floor). Lands on the deterministic node-count basis (like onecast/preamble/addnorm), NOT a wall win; combined headline UNCHANGED **36.97 tok/s = 86.0%**. The ranking confirms the byte-exact decode-tail fold tier is now essentially EXHAUSTED (residual tail = already-folded norms + attention compute + cuBLAS-adjacent router/topk + ported-Marlin `MoeAlign`/`SiluAndMul`/`MoeCombine`); the gap to vLLM 43 is genuine device compute at the practical ceiling. See BENCHMARKS.md `CLAIM-LAGUNA-TAIL-FUSED`. **KERNEL-EFFICIENCY tier (2026-08-03, `VT_LAGUNA_FAST_NORM` default ON + f32 ext of `VT_RMSNORM_DECODE_FAST`):** the fold tier was exhausted but the residual-stream norm KERNELS were still under-occupied — `ncu` on the shipped `<<<1,256>>>` `AddAdd2RmsNormStdBf16`/`RmsNormRow` decode norms: `launch__waves_per_multiprocessor≈0.00`, `sm__throughput≈0.06%` (one 256-thread block on 1 SM of ~100+, latency-bound). Porting the PROVEN bit-identical `RmsNormRowFastKernel` structure (1024-thread float4 memory passes; 256-strided-partial + tree reduction reproduced byte-for-byte) to the f32 kernels cut each **286→~155 µs/tok (1.85×)**, **byte-exact** (160-tok ids identical `=1`vs`=0`; the f32 fix vs the bf16 sibling: store `v` not `v²` and square in the reduction so nvcc emits shipped's `acc += v*v` **fma** — a pre-squared f32 `v²` is not exact and flipped an XS near-tie at tok 108). **−0.81% decode-step GPU time** (paging-immune 70-vs-20 2-length diff, 26192→25980 µs/step); wall-clock ON/OFF overlap (noise floor). Residual: the byte-exact 256-strided reduction can't reach vLLM's per-kernel norm floor (~2.4× vLLM) without breaking byte-exactness → that remainder is byte-exactness-BLOCKED. See BENCHMARKS.md `CLAIM-LAGUNA-FAST-NORM`. **Router top-k warp-shuffle LANDED (2026-08-03, `CLAIM-LAGUNA-TOPK-SHFL`, default-ON `VT_LAGUNA_TOPK_SHFL`, `=0` A/B): BYTE-EXACT** — an nsys 2-length rank of the remaining small kernels (past the at-parity `gemvx` projection GEMVs ~69% of step + Marlin MoE) put the router `SigmoidTopKKernel` top (415 µs/step); `ncu` showed it `<<<1,256>>>` at `waves≈0.000`/`sm≈0.2%` — pure latency (8 serially-dependent rounds × a ~10-sync `sh[256]` argmax tree). New `SigmoidTopKShflKernel` reduces each round by warp-shuffle argmax (2 syncs/round; argmax over the total order is associative ⇒ SAME winner) → **`SigmoidTopK` 414.6→248.8 µs/step (1.67×)**, decode-step GPU **−0.57%** (26.018→25.869 ms/step), 37.39→37.49 tok/s decode_hp (**87.2% of vLLM-NVFP4 43**); 160-id stream byte-identical `=1`vs`=0`. **NOT landed — norm warp-shuffle (`VT_LAGUNA_NORM_SHFL`):** a near-tie register-accumulate+shuffle reduce for the Laguna `AddAdd2RmsNormStd{,Bf16}Fast` norms PASSED the distributional gate (coherent, in-set 38/40 = baseline, one near-tie fork at pos 37) and was −19.3% per-kernel (`AddAdd2RmsNormStdBf16` 150.3→121.3 µs/step) BUT washed at whole-step (0.6% of step; +0.02% within noise) — a near-tie fork isn't justified by a below-noise gain, so it was dropped. The small-kernel norm tail is at its occupancy floor; the decode step is dominated by the at-parity projection GEMVs. See BENCHMARKS.md `CLAIM-LAGUNA-TOPK-SHFL`. **Shared-expert 2-stream overlap LANDED (2026-08-03, `CLAIM-LAGUNA-SHARED-AUX`, default-ON `VT_LAGUNA_SHARED_AUX`, `=0` A/B):** mirror of vLLM's `MULTI_STREAM_OVERLAPPED` — in `LagunaGraph::RunChain` the fp4-shared arm's shared expert is EARLY-forked onto a second CUDA stream from the post-attn hidden `hn` BEFORE the router GEMV (aux reads `hn` f32 + does its own byte-identical cast; scratch from `AuxPool`), overlapping router+`sigmoid_topk`+routed grouped GEMM, joined before the combine — the SAME machinery the 35B ships default-ON (ENG-MOE-SHARED-AUX, runs inside the captured graph). This is the EARLY fork the prior fused-`router_shared_gu` attempt (`89e0d074`, −0.35% wash) could not reach. Capture-safe (aux stream+2 events in the ctor; gstate-0 warm-run builds residents + warms `AuxPool`). **BYTE-EXACT** (`=1`vs`=0` byte-identical 63-tok ids). REAL concurrency: nsys `--cuda-graph-trace=node` 20↔70 sum-vs-union → OVERLAP **2.34 ms/step** (SUM/UNION 1.092) vs `=0`'s 0.0004 ms; net GPU-busy wall **26.213→25.467 ms/step (−2.9%, 38.15→39.27 tok/s)**, wall @200 37.08→37.93 (+2.3%). Net"}}, +]}]) +``` + +Accepted part types (`src/vllm/entrypoints/openai/chat_mm.cpp`): + +| part type | modality | +|---|---| +| `image_url` | image | +| `video_url` | video | +| `input_audio` / `audio_url` | audio | + +## MiniMax-H3: video + audio generation + + +Renders an MP4 with a stereo track. Weights: a GGUF DiT (use **Q4_K_M**), the Qwen3-VL-32B +encoder, and both VAEs. + +```sh +build/examples/minimax-h3-gen \ + --dit MiniMax-H3-FL2VA-Q4_K_M.gguf --dequant-bf16 \ + --encoder qwen3vl-32B-MiniMax-H3-Q4_K_M.gguf --tokenizer tokenizer.json \ + --prompt "A golden retriever runs across a sunlit beach, waves crashing behind it" \ + --video-vae video_vae.safetensors --video-vae-config video_vae_config.json \ + --audio-vae audio_vae.safetensors --audio-vae-config audio_vae_config.json \ + --frames 124 --height 480 --width 864 --steps 50 \ + --device cuda --out out.mp4 --workdir /tmp/h3 +``` + +Conditioning modes, all optional and mutually exclusive where noted: + +```sh +--first-frame start.ppm --last-frame end.ppm # pin the first and/or last frame (fl2va) +--ref-image subject.ppm # reference image, repeatable (ref2va) +--ref-video prev_workdir/ # reference clip, reads frame_%06d.ppm +--ref-audio voice.wav # reference audio +--noise-aug 0.9 # how hard a keyframe is pinned (1.0 = exact) +``` + +Reference frames are binary PPM, which is what this tool also **writes**, so one run's `--workdir` +feeds straight back in as `--ref-video` and clips chain. Convert anything else with +`ffmpeg -i in.png -pix_fmt rgb24 out.ppm`. + +Useful for measurement: `--denoise-only` times the DiT loop without loading the VAEs, +`--dump-params` prints the geometry a checkpoint implies (manifest only, no weights), +`--save-embeds` writes the text conditioning so a second run can replay it with +`--prompt-embeds` and compare checkpoints on identical conditioning. + +Served over HTTP too: pass `--video-dit` (plus the VAEs and configs) to `examples/server` and +`POST /v1/videos`, `POST /v1/videos/sync` and `GET /v1/videos/{id}` register. Without it the +routes stay unregistered. + diff --git a/src/vllm/model_executor/models/minimax_h3_pipeline.cpp b/src/vllm/model_executor/models/minimax_h3_pipeline.cpp index ccd03e98..20026fdf 100644 --- a/src/vllm/model_executor/models/minimax_h3_pipeline.cpp +++ b/src/vllm/model_executor/models/minimax_h3_pipeline.cpp @@ -436,7 +436,12 @@ MiniMaxH3T2vaResult MiniMaxH3GenerateT2va(vt::Device device, const MiniMaxH3T2va const int64_t video_per_channel = request.latent_t * request.latent_h * request.latent_w; denormalize(video_latent, dit_params.latents_dim, video_per_channel, request.video_latents_mean, request.video_latents_std); - const int64_t audio_steps = request.audio_t / request.audio_channel; + // `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). Dividing by the channel count here halved + // the decoded audio, and because the muxer passes `-shortest` that silently + // truncated the VIDEO to half its frames too: a 124-frame render muxed 61. + const int64_t audio_steps = request.audio_t; denormalize(audio_latent, dit_params.audio_latents_dim, audio_steps * request.audio_channel, request.audio_latents_mean, request.audio_latents_std); diff --git a/tests/vllm/models/test_minimax_h3.cpp b/tests/vllm/models/test_minimax_h3.cpp index 3d28ac7d..f5361a5a 100644 --- a/tests/vllm/models/test_minimax_h3.cpp +++ b/tests/vllm/models/test_minimax_h3.cpp @@ -2805,6 +2805,38 @@ TEST_CASE("minimax_h3: the encoder text tower matches upstream, with all three H // ImageNet-normalized values to a writer that expects [-1, 1]: it casts colour (the // per-channel means differ) and compresses the dynamic range ~4.4x (std ~0.22), // which is what "dark and washed out" looks like. +// The decoded AUDIO must last as long as the decoded VIDEO. Nothing checked this, +// and it was wrong: `audio_t` is the PER-CHANNEL latent length (planner: 40 Hz * +// duration) but the pipeline divided it by the channel count, halving the audio. +// Because the muxer passes `-shortest`, that silently truncated the VIDEO too -- +// a 124-frame render produced a 61-frame MP4. Every structural check passed +// throughout: shapes were self-consistent, just half as long as intended. +// +// Found by RENDERING, not by the suite, which is the part worth recording: the +// existing gates all assert shape SELF-CONSISTENCY, and a uniformly halved +// pipeline is self-consistent. The invariant below is the one that is not. +TEST_CASE("minimax_h3: decoded audio spans the same duration as the video") { + // Planner geometry: audio latents run at 40 Hz over num_frames / fps seconds. + const int64_t num_frames = 124; + const double seconds = static_cast(num_frames) / vllm::kMiniMaxH3Fps; + const int64_t audio_t = vllm::MiniMaxH3AudioLatentT(seconds); + INFO("num_frames=" << num_frames << " seconds=" << seconds << " audio_t=" << audio_t); + + // 40 Hz * 5.1667 s = 207 latent steps PER CHANNEL, not 103. + CHECK(audio_t == 207); + + // The invariant the bug broke: latent steps / 40 Hz must equal the video + // duration, to within one latent step. + const double audio_seconds = static_cast(audio_t) / 40.0; + INFO("audio " << audio_seconds << " s vs video " << seconds << " s"); + CHECK(std::abs(audio_seconds - seconds) <= 1.0 / 40.0); + + // And the halving specifically: dividing by the 2 channels would give ~half the + // duration, which is what shipped. + const double halved = static_cast(audio_t / vllm::kMiniMaxH3AudioChannels) / 40.0; + CHECK(std::abs(halved - seconds) > 1.0); // the wrong value is off by ~2.5 s +} + TEST_CASE("minimax_h3: ImageNet pixel de/normalization matches upstream's wrapper") { const int64_t n = 5; // Round trip: normalize then de-normalize must return the original, for values @@ -3474,6 +3506,34 @@ TEST_CASE("minimax_h3: the WHOLE t2va path composes end to end") { CHECK(out.audio_channels == request.audio_channel); CHECK(out.sample_rate == vllm::kMiniMaxH3AudioSampleRate); CHECK(out.audio_samples_per_channel > 0); + + // DURATION, not merely self-consistency. `audio_t` is the PER-CHANNEL latent + // length; the pipeline used to divide it by the channel count, so the decoded + // audio ran half the video's length and `-shortest` truncated the MP4 to match + // (a 124-frame render muxed as 61). Every shape check still passed, because a + // uniformly halved pipeline is self-consistent. The decoder is linear in latent + // steps, so decoding `audio_t` steps INDEPENDENTLY and requiring the pipeline to + // have produced exactly that many samples is the check that is not + // self-consistency: with the bug this is off by the channel count. + { + const std::vector probe( + static_cast(p.audio_latents_dim * request.audio_t), 0.0f); + int64_t expected_samples = 0; + vllm::MiniMaxH3AudioVaeDecode(audio_config, audio_weights, probe, request.audio_t, + &expected_samples); + INFO("audio_t=" << request.audio_t << " channels=" << request.audio_channel + << " expected=" << expected_samples + << " got=" << out.audio_samples_per_channel); + CHECK(out.audio_samples_per_channel == expected_samples); + // And the halved value is genuinely different, so the check has teeth. + int64_t halved_samples = 0; + const std::vector halved_probe( + static_cast(p.audio_latents_dim * (request.audio_t / request.audio_channel)), + 0.0f); + vllm::MiniMaxH3AudioVaeDecode(audio_config, audio_weights, halved_probe, + request.audio_t / request.audio_channel, &halved_samples); + CHECK(halved_samples != expected_samples); + } CHECK(static_cast(out.waveform.size()) == out.audio_channels * out.audio_samples_per_channel); for (float v : out.waveform) {