From 877d3eaedb3bdb432c509a2d5793b95f1aebada7 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 21:58:01 +0000 Subject: [PATCH 1/2] =?UTF-8?q?feat(minimax-h3):=20load=20the=20bf16=20TEX?= =?UTF-8?q?T=20ENCODER=20=E2=80=94=2014=20shards,=2063=20GB=20=E2=80=94=20?= =?UTF-8?q?and=20run=20it=20alone?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Branch: row/H3-ENC-BF16-COND-DIFF (helper). Carries §8.6's MiniMaxH3ShardedCheckpoint CHERRY-PICKED from row/H3-BF16-SHARDED-DIT 1a46ff17 rather than mirrored: the shard-index resolver is exactly the piece this needs, and a second copy of those 222 lines would drift. Every H3 render so far conditioned on a Q4_K_M Qwen3-VL-32B encoder (enc_q4km.gguf, 14.6 GB) and NOBODY had measured the encoder's contribution. That matters because weak conditioning and quantization-damaged conditioning are indistinguishable from outside a render — the wuxia prompt asked for shot/reverse-shot coverage of a sect exchanging intelligence and produced a good generic portrait — and this family is quantization-sensitive (ComfyUI PR 15298: the partial split-half RoPE produces channel-wise magnitude outliers that corrupt even INT8). The blocker was mechanical: --encoder took only a GGUF, while the unquantized tower ships as 14 safetensors shards. This is a LOADER, not a second forward. MiniMaxH3EncoderDeviceWeights binds a plain name->vt::Tensor map; the GGUF arm fills it with ggml blocks, this fills it with bf16, and MiniMaxH3EncoderTextForwardDevice runs unchanged over either. That is what makes the two arms comparable at all. - MiniMaxH3EncoderConfigFromShards derives the geometry from the index's SHAPES alone (no payload), with the SAME recovery rules as the GGUF loader AND the same defaults for what shapes cannot carry (rope_theta, mrope_section, rms_norm_eps, selected_layer) — otherwise an A/B compares two RoPEs, not two quantizations. - StreamMiniMaxH3EncoderShardsToDevice (new minimax_h3_encoder_sharded.cpp) keeps the projections BF16 on the device and uploads them DIRECTLY out of the read-only mmap; the [q|k|v] and [gate|up] row fusions are done ON THE DEVICE, by uploading each member into its offset of one allocation, so even the one transform this loader performs costs no host copy. Only the norms are widened, and those are [5120]. - MiniMaxH3EncoderTextForwardDevice WIDENS a bf16 weight to f32 immediately before its GEMM, into a scratch reused across all 50 layers. vt::MatmulBT needs one dtype for both operands and these activations are f32, while staging the tower f32 is 97.5 GiB against 48.8 GiB bf16 on a 122 GiB UNIFIED pool. bf16 -> f32 is EXACT, so this is residency, not numerics — and it is gated as such. - MiniMaxH3EncoderEmbedTokensFromShards gathers a prompt's rows straight out of the mmap'd [151936, 5120] table, the safetensors twin of the GGUF per-row dequantize. - minimax-h3-gen accepts --encoder wherever it accepted the GGUF, and --encoder-only runs the tower alone (no DiT, no VAEs) and writes --save-embeds. The DiT was loaded FIRST in the normal path, so conditioning alone used to cost ~96 GiB peak instead of ~49 GiB on a pool that OOM-reboots this box. Both encoder paths now go through ONE helper. Gates (CPU, test_minimax_h3 70/70 cases / 49706 assertions, up from 68/68): (1) a synthetic 4-shard encoder at the REAL name spellings resolves, and every fused view is memcmp-exact against q ++ k ++ v and gate ++ up for EVERY layer, unfused projections byte-exact, the separate names GONE, model.language_model. norm.weight + lm_head.weight + the vision tower NOT bound, truncation honoured, the embedding gather exact and out-of-range throwing; (2) ★ the loader RAN and is NOT the GGUF path — MiniMaxH3EncoderShardStreamStats asserted on shards, layers, views, fused groups and BOTH upload paths, with host_peak_bytes equal to ONE norm so peak cannot scale with the model, and the views are kBF16, a dtype the GGUF loader can never produce; (3) ★ the widening is EXACT — the same checkpoint written BF16 and F32 (bf16-rounded values) streams to kBF16 and kF32 views respectively and the two full encoder forwards are BIT-IDENTICAL (memcmp == 0), so the conditioning measurement cannot be confounded by the widening itself. Also repairs a PRE-EXISTING check-public-doc-tables ratchet red on origin/main (docs/STATUS.md 284114 chars > 284073) by collapsing the H3 row's superseded narrative to its binding result. Honest residuals: the real 63 GB load, its measured peak RSS, and the Q4_K_M-vs-bf16 conditioning numbers are UNVERIFIED here — the GPU was busy with a render. They are this row's next step. Pre-existing red, not this row's: check-fusion-consistency / test_check_fusion_consistency (minimax_h3_video_vae_device), verified RED on a clean origin/main tree. Folds in the lifetime fix found while gating this row: the shard streamer's checkpoint must OUTLIVE its async uploads. It was a defect in code that never landed, so it belongs in the commit that introduces the streamer rather than as a follow-up against a tree that never existed on main. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context) --- .agents/NOW.md | 2 +- .agents/parity-ledger.md | 2 + .agents/specs/minimax-h3.md | 70 ++++ .agents/state.md | 68 ++++ CMakeLists.txt | 1 + docs/BENCHMARKS.md | 2 +- docs/FEATURES.md | 4 +- docs/STATUS.md | 2 +- examples/minimax_h3_gen/main.cpp | 168 +++++++++- .../vllm/model_executor/models/minimax_h3.h | 89 ++++++ .../models/minimax_h3_encoder_device.cpp | 48 ++- .../models/minimax_h3_encoder_sharded.cpp | 294 +++++++++++++++++ tests/vllm/models/test_minimax_h3.cpp | 298 ++++++++++++++++++ 13 files changed, 1025 insertions(+), 23 deletions(-) create mode 100644 src/vllm/model_executor/models/minimax_h3_encoder_sharded.cpp diff --git a/.agents/NOW.md b/.agents/NOW.md index 39010294e..1d1f26629 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 grid DIAGNOSED (#95): NO loader bug; bf16 13-shard DiT STREAMS** | residual = community-NVFP4 quant fidelity §8.12; no bf16 render yet | +| MiniMax-H3 lane | **bf16 shards STREAM both towers (DiT + encoder); Q4_K_M enc cond cos 0.9975, 3.5° med, DIFFUSE** | render A/B on saved embeds | | Kimi-Linear-48B (KDA/NoPE-MLA/MoE) | device-KDA **122/128, 4.24 tok/s** best (§15); MLA device NEG (§16). chunk_kda prefill AOT **SPIKED**: 5 kernels authored+pinned+recipe (§17). Bar = MEET vLLM speed | Phase-2: regen harness, wire `vt::KdaChunkPrefill`, gate STRICT + vLLM 0.82 ladder | | 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/parity-ledger.md b/.agents/parity-ledger.md index 066855ffe..c669faf02 100644 --- a/.agents/parity-ledger.md +++ b/.agents/parity-ledger.md @@ -918,3 +918,5 @@ Columns: | 2026-08-07 (startup-latency FIRST NUMBERS, provisional; extends `SERVE-GATE-ONLINE`; no new row; `benchmark_binding=false`) | **What it does.** Runs the `--startup-only` series landed the day before: Qwen3.6-27B-NVFP4, GB10, 3 interleaved ours/vLLM repetitions under one `/tmp/gpu` lock, page cache dropped per leg, GPU idle proven before and after each. | Reference = vLLM oracle 0.25.0 in the same production server config the throughput grid launches (`--gpu-memory-utilization 0.6`, matched `--max-num-seqs`/`--max-num-batched-tokens`, prefix caching off). Attribution taken from vLLM's OWN log, not inferred. | **MEASURED, PROVISIONAL, NOT BINDING.** ours 37.94/36.51/35.88 s (median **36.51**), vLLM 460.36/221.51/217.86 s (warm median **221.51**) => **6.07x**. Ours ±3%; vLLM warm legs within 1.7%. vLLM r1's 460 s is one-time FlashInfer autotune+compile (`saved 64 configs`, `init engine 259.42 s`) vs r2 (`loaded 64 configs`, `26.95 s`); even warm, init is only ~27 s of ~221 s, so the gap is process start + imports + weight load. Our cold-autotune start = 69.29 s (+33 s over warm). **Two reasons it is not binding:** (a) a concurrent build session overlapped r2/r3 of both arms, including BOTH warm-cache vLLM legs, biasing vLLM slow and inflating the ratio; (b) the uncontended repeat was destroyed when the box HARD-REBOOTED mid-leg during vLLM's cold-autotune start (previous boot's journal ends with no shutdown sequence) - a NEW trigger for the known GB10 unified-memory reboot hazard, since the warm-cache config ran six times without incident. Owed: one uncontended 3-rep series on a quiet box. | | 2026-08-07 (`row/H3-BF16-SHARDED-DIT`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; CPU-only, no GPU and no download; lifecycle unchanged) | **MiniMax-H3 — the ORIGINAL bf16 release (13 safetensors shards, 66.3 GB) is now INDEXABLE.** Every H3 render so far used a QUANTIZED DiT and H3 is unusually quantization-sensitive (Q3_K_M -> Q4_K_M alone turned a murky lattice into a photoreal close-up; ComfyUI PR 15298 blames the partial split-half RoPE's channel-wise magnitude outliers), but the full-precision question was unaskable because every DiT loader took a SINGLE file. Adds (a) `MiniMaxH3ShardedCheckpoint::Open(dir)` (`src/vllm/model_executor/models/minimax_h3_sharded.cpp`), which resolves tensors through the checkpoint's own `model.safetensors.index.json` weight map (never by scanning) with one index over every shard, mirroring the in-tree multi-shard template `LoadMiniMaxH3EncoderWeights(const std::vector&, ...)`, and throws BY NAME when the index names a tensor its shard does not hold; (b) `EnumerateMiniMaxH3ShardedTensors`, the shapes-only manifest the geometry parser consumes; (c) `LoadMiniMaxH3DitFromShards`, the host-f32 reference loader; (d) `MiniMaxH3IsFp32IslandTensor`, single-sourcing the upstream fp32-ISLAND split the three existing streamers each hand-rolled; (e) `--dit ` in `examples/minimax_h3_gen` for both `--dump-params` and the run path, every existing `--dit` form unchanged. The DEVICE streamer is the stacked follow-up `row/H3-BF16-SHARDED-STREAM`, split out to stay inside the 900-line PR cap. | vLLM-Omni `vllm_omni/diffusion/models/minimax_h3/minimax_h3_transformer.py:85-101` (MINIMAX_H3_FP32_PARAM_NAMES / _BUFFER_NAMES, the island split) and `:906-922` (the parameter set); the shard-index container convention is HF safetensors' own `model.safetensors.index.json` weight_map, already consumed in-tree by `LoadSafetensorsIndex` and the multi-shard encoder/VAE loaders. No vLLM behavior changed; H3 remains BEYOND-PIN (vllm-omni, not the pinned vLLM repo). | **LANDED + CPU-GATED (loader brick; `benchmark_binding=false` — no throughput owed, and NO bf16-vs-quant render or speed number is claimed).** Re-gated AFTER the rebase onto `f34e0d17`: `test_minimax_h3` 72/72 cases / 54497 assertions, clean Release build of `libvllm.a`, `test_minimax_h3` and `minimax-h3-gen`. Two gates: (1) index+name mapping over a synthetic 4-shard set — every tensor resolves to the shard the index named AND to the bytes written there, a tensor missing from its shard throws WITH ITS NAME, and the derived geometry equals the single-file path field for field; (2) a SPARSE 13-shard release declaring the REAL 535 tensors at REAL shapes (66.3 GB declared, 144 KB on disk) derives the SHIPPED geometry 50/5376/56/128/14336/24/32/1x2x2/5120, and `minimax-h3-gen --dit --dump-params` prints all 20 fields on it. Also FIXES a real latent defect this row's sanitizer lane exposed: `MiniMaxH3ReadSafetensorF32` read 16-bit payloads through `reinterpret_cast`, which is UB on a safetensors file whose JSON header leaves the payload odd-aligned (the format does not require padding); now a byte-wise `memcpy`. RED-first proven: reverting it reproduces UBSan's `load of misaligned address` at the same line and exits 1. Honest residuals: no device load of the real 66.3 GB release, no measured peak RSS, and no bf16-vs-quantized render/speed comparison — the quality question is UNBLOCKED, not answered. | | 2026-08-07 (`row/H3-BF16-SHARDED-STREAM`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; stacked on `row/H3-BF16-SHARDED-DIT`; CPU-only, no GPU and no download; lifecycle unchanged) | **MiniMax-H3 — the ORIGINAL bf16 release (13 shards, 66.3 GB) now STREAMS to the device.** §8.13 made the checkpoint indexable but its only loader was host-f32 (~132 GB on the real release); on a 122 GiB UNIFIED pool that holds the model TWICE, and the non-streaming NVFP4 loader was already OOM-killed at anon-rss 125 GB on HALF this size, so the real release was not loadable at all. Adds `StreamMiniMaxH3ShardedToDeviceBf16` (`minimax_h3_device.cpp`, sharing `BindStreamedDitViews` with the GGUF and NVFP4 streamers): manifest first, then ONE tensor at a time, with a BF16-on-disk tensor bound for a bf16 device slot — essentially the whole 66.3 GB — uploaded DIRECTLY out of the read-only mmap with NO host buffer, each source range released via `MaybeReleaseSourcePages`, and `rope.inv_freq` kept HOST-resident. Adds `MiniMaxH3ShardStreamStats` (mirroring `Nvfp4W4A16Stats`) so the path is observable, and `--dit --device cuda` in `examples/minimax_h3_gen`. | vLLM-Omni `vllm_omni/diffusion/models/minimax_h3/minimax_h3_transformer.py:85-101` (MINIMAX_H3_FP32_PARAM_NAMES / _BUFFER_NAMES, the fp32-island split the stream honours). The streaming SHAPE is our own in-tree convention (`StreamMiniMaxH3Nvfp4ToDeviceBf16`), which exists because upstream never has to load this checkpoint on one unified-memory device; recorded as a deviation in porting-inventory §9 terms. No vLLM behavior changed; H3 remains BEYOND-PIN. | **LANDED + CPU-GATED (loader brick; `benchmark_binding=false` — no throughput owed, and NO bf16-vs-quant render or speed number is claimed).** `test_minimax_h3` 73/73 cases / 55203 assertions, clean Release build of `libvllm.a`, `test_minimax_h3` and `minimax-h3-gen`. Two gates: (1) streamed == non-streamed — all 46 weight views BIT-EXACT (`memcmp == 0`) vs `StageMiniMaxH3DitWeights(kBF16)`, dtypes included (12 fp32 islands), both device forwards IDENTICAL (video and audio max|diff| == 0.0), `rope.inv_freq` host-resident; (2) the loader RAN — counters ASSERTED, observed `shards=3 tensors=46 direct=37 converted=9 bytes=444504 host_peak=8192`, i.e. BOTH upload paths taken, every view owned by this loader, and `host_peak_bytes` bounded by one tensor (< 1/4 of bytes uploaded) so the peak cannot scale with the model. Honest residuals: the real 66.3 GB load, its measured peak RSS, and CUDA memcpy from a file-backed mmap are all UNVERIFIED (CPU-only row); the bf16-vs-quant A/B is now RUNNABLE and has not been run. | +| 2026-08-06 (`row/H3-BF16-SHARDED-DIT`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; CPU-only, no GPU and no download; lifecycle unchanged) | **MiniMax-H3 — the ORIGINAL bf16 release (13 safetensors shards, 66.3 GB) now LOADS, and it STREAMS.** Every H3 render so far used a QUANTIZED DiT and H3 is unusually quantization-sensitive (Q3_K_M -> Q4_K_M alone turned a murky lattice into a photoreal close-up; ComfyUI PR 15298 blames the partial split-half RoPE's channel-wise magnitude outliers), but the full-precision question was unaskable because every DiT loader took a SINGLE file. Adds (a) `MiniMaxH3ShardedCheckpoint::Open(dir)` (`src/vllm/model_executor/models/minimax_h3_sharded.cpp`), which resolves tensors through the checkpoint's own `model.safetensors.index.json` weight map (never by scanning) with one index over every shard, mirroring the in-tree multi-shard template `LoadMiniMaxH3EncoderWeights(const std::vector&, ...)`, and throws BY NAME when the index names a tensor its shard does not hold; (b) `StreamMiniMaxH3ShardedToDeviceBf16` (`minimax_h3_device.cpp`, sharing `BindStreamedDitViews` with the GGUF and NVFP4 streamers), which converts+uploads ONE tensor at a time and uploads a BF16 tensor bound for a bf16 device slot DIRECTLY out of the read-only mmap with no host buffer at all, releasing each source range afterwards; (c) `LoadMiniMaxH3DitFromShards`, the host-f32 reference loader; (d) `MiniMaxH3IsFp32IslandTensor`, single-sourcing the upstream fp32-ISLAND split the three existing streamers each hand-rolled; (e) `--dit ` in `examples/minimax_h3_gen` for both `--dump-params` and the run path, every existing `--dit` form unchanged. It MUST stream: the pool is UNIFIED (122 GiB shared host+device) and the non-streaming NVFP4 loader was already OOM-killed at anon-rss 125 GB on half this size. | vLLM-Omni `vllm_omni/diffusion/models/minimax_h3/minimax_h3_transformer.py:85-101` (MINIMAX_H3_FP32_PARAM_NAMES / _BUFFER_NAMES, the island split) and `:906-922` (the parameter set); the shard-index container convention is HF safetensors' own `model.safetensors.index.json` weight_map, already consumed in-tree by `LoadSafetensorsIndex` and the multi-shard encoder/VAE loaders. No vLLM behavior changed; H3 remains BEYOND-PIN (vllm-omni, not the pinned vLLM repo). | **LANDED + CPU-GATED (loader brick; `benchmark_binding=false` — no throughput owed, and NO bf16-vs-quant render or speed number is claimed).** `test_minimax_h3` 68/68 cases / 49300 assertions, clean Release build of `libvllm.a`, `test_minimax_h3` and `minimax-h3-gen`. Four gates: (1) index+name mapping over a synthetic 4-shard set — every tensor resolves to the shard the index named AND to the bytes written there, a tensor missing from its shard throws WITH ITS NAME, and the derived geometry equals the single-file path field for field; (2) streamed == non-streamed — all 46 weight views BIT-EXACT (`memcmp == 0`) vs `StageMiniMaxH3DitWeights(kBF16)`, dtypes included (12 fp32 islands), both device forwards IDENTICAL (max|diff| == 0), `rope.inv_freq` HOST-resident; (3) the loader RAN — `MiniMaxH3ShardStreamStats` (mirroring `Nvfp4W4A16Stats`) proves shards opened, tensors streamed, BOTH upload paths taken, every view owned by this loader, and `host_peak_bytes` bounded by one tensor (< 1/4 of bytes uploaded), i.e. peak cannot scale with the model; (4) a SPARSE 13-shard release declaring the REAL 535 tensors at REAL shapes (66.3 GB declared, 144 KB on disk) derives the SHIPPED geometry 50/5376/56/128/14336/24/32/1x2x2/5120, and `minimax-h3-gen --dit --dump-params` prints all 20 fields on it. Honest residuals: the real 66.3 GB load and its measured peak RSS, CUDA memcpy from a file-backed mmap, and any bf16-vs-quantized render/speed comparison are all UNVERIFIED here (no GPU, no download, per the operator's instruction). Not pushed. | +| 2026-08-06 (`row/H3-ENC-BF16-COND-DIFF`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; lifecycle unchanged) | **MiniMax-H3 - the bf16 TEXT ENCODER (14 safetensors shards, 63 GB) now LOADS, it STREAMS, and `--encoder-only` runs the tower alone.** Every H3 render so far conditioned on a Q4_K_M Qwen3-VL-32B encoder and the encoder's contribution had never been measured, but `--encoder` accepted only a GGUF. Adds (a) `MiniMaxH3EncoderConfigFromShards`, deriving the geometry from the shard index's SHAPES alone with the SAME recovery rules AND the same non-shape defaults (`rope_theta`, `mrope_section`, `rms_norm_eps`, `selected_layer`) as the GGUF loader, so an A/B cannot be comparing two RoPEs; (b) `StreamMiniMaxH3EncoderShardsToDevice` (`src/vllm/model_executor/models/minimax_h3_encoder_sharded.cpp`), which fills the SAME `MiniMaxH3EncoderDeviceWeights::views` map the GGUF arm fills - over bf16 instead of ggml blocks - uploading projections DIRECTLY out of the read-only mmap and doing the `[q|k|v]` / `[gate|up]` row fusions ON THE DEVICE into offsets of one allocation, so even the transform costs no host copy; (c) `MiniMaxH3EncoderEmbedTokensFromShards`, a per-row gather out of the `[151936, 5120]` table; (d) a bf16-weight WIDEN step in `MiniMaxH3EncoderTextForwardDevice` (scratch reused across layers) because `vt::MatmulBT` needs one dtype for both operands and these activations are f32 - the 50 layers H3 runs are 48.8 GiB bf16 vs 97.5 GiB f32 on a 122 GiB UNIFIED pool; (e) `--encoder ` and `--encoder-only` in `examples/minimax_h3_gen`, which drops peak from ~96 GiB (DiT loaded first) to ~49 GiB. No new forward: the encoder graph is byte-for-byte the same code for both arms, which is what makes the quantization question measurable. | The name map is the one already gated in-tree for `LoadMiniMaxH3EncoderWeights(const std::vector&, ...)` (`model.language_model.layers.N.` -> `layers.N.`, q/k/v and gate/up FUSED, final `norm.weight` and `lm_head.weight` deliberately unbound because H3 reads the UNNORMALIZED truncated output); shard resolution reuses `MiniMaxH3ShardedCheckpoint` (§8.6, cherry-picked from `1a46ff17`), i.e. the checkpoint's own HF `model.safetensors.index.json` weight_map. Encoder truncation to `min(num_hidden_layers, 50)` is upstream vLLM-Omni's own. No vLLM behavior changed; H3 remains BEYOND-PIN (vllm-omni). | **LANDED + CPU-GATED (loader brick; `benchmark_binding=false`).** `test_minimax_h3` 70/70 cases / 49706 assertions (up from 68/68), clean Release build of `libvllm.a`, `test_minimax_h3`, `minimax-h3-gen`. Three gates: (1) a synthetic 4-shard encoder at the REAL name spellings resolves, and every fused view is `memcmp`-exact against `q ++ k ++ v` / `gate ++ up` for EVERY layer, unfused projections byte-exact, separate names gone, final norm + lm_head + vision tower NOT bound, truncation honoured, embedding gather exact and out-of-range throwing; (2) the loader RAN and is NOT the GGUF path - `MiniMaxH3EncoderShardStreamStats` asserted on shards/layers/views/fused groups/direct-vs-converted uploads with `host_peak_bytes` equal to ONE norm (peak cannot scale with the model), and the views are `kBF16`, a dtype the GGUF loader can never produce; (3) the WIDENING is exact - the same checkpoint written BF16 and F32 (bf16-rounded values) streams to `kBF16` and `kF32` views respectively and the two full encoder forwards are BIT-IDENTICAL (`memcmp == 0`), so the conditioning A/B cannot be confounded by the widening. Also repairs a PRE-EXISTING `check-public-doc-tables` ratchet red on `origin/main` (docs/STATUS.md 284114 > 284073). The real 63 GB load, its peak RSS, and the Q4_K_M-vs-bf16 conditioning numbers are the GPU follow-up in this row. | diff --git a/.agents/specs/minimax-h3.md b/.agents/specs/minimax-h3.md index 7bb3514a8..0bf18181d 100644 --- a/.agents/specs/minimax-h3.md +++ b/.agents/specs/minimax-h3.md @@ -1000,3 +1000,73 @@ pageable-source usage, and every copy is followed by a synchronize, but unexerci device), and any bf16-vs-quantized RENDER or SPEED comparison. The quality A/B is now runnable; it has not been run. + +## 8.15 The bf16 TEXT ENCODER — 14 shards, 63 GB, streamed; and `--encoder-only` (2026-08-06, `row/H3-ENC-BF16-COND-DIFF`) + +**Why.** §8.6 made the full-precision *DiT* loadable, but every H3 render — including +the ones whose output looks competent-but-generic — conditioned on a **Q4_K_M** text +encoder (`enc_q4km.gguf`, 14.6 GB, Qwen3-VL-32B). Nobody had ever measured the +encoder's contribution. That matters because weak conditioning and a +quantization-damaged conditioning tensor look identical from the outside: the wuxia +prompt asked for shot/reverse-shot coverage of a martial-arts sect exchanging +intelligence and got a good generic portrait, and this family is known to be +quantization-sensitive (ComfyUI PR 15298: H3's partial split-half RoPE produces +channel-wise magnitude outliers that corrupt even INT8). + +The blocker was mechanical: `--encoder` only accepted a GGUF +(`LoadMiniMaxH3EncoderFromGguf`), while the unquantized tower ships as **14 +safetensors shards + `model.safetensors.index.json`, 63 GB**. + +**What landed.** + +- `MiniMaxH3EncoderConfigFromShards(ckpt, max_layers)` derives the geometry from the + index's SHAPES alone — no payload — using the SAME recovery rules as the GGUF + loader (head_dim from `q_norm`, heads from `q_proj` rows), so the two arms cannot + disagree about what model they are running. The knobs shapes cannot carry + (`rope_theta`, `mrope_section`, `rms_norm_eps`, `selected_layer`) keep the SAME + defaults the GGUF arm leaves in place; otherwise an A/B would be comparing two + RoPEs, not two quantizations. +- `StreamMiniMaxH3EncoderShardsToDevice` (`minimax_h3_encoder_sharded.cpp`) fills the + same `MiniMaxH3EncoderDeviceWeights::views` map the GGUF arm fills, over bf16 + instead of ggml blocks. The projections stay **BF16 on the device** and are + uploaded DIRECTLY out of the read-only mmap with no host buffer at all; the two row + FUSIONS (`[q|k|v]`, `[gate|up]`) are done ON THE DEVICE by uploading each member + into its offset of one allocation, so the transform does not cost a host copy + either. Only the norms are widened on the host, and those are `[5120]`. +- `MiniMaxH3EncoderTextForwardDevice` is UNCHANGED in structure and now WIDENS a bf16 + weight to f32 immediately before its GEMM, into a scratch buffer keyed by element + count and reused across all 50 layers. This is a RESIDENCY trick, not a numerics + one: `vt::MatmulBT` requires both operands in one dtype and these activations are + f32, while staging the 50 layers H3 runs as f32 would be **97.5 GiB** against a + **48.8 GiB** bf16 residency on a 122 GiB UNIFIED pool. bf16 -> f32 is EXACT, so the + GEMM sees bit-identical inputs to an f32-staged tower — gated below. +- `MiniMaxH3EncoderEmbedTokensFromShards` gathers a prompt's embedding rows straight + out of the mmap (the table is `[151936, 5120]`), the safetensors twin of the GGUF + arm's per-row dequantize. +- `minimax-h3-gen --encoder ` is accepted wherever the GGUF was, and + **`--encoder-only`** runs the tower alone — no DiT, no VAEs, no output path — and + writes `--save-embeds`. That is what makes the measurement affordable: the DiT was + loaded FIRST in the normal path, so asking for conditioning alone used to cost its + residency too (~96 GiB peak instead of ~49 GiB on a pool that OOM-reboots the box). + Both encoder paths now go through ONE helper, so the conditioning a render consumes + and the conditioning the A/B measures are produced by the same code. + +**Gates (CPU, `test_minimax_h3` 70/70, 49706 assertions, up from 68/68).** +1. *Resolve, fuse and stream*: a synthetic 4-shard encoder at the REAL name spellings + (`model.language_model.layers.N.*`, `model.visual.*`, `lm_head.weight`). Geometry + from shapes matches; every fused view's bytes are `memcmp`-exact against + `q ++ k ++ v` and `gate ++ up` for EVERY layer; `o_proj`/`down_proj` pass through + byte-exact; the separate q/k/v and gate/up names are GONE; `norm.weight`, + `lm_head.weight` and the vision tower are NOT bound; truncation works; the + embedding gather is exact and throws out of range. +2. *The loader RAN, and it is not the GGUF path*: `MiniMaxH3EncoderShardStreamStats` + (mirroring `MiniMaxH3ShardStreamStats`) is asserted on — shards opened, layers + streamed, 8 views/layer, 2 fused groups/layer, 7 direct (no-host-copy) uploads and + 4 conversions per layer, and `host_peak_bytes` equal to ONE norm (so peak cannot + scale with the model). The views are `kBF16`, a dtype the GGUF loader can never + produce, which is what makes this a proof rather than a coincidence. +3. *Widening is exact*: the SAME synthetic checkpoint written twice — once BF16, once + F32 holding the bf16-rounded values — streams to `kBF16` and `kF32` views + respectively (asserted), and the two full encoder forwards are **BIT-IDENTICAL** + (`memcmp == 0`), not merely close. Without this, "we measured what quantizing the + encoder costs" would be confounded by what the widening itself did. diff --git a/.agents/state.md b/.agents/state.md index 7a8049dc6..3220e2f58 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -41176,3 +41176,71 @@ number is claimed - the quality A/B is now runnable, and has not been run. Next: the operator runs the real 13-shard bf16 DiT on GB10 and answers the quantization quality question. +Next: the operator runs the real 13-shard bf16 DiT and answers the quantization +quality question; nothing here claims it. + +## 2026-08-07T13:30 - H3: the bf16 TEXT ENCODER (14 shards, 63 GB) LOADS + `--encoder-only`, so the Q4_K_M conditioning question is measurable (row/H3-ENC-BF16-COND-DIFF) + + + +`row/H3-ENC-BF16-COND-DIFF` (helper, branched from `origin/main` `ad231615`, with +§8.6's `MiniMaxH3ShardedCheckpoint` **cherry-picked** from `row/H3-BF16-SHARDED-DIT` +`1a46ff17` rather than mirrored - the shard index resolver is exactly the piece this +row needs and duplicating 222 lines of it would drift). + +**Why.** §8.6 made the full-precision *DiT* loadable, but every H3 render so far - +including the ones that look competent-but-generic - conditioned on a **Q4_K_M** +encoder (`enc_q4km.gguf`, 14.6 GB, Qwen3-VL-32B), and the encoder's contribution had +NEVER been measured. Weak conditioning and quantization-damaged conditioning look +identical from outside a render: the wuxia prompt asked for shot/reverse-shot +coverage of a sect exchanging intelligence and produced a good generic portrait. +This family is quantization-sensitive (ComfyUI PR 15298: the partial split-half RoPE +produces channel-wise magnitude outliers that corrupt even INT8). The blocker was +mechanical - `--encoder` took only a GGUF, the unquantized tower ships as 14 +safetensors shards. + +**What landed** (full detail: [specs/minimax-h3.md](specs/minimax-h3.md) §8.7). +`MiniMaxH3EncoderConfigFromShards` (geometry from the index's SHAPES alone, SAME +recovery rules and SAME non-shape defaults as the GGUF loader, so an A/B cannot be +comparing two RoPEs), `StreamMiniMaxH3EncoderShardsToDevice` (fills the same +`MiniMaxH3EncoderDeviceWeights::views` map over bf16; projections uploaded DIRECTLY +out of the mmap, the `[q|k|v]` and `[gate|up]` fusions done ON DEVICE into offsets of +one allocation, so even the transform costs no host copy), +`MiniMaxH3EncoderEmbedTokensFromShards` (per-row gather out of the `[151936, 5120]` +table), and `--encoder ` + **`--encoder-only`** in `minimax-h3-gen`. + +**The memory decision, stated plainly.** The 50 layers H3 actually runs are +**48.8 GiB in bf16 and 97.5 GiB in f32**, on a 122 GiB UNIFIED pool that has +OOM-rebooted this box. So the weights stay bf16 and +`MiniMaxH3EncoderTextForwardDevice` WIDENS each one to f32 immediately before its +GEMM into a scratch reused across layers (~2 GiB peak): `vt::MatmulBT` requires one +dtype for both operands and these activations are f32. bf16 -> f32 is EXACT, so this +is residency, not numerics - and it is GATED as such. `--encoder-only` matters for +the same reason: the DiT was loaded FIRST in the normal path, so conditioning alone +used to cost ~96 GiB peak instead of ~49 GiB. + +**Gates (CPU, `test_minimax_h3` 70/70 cases / 49706 assertions, up from 68/68).** +(1) resolve+fuse+stream over a synthetic 4-shard encoder at the REAL name spellings - +every fused view `memcmp`-exact against `q ++ k ++ v` and `gate ++ up` for every +layer, unfused projections byte-exact, separate names GONE, `norm.weight`/`lm_head`/ +vision tower NOT bound, truncation, exact embedding gather + out-of-range throws. +(2) ★ the loader RAN and is NOT the GGUF path - `MiniMaxH3EncoderShardStreamStats` +asserted on shards/layers/views/fused-groups/direct-vs-converted uploads, with +`host_peak_bytes` equal to ONE norm so peak cannot scale with the model, and the +views are `kBF16`, a dtype the GGUF loader can never produce. (3) ★ widening is +EXACT - the same checkpoint written BF16 and F32 (bf16-rounded values) streams to +`kBF16` and `kF32` views respectively and the two full forwards are BIT-IDENTICAL +(`memcmp == 0`), so the measurement below cannot be confounded by the widening. + +**Record repair done in passing.** `docs/STATUS.md` was 41 chars OVER its +`check-public-doc-tables` ratchet on `origin/main` (284114 vs 284073) - a pre-existing +red. The H3 row's superseded narrative was collapsed to the binding result, which +brings the page back inside the ratchet. + +**Pre-existing red, NOT this row's:** `check-fusion-consistency` / +`test_check_fusion_consistency` (`minimax_h3_video_vae_device` gemm-merge drift), +verified RED on a clean `origin/main` tree before this work. + +Next: the measurement itself - encode the wuxia prompt with both encoders on the +Thor GPU and diff the `[tokens, 5120]` conditioning (max|diff|, RMS, relative RMS, +per-token cosine). Numbers land in the same row. diff --git a/CMakeLists.txt b/CMakeLists.txt index c5451a098..4bcc5d173 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -580,6 +580,7 @@ add_library(vllm STATIC src/vllm/model_executor/models/minimax_h3_device.cpp src/vllm/model_executor/models/minimax_h3_vae_loader.cpp src/vllm/model_executor/models/minimax_h3_encoder_gguf.cpp + src/vllm/model_executor/models/minimax_h3_encoder_sharded.cpp src/vllm/model_executor/models/minimax_h3_encoder_device.cpp src/vllm/model_executor/models/minimax_h3_vision_gguf.cpp src/vllm/entrypoints/openai/video_api.cpp diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index bf6dc9721..c827a51fd 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -313,7 +313,7 @@ built on it rather than keeping the flattering one. | Qwen3-dense decode CUDA-graph | Token-exact pass, ~4.3% e2e directional | Steady-state per-step tok/s | | Kimi-Linear-48B-A3B (KDA+MLA+MoE) | e2e RUNS (bf16-resident §13); KDA device op `vt::KdaGatedDeltaRule` GB10 **106→122/128 + 4.24 tok/s (3.1×)**, NOT STRICT, default OFF; NoPE-MLA `VT_KIMI_DEVICE_MLA` NEG 122→109 (§16); chunk_kda prefill AOT SPIKED (§17) | device-KDA = vLLM's actual GPU recurrence; MLA vt::Attention f32-softmax ≠ FA2 order coin-flips (§16). p7 near-tie; STRICT residual = chunk_kda prefill (5-kernel AOT recipe §17) + paged FA2 MLA + incremental | | vLLM 0.26 re-benchmark | Pending | Re-run the binding grids on the advanced pin | -| 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. bf16-vs-quant A/B UNBLOCKED (the 13-shard bf16 DiT now STREAMS to device) but NOT MEASURED: no bf16 render exists. Detail: benchmark-record + spec §8 | +| 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. bf16-vs-quant A/B: ENCODER half MEASURED (§8.15), DiT half NOT (no bf16 render exists). 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) | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 82cc1d7ac..dbf7c28e6 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -133,7 +133,7 @@ they sit outside the gated list above. |---|---|---|---| | Voxtral audio (`VoxtralForConditionalGeneration`) | Voxtral-Mini-3B-2507 | near-tie-robust 16/16 vs vLLM 0.25.0 | decode 0.97x (beats vLLM); encoder TTFT ~17x, pending | | Whisper audio encoder | openai/whisper-small; whisper-large-v3 (Voxtral cfg) | encoder tower 77/77; large-v3 tower 203/203 | pending | -| MiniMax-H3 DiT (`MiniMaxH3DiTModel`, vllm-omni lane) | MiniMax-H3 (33.1B video+audio) | portable 72/72; t2va+fl2va COHERENT; ref2va NVFP4 grid = the community checkpoint's own quant fidelity, NO loader bug (§8.12); loads GGUF + NVFP4, STREAMS the bf16 13-shard release to device | FP4/Marlin landed; ref2va NVFP4 render blocked on checkpoint quant (needs official modelopt NVFP4), speed pending; no bf16 render yet | +| MiniMax-H3 DiT (`MiniMaxH3DiTModel`, vllm-omni lane) | MiniMax-H3 (33.1B video+audio) | portable 72/72; t2va+fl2va COHERENT; ref2va NVFP4 grid = the community checkpoint's own quant fidelity, NO loader bug (§8.12); loads GGUF + NVFP4, STREAMS the bf16 13-shard DiT and the 14-shard bf16 text encoder | FP4/Marlin landed; ref2va NVFP4 render blocked on checkpoint quant (needs official modelopt NVFP4), speed pending; no bf16 render yet | | MTP speculator | Qwen3.6-27B, Qwen3.6-35B-A3B | token-identical to vLLM `mtp` at c1 | ~4% faster c1; +16% output tput (MoE) | | DFlash block-diffusion | Qwen3 (DFlash draft) | near-tie e2e 27/27 vs vLLM | 2.9x over spec-off, 1.003x vs vLLM DFlash-on | | DeepSeek-V4 MTP | DeepSeek-V4-Flash (nextn head) | lossless 5/5; real-model weight-blocked | pending | @@ -161,7 +161,7 @@ model architecture is wired. | Image | ✅ correctness-gated | ✅ | ✅ | ◐ | | Video | ✅ correctness-gated | ✅ | ✅ | ☐ | | Audio | ✅ correctness-gated | ✅ | ◐ | ◐ | -| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va+fl2va COHERENT on GB10; ref2va NVFP4 grid = the community checkpoint's own quant fidelity, NO loader bug (§8.12); DiT loads GGUF or NVFP4, and streams the bf16 13-shard release to device | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | +| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ t2va+fl2va COHERENT on GB10; ref2va NVFP4 grid = the community ckpt's own quant fidelity, NO loader bug (§8.12); DiT loads GGUF/NVFP4/bf16-13-shard, encoder loads GGUF or bf16-14-shard | ✅ (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 ca1c3895a..308b82229 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 +// # a DIRECTORY of the original bf16 encoder release's shards +// # plus model.safetensors.index.json is accepted wherever the +// # Q4_K_M GGUF is. +// --prompt --tokenizer +// --save-embeds [--encoder-max-layers N] +// // PROMPT EMBEDDINGS are taken as a file rather than computed here, deliberately: // the encoder tower needs a tokenizer + a 32B forward, which is its own driver. // This keeps the assembly question ("do the checkpoints compose into a video?") -// separable from the encoding question. +// separable from the encoding question. `--encoder-only` is the other side of +// that seam: run the tower alone, write the conditioning, exit — no DiT, no VAEs. +// On a 122 GiB UNIFIED pool that is what makes encoding the same prompt with two +// different encoders affordable, which is how "what does quantizing the encoder +// cost?" gets a number instead of an opinion. #include #include @@ -193,13 +205,126 @@ std::string Need(int argc, char** argv, int i, const std::string& flag) { return argv[i]; } +// Encode `prompt` with the H3 text encoder; returns the [seq, hidden] f32 +// conditioning the DiT consumes. +// +// `encoder_path` is EITHER a ComfyUI-format GGUF (the shipped Q4_K_M tower) or a +// DIRECTORY holding the ORIGINAL bf16 release's safetensors shards plus +// model.safetensors.index.json. Everything after the weight bytes is shared: same +// tokenizer, same text-only M-RoPE positions, same +// MiniMaxH3EncoderTextForwardDevice, same f32 activations. That is exactly what +// makes "how much does quantizing the encoder change the conditioning?" a +// measurable question rather than an opinion — run both, diff the output. +// +// TEXT-ONLY, on purpose. This backs `--encoder-only`, which exists to produce a +// conditioning tensor for that A/B and nothing else. The normal run path keeps its +// own inline encoder block because that one also carries the VISION path +// (`--cond-image`: merged features masked_scatter'd into inputs_embeds plus the 3 +// DeepStack taps), which this helper deliberately does not duplicate. +std::vector EncodeH3Prompt(const std::string& encoder_path, const std::string& prompt, + const std::string& tokenizer_path, int64_t encoder_max_layers, + const std::string& device_name, int64_t* out_seq, + int64_t* out_hidden) { + const bool sharded = vllm::MiniMaxH3ShardedCheckpoint::IsShardedDir(encoder_path); + std::cerr << "loading encoder " << encoder_path + << (sharded ? " (bf16 shards)" : " (keep-quant GGUF)") << "\n"; + + // The queue is created BEFORE the weights are read: on a unified-memory box the + // CUDA context must exist first, or the driver's reservation lands on top of a + // pool the weights already filled. + vt::Device enc_dev{}; + if (device_name == "cuda") { + enc_dev = vt::GetBackend(vt::DeviceType::kCUDA).CreateQueue().device; + } + vt::Queue eq{enc_dev, nullptr}; + vt::Backend& eb = vt::GetBackend(enc_dev.type); + if (enc_dev.type != vt::DeviceType::kCPU) eq = eb.CreateQueue(); + + vllm::MiniMaxH3EncoderConfig ec; + std::vector ids; + std::vector embeds; + vllm::MiniMaxH3EncoderDeviceWeights staged; + const auto t0 = std::chrono::steady_clock::now(); + + if (sharded) { + const vllm::MiniMaxH3ShardedCheckpoint ckpt = + vllm::MiniMaxH3ShardedCheckpoint::Open(encoder_path); + std::cerr << " shards=" << ckpt.ShardCount() << " tensors=" << ckpt.Names().size() << "\n"; + ec = vllm::MiniMaxH3EncoderConfigFromShards(ckpt, encoder_max_layers); + if (tokenizer_path.empty()) { + throw std::runtime_error("--tokenizer is required with a safetensors-shard --encoder"); + } + const vllm::tok::Tokenizer tokenizer = vllm::tok::Tokenizer::FromHfJson(tokenizer_path); + ids = tokenizer.Encode(prompt); + VT_CHECK(!ids.empty(), "minimax-h3-gen: the prompt tokenized to nothing"); + // Gathered straight out of the mmap'd shard: the table is [151936, hidden] and + // the prompt touches a few dozen rows, so nothing is materialized. + embeds = vllm::MiniMaxH3EncoderEmbedTokensFromShards(ckpt, ids); + staged = vllm::StreamMiniMaxH3EncoderShardsToDevice(eq, ckpt, encoder_max_layers, &ec); + const vllm::MiniMaxH3EncoderShardStreamStats st = vllm::GetMiniMaxH3EncoderShardStreamStats(); + std::cerr << " streamed bf16 encoder -> device: layers=" << st.layers_streamed + << " tensors=" << st.tensors_streamed << " direct=" << st.direct_uploads + << " converted=" << st.converted_uploads << " fused=" << st.fused_groups + << " uploaded=" << (st.bytes_uploaded / (1024.0 * 1024.0 * 1024.0)) + << " GiB host_peak=" << (st.host_peak_bytes / (1024.0 * 1024.0)) << " MiB\n"; + } else { + const vllm::GgufFile ef = vllm::GgufFile::Open(encoder_path); + const vllm::MiniMaxH3EncoderQuantWeights enc = + vllm::LoadMiniMaxH3EncoderFromGguf(ef, encoder_max_layers); + ec = enc.config; + size_t quant_bytes = 0; + for (const auto& kv : enc.quant_storage) quant_bytes += kv.second.size(); + std::cerr << " encoder resident (keep-quant) = " + << (quant_bytes / (1024.0 * 1024.0 * 1024.0)) << " GiB\n"; + // The ComfyUI-style encoder GGUF is WEIGHTS ONLY — it carries no + // `tokenizer.ggml.*` metadata, unlike a llama.cpp export — so the vocab comes + // from the checkpoint's own tokenizer.json. + const vllm::tok::Tokenizer tokenizer = tokenizer_path.empty() + ? vllm::tok::Tokenizer::FromGguf(ef) + : vllm::tok::Tokenizer::FromHfJson(tokenizer_path); + ids = tokenizer.Encode(prompt); + VT_CHECK(!ids.empty(), "minimax-h3-gen: the prompt tokenized to nothing"); + embeds = vllm::MiniMaxH3EncoderEmbedTokens(enc, ids); + staged = vllm::StageMiniMaxH3EncoderWeights(eq, enc); + } + + std::cerr << " encoder layers=" << ec.num_hidden_layers << " hidden=" << ec.hidden_size + << " heads=" << ec.num_attention_heads << " kv_heads=" << ec.num_key_value_heads + << " head_dim=" << ec.head_dim << " ffn=" << ec.intermediate_size << "\n"; + std::cerr << " prompt tokens = " << ids.size() << " (load " + << std::chrono::duration(std::chrono::steady_clock::now() - t0).count() + << " s)\n"; + + // Text-only: all three M-RoPE axes are the token index. + const int64_t seq = static_cast(ids.size()); + std::vector pos(static_cast(3 * seq)); + for (int64_t a = 0; a < 3; ++a) { + for (int64_t s = 0; s < seq; ++s) pos[static_cast(a * seq + s)] = s; + } + std::cerr << " encoding prompt...\n"; + std::vector conditioning = + vllm::MiniMaxH3EncoderTextForwardDevice(eq, ec, staged, embeds, pos.data(), seq); + std::cerr << " conditioning = [" << seq << ", " << ec.hidden_size << "]\n"; + if (out_seq != nullptr) *out_seq = seq; + if (out_hidden != nullptr) *out_hidden = ec.hidden_size; + return conditioning; +} + +void WriteEmbeds(const std::string& path, const std::vector& values) { + std::ofstream out(path, std::ios::binary); + if (!out) throw std::runtime_error("cannot write " + path); + out.write(reinterpret_cast(values.data()), + static_cast(values.size() * sizeof(float))); + std::cerr << " saved conditioning -> " << path << "\n"; +} + } // namespace int main(int argc, char** argv) { std::string dit_path, video_vae_path, video_cfg_path, audio_vae_path, audio_cfg_path; std::string embeds_path, out_path, workdir = "/tmp/minimax_h3_gen", ffmpeg = "ffmpeg"; bool keep_quant = false, dry_run = false, dequant_bf16 = false, denoise_only = false; - bool dump_params = false, fp4_resident = false; + bool dump_params = false, fp4_resident = false, encoder_only = false; std::string device_name = "cpu"; std::string encoder_path, prompt, tokenizer_path, save_embeds_path; std::string first_frame_path, last_frame_path; @@ -237,6 +362,7 @@ int main(int argc, char** argv) { else if (f == "--dry-run") dry_run = true; else if (f == "--denoise-only") denoise_only = true; else if (f == "--dump-params") dump_params = true; + else if (f == "--encoder-only") encoder_only = true; else if (f == "--decode-latent") decode_latent_path = Need(argc, argv, ++i, f); else if (f == "--roundtrip") roundtrip_path = Need(argc, argv, ++i, f); else if (f == "--prompt-image") prompt_image_path = Need(argc, argv, ++i, f); @@ -270,12 +396,17 @@ int main(int argc, char** argv) { // --prompt-image runs the vision tower ONLY (from --encoder); it needs no DiT/VAE/out. const bool vision_probe = !prompt_image_path.empty(); const bool diag_vae_only = !decode_latent_path.empty() || !roundtrip_path.empty(); - const bool need_vaes = !denoise_only && !dump_params && !diag_vae_only; - const bool need_cond = !dump_params && !diag_vae_only; + const bool need_vaes = !denoise_only && !dump_params && !encoder_only && !diag_vae_only; + const bool need_cond = !dump_params && !encoder_only && !diag_vae_only; // --decode-latent / --roundtrip / --prompt-image need NO DiT and NO conditioning // (their own blocks validate their inputs); the shared check below would otherwise // reject --dit. - if (!diag_vae_only && !vision_probe && + // --encoder-only needs NO DiT either, and that is the point: the DiT is loaded + // FIRST in the normal path, so asking for conditioning alone used to cost the + // DiT's residency on top of the tower's. On a 122 GiB UNIFIED pool that is the + // difference between ~49 GiB and ~96 GiB peak — i.e. between a run and an OOM + // reboot — and it is faster besides. + if (!diag_vae_only && !vision_probe && !encoder_only && (dit_path.empty() || (need_vaes && (video_vae_path.empty() || audio_vae_path.empty())) || (need_vaes && out_path.empty()) || (need_cond && embeds_path.empty() && (encoder_path.empty() || prompt.empty())))) { @@ -286,10 +417,29 @@ int main(int argc, char** argv) { "[--dry-run] [--denoise-only] [--dump-params] " "[--first-frame f.ppm] [--last-frame f.ppm] [--noise-aug A] " "[--ref-image f.ppm ...] [--ref-video DIR] [--ref-audio f.wav] " - "[--partition fl2va|ref2va]\n"; + "[--partition fl2va|ref2va]\n" + " or: minimax-h3-gen --encoder-only --encoder " + "--prompt --tokenizer --save-embeds " + "[--encoder-max-layers N] [--device cpu|cuda]\n"; return 2; } + // --encoder-only: run the text tower, write its conditioning, exit. No DiT, no + // VAEs, no output path. This is what makes the encoder A/B affordable — and it + // is the tool for "produce conditioning once, reuse it across renders". + if (encoder_only) { + if (encoder_path.empty() || prompt.empty() || save_embeds_path.empty()) { + throw std::runtime_error( + "--encoder-only needs --encoder, --prompt and --save-embeds"); + } + int64_t seq = 0, hidden = 0; + const std::vector conditioning = EncodeH3Prompt( + encoder_path, prompt, tokenizer_path, encoder_max_layers, device_name, &seq, &hidden); + WriteEmbeds(save_embeds_path, conditioning); + std::cout << "tokens=" << seq << "\nhidden=" << hidden << "\n"; + return 0; + } + // --dump-params reads the MANIFEST ONLY -- names and shapes, no payload -- and // prints the geometry those shapes imply. That makes it safe on a checkpoint // whose weights do not fit (the NVFP4 reference loader is ~132 GB of host f32), @@ -616,9 +766,9 @@ int main(int argc, char** argv) { std::cerr << " layers=" << dit.params.num_layers << " hidden=" << dit.params.hidden_size << " heads=" << dit.params.num_attention_heads << "\n"; - // --- 1b. optional encoder probe. Loading the 32B tower keep-quant is the - // precondition for real text conditioning; this reports the geometry it - // recovered so the loader can be validated against the REAL file. --- + // --- 1b. optional encoder run. Same helper --encoder-only uses, so the + // conditioning a render consumes and the conditioning the A/B measures come + // out of ONE code path. --- std::vector encoded_prompt; if (!encoder_path.empty()) { std::cerr << "loading encoder " << encoder_path << " (keep-quant)\n"; diff --git a/include/vllm/model_executor/models/minimax_h3.h b/include/vllm/model_executor/models/minimax_h3.h index 2dbfb8f42..0c4f760ac 100644 --- a/include/vllm/model_executor/models/minimax_h3.h +++ b/include/vllm/model_executor/models/minimax_h3.h @@ -1076,6 +1076,16 @@ MiniMaxH3EncoderDeviceWeights StageMiniMaxH3EncoderWeights( // MERGED-feature masked_scatter into `inputs_embeds` is the CALLER's job (upstream // `_encode` does it on inputs_embeds before the tower runs); this forward consumes // the already-scattered stream, exactly like the host reference. +// +// A BF16 projection is WIDENED to f32 on the device immediately before its GEMM, +// into a scratch buffer reused across layers. That is not a precision choice — +// bf16 -> f32 is EXACT — it is what makes the unquantized (bf16 safetensors) arm +// runnable at all: the activations here are f32 and `vt::MatmulBT` rejects a mixed +// (f32 activation, bf16 weight) pair, while staging the tower as f32 would double +// a 48.8 GiB residency to 97.5 GiB on a 122 GiB UNIFIED pool. Widening per layer +// costs ONE layer's worth of scratch (~2 GiB) and leaves the GEMM inputs +// bit-identical to what an f32-staged tower would have fed it — so the Q4_K_M and +// bf16 arms differ in their WEIGHT BYTES and nothing else. std::vector MiniMaxH3EncoderTextForwardDevice( vt::Queue& queue, const MiniMaxH3EncoderConfig& config, const MiniMaxH3EncoderDeviceWeights& weights, const std::vector& inputs_embeds, @@ -1114,6 +1124,85 @@ multimodal::Qwen3VLVisionConfig MiniMaxH3EncoderVisionConfig(); multimodal::Qwen3VLVisionWeights LoadQwen3VLVisionFromGguf( const GgufFile& file, const multimodal::Qwen3VLVisionConfig& cfg); +// --------------------------------------------------------------------------- +// H3-Encoder from the ORIGINAL bf16 release — 14 safetensors shards, 63 GB +// (minimax_h3_encoder_sharded.cpp) +// +// WHY: every H3 render so far conditioned on a Q4_K_M encoder, and nobody had +// measured what that quantization does to the conditioning tensor. Answering it +// needs the SAME prompt encoded by the unquantized tower, which ships as 14 +// shards — and `--encoder` only ever accepted a GGUF. +// +// This is a LOADER, not a second forward: it fills the same `views` map +// `MiniMaxH3EncoderDeviceWeights` already binds, over bf16 data instead of ggml +// blocks, and `MiniMaxH3EncoderTextForwardDevice` runs unchanged. +// +// The name map is the one already gated for `LoadMiniMaxH3EncoderWeights`: +// `model.language_model.layers.N.` -> `layers.N.`, with q/k/v and gate/up FUSED +// by row concatenation ([q|k|v], [gate|up]) because the forward slices them that +// way. `model.language_model.norm.weight` and `lm_head.weight` are deliberately +// NOT bound — H3 reads the UNNORMALIZED truncated output. +// --------------------------------------------------------------------------- +class MiniMaxH3ShardedCheckpoint; + +// The encoder geometry implied by the shard index's SHAPES alone — no payload is +// read, so this is safe on a checkpoint far larger than RAM and is the answer to +// "do the GGUF and bf16 arms agree on geometry?" without loading either. +// `max_layers` truncates the text tower exactly as the GGUF loader's does. +MiniMaxH3EncoderConfig MiniMaxH3EncoderConfigFromShards(const MiniMaxH3ShardedCheckpoint& ckpt, + int64_t max_layers = 0); + +// Gather `ids`' embedding rows STRAIGHT out of the mmap'd shard holding +// `model.language_model.embed_tokens.weight`. The table is [151936, 5120] — 1.6 GB +// even in bf16 — and a prompt touches a few dozen rows, so nothing is materialized: +// this is the safetensors twin of MiniMaxH3EncoderEmbedTokens' per-row dequantize. +// Returns [ids.size(), hidden] f32. +std::vector MiniMaxH3EncoderEmbedTokensFromShards(const MiniMaxH3ShardedCheckpoint& ckpt, + const std::vector& ids); + +// "This loader actually RAN" counters. A green suite over a path that silently +// fell back to the GGUF loader is a failure mode this codebase has hit before, so +// the streamer is OBSERVABLE and the gate asserts on it. Mirrors +// MiniMaxH3ShardStreamStats. +struct MiniMaxH3EncoderShardStreamStats { + uint64_t shards_opened = 0; // shards the checkpoint resolved to + uint64_t layers_streamed = 0; // text-tower layers bound + uint64_t tensors_streamed = 0; // device views produced + uint64_t direct_uploads = 0; // uploaded straight from the mmap, NO host copy + uint64_t converted_uploads = 0; // needed one host dtype conversion first + uint64_t fused_groups = 0; // qkv / gate_up concatenations done ON DEVICE + uint64_t bytes_uploaded = 0; // total device bytes staged + uint64_t host_peak_bytes = 0; // largest host conversion buffer alive at once +}; + +inline MiniMaxH3EncoderShardStreamStats& MutableMiniMaxH3EncoderShardStreamStats() { + static MiniMaxH3EncoderShardStreamStats s; + return s; +} +inline MiniMaxH3EncoderShardStreamStats GetMiniMaxH3EncoderShardStreamStats() { + return MutableMiniMaxH3EncoderShardStreamStats(); +} +inline void ResetMiniMaxH3EncoderShardStreamStats() { + MutableMiniMaxH3EncoderShardStreamStats() = MiniMaxH3EncoderShardStreamStats{}; +} + +// ★ Stream the bf16 tower STRAIGHT ONTO THE DEVICE, one tensor at a time. +// +// It MUST stream. The box has 122 GiB of UNIFIED memory (host and device share +// ONE pool) and a previous non-streaming H3 loader was OOM-KILLED at anon-rss +// 125 GB. Here the projections stay BF16 on the device (~48.8 GiB for the 50 +// layers H3 actually runs, against 97.5 GiB as f32) and are uploaded DIRECTLY out +// of the read-only mmap — a bf16 shard tensor bound for a bf16 device slot needs +// no host buffer at all. Only the norms are widened, and those are [5120] each. +// Each source range goes to MaybeReleaseSourcePages the moment its copy returns. +// +// The FUSIONS are done on the DEVICE: one allocation per fused group, with q, k +// and v uploaded into its row offsets. That keeps the "no host copy" property +// through the one transform this loader performs. +MiniMaxH3EncoderDeviceWeights StreamMiniMaxH3EncoderShardsToDevice( + vt::Queue& queue, const MiniMaxH3ShardedCheckpoint& ckpt, int64_t max_layers = 0, + MiniMaxH3EncoderConfig* out_config = nullptr); + // Materialize the H3-Encoder (FL2VA/text_encoder, 14 shards / 1058 tensors) into // the name map both encoder forwards read. // diff --git a/src/vllm/model_executor/models/minimax_h3_encoder_device.cpp b/src/vllm/model_executor/models/minimax_h3_encoder_device.cpp index 00ef5ce0b..8824c4672 100644 --- a/src/vllm/model_executor/models/minimax_h3_encoder_device.cpp +++ b/src/vllm/model_executor/models/minimax_h3_encoder_device.cpp @@ -21,6 +21,7 @@ #include #include +#include #include #include #include @@ -136,6 +137,35 @@ std::vector MiniMaxH3EncoderTextForwardDevice( vt::RmsNormArgs norm_args; norm_args.eps = static_cast(config.rms_norm_eps); + // GEMM WEIGHT ACCESS. A block-quant weight (the GGUF arm) goes to MatmulBT + // untouched — it dispatches kMatmulBTQuant, which takes the f32 activation as + // it is. A BF16 weight (the unquantized safetensors arm) cannot: MatmulBT + // requires BOTH operands in the same dtype and these activations are f32. + // + // So a bf16 weight is WIDENED here, immediately before its GEMM, into a scratch + // buffer keyed by element count and reused across all 50 layers. bf16 -> f32 is + // EXACT, so the GEMM sees bit-identical inputs to what an f32-staged tower would + // have given it — the widening is a residency trick, not a numerics one. It has + // to be: the 50 layers H3 runs are 48.8 GiB in bf16 and 97.5 GiB in f32, and the + // pool is 122 GiB shared with the host. Peak cost is ONE layer's projections + // (~2 GiB), not the model. + std::map widen_scratch; + auto weight = [&](const std::string& name) -> Tensor { + const Tensor& w = weights.Get(name); + if (w.dtype != DType::kBF16) return w; + const int64_t numel = w.Numel(); + auto it = widen_scratch.find(numel); + if (it == widen_scratch.end()) { + it = widen_scratch.emplace(numel, DBuf(d, DType::kF32, {numel})).first; + } + // Same allocation every layer, and the casts and GEMMs are enqueued on ONE + // stream, so layer L's GEMM has consumed it before layer L+1's cast writes it. + Tensor flat_src = dense_attn::Reshape(w, {numel}); + vt::CastF32(d.q, it->second.t(), flat_src); + return dense_attn::Reshape(it->second.t(), + std::vector(w.shape, w.shape + w.rank)); + }; + const int64_t num_layers = MiniMaxH3EncoderNumLayers(config.num_hidden_layers, config.selected_layer); for (int64_t layer = 0; layer < num_layers; ++layer) { @@ -148,15 +178,15 @@ std::vector MiniMaxH3EncoderTextForwardDevice( if (weights.Has(p + "self_attn.qkv_proj.weight")) { // Uniform-encoding checkpoint: one GEMM then split. DBuf qkv(d, DType::kF32, {seq, q_width + 2 * kv_width}); - vt::MatmulBT(d.q, qkv.t(), normed.t(), weights.Get(p + "self_attn.qkv_proj.weight")); + vt::MatmulBT(d.q, qkv.t(), normed.t(), weight(p + "self_attn.qkv_proj.weight")); vt::QkvSplit(d.q, q.t(), k.t(), v.t(), qkv.t()); } else { // MIXED-encoding checkpoint (the shipped Q4_K_M keeps v_proj at Q6_K), so the // group was never fused: three GEMMs, each in its own encoding. Costs launches, // not precision. - vt::MatmulBT(d.q, q.t(), normed.t(), weights.Get(p + "self_attn.q_proj.weight")); - vt::MatmulBT(d.q, k.t(), normed.t(), weights.Get(p + "self_attn.k_proj.weight")); - vt::MatmulBT(d.q, v.t(), normed.t(), weights.Get(p + "self_attn.v_proj.weight")); + vt::MatmulBT(d.q, q.t(), normed.t(), weight(p + "self_attn.q_proj.weight")); + vt::MatmulBT(d.q, k.t(), normed.t(), weight(p + "self_attn.k_proj.weight")); + vt::MatmulBT(d.q, v.t(), normed.t(), weight(p + "self_attn.v_proj.weight")); } // Per-head q/k RMSNorm over head_dim, THEN RoPE — that order is upstream's. @@ -185,7 +215,7 @@ std::vector MiniMaxH3EncoderTextForwardDevice( vt::DFlashBlockAttention(d.q, attn.t(), q3, k3, tv, args); Tensor flat = dense_attn::Reshape(attn.t(), {seq, q_width}); - vt::MatmulBT(d.q, attn_out.t(), flat, weights.Get(p + "self_attn.o_proj.weight")); + vt::MatmulBT(d.q, attn_out.t(), flat, weight(p + "self_attn.o_proj.weight")); vt::Add(d.q, h.t(), h.t(), attn_out.t()); vt::RmsNorm(d.q, normed.t(), h.t(), weights.Get(p + "post_attention_layernorm.weight"), @@ -194,13 +224,13 @@ std::vector MiniMaxH3EncoderTextForwardDevice( DBuf act(d, DType::kF32, {seq, ffn}); if (weights.Has(p + "mlp.gate_up_proj.weight")) { DBuf gate_up(d, DType::kF32, {seq, 2 * ffn}); - vt::MatmulBT(d.q, gate_up.t(), normed.t(), weights.Get(p + "mlp.gate_up_proj.weight")); + vt::MatmulBT(d.q, gate_up.t(), normed.t(), weight(p + "mlp.gate_up_proj.weight")); vt::SiluAndMul(d.q, act.t(), gate_up.t()); } else { DBuf gate(d, DType::kF32, {seq, ffn}); DBuf up(d, DType::kF32, {seq, ffn}); - vt::MatmulBT(d.q, gate.t(), normed.t(), weights.Get(p + "mlp.gate_proj.weight")); - vt::MatmulBT(d.q, up.t(), normed.t(), weights.Get(p + "mlp.up_proj.weight")); + vt::MatmulBT(d.q, gate.t(), normed.t(), weight(p + "mlp.gate_proj.weight")); + vt::MatmulBT(d.q, up.t(), normed.t(), weight(p + "mlp.up_proj.weight")); // SiluAndMul wants [gate | up] contiguous, so stage the pair once. DBuf gate_up(d, DType::kF32, {seq, 2 * ffn}); for (int64_t r = 0; r < seq; ++r) { @@ -213,7 +243,7 @@ std::vector MiniMaxH3EncoderTextForwardDevice( } vt::SiluAndMul(d.q, act.t(), gate_up.t()); } - vt::MatmulBT(d.q, attn_out.t(), act.t(), weights.Get(p + "mlp.down_proj.weight")); + vt::MatmulBT(d.q, attn_out.t(), act.t(), weight(p + "mlp.down_proj.weight")); vt::Add(d.q, h.t(), h.t(), attn_out.t()); // DeepStack: ADD the visual features into the visual-token rows, for the FIRST diff --git a/src/vllm/model_executor/models/minimax_h3_encoder_sharded.cpp b/src/vllm/model_executor/models/minimax_h3_encoder_sharded.cpp new file mode 100644 index 000000000..3798d939b --- /dev/null +++ b/src/vllm/model_executor/models/minimax_h3_encoder_sharded.cpp @@ -0,0 +1,294 @@ +// H3-Encoder — the ORIGINAL bf16 Qwen3-VL-32B tower, 14 safetensors shards, 63 GB. +// +// WHY THIS EXISTS. Every H3 render so far conditioned on the Q4_K_M encoder, and +// nobody had ever measured what that quantization does to the conditioning tensor +// the DiT actually consumes. The question needs the SAME prompt encoded by the +// UNQUANTIZED tower — and `--encoder` only ever accepted a single GGUF, while the +// bf16 release ships 14 shards. This file is the missing half. +// +// It is a LOADER, not a second forward. `MiniMaxH3EncoderDeviceWeights` binds a +// plain `std::map`; the GGUF arm fills it with ggml +// blocks, this one fills it with bf16, and `MiniMaxH3EncoderTextForwardDevice` +// runs unchanged over either. Everything downstream of the weight bytes — the +// M-RoPE, the per-head q/k norms, the causal GQA attention, the f32 activations, +// the truncation to min(num_hidden_layers, 50), the UNNORMALIZED output — is +// literally the same code, which is what makes the two arms comparable at all. +// +// THE NAME MAP IS NOT RE-DERIVED. It is the one already gated for the in-tree +// `LoadMiniMaxH3EncoderWeights(const std::vector&, ...)`: +// model.language_model.layers.N.* -> layers.N.* +// self_attn.{q,k,v}_proj -> self_attn.qkv_proj ([q|k|v] rows) +// mlp.{gate,up}_proj -> mlp.gate_up_proj ([gate|up] rows) +// `model.language_model.norm.weight` and `lm_head.weight` are deliberately NOT +// bound: H3 reads the UNNORMALIZED layer-49 output, and binding the final norm +// would imply it is applied. Shard resolution goes through +// `MiniMaxH3ShardedCheckpoint`, i.e. the checkpoint's own index, never a scan. +// +// MEMORY IS THE DESIGN CONSTRAINT. The box has 122 GiB of UNIFIED memory (host +// and device share ONE pool) and a previous non-streaming H3 loader was +// OOM-KILLED at anon-rss 125 GB. So: the projections stay BF16 on the device +// (48.8 GiB for the 50 layers H3 runs, against 97.5 GiB as f32) and are uploaded +// DIRECTLY out of the read-only mmap with no host buffer at all; the row +// concatenations are done ON THE DEVICE by uploading each member into its offset +// of one allocation; and every source range is released the moment its copy +// returns. Only the norms — [5120] each — are widened on the host, because +// vt::RmsNorm takes an f32 weight. +#include +#include +#include +#include +#include +#include + +#include "vllm/model_executor/model_loader/safetensors_reader.h" +#include "vllm/model_executor/models/dense_device_glue.h" +#include "vllm/model_executor/models/minimax_h3.h" +#include "vt/backend.h" +#include "vt/dtype.h" + +namespace vllm { +namespace { + +const char* const kLmPrefix = "model.language_model."; + +// One destination view and the safetensors tensors that make it. `srcs.size() > 1` +// means a ROW CONCATENATION, and the order is load-bearing: the forward slices +// qkv_proj at [0,q) / [q,q+kv) / [q+kv,...), so any other order silently feeds +// keys into the query path — which still runs, and is still wrong. +struct EncoderShardGroup { + std::string dst; + std::vector srcs; + bool norm = false; // f32 on the device (vt::RmsNorm takes an f32 weight) +}; + +std::string LayerSrc(int64_t layer, const std::string& leaf) { + return std::string(kLmPrefix) + "layers." + std::to_string(layer) + "." + leaf; +} + +// The whole name map in ONE place, so the geometry deriver, the streamer and the +// gate all read the same plan. `max_layers` truncates the text tower, which is +// H3's own behaviour (min(num_hidden_layers, 50) — the release ships 64); 0 keeps +// every layer the index names. +std::vector PlanEncoderShardGroups(const MiniMaxH3ShardedCheckpoint& ckpt, + int64_t max_layers, int64_t* out_layers) { + std::vector plan; + int64_t layers = 0; + for (int64_t layer = 0;; ++layer) { + if (!ckpt.Has(LayerSrc(layer, "input_layernorm.weight"))) break; + if (max_layers > 0 && layer >= max_layers) break; + const std::string dst = "layers." + std::to_string(layer) + "."; + plan.push_back({dst + "input_layernorm.weight", {LayerSrc(layer, "input_layernorm.weight")}, + true}); + plan.push_back({dst + "post_attention_layernorm.weight", + {LayerSrc(layer, "post_attention_layernorm.weight")}, true}); + plan.push_back({dst + "self_attn.q_norm.weight", {LayerSrc(layer, "self_attn.q_norm.weight")}, + true}); + plan.push_back({dst + "self_attn.k_norm.weight", {LayerSrc(layer, "self_attn.k_norm.weight")}, + true}); + plan.push_back({dst + "self_attn.qkv_proj.weight", + {LayerSrc(layer, "self_attn.q_proj.weight"), + LayerSrc(layer, "self_attn.k_proj.weight"), + LayerSrc(layer, "self_attn.v_proj.weight")}, + false}); + plan.push_back({dst + "self_attn.o_proj.weight", {LayerSrc(layer, "self_attn.o_proj.weight")}, + false}); + plan.push_back({dst + "mlp.gate_up_proj.weight", + {LayerSrc(layer, "mlp.gate_proj.weight"), LayerSrc(layer, "mlp.up_proj.weight")}, + false}); + plan.push_back({dst + "mlp.down_proj.weight", {LayerSrc(layer, "mlp.down_proj.weight")}, false}); + ++layers; + } + VT_CHECK(layers > 0, + "minimax_h3 encoder shards: no text-tower layers were found (expected " + "model.language_model.layers.0.input_layernorm.weight)"); + if (out_layers != nullptr) *out_layers = layers; + return plan; +} + +// A rank-2 [rows, K] projection's shape, checked. +void ProjShape(const MiniMaxH3ShardedCheckpoint& ckpt, const std::string& name, int64_t* rows, + int64_t* k) { + const StTensor& t = ckpt.Get(name); + VT_CHECK(t.shape.size() == 2, + "minimax_h3 encoder shards: '" + name + "' is not a rank-2 projection"); + *rows = t.shape[0]; + *k = t.shape[1]; +} + +} // namespace + +MiniMaxH3EncoderConfig MiniMaxH3EncoderConfigFromShards(const MiniMaxH3ShardedCheckpoint& ckpt, + int64_t max_layers) { + int64_t layers = 0; + (void)PlanEncoderShardGroups(ckpt, max_layers, &layers); + + MiniMaxH3EncoderConfig config; + config.num_hidden_layers = layers; + + // head_dim comes from q_norm, which is [head_dim] — the same recovery the GGUF + // loader does, so the two arms cannot disagree on geometry by construction. + const StTensor& qn = ckpt.Get(LayerSrc(0, "self_attn.q_norm.weight")); + VT_CHECK(qn.shape.size() == 1, "minimax_h3 encoder shards: q_norm must be rank-1 [head_dim]"); + config.head_dim = qn.shape[0]; + + int64_t q_rows = 0, k_rows = 0, hidden = 0, kv_k = 0; + ProjShape(ckpt, LayerSrc(0, "self_attn.q_proj.weight"), &q_rows, &hidden); + ProjShape(ckpt, LayerSrc(0, "self_attn.k_proj.weight"), &k_rows, &kv_k); + VT_CHECK(kv_k == hidden, "minimax_h3 encoder shards: q_proj and k_proj disagree on K"); + config.hidden_size = hidden; + VT_CHECK(config.head_dim > 0 && q_rows % config.head_dim == 0 && k_rows % config.head_dim == 0, + "minimax_h3 encoder shards: projection rows are not a multiple of head_dim"); + config.num_attention_heads = q_rows / config.head_dim; + config.num_key_value_heads = k_rows / config.head_dim; + + int64_t ffn = 0, gate_k = 0; + ProjShape(ckpt, LayerSrc(0, "mlp.gate_proj.weight"), &ffn, &gate_k); + VT_CHECK(gate_k == hidden, "minimax_h3 encoder shards: gate_proj K is not hidden_size"); + config.intermediate_size = ffn; + return config; +} + +std::vector MiniMaxH3EncoderEmbedTokensFromShards(const MiniMaxH3ShardedCheckpoint& ckpt, + const std::vector& ids) { + const std::string name = std::string(kLmPrefix) + "embed_tokens.weight"; + VT_CHECK(ckpt.Has(name), + "minimax_h3 encoder shards: checkpoint has no " + name); + const StTensor& table = ckpt.Get(name); + VT_CHECK(table.shape.size() == 2, + "minimax_h3 encoder shards: embed_tokens must be [vocab, hidden]"); + const int64_t vocab = table.shape[0], hidden = table.shape[1]; + + VT_CHECK(vocab > 0 && table.nbytes % static_cast(vocab) == 0, + "minimax_h3 encoder shards: embed_tokens byte span does not divide by its rows"); + const size_t row_bytes = table.nbytes / static_cast(vocab); + + // ROW AT A TIME out of the mmap. The table is the single largest tensor in the + // checkpoint ([151936, 5120]) and a prompt touches a few dozen rows, so + // materializing it to gather them would be the expensive way to the same answer. + // Each row is handed to the SHARED, already-gated dtype converter as a one-row + // view rather than getting its own bf16/f16 decode here. + std::vector out(ids.size() * static_cast(hidden)); + for (size_t i = 0; i < ids.size(); ++i) { + const int64_t id = ids[i]; + VT_CHECK(id >= 0 && id < vocab, + "minimax_h3 encoder shards: token id out of vocabulary range"); + StTensor row; + row.dtype = table.dtype; + row.shape = {hidden}; + row.data = table.data + static_cast(id) * row_bytes; + row.nbytes = row_bytes; + const std::vector values = MiniMaxH3ReadSafetensorF32(row); + VT_CHECK(static_cast(values.size()) == hidden, + "minimax_h3 encoder shards: an embedding row decoded to the wrong width"); + std::memcpy(out.data() + i * static_cast(hidden), values.data(), + static_cast(hidden) * sizeof(float)); + } + return out; +} + +MiniMaxH3EncoderDeviceWeights StreamMiniMaxH3EncoderShardsToDevice( + vt::Queue& queue, const MiniMaxH3ShardedCheckpoint& ckpt, int64_t max_layers, + MiniMaxH3EncoderConfig* out_config) { + vt::Backend& backend = vt::GetBackend(queue.device.type); + int64_t layers = 0; + const std::vector plan = PlanEncoderShardGroups(ckpt, max_layers, &layers); + if (out_config != nullptr) *out_config = MiniMaxH3EncoderConfigFromShards(ckpt, max_layers); + + MiniMaxH3EncoderShardStreamStats& stats = MutableMiniMaxH3EncoderShardStreamStats(); + stats = MiniMaxH3EncoderShardStreamStats{}; + stats.shards_opened = static_cast(ckpt.ShardCount()); + stats.layers_streamed = static_cast(layers); + + MiniMaxH3EncoderDeviceWeights out; + for (const EncoderShardGroup& group : plan) { + // Shape of the destination: rows are SUMMED over the group, K is shared. + int64_t rows = 0, k = -1; + bool all_bf16 = true; + std::vector member_rows; + member_rows.reserve(group.srcs.size()); + for (const std::string& src : group.srcs) { + const StTensor& t = ckpt.Get(src); + VT_CHECK(!t.shape.empty(), "minimax_h3 encoder shards: '" + src + "' has no shape"); + const int64_t r = t.shape[0]; + const int64_t kk = t.shape.size() == 2 ? t.shape[1] : -1; + if (k == -1) { + k = kk; + } else { + VT_CHECK(kk == k, "minimax_h3 encoder shards: fused group '" + group.dst + + "' disagrees on K"); + } + if (t.dtype != "BF16") all_bf16 = false; + member_rows.push_back(r); + rows += r; + } + + // Norms go to f32 (vt::RmsNorm's contract). Projections stay BF16 when the + // shards store bf16 — the whole point, since that is the 48.8 GiB residency + // and the no-host-copy upload — and are widened to f32 only if the shards + // store something else, so no precision is ever ROUNDED AWAY by this loader. + const vt::DType dtype = (group.norm || !all_bf16) ? vt::DType::kF32 : vt::DType::kBF16; + const size_t elem = vt::SizeOf(dtype); + const int64_t numel = k >= 0 ? rows * k : rows; + const size_t bytes = static_cast(numel) * elem; + + void* base = backend.Alloc(bytes); + std::shared_ptr owner(base, [&backend](void* p) { backend.Free(p); }); + + size_t offset = 0; + for (size_t m = 0; m < group.srcs.size(); ++m) { + const StTensor& t = ckpt.Get(group.srcs[m]); + const size_t member_numel = + static_cast(member_rows[m]) * static_cast(k >= 0 ? k : 1); + const size_t member_bytes = member_numel * elem; + VT_CHECK(offset + member_bytes <= bytes, + "minimax_h3 encoder shards: '" + group.dst + "' member overruns its allocation"); + // DECLARED OUTSIDE the branch on purpose: vt::Backend::Copy is + // cudaMemcpyAsync, so a conversion buffer that dies at the end of its own + // block is a use-after-free the stream may or may not have read yet. This + // codebase has been bitten by exactly that (uploads from function-local + // temporaries), so the buffer outlives the Synchronize below. + std::vector host; + if (dtype == vt::DType::kBF16) { + // The on-disk bytes ARE the device bytes: no host buffer exists at any + // point, which is what keeps peak host memory flat across a 63 GB load. + VT_CHECK(t.nbytes == member_bytes, + "minimax_h3 encoder shards: '" + group.srcs[m] + + "' byte span does not match its shape"); + backend.Copy(queue, static_cast(base) + offset, t.data, member_bytes); + ++stats.direct_uploads; + } else { + host = MiniMaxH3ReadSafetensorF32(t); + VT_CHECK(host.size() == member_numel, + "minimax_h3 encoder shards: '" + group.srcs[m] + + "' read produced the wrong element count"); + backend.Copy(queue, static_cast(base) + offset, host.data(), member_bytes); + ++stats.converted_uploads; + const size_t host_bytes = host.size() * sizeof(float); + if (host_bytes > stats.host_peak_bytes) stats.host_peak_bytes = host_bytes; + } + // The SOURCE — mmap range or conversion buffer — must stay live until the + // copy has landed, and `host` must not be freed before this returns. + backend.Synchronize(queue); + MaybeReleaseSourcePages(t.data, t.nbytes); + offset += member_bytes; + stats.bytes_uploaded += member_bytes; + } + VT_CHECK(offset == bytes, + "minimax_h3 encoder shards: '" + group.dst + "' did not fill its allocation"); + if (group.srcs.size() > 1) ++stats.fused_groups; + + std::vector shape; + if (k >= 0) { + shape = {rows, k}; + } else { + shape = {rows}; + } + out.views[group.dst] = dense_attn::MakeTensor(base, dtype, queue.device, shape); + out.storage.push_back(std::move(owner)); + ++stats.tensors_streamed; + } + backend.Synchronize(queue); + return out; +} + +} // namespace vllm diff --git a/tests/vllm/models/test_minimax_h3.cpp b/tests/vllm/models/test_minimax_h3.cpp index c8d406307..58e91a0bb 100644 --- a/tests/vllm/models/test_minimax_h3.cpp +++ b/tests/vllm/models/test_minimax_h3.cpp @@ -7395,3 +7395,301 @@ TEST_CASE("minimax_h3: the embedding gather decodes ONLY the rows it needs, exac CHECK_THROWS(vllm::MiniMaxH3EncoderEmbedTokens(w, {VOCAB})); CHECK_THROWS(vllm::MiniMaxH3EncoderEmbedTokens(w, {-1})); } + +// --------------------------------------------------------------------------- +// The ORIGINAL bf16 H3-Encoder: 14 safetensors shards, 63 GB +// --------------------------------------------------------------------------- +// Every H3 render so far conditioned on the Q4_K_M encoder, and nobody had ever +// measured what that quantization does to the conditioning tensor. Asking needs +// the SAME prompt encoded by the unquantized tower — which ships as 14 shards, +// while `--encoder` only ever accepted a GGUF. These gates cover the loader that +// closes that gap: the name map, the two row fusions, the bf16 residency, and — +// the one that has bitten this codebase before — that the new path actually RAN +// rather than silently falling back to the GGUF one. + +namespace { + +// The reduced geometry the encoder-shard gates run at, at the REAL layout: +// q/k/v and gate/up SEPARATE on disk, GQA (heads > kv_heads), head_dim dividing +// both projection row counts. +struct EncShardGeometry { + int64_t layers = 3; + int64_t hidden = 32; + int64_t heads = 4; + int64_t kv_heads = 2; + int64_t head_dim = 8; + int64_t ffn = 48; + int64_t vocab = 24; +}; + +// Round through bf16, so an "F32" copy of a checkpoint holds values the BF16 copy +// can represent EXACTLY. That is what lets the widen-vs-native comparison below +// demand BIT-IDENTICAL outputs rather than a tolerance. +std::vector Bf16RoundTrip(const std::vector& v) { + std::vector out(v.size()); + for (size_t i = 0; i < v.size(); ++i) { + uint32_t bits; + std::memcpy(&bits, &v[i], sizeof(bits)); + const uint32_t rounded = bits + 0x7FFFu + ((bits >> 16) & 1u); + const uint32_t back = rounded & 0xFFFF0000u; + std::memcpy(&out[i], &back, sizeof(back)); + } + return out; +} + +// The REAL encoder release's tensor names, at reduced dims. Deliberately includes +// the three things the loader must NOT bind (`model.language_model.norm.weight`, +// `lm_head.weight`, the whole `model.visual.` tower) — a loader that grabs them +// would still "work", and would apply a final RMSNorm H3 does not. +std::vector BuildEncoderShardEntries(const EncShardGeometry& g, bool bf16) { + std::vector entries; + auto add = [&](const std::string& name, const std::vector& shape) { + int64_t n = 1; + for (int64_t d : shape) n *= d; + const std::vector values = Bf16RoundTrip(MakeParam("encsh." + name, n, 0.2)); + entries.push_back({name, bf16 ? "BF16" : "F32", shape, + bf16 ? PackBf16(values) : PackF32(values)}); + }; + const std::string lm = "model.language_model."; + add(lm + "embed_tokens.weight", {g.vocab, g.hidden}); + for (int64_t l = 0; l < g.layers; ++l) { + const std::string p = lm + "layers." + std::to_string(l) + "."; + add(p + "input_layernorm.weight", {g.hidden}); + add(p + "post_attention_layernorm.weight", {g.hidden}); + add(p + "self_attn.q_norm.weight", {g.head_dim}); + add(p + "self_attn.k_norm.weight", {g.head_dim}); + add(p + "self_attn.q_proj.weight", {g.heads * g.head_dim, g.hidden}); + add(p + "self_attn.k_proj.weight", {g.kv_heads * g.head_dim, g.hidden}); + add(p + "self_attn.v_proj.weight", {g.kv_heads * g.head_dim, g.hidden}); + add(p + "self_attn.o_proj.weight", {g.hidden, g.heads * g.head_dim}); + add(p + "mlp.gate_proj.weight", {g.ffn, g.hidden}); + add(p + "mlp.up_proj.weight", {g.ffn, g.hidden}); + add(p + "mlp.down_proj.weight", {g.hidden, g.ffn}); + } + // MUST NOT be bound: H3 reads the UNNORMALIZED truncated output, never a head. + add(lm + "norm.weight", {g.hidden}); + add("lm_head.weight", {g.vocab, g.hidden}); + // The vision tower shares the release and is not part of a text-only encode. + add("model.visual.blocks.0.attn.qkv.weight", {3 * g.hidden, g.hidden}); + add("model.visual.merger.norm.weight", {g.hidden}); + return entries; +} + +const H3StEntry& FindEntry(const std::vector& entries, const std::string& name) { + for (const H3StEntry& e : entries) { + if (e.name == name) return e; + } + REQUIRE_MESSAGE(false, "synthetic encoder checkpoint has no tensor named " << name); + return entries.front(); +} + +} // namespace + +TEST_CASE("minimax_h3: the bf16 encoder shards resolve, FUSE and STREAM to the device") { + const EncShardGeometry g; + const std::vector entries = BuildEncoderShardEntries(g, /*bf16=*/true); + const std::string dir = "/tmp/minimax_h3_enc_shards"; + const size_t kShards = 4; + const std::map promised = + WriteMiniMaxH3ShardedDit(entries, dir, kShards); + const vllm::MiniMaxH3ShardedCheckpoint ckpt = vllm::MiniMaxH3ShardedCheckpoint::Open(dir); + CHECK(ckpt.ShardCount() == kShards); + CHECK(ckpt.Names().size() == entries.size()); + + // GEOMETRY FROM SHAPES ALONE — no payload read. The recovery rules are the GGUF + // loader's (head_dim from q_norm, heads from q_proj rows), so the two arms + // cannot disagree about what model they are running. + const vllm::MiniMaxH3EncoderConfig cfg = vllm::MiniMaxH3EncoderConfigFromShards(ckpt); + CHECK(cfg.num_hidden_layers == g.layers); + CHECK(cfg.hidden_size == g.hidden); + CHECK(cfg.num_attention_heads == g.heads); + CHECK(cfg.num_key_value_heads == g.kv_heads); + CHECK(cfg.head_dim == g.head_dim); + CHECK(cfg.intermediate_size == g.ffn); + // The knobs the shapes cannot carry keep their defaults, and they are the SAME + // defaults the GGUF arm leaves in place — otherwise an A/B would be comparing + // two different RoPEs, not two quantizations. + CHECK(cfg.selected_layer == vllm::kMiniMaxH3EncoderSelectedLayer); + CHECK(cfg.rope_theta == doctest::Approx(5000000.0)); + CHECK(cfg.mrope_section == std::vector{24, 20, 20}); + + vllm::ResetMiniMaxH3EncoderShardStreamStats(); + vt::Queue q{Cpu(), nullptr}; + vllm::MiniMaxH3EncoderConfig streamed_cfg; + const vllm::MiniMaxH3EncoderDeviceWeights w = + vllm::StreamMiniMaxH3EncoderShardsToDevice(q, ckpt, /*max_layers=*/0, &streamed_cfg); + CHECK(streamed_cfg.num_hidden_layers == g.layers); + CHECK(streamed_cfg.intermediate_size == g.ffn); + + // ★ THE LOADER RAN. A green suite over a path that silently fell back to the + // GGUF loader is a failure mode this codebase has hit; the counters make the + // bf16 path OBSERVABLE and this asserts on them rather than on a comment. + const vllm::MiniMaxH3EncoderShardStreamStats st = + vllm::GetMiniMaxH3EncoderShardStreamStats(); + CHECK(st.shards_opened == kShards); + CHECK(st.layers_streamed == static_cast(g.layers)); + CHECK(st.tensors_streamed == static_cast(8 * g.layers)); + CHECK(st.fused_groups == static_cast(2 * g.layers)); // qkv + gate_up per layer + // Every PROJECTION took the no-host-copy path (mmap -> device); only the four + // norms per layer needed a conversion, and those are [hidden]/[head_dim]. + CHECK(st.direct_uploads == static_cast(7 * g.layers)); + CHECK(st.converted_uploads == static_cast(4 * g.layers)); + CHECK(st.bytes_uploaded > 0); + // Peak host memory is bounded by the largest CONVERSION, i.e. one norm — it + // cannot scale with the model. On the real 63 GB release that is the difference + // between a run and the OOM-kill a previous H3 loader took at anon-rss 125 GB. + CHECK(st.host_peak_bytes == static_cast(g.hidden) * sizeof(float)); + CHECK(st.host_peak_bytes * 16 < st.bytes_uploaded); + + // ★ These views are BF16 — a dtype the GGUF loader can never produce (it binds + // block-quant projections and f32 norms). So this is not the GGUF path wearing + // a different name. + const vt::Tensor& qkv = w.Get("layers.0.self_attn.qkv_proj.weight"); + CHECK(qkv.dtype == vt::DType::kBF16); + CHECK(w.Get("layers.0.mlp.gate_up_proj.weight").dtype == vt::DType::kBF16); + CHECK(w.Get("layers.0.self_attn.o_proj.weight").dtype == vt::DType::kBF16); + CHECK(w.Get("layers.0.mlp.down_proj.weight").dtype == vt::DType::kBF16); + // Norms stay f32: vt::RmsNorm takes an f32 weight. + CHECK(w.Get("layers.0.input_layernorm.weight").dtype == vt::DType::kF32); + CHECK(w.Get("layers.0.self_attn.q_norm.weight").dtype == vt::DType::kF32); + + // SHAPES: the fusions are row concatenations, and the forward slices them. + CHECK(qkv.shape[0] == (g.heads + 2 * g.kv_heads) * g.head_dim); + CHECK(qkv.shape[1] == g.hidden); + const vt::Tensor& gu = w.Get("layers.0.mlp.gate_up_proj.weight"); + CHECK(gu.shape[0] == 2 * g.ffn); + CHECK(gu.shape[1] == g.hidden); + CHECK(w.Get("layers.0.self_attn.q_norm.weight").shape[0] == g.head_dim); + + // ★ AND THE ORDER IS RIGHT, byte for byte. [q|k|v] and [gate|up] are what the + // forward assumes; any other order still runs and silently feeds keys into the + // query path. + for (int64_t l = 0; l < g.layers; ++l) { + const std::string src = "model.language_model.layers." + std::to_string(l) + "."; + const std::string dst = "layers." + std::to_string(l) + "."; + const std::string want_qkv = FindEntry(entries, src + "self_attn.q_proj.weight").bytes + + FindEntry(entries, src + "self_attn.k_proj.weight").bytes + + FindEntry(entries, src + "self_attn.v_proj.weight").bytes; + const vt::Tensor& got_qkv = w.Get(dst + "self_attn.qkv_proj.weight"); + REQUIRE(static_cast(got_qkv.Numel()) * 2 == want_qkv.size()); + CHECK(std::memcmp(got_qkv.data, want_qkv.data(), want_qkv.size()) == 0); + + const std::string want_gu = FindEntry(entries, src + "mlp.gate_proj.weight").bytes + + FindEntry(entries, src + "mlp.up_proj.weight").bytes; + const vt::Tensor& got_gu = w.Get(dst + "mlp.gate_up_proj.weight"); + REQUIRE(static_cast(got_gu.Numel()) * 2 == want_gu.size()); + CHECK(std::memcmp(got_gu.data, want_gu.data(), want_gu.size()) == 0); + + // The unfused projections pass through untouched. + const std::string want_o = FindEntry(entries, src + "self_attn.o_proj.weight").bytes; + const vt::Tensor& got_o = w.Get(dst + "self_attn.o_proj.weight"); + CHECK(std::memcmp(got_o.data, want_o.data(), want_o.size()) == 0); + } + + // The SEPARATE names must be gone — leaving them would let a forward read an + // unfused tensor and take the mixed-encoding branch by accident. + CHECK_FALSE(w.Has("layers.0.self_attn.q_proj.weight")); + CHECK_FALSE(w.Has("layers.0.mlp.gate_proj.weight")); + // H3 deltas: no final norm, no lm_head, no vision tower on a text-only encode. + CHECK_FALSE(w.Has("norm.weight")); + CHECK_FALSE(w.Has("lm_head.weight")); + CHECK_FALSE(w.Has("blocks.0.attn.qkv.weight")); + + // TRUNCATION — H3 runs min(num_hidden_layers, 50) and the release ships 64, so + // the loader must be able to stop early. That is not cosmetic: it is 14 layers + // of a 48.8 GiB residency. + CHECK(w.Has("layers." + std::to_string(g.layers - 1) + ".self_attn.qkv_proj.weight")); + vllm::ResetMiniMaxH3EncoderShardStreamStats(); + const vllm::MiniMaxH3EncoderDeviceWeights trunc = + vllm::StreamMiniMaxH3EncoderShardsToDevice(q, ckpt, /*max_layers=*/1, nullptr); + CHECK(trunc.Has("layers.0.self_attn.qkv_proj.weight")); + CHECK_FALSE(trunc.Has("layers.1.self_attn.qkv_proj.weight")); + CHECK(vllm::GetMiniMaxH3EncoderShardStreamStats().layers_streamed == 1); + CHECK(vllm::MiniMaxH3EncoderConfigFromShards(ckpt, /*max_layers=*/1).num_hidden_layers == 1); + + // The EMBEDDING gather reads rows out of the mmap without materializing the + // table, and must be exact — an off-by-one row is a different prompt. + const std::vector ids = {5, 0, static_cast(g.vocab - 1), 5}; + const std::vector got_rows = vllm::MiniMaxH3EncoderEmbedTokensFromShards(ckpt, ids); + REQUIRE(got_rows.size() == ids.size() * static_cast(g.hidden)); + const std::vector table = + Bf16RoundTrip(MakeParam("encsh.model.language_model.embed_tokens.weight", + g.vocab * g.hidden, 0.2)); + for (size_t i = 0; i < ids.size(); ++i) { + for (int64_t c = 0; c < g.hidden; ++c) { + CHECK(got_rows[i * static_cast(g.hidden) + c] == + table[static_cast(ids[i]) * static_cast(g.hidden) + c]); + } + } + CHECK_THROWS(vllm::MiniMaxH3EncoderEmbedTokensFromShards(ckpt, {static_cast(g.vocab)})); + CHECK_THROWS(vllm::MiniMaxH3EncoderEmbedTokensFromShards(ckpt, {-1})); + + RemoveShardedDit(dir, kShards); +} + +TEST_CASE("minimax_h3: the bf16 encoder runs the SAME forward as an f32-staged tower") { + // THE CLAIM UNDER TEST. The unquantized arm keeps its projections BF16 on the + // device (48.8 GiB for the 50 layers H3 runs; 97.5 GiB as f32, against a 122 GiB + // UNIFIED pool) and widens each one to f32 immediately before its GEMM. That is + // a RESIDENCY trick, and this asserts it is nothing more: bf16 -> f32 is exact, + // so a widened tower and an f32-staged tower holding the same values must agree + // BIT FOR BIT — not to a tolerance. + // + // Without this, "we measured what quantizing the encoder costs" would be + // confounded by whatever the widening itself did. + const EncShardGeometry g; + const int64_t SEQ = 6; + + const std::string bf16_dir = "/tmp/minimax_h3_enc_shards_bf16"; + const std::string f32_dir = "/tmp/minimax_h3_enc_shards_f32"; + const size_t kShards = 3; + const std::vector bf16_entries = BuildEncoderShardEntries(g, /*bf16=*/true); + const std::vector f32_entries = BuildEncoderShardEntries(g, /*bf16=*/false); + WriteMiniMaxH3ShardedDit(bf16_entries, bf16_dir, kShards); + WriteMiniMaxH3ShardedDit(f32_entries, f32_dir, kShards); + + const vllm::MiniMaxH3ShardedCheckpoint bf16_ckpt = + vllm::MiniMaxH3ShardedCheckpoint::Open(bf16_dir); + const vllm::MiniMaxH3ShardedCheckpoint f32_ckpt = + vllm::MiniMaxH3ShardedCheckpoint::Open(f32_dir); + + vt::Queue q{Cpu(), nullptr}; + const vllm::MiniMaxH3EncoderDeviceWeights bf16_w = + vllm::StreamMiniMaxH3EncoderShardsToDevice(q, bf16_ckpt); + const vllm::MiniMaxH3EncoderDeviceWeights f32_w = + vllm::StreamMiniMaxH3EncoderShardsToDevice(q, f32_ckpt); + // The two really are staged differently — otherwise this compares a path to + // itself, which is the trap the counters exist to catch. + CHECK(bf16_w.Get("layers.0.self_attn.qkv_proj.weight").dtype == vt::DType::kBF16); + CHECK(f32_w.Get("layers.0.self_attn.qkv_proj.weight").dtype == vt::DType::kF32); + + vllm::MiniMaxH3EncoderConfig cfg = vllm::MiniMaxH3EncoderConfigFromShards(bf16_ckpt); + cfg.selected_layer = g.layers; + cfg.mrope_section = {2, 1, 1}; + cfg.rope_theta = 10000.0; + + const std::vector embeds = MakeParam("encsh.embeds", SEQ * g.hidden, 1.0); + std::vector pos(static_cast(3 * SEQ)); + for (int64_t a = 0; a < 3; ++a) { + for (int64_t s = 0; s < SEQ; ++s) pos[static_cast(a * SEQ + s)] = s; + } + + const std::vector from_bf16 = + vllm::MiniMaxH3EncoderTextForwardDevice(q, cfg, bf16_w, embeds, pos.data(), SEQ); + const std::vector from_f32 = + vllm::MiniMaxH3EncoderTextForwardDevice(q, cfg, f32_w, embeds, pos.data(), SEQ); + + REQUIRE(from_bf16.size() == static_cast(SEQ * g.hidden)); + REQUIRE(from_f32.size() == from_bf16.size()); + CHECK(std::memcmp(from_bf16.data(), from_f32.data(), from_bf16.size() * sizeof(float)) == 0); + // ...and it produced a tower's output, not zeros. + double mag = 0.0; + for (float v : from_bf16) { + REQUIRE(std::isfinite(v)); + mag = std::max(mag, std::abs(static_cast(v))); + } + CHECK(mag > 1e-3); + + RemoveShardedDit(bf16_dir, kShards); + RemoveShardedDit(f32_dir, kShards); +} From 8d392a80f60c198e8a9e6ad505a44bbe4074c145 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 23:06:42 +0000 Subject: [PATCH 2/2] =?UTF-8?q?bench(minimax-h3):=20THE=20NUMBER=20?= =?UTF-8?q?=E2=80=94=20Q4=5FK=5FM=20encoder=20moves=20the=20conditioning?= =?UTF-8?q?=20as=20much=20as=20a=20one-word=20prompt=20edit,=20but=20DIFFU?= =?UTF-8?q?SELY?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Branch: row/H3-ENC-BF16-COND-DIFF (helper). Records only. Measured on Thor sm_110, CUDA 13.0.1 container, build d1085374 (built and measured as d3861693, amended only for the row-branch trailer; IDENTICAL tree dd9283cf, so the measurement binary IS this commit). GPU idle — the LocalAI render finished on its own and was never touched. CONTROLLED THE WAY IT HAS TO BE: same prompt (wuxia.txt, 233 tokens), same tokenizer, same 50-layer truncation, same MiniMaxH3EncoderTextForwardDevice, same f32 activations — only the weight bytes differ. Both arms self-report IDENTICAL geometry (50 / 5120 / 64 / 8 / 128 / 25600), which is what proves they are the same model rather than two checkpoints sharing a name. A CALIBRATION ARM, because a cosine means nothing without a yardstick: the bf16 encoder also encoded a ONE-WORD edit of the same prompt ("at night" -> "at dawn", also 233 tokens). That is the scale the quantization number is read on. relRMS relRMS(excl sink) cos mean cos med ang med Q4_K_M vs bf16 0.03403 0.06849 0.99745 0.99810 3.535 deg bf16, one-word edit 0.01897 0.06666 0.99769 0.99963 1.565 deg Q4_K_M vs bf16: max|diff| 154.0, RMS 0.5045, cosine min 0.90916 (token 69), max rotation 24.61 deg, 232 of 233 tokens below cosine 0.999. THE READ. 1. NOT a scale change. Q4 conditioning is uniformly ~1% smaller (norm ratio 0.99010) but the best single global rescale removes almost none of the difference (0.03403 -> 0.03280). It is DIRECTIONAL — the kind that matters. 2. Its total energy is ON PAR with a one-word prompt edit (6.85% vs 6.67% excluding the attention-sink token). Quantizing the encoder moves the conditioning about as much as rewriting a word of the prompt. 3. The SHAPE is opposite, and that is the interesting part. The edit is SPARSE: 172 of 233 tokens stay above cosine 0.999 (median rotation 0.16 deg), the change concentrating on ~6 tokens, the largest at 32 deg. Quantization is DIFFUSE: 232 of 233 tokens fall below 0.999, EVERY token rotates a few degrees, one by 24.6 deg. A smear applied everywhere, not a different prompt. 4. max|diff| 154 is the ATTENTION SINK, not corruption. Token 0 has norm 15,522 against a 366 mean (42x) and carries 68% of the total squared error while its DIRECTION is intact (cosine 0.99962) — ComfyUI PR 15298's channel-wise magnitude outliers showing up concretely, and the reason the sink-excluded column is the honest aggregate. VERDICT: Q4_K_M does real, measurable, directional damage to the conditioning, comparable in magnitude to editing the prompt, but diffusely. A uniform few-degree rotation of every token is the signature that blunts fine-grained compositional instruction (coverage, blocking, staging) toward a prompt's average semantics — which is exactly the "competent but generic" symptom. The bf16 encoder is worth a render A/B. EXPLICITLY NOT ESTABLISHED: that the RENDER changes. Nothing here measures the DiT's sensitivity to a 3.5-degree median rotation. The owed follow-up is a byte-identical-everything-else render A/B — same DiT, seed and steps, with --prompt-embeds cond_q4km.bin vs cond_bf16.bin (both saved on the box) — which is exactly what the save/replay seam makes controllable. REPRODUCED: both arms re-run from scratch produced BYTE-IDENTICAL files (cond_q4km.bin md5 a331232096ef1da2628f885950b2fc55, cond_bf16.bin md5 9c096b63b9bd07f604daebb2fc090f46 on both runs), so these are deterministic numbers, not samples. Cost: Q4_K_M arm 40 s / 18.0 GiB peak; bf16 arm 40 s (35 s streaming) / 45.41 GiB uploaded / host conversion peak 0.0195 MiB (ONE norm — the projections never touch a host buffer) / 51.95 GiB total peak of the 122 GiB unified pool. The real-checkpoint streamer counters (layers=50 tensors=400 direct=350 converted=200 fused=100) confirm the shard path RAN and that every projection took the no-host-copy upload. benchmark_binding=false; no throughput claimed. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Opus 5 (1M context) --- .agents/benchmark-record.md | 78 ++++++++++++++++++++++++++++++ .agents/model-matrix.md | 2 +- .agents/parity-ledger.md | 3 +- .agents/roadmap_v1.md | 2 +- .agents/specs/minimax-h3.md | 63 +++++++++++++++++++++++- .agents/state.md | 95 ++++++++++++++++++++++++++++++++----- docs/BENCHMARKS.md | 1 + docs/FEATURES.md | 2 +- docs/STATUS.md | 2 +- 9 files changed, 229 insertions(+), 19 deletions(-) diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index dd2adb85d..b323f04a2 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -14687,3 +14687,81 @@ DELIVERED: 5 harness bodies authored (verbatim FLA ports, AOT-adapted) STAGED in NOT YET (Phase-2, coupled — the harness signatures depend on the op's confirmed buffer dtypes so regen is premature until the op exists): move harness to `triton_kernels/`, add the §17.3 declarations, regen sm_121a cubins (`scripts/regen-triton-aot.sh`), wire `vt::KdaChunkPrefill`, run the RED-first unit + FLA golden + 48.9B STRICT gate + tok/s/TTFT ladder. Row STAYS ACTIVE. No default flips (nothing measured). **USER 2026-08-07 (mid-flight): the Kimi-Linear success bar is MEET vLLM SPEED — the §14/§16/#107 "HW-forced-indirect" framing is SUPERSEDED.** vLLM demonstrably RUNS Kimi-Linear-48B on ONE GB10: the §12 STRICT golden capture used it at `gpu_memory_utilization=0.82`, single-seq, eager. So the Phase-2 speed ladder MUST include a vLLM arm at that EXACT recipe (single-seq, eager, util 0.82, the §12 launch config) measuring steady decode tok/s + prefill TTFT on the SAME prompts as our arm — SEQUENTIAL after our runs, `local-ai-worker` PARKED, `drop_caches` before wall-clock, and PRE-WARM FlashInfer's autotune in a throwaway start at TINY util FIRST (cold autotune at util 0.82 with 91.5 GiB weights = the tightest vLLM config ever run on this box = a recorded OOM-reboot trigger; memory monitor mandatory, ONE attempt, if it OOMs record honestly and do NOT retry higher). The tok/s ladder then reads ours-vs-vLLM matched-config: the lane's distance-to-bar becomes a MEASURED number. Recorded in spec §17.5; below vLLM on any axis is an open gap, not done. +## MiniMax-H3 — what quantizing the TEXT ENCODER to Q4_K_M does to the conditioning (2026-08-06, `row/H3-ENC-BF16-COND-DIFF`, Thor sm_110) + +**The question.** Every H3 render so far conditioned on a Q4_K_M Qwen3-VL-32B text +encoder (`enc_q4km.gguf`, 14.6 GB), and the encoder's contribution had never been +measured. It mattered because weak conditioning and quantization-damaged +conditioning are indistinguishable from outside a render: the wuxia prompt asked +for measured shot/reverse-shot coverage of a martial-arts sect exchanging +intelligence and produced a good but generic portrait. + +**Method.** The SAME prompt (`wuxia.txt`, 233 tokens), the SAME tokenizer, the SAME +50-layer truncation, the SAME `MiniMaxH3EncoderTextForwardDevice`, the SAME f32 +activations — only the weight bytes differ (Q4_K_M ggml blocks vs the original +bf16 14-shard release). Both arms self-report identical geometry +(`layers=50 hidden=5120 heads=64 kv_heads=8 head_dim=128 ffn=25600`), which is what +establishes they are the same model. Conditioning is `[233, 5120]` f32, written by +`minimax-h3-gen --encoder-only --save-embeds`. Build `d3861693d51a`, CUDA 13.0.1 +container, Thor sm_110, GPU idle (the LocalAI render had finished). + +**A CALIBRATION arm, because a cosine is meaningless without a scale.** The bf16 +encoder also encoded a one-word edit of the same prompt (`bamboo forest at night` +-> `at dawn`, also 233 tokens). That is a real semantic change to the scene, and it +is the yardstick the quantization number is read against. + +| | max\|diff\| | RMS | rel RMS | rel RMS excl. sink tok | cos min | cos mean | cos median | angle mean | angle max | +|---|---:|---:|---:|---:|---:|---:|---:|---:|---:| +| **Q4_K_M vs bf16** (quantization) | 154.0 | 0.5045 | **0.03403** | **0.06849** | 0.90916 | **0.99745** | 0.99810 | 3.793° | 24.61° | +| bf16 `night`->`dawn` (one-word edit) | 31.1 | 0.2812 | 0.01897 | 0.06666 | 0.84736 | 0.99769 | 0.99963 | 2.228° | 32.07° | + +**Reading it.** + +1. **It is NOT a scale change.** The Q4 conditioning is uniformly ~1% smaller + (per-token norm ratio mean 0.99010), but the best single global rescale removes + almost none of the difference (0.03403 -> 0.03280). The change is DIRECTIONAL, + which is the kind that matters for conditioning. +2. **Its total energy is on par with a one-word prompt edit.** Excluding token 0, + the perturbations are 6.85% (quantization) vs 6.67% (the edit) of the + conditioning norm. Quantizing the encoder moves the conditioning about as much + as rewriting a word of the prompt. +3. **But the SHAPE is opposite, and this is the interesting part.** The word edit + is SPARSE: 172 of 233 tokens stay above cosine 0.999 (median rotation 0.16°) and + the change concentrates on ~6 tokens, the biggest at 32°. Quantization is + DIFFUSE: 232 of 233 tokens fall below cosine 0.999, EVERY token rotates by a few + degrees (median 3.5°), and one token (69) rotates 24.6°. That is a smear applied + everywhere, not a different prompt. +4. **`max|diff|` = 154 is the attention sink, not corruption.** Token 0 has norm + 15,522 against a 366 mean (42x) and carries 68% of the total squared error, yet + its DIRECTION is nearly untouched (cosine 0.99962). This is the channel-wise + magnitude-outlier behaviour ComfyUI PR 15298 attributes to H3's partial + split-half RoPE, showing up concretely: the outlier dominates every norm-weighted + aggregate, which is why the sink-excluded column is the honest one. + +**Verdict.** Q4_K_M is doing real, measurable damage to the conditioning — not a +rounding artifact, and comparable in magnitude to editing the prompt — but it +damages it DIFFUSELY. A uniform few-degree rotation of every token is the signature +that blunts fine-grained compositional instruction (coverage, blocking, staging) +toward a prompt's average semantics, which is exactly the "competent but generic" +symptom. So the bf16 encoder is worth a render A/B. + +**What this does NOT establish.** It does not prove the render changes. The DiT +consumes conditioning through a token refiner and cross-attention, and nothing here +measures that path's sensitivity to a 3.5°-median rotation. The owed next +measurement is a byte-identical-everything-else render A/B: same DiT, same seed, +same steps, `--prompt-embeds cond_q4km.bin` vs `cond_bf16.bin` (which is exactly +what `--save-embeds` / `--prompt-embeds` make controllable). + +**REPRODUCED.** Both arms were re-run from scratch (fresh process, fresh +load of the checkpoint) and each produced a BYTE-IDENTICAL conditioning file: +`cond_q4km.bin` md5 `a331232096ef1da2628f885950b2fc55` and `cond_bf16.bin` md5 +`9c096b63b9bd07f604daebb2fc090f46` on both runs. So every number above is +deterministic, not a sample — there is no noise band to argue about, and a +future change to either path shows up as an md5 change. + +**Cost, for the next person.** Q4_K_M arm: 40 s wall, host+device peak **18.0 GiB**. +bf16 arm: 40 s wall (35 s of it streaming), **45.41 GiB uploaded** to the device, +host conversion peak **0.0195 MiB** (one norm — the projections never touch a host +buffer), total host+device peak **51.95 GiB** on the 122 GiB UNIFIED pool. The +streamer's counters confirm the path ran: `layers=50 tensors=400 direct=350 +converted=200 fused=100`. diff --git a/.agents/model-matrix.md b/.agents/model-matrix.md index 3acd3eb9a..585fdb00d 100644 --- a/.agents/model-matrix.md +++ b/.agents/model-matrix.md @@ -81,7 +81,7 @@ Engaged architectures (the 45 non-`INVENTORIED` rows): | ✅ | `Glm4MoeLiteForCausalLM` | GLM-4.7-Flash (31.2B MLA + GLM MoE) | SACRED gate 8/8 vs vLLM 0.25.0 (STRICT token-exact 1/8 + near-tie-band 7/8, 69/128 tokens strictly exact, max teacher-forced gap 0.0 nats, 0 forward-divergent; vLLM K=5 self-deterministic → STRICT bar); FIRST e2e coverage of the q_lora query branch AND the noaux_tc sigmoid router (closes the MLA campaign's two gaps, C2); speed pending | `MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm` | | 🚧 | `KimiLinearForCausalLM` | Kimi-Linear-48B-A3B | **DEVICE-KDA GB10 122/128 + 4.24 tok/s (§15, #104); device NoPE-MLA lever MEASURED-NEGATIVE (2026-08-07, §16, `row/KIMI-STRICT-CLOSE` #107):** the per-channel-decay device recurrence `vt::KdaGatedDeltaRule` moves 106→**122/128** (p0-p6 16/16; sole p7 pos-6 comma near-tie) AND **1.35→4.24 tok/s (3.1×)** — vLLM's ACTUAL f32-on-bf16 arithmetic, beats §14's host-precision 120. The §15 residual (d) was attempted in device-COMPUTE form: `VT_KIMI_DEVICE_MLA` routes the 7 NoPE-MLA layers' softmax core through `vt::Attention` (pad-V: value zero-padded qk_nope+qk_rope=192 vs v=128, `out[:,:,:v]` byte-exact). CPU RED-first gate GREEN (`test_kimi_linear_forward` **14/14·825**, pad-V==f64 ref rtol 3e-3; perturbation fails 108). Full 48.9B GB10 gate (single-load, flock, min-avail 21 GiB, no reboot): control device-KDA reproduces **122/128, 4.24 tok/s** EXACTLY; **+device-MLA REGRESSES to 109/128 AND 3.89 tok/s** — `vt::Attention`'s f32 online-softmax is the right math but a DIFFERENT reduction order than vLLM's FA2, so it coin-flips near-ties (breaks p3 16→3 into §14's `163586×` repeat) and the per-(t,h) build slows the O(n²) recompute. `VT_KIMI_DEVICE_MLA` STAYS OFF, kept as a documented-MEASURED-NEGATIVE A/B knob (§14 `ISLAND_F32ACC` precedent). MLA dims VERIFIED from the real config (nah=32, qk_nope=128, qk_rope=64, v=128, kv_lora=512, q_lora=None; 7 full-attn/20 KDA). Both device knobs default OFF (122 ≠ STRICT, K=3-deterministic golden). STRICT residual, sharpened: needs vLLM's ACTUAL kernels — (c) chunk_kda prefill family (Triton-AOT regen for sm_121a) + (d) paged FA2 `mla::ForwardMlaAttentionBlock` (NOT the vt::Attention approximation) + (e) paged-incremental decode (needs a decode/paged-attn op, query_len≠key_len; kills the O(n²)) — each a substantial multi-kernel brick (§16). Row STAYS 🚧. **FULL-MODEL GB10 e2e RUNS — NEAR-TIE 106/128 (2026-08-06, `row/MODEL-KIMI-LINEAR-BF16`):** the bf16-resident path CLEARS the f32-loader block — the full 48.9B model now runs e2e on one GB10. dgx CUDA build (`-Werror` clean, 14 GDN AOT symbols nm-linked, `test_kimi_linear_forward` 13/13·656 in the CUDA binary); `kimi-linear-gen --gpu` greedy-decodes the §12 8-prompt battery x16 vs `greedy_ids.npy`. MEMORY: load 117.6s, host RSS PEAK **1.7 GiB** (stage-then-ReleaseHost), device peak 98.5 GiB, min-avail **21.6 GiB** (above the 15 GiB floor, matches the ~25 GiB pool-math headroom), NO OOM/reboot. TOKEN gate **NEAR-TIE 106/128 (82.8%)** — prompts 0,1,3,4,5,6 are 16/16 token-exact, p2/p7 diverge at punctuation/word near-ties; 96 consecutive exact tokens across 6 prompts prove the WIRING (a wiring bug can't). Root cause (honest): the f32 residual stream + host-f64 islands are MORE precise than vLLM's bf16 device kernels, so they flip the argmax where vLLM's deterministic bf16 top-1 has a small margin. STRICT path = the named W7-speed residuals (device GDN/MLA islands -> bf16 stream matching vLLM's rounding). 1.59 tok/s (recompute+island rate). `VT_KIMI_DEVICE_COMPUTE` STAYS OFF (parity-enablers: near-tie != token-exact). Row STAYS 🚧. **bf16-RESIDENT loader/forward IMPLEMENTED + CPU-gated (2026-08-06, `row/MODEL-KIMI-LINEAR-BF16`):** the §13 design is coded — `LoadKimiLinearResidentBf16Weights`/`StageKimiResidentBf16`/`BuildKimiResidentFromHost` (`kimi_linear_weights.cpp`; `LoadBf16Direct` -> `OwnedTensor`, per-tensor stage-to-`d_dev` + `ReleaseHost`, tiny vectors host f32), `KimiLinearResidentWeights` (`kimi_linear.h`), bf16 device forward `DeviceForwardBodyBf16` + `Gemm Bf16` cast-act at ~20 GEMM sites with the two host-fallback islands EXTRACTED+shared (`kimi_linear_device.cpp`), `ForwardDevice` resident-path dispatch (`kimi_linear.cpp`), and the `kimi-linear-gen` e2e harness. CPU **13/13·656** (12/12·614 f32 path UNTOUCHED + NEW tiny-config bf16-vs-f32 gate). PENDING: dgx CUDA build + full-model GB10 e2e vs the STRICT golden. Row STAYS 🚧. **bf16-RESIDENT brick POOL-MATH+DESIGN (2026-08-06, `row/MODEL-KIMI-LINEAR-BF16`):** pool math CLOSES (91.5 GiB bf16 device-resident + ~2.4 GiB act/norms/ctx ≈ 94 GiB, ~25 GiB headroom); design grounded §13 (Laguna `GemmBf16` cast-act + `OwnedTensor::d_dev`, `LoadBf16Direct`, f32 `MaterializeHost` kept for the unit gate). Impl (loader/forward rewrite + gate + e2e) pending. Row STAYS 🚧. **§8 GOLDEN CAPTURED — STRICT (2026-08-06, `row/MODEL-KIMI-LINEAR-E2E`):** the §8 SACRED oracle golden is captured on GB10 (0.25.0-stage, util 0.82, moe=triton, min 15 GiB avail, NO reboot), **8/8 prompts DETERMINISTIC over K=3 → STRICT gate**, committed at `tests/parity/goldens/kimi_linear_greedy/`. Full our-engine e2e BLOCKED on OUR f32 loader (materializes ~183 GiB > 119 pool), the bf16-residency residual; row STAYS 🚧. **W7 GPU-VERIFY (2026-08-06, branch `row/MODEL-KIMI-LINEAR-GPU`):** the device compute runs **12/12·614 GREEN on GB10 sm_121a CUDA build**, BOTH arms (`VT_KIMI_DEVICE_COMPUTE=1` + host-ref); prod stack (CUTLASS-NVFP4 GEMM + FA2 ENABLED + Triton-AOT GDN, 14 cubins nm-verified); f32 device==W2 ref, no divergence, no DeepSeek-class trap. Oracle gateability re-confirmed (0.25.0-stage registers `KimiLinearForCausalLM`). e2e §8 SACRED golden STILL disk-blocked (91.5 GiB checkpoint absent, dgx root 100% full, 34G free). Row STAYS 🚧. **W7 DBuf-resident device COMPUTE landed, CPU-gated** (`CLAIM-KIMI-LINEAR-W7`): the real device compute (`ForwardDeviceCompute`, `kimi_linear_device.cpp`) composes the whole 27-layer KDA/NoPE-MLA + 256-expert-MoE hybrid over pooled f32 `DBuf`s through the SHARED `vt::` ops (embed/`FusedChain` add+RMSNorm/`MatmulBT` projections/`CausalConv1dFwd` convs/`L2Norm`/`RmsNormGated`/`MoeRouterTopK` sigmoid-`noaux_tc`/`MoeSiluMul`/`MoeCombine`/lm_head), returning DEVICE-RESIDENT logits; 2 documented HOST-FALLBACK islands (the KDA per-k-channel gated-delta recurrence + its exp/softplus decay gate — `vt::GdnDecode` carries only a per-HEAD scalar decay; the NoPE-MLA softmax core — the paged `mla::ForwardMlaAttentionBlock` device path is born-on-runner) are the W7-speed residuals. CPU-gated vs the W2 host reference (the CPU backend runs the SAME `vt::` dispatch): `test_kimi_linear_forward` **12/12·614** (per-op KDA/NoPE-MLA/MoE/dense device==ref within f32-accumulation tolerance; the whole `ForwardDeviceCompute` == ref logits + greedy-token-identical + device-resident). Runner opt-in via `VT_KIMI_DEVICE_COMPUTE=1` (default OFF keeps the CPU-verified W6 host-ref compose). GPU numerics (bf16 activations, GDN Triton-AOT cubins, paged het-KV, grouped-MoE slabs) + the e2e SACRED golden stay a NAMED pending (box down) — row STAYS 🚧. ON TOP OF **W6 DEVICE forward SEAM** (`CLAIM-KIMI-LINEAR-W6`): the born-on-the-runner `ForwardDevice` (the DEFAULT `gather_logits` runner path) no longer refuses — it composes the `[rows,vocab]` logits via the CPU reference and hands them back DEVICE-RESIDENT (a pooled `DBuf`, wrapped like deepseek_v2 `WrapDeviceLogits`; `on_device()==true` on CPU+CUDA) so the on-GPU sampler consumes them with NO host download. Kimi-Linear now ROUTES device-resident (`check-runner-routing-consistency` reclassifies it, refuse-skipped stubs 2→1, NO allowlist; `check-fusion-consistency` green); `test_kimi_linear_forward` **7/7·300** (adds the `ForwardDevice`==host-ref device-resident gate). The DBuf-resident device COMPUTE (KDA via the GDN family, NoPE-MLA via `mla::ForwardMlaAttentionBlock`, DeepSeek-V2 grouped-MoE over the paged het-KV; full plan in `kimi_linear.cpp`) is the GPU-verify-pending W7 residual. ON TOP OF **W2-W6 CPU REFERENCE forward** (`CLAIM-KIMI-LINEAR-W2`): the real host `KimiLinearModel::Forward` composes the whole 27-layer hybrid from the landed primitives (KDA layer via `vllm::kimi_kda` refs + the gated-delta recurrence; NoPE-MLA materialized-MHA ref; sigmoid `noaux_tc` MoE + shared expert; dense SwiGLU); loader now materializes host float weights; `test_kimi_linear_forward` 6/6·246 (per-op gates + finite whole forward + greedy decode). ON TOP OF **W1 scaffolding** (registry + `ParseKimiLinearParams` 20 KDA + 7 NoPE-MLA + index-verified name-map + het-KV spec). e2e-gateable (FITS one GB10, 0.77× pool). RESIDUAL = the DEVICE born-on-runner forward (KDA kernel/absorbed-MLA/grouped-MoE slabs) + the W0/W7 e2e SACRED golden. Row → `ACTIVE` (device SEAM wired; the DBuf device compute + e2e SACRED golden pending) | `MODEL-TEXT-kimi-linear-kimi-linear-for-causal-lm` | | 📋 | `KimiK3ForConditionalGeneration` | Kimi K3 (2.8T MoE + MoonViT-V2, DERIVE-AND-SHIP) | **W2/W5 CPU scaffolding landed** (registry stub + nested text/vision/quant config descent + text-backbone structural name-map + REFUSE-by-name forward + MXFP4-refuse loader; clean CPU build, scaffold gate 6/6). text backbone IS `KimiLinearForCausalLM` (KDA+MLA+MoE hybrid, HEAVY reuse); **does NOT fit GB10 (~1.56 TB MXFP4, ~12×)** and NOT in the pinned oracle ⇒ no on-box golden — DERIVED, proxy-gated on Kimi-Linear-48B; forward + MXFP4 + KDA delta + MoonViT-V2 not implemented (NOT-YET-BUILDABLE) | `MODEL-MM-kimi-k3-kimi-k3-for-conditional-generation` | -| 🚧 | `MiniMaxH3DiTModel` | MiniMax-H3 (33.1B omni-modal video+audio DiT, DERIVE-AND-SHIP) | **W1/W2 landed**: packed layout (fl2va + ref2va, fp64 position grid BIT-EXACT), latent packing, euler-ancestral eta0 scheduler, and the full DiT forward all parity-gated against the UPSTREAM vLLM-Omni modules executed at reduced dimensions (**max abs diff 1.6e-7**, 10/10 cases / 2539 assertions). NOT autoregressive (no KV cache, no sampler, no logits) and **e2e HW-BLOCKED** (~354 GB checkpoint, ~133 GB/rank on 4x B300 vs 119 GiB unified); bf16 production stream + request planning + the ComfyUI-GGUF arm also landed (535 REAL tensors resolve onto our contract, geometry from shapes alone). **HW verdict CORRECTED: quantized arms FIT (~41 GB in 119 GiB)**, so e2e + speed are reachable; encoder/VAEs/audio VAE DONE (4.2e-9 vs the checkpoint's remote code); NVFP4 layout GATED as identical to ours (speed path is loader wiring); BOTH VAE DECODERS done (audio 4.2e-9, video ViT3D 8.9e-8); video tiling + 3D-CNN encoder (conditioning only) pending; encoder TEXT tower done (1.2e-7); **serving `/v1/videos` DONE and the DEVICE-RESIDENT forward (W2b, f32) LANDED + GPU-VERIFIED on Thor sm_110 at video 1.49e-7 / audio 8.94e-8**; bf16 stream + fusion folds + the FP4 path (needs sm_121a) + a real-checkpoint run pending. **2026-08-05: the AUDIO-VAE ENCODER is ported** (DAC analysis stack + `pre_block` AttnProjection + `mean_proj`, gated stage by stage vs the checkpoint's own remote code at 2.98e-8 / 1.64e-7 / 1.86e-8) with its own checkpoint loader gated on the real 1087-tensor manifest — so **ref2va AUDIO and VIDEO+AUDIO references are now WIRED** (audio rows move by 0.51 / 0.71; a different waveform still moves them by 7.1e-4). Both VAEs are now complete in both directions. **bf16 13-SHARD RELEASE INDEXES 2026-08-07 (`row/H3-BF16-SHARDED-DIT`)**: `MiniMaxH3ShardedCheckpoint` resolves the ORIGINAL 66.3 GB release through its own `model.safetensors.index.json` (a tensor named in the index but missing from its shard throws BY NAME), `EnumerateMiniMaxH3ShardedTensors` feeds the shared shapes-only geometry parser, and `LoadMiniMaxH3DitFromShards` is the host-f32 reference loader. Gated CPU-only at 72/72/54497 (post-rebase): every tensor resolves to the shard the index named AND to the bytes written there, the derived geometry equals the single-file path field for field, and a SPARSE 13-shard release with the REAL 535 tensors at REAL shapes (66.3 GB declared, 144 KB on disk) derives the SHIPPED geometry (50/5376/56/128/14336/24/32/1x2x2/5120). **STREAMS 2026-08-07 (`row/H3-BF16-SHARDED-STREAM`)**: `StreamMiniMaxH3ShardedToDeviceBf16` uploads it one tensor at a time — a BF16 tensor bound for a bf16 device slot goes straight from the mmap with ZERO host buffer, so peak host is bounded by ONE tensor (observed `host_peak=8192`, `direct=37 converted=9`); bit-exact vs the non-streamed `StageMiniMaxH3DitWeights` reference over all 46 views with identical logits, rope.inv_freq host-resident, 73/73/55203. Spec §8.14. This UNBLOCKS the quantization-quality question; no bf16-vs-quant render or speed number is claimed. Spec §8.13. **W-FP4a LANDED (CPU) 2026-08-06 (`row/H3-FP4-SPEED`)**: the device DiT forward now routes the NVFP4 projections through the shared Marlin W4A16 dispatcher (fp4 kept packed; no new quant code), fp4-vs-bf16 wiring gate GREEN (62/62·30039). **W-FP4a GB10 leg LANDED 2026-08-06 (`row/H3-FP4-GPU-E2E`, PR #64):** on sm_121a the Marlin W4A16 path RAN for all 11 projections (`dense_gemms==11` default — VT_MARLIN_DENSE is default-ON → vLLM's own DENSE Marlin GEMM, not the grouped route; `marlin_gemms==11` under VT_MARLIN_DENSE=0; `fallback_gemms==0`), fp4-vs-bf16 BYTE-EXACT (max\|diff\|=0), and the fp4 arm is a MEMORY win not a diffusion-forward speed win (per-forward bf16/fp4 3.47× @seq64 → 0.79–0.83× @seq4224–7040; ~16 vs ~66 GB device). Real-checkpoint fp4-resident t2va e2e RUNS (real 18.75 GB NVFP4 DiT + VAEs + GGUF Qwen3-VL-32B encoder → valid mp4/wav; DiT s/step 5.45/20.0/209 s @512/768/REF-209f) but frames are a non-scene patch-grid at 12/20/50 steps → OPEN render bug (device VAE/denoise). vLLM-Omni has no quantized H3 arm (BF16-only) so any comparison is HW/loader-forced-indirect — spec §8 | `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit` | +| 🚧 | `MiniMaxH3DiTModel` | MiniMax-H3 (33.1B omni-modal video+audio DiT, DERIVE-AND-SHIP) | **W1/W2 landed**: packed layout (fl2va + ref2va, fp64 position grid BIT-EXACT), latent packing, euler-ancestral eta0 scheduler, and the full DiT forward all parity-gated against the UPSTREAM vLLM-Omni modules executed at reduced dimensions (**max abs diff 1.6e-7**, 10/10 cases / 2539 assertions). NOT autoregressive (no KV cache, no sampler, no logits) and **e2e HW-BLOCKED** (~354 GB checkpoint, ~133 GB/rank on 4x B300 vs 119 GiB unified); bf16 production stream + request planning + the ComfyUI-GGUF arm also landed (535 REAL tensors resolve onto our contract, geometry from shapes alone). **HW verdict CORRECTED: quantized arms FIT (~41 GB in 119 GiB)**, so e2e + speed are reachable; encoder/VAEs/audio VAE DONE (4.2e-9 vs the checkpoint's remote code); NVFP4 layout GATED as identical to ours (speed path is loader wiring); BOTH VAE DECODERS done (audio 4.2e-9, video ViT3D 8.9e-8); video tiling + 3D-CNN encoder (conditioning only) pending; encoder TEXT tower done (1.2e-7); **serving `/v1/videos` DONE and the DEVICE-RESIDENT forward (W2b, f32) LANDED + GPU-VERIFIED on Thor sm_110 at video 1.49e-7 / audio 8.94e-8**; bf16 stream + fusion folds + the FP4 path (needs sm_121a) + a real-checkpoint run pending. **2026-08-05: the AUDIO-VAE ENCODER is ported** (DAC analysis stack + `pre_block` AttnProjection + `mean_proj`, gated stage by stage vs the checkpoint's own remote code at 2.98e-8 / 1.64e-7 / 1.86e-8) with its own checkpoint loader gated on the real 1087-tensor manifest — so **ref2va AUDIO and VIDEO+AUDIO references are now WIRED** (audio rows move by 0.51 / 0.71; a different waveform still moves them by 7.1e-4). Both VAEs are now complete in both directions. **bf16 13-SHARD RELEASE INDEXES 2026-08-07 (`row/H3-BF16-SHARDED-DIT`)**: `MiniMaxH3ShardedCheckpoint` resolves the ORIGINAL 66.3 GB release through its own `model.safetensors.index.json` (a tensor named in the index but missing from its shard throws BY NAME), `EnumerateMiniMaxH3ShardedTensors` feeds the shared shapes-only geometry parser, and `LoadMiniMaxH3DitFromShards` is the host-f32 reference loader. Gated CPU-only at 72/72/54497 (post-rebase): every tensor resolves to the shard the index named AND to the bytes written there, the derived geometry equals the single-file path field for field, and a SPARSE 13-shard release with the REAL 535 tensors at REAL shapes (66.3 GB declared, 144 KB on disk) derives the SHIPPED geometry (50/5376/56/128/14336/24/32/1x2x2/5120). **STREAMS 2026-08-07 (`row/H3-BF16-SHARDED-STREAM`)**: `StreamMiniMaxH3ShardedToDeviceBf16` uploads it one tensor at a time — a BF16 tensor bound for a bf16 device slot goes straight from the mmap with ZERO host buffer, so peak host is bounded by ONE tensor (observed `host_peak=8192`, `direct=37 converted=9`); bit-exact vs the non-streamed `StageMiniMaxH3DitWeights` reference over all 46 views with identical logits, rope.inv_freq host-resident, 73/73/55203. Spec §8.14. **bf16 TEXT ENCODER + THE CONDITIONING NUMBER 2026-08-07 (`row/H3-ENC-BF16-COND-DIFF`)**: the 14-shard 63 GB bf16 Qwen3-VL-32B encoder streams to device too (`StreamMiniMaxH3EncoderShardsToDevice`, q/k/v and gate/up fused ON DEVICE), `--encoder-only` runs the tower alone (peak ~96 -> ~49 GiB by not loading the DiT first), and the widening is gated BIT-IDENTICAL vs an f32-staged tower so the A/B cannot be confounded. MEASURED on Thor over 233 tokens: Q4_K_M vs bf16 conditioning is cos 0.99745 mean / 0.909 min, rel RMS 6.85% excluding the attention sink, median rotation 3.5 deg — same energy as a ONE-WORD prompt edit but DIFFUSE (232/233 tokens rotate vs 172/233). Whether the RENDER changes is NOT established. 75/75/55609. Spec §8.15. This UNBLOCKS the quantization-quality question; no bf16-vs-quant render or speed number is claimed. Spec §8.13. **W-FP4a LANDED (CPU) 2026-08-06 (`row/H3-FP4-SPEED`)**: the device DiT forward now routes the NVFP4 projections through the shared Marlin W4A16 dispatcher (fp4 kept packed; no new quant code), fp4-vs-bf16 wiring gate GREEN (62/62·30039). **W-FP4a GB10 leg LANDED 2026-08-06 (`row/H3-FP4-GPU-E2E`, PR #64):** on sm_121a the Marlin W4A16 path RAN for all 11 projections (`dense_gemms==11` default — VT_MARLIN_DENSE is default-ON → vLLM's own DENSE Marlin GEMM, not the grouped route; `marlin_gemms==11` under VT_MARLIN_DENSE=0; `fallback_gemms==0`), fp4-vs-bf16 BYTE-EXACT (max\|diff\|=0), and the fp4 arm is a MEMORY win not a diffusion-forward speed win (per-forward bf16/fp4 3.47× @seq64 → 0.79–0.83× @seq4224–7040; ~16 vs ~66 GB device). Real-checkpoint fp4-resident t2va e2e RUNS (real 18.75 GB NVFP4 DiT + VAEs + GGUF Qwen3-VL-32B encoder → valid mp4/wav; DiT s/step 5.45/20.0/209 s @512/768/REF-209f) but frames are a non-scene patch-grid at 12/20/50 steps → OPEN render bug (device VAE/denoise). vLLM-Omni has no quantized H3 arm (BF16-only) so any comparison is HW/loader-forced-indirect — spec §8 | `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit` | | ✅ | `LagunaForCausalLM` | Poolside Laguna-S-2.1 (118B/8B MoE) | **LONG-CTX DECODE LEVERS LANDED + MEASURED (2026-08-03, `CLAIM-LAGUNA-LONGCTX-LEVERS`): window-bounded SWA reads (`VT_LAGUNA_SWA_WINDOW`, default-ON, BYTE-EXACT) bound the four `DecodeAttnGqa*` kernels' read to the ~512 sliding window (vLLM `laguna.py:412`) — GB10 A/B token-IDENTICAL `=1` vs `=0` at 520-token context (truncation active), MEASURED −0.30 ms/step at ~2k (~0 at ≤512, grows linearly). bf16 paged KV (`VT_LAGUNA_KV_BF16`, default-OFF opt-in) a distributional near-tie left UNRATIFIED. See BENCHMARKS `CLAIM-LAGUNA-LONGCTX-LEVERS`.** — **NVFP4 W4A4 ARM RAN on GB10 (N4, 2026-08-01, `CLAIM-LAGUNA-NVFP4-N4`): the additive safetensors NVFP4 arm (N1a/N1b/N2/N3 — `Nvfp4Weight` expert fields + `LoadLagunaForCausalLMWeights` + `LqGemmNvfp4Fp4` per-expert TRUE-W4A4 + `LagunaFfnBlock` `fp4` branch + `laguna_gen` dir-autodetect; CPU-gated `test_laguna_nvfp4_loader` 3/3·61, GGUF path byte-identical) generates COHERENTLY on the real 67 GiB `poolside/Laguna-S-2.1-NVFP4`. vs the vLLM MARLIN golden (vLLM's exact prompt ids injected): FIRST 2 TOKENS MATCH exactly, then near-tie divergence (our TRUE-W4A4 fp4-activations vs the MARLIN golden's W4A16 bf16-activations — different precision, EXPECTED; shares golden vocab). SPEED (N5, trace-driven, 2026-08-01): 0.16 → ~4.5 tok/s (~28× THIS SESSION), now ~4× from vLLM 18.8. **Lever #2** (nsys found the bf16 tower running host `MatmulNK` on the CUDA queue): route it to the GPU (`LqGemm` bf16 → `CastBf16` + `MatmulBT`, weight stays bf16) → 6.34 → 0.39 s/tok (16×). **Lever #1** (nsys found the emulation expert GEMM at 92%, GPU 87% busy): the engine's native sm120a fp4 tensor-core MMA (`MatmulNvfp4Fp4Native`) reads the SAME linear scales — it was gated OFF behind `VT_NVFP4_FP4_NATIVE`; default it ON in the driver → 0.39 → ~0.20-0.24 s/tok (~2×). Both coherent + near-tie (byte-identical ids to emulation; first token matches golden). Two GB10 memory fixes landed to run (shard-release + context-before-load). OPEN #234 (remaining ~4×): grouped W4A4 MoE (top_k×3 launches → 3), `ResidentNvfp4`, decode CUDA-graph + on-GPU sampling (the host-orchestration tail). Spec `.agents/specs/laguna-nvfp4-arm-2026-07-31.md` §N4/§N5. The GGUF-Q4_K track (below) is the separate keep-quant vehicle.** Prior **FASTER DECODE (W9, 2026-07-31, `CLAIM-LAGUNA-W9-GROUPED`): the 30 un-grouped per-expert keep-quant GEMV launches/step (top_k × {gate,up,down} `LqGemmRowSlice`) fold onto the SHARED `vt::MatmulBTQuantGrouped` op — per token, Pk experts' gate/up/down each collapse to ONE grouped launch over the already-stacked `[E*N,H]` tower (no loader change). Same-binary A/B on real UD-Q4_K_XL (GB10, `--gpu`, drop_caches cold, 24 tok): grouped (`VT_LAGUNA_GROUPED_MOE=1`, default) == per-expert (`=0`) BYTE-IDENTICAL (md5 `754728c6`, both == W6 golden) + decode 0.18 → 0.13 s/tok (1.38×). Routes through the shared vt op (fold policy). Cumulative with W8: decode 0.66 → 0.13 s/tok (5.1×; 1.5 → 7.7 tok/s; 18× → 3.6× vs llama.cpp 27.8). Next lever: device-resident decode (#1). See spec §W9.** Prior **FASTER DECODE (W8, 2026-07-31, `CLAIM-LAGUNA-W8-EMBED`): `LagunaEmbed` no longer converts the whole 1.23 GB embed table to f32 every token (it gathered T rows out of the whole [Vsz,H] table via `ReadF32` — ~311M host element-converts/token, the DOMINANT decode cost the W7 profile under-filed as "#5"); now gathers only the T needed rows directly (BIT-IDENTICAL — same per-element conversion, same rows). GATED on the real 3-shard UD-Q4_K_XL GGUF (GB10, `--gpu`, W6 cached, drop_caches cold, 24 tok): TOKEN-IDENTICAL to the W5/W6 golden (`22345 83 350 785 …`, coherent " Paris.") + decode 0.66 → 0.17 s/tok = 3.9× (1.5 → 5.9 tok/s; 18× → 4.7× vs llama.cpp 27.8). See `.agents/specs/laguna-s21-w7-speed-2026-07-31.md` §W8. Next: grouped-expert GEMM (=A3) then device-resident decode.** Prior **DECODE-SPEED ATTRIBUTED (W7 profile-only, 2026-07-31, `CLAIM-LAGUNA-W7-SPEED`): `nsys` of the W6 decode (real UD-Q4_K_XL GGUF, GB10) attributes the 0.66 s/tok (~1.5 tok/s vs llama.cpp 27.8 on identical bytes, ~15-18x) to HOST-ORCHESTRATION, not kernel compute — GPU active only 32.7% of the step, 67.3% host/idle; 22,115 `cudaStreamSynchronize` (~2,764/step, zero GPU overlap) from the ~1,795 per-GEMM `DrainQueue` in `LagunaForwardGgufCached` + scalar host glue; 39.4% of GPU time is `QuantizeQ8K` activation-quant (per-GEMM), weight GEMVs un-grouped at ~22% of the 240 GB/s peak (llama.cpp ~76%); no H2D/D2H (unified memory). Ranked levers (all in-tree from ds4): device-resident decode 1.5->~5-7 tok/s, grouped-expert GEMM (`MatmulBTQuantGrouped`) +1.5-2x + dedupes the activation-quant, decode CUDA-graph, tuned MMVQ; + free host cleanups (`LagunaEmbed` copies the whole 1.23 GB embed table/token, per-token RoPE-cache rebuild). Honest reachable ~13-20 tok/s, 27.8 a stretch. NO code changed. See `.agents/specs/laguna-s21-w7-speed-2026-07-31.md`. Prior RUNNABLE + FAST DECODE (W6, 2026-07-31): a per-layer K/V cache + single-token incremental decode replaces W5's O(n²) STATELESS recompute — TOKEN-IDENTICAL (byte-equal ids, md5 `754728c6…` match, == the W5 golden) and 5.05× faster per token: decode 3.33 → 0.66 s/tok on the real UD-Q4_K_XL GGUF (GB10, `--gpu`, keep-quant), same " Paris.…" text. `LagunaKvCache` (mirrors `DeepseekV4KvCache`, MLA-latent → GQA multi-head K/V; caches post-QK-RMSNorm/post-RoPE K + raw V at f32 — bit-exact since RoPE/QK-norm are position-only and attention is causal), MIXED attention per-layer: 12 GLOBAL layers grow unbounded + 36 SLIDING-WINDOW-512 layers EVICT rows beyond the 512 window (gemma2/3 `is_sliding`); `LagunaForwardGgufCached` + shared `LagunaAttention`/`LagunaFfnBlock` helpers used by BOTH forwards (identical float ops; recompute ids unchanged after refactor), `examples/laguna_gen --stateless` A/B flag. No cache bug (bit-exact 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) — "The capital of France is" → " Paris. …", first token "Paris." matches the llama.cpp-Poolside reference. Multi-shard GGUF reader + keep-quant tower (`LoadLagunaFromGgufShards`) + `LagunaForwardGguf` (ds4 keep-quant Gemm/GemmRowSlice) + `examples/laguna_gen`; load 20.6s, peak 71 GiB, 3.27 s/tok stateless recompute (speed=W6).** Prior W3: **W3 REAL forward + 3 new ops landed** (`laguna_ops.cpp`: per-head softplus attn out-gate + ungrouped sigmoid-noaux router + dual per-layer RoPE cos/sin builders; `LagunaModel::Forward` 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 W1/W2 `VT_CHECK(false)` stub; CPU `-Werror` full-library build clean; `test_laguna_scaffold` **8/8·166** incl. softplus math, router selection+tie-break RED-first, dual-RoPE cos/sin bit-match, variable-Q-head shapes, forward composition on synthetic weights; `test_model_registry` 24/24). W1 oracle DECISION: vLLM native `laguna.py` in pin ⇒ config constructs; dual-oracle = vLLM-NVFP4/-FP8 (fits GB10, BF16 235 GiB does NOT) + llama.cpp-Q4_K token-exact. DEFERRED to W4 (needs 73 GB checkpoint): GGUF keep-quant tower materialization + device/paged production forward + strict dual-oracle greedy gate. ~85-90% reuse (ds4-MoE + gemma-sliding + olmo3-dual-rope + landed Q4_K keep-quant); NEW = the 3 landed host ops + name-map + variable-Q-head device runner. **W4 (2026-07-31, `CLAIM-LAGUNA-W4`, in progress):** the UD-Q4_K_XL GGUF (73.4 GiB, 3 shards) FETCHED to dgx + its metadata/tensor-map READ AUTHORITATIVELY (814 tensors, arch `laguna`, `expert_gating_func=2` sigmoid, `leading_dense_block_count=1`, `expert_weights_scale=2.5`). Three CPU-verified FIDELITY corrections the W1-W3 scaffold got wrong, each grounded in the real GGUF + llama.cpp: (1) **per-head QK-RMSNorm** (`attn_q_norm`/`attn_k_norm` F32[128]) added to params+forward — the scope MISSED it (surfaces only in the tensor map); (2) **dual-RoPE mscale** now uses llama.cpp's `yarn_attn_factor·(1+0.1·ln(factor))` off the GGUF-authoritative `factor=32`/`yarn_attn_factor=1.0` (256K-ctx build, NOT HF's factor-128/1.4852 1M-ctx scalar) — resolves the numerics-delicate residual; (3) **separate** `ffn_gate_exps`/`ffn_up_exps` (Q4_K) + `ffn_down_exps` (Q5_K) + Q8_0 shared/attn (the scaffold assumed merged gate_up). GGUF keep-quant tower materialization (`Mw`/`Sew` mirror of ds4) + keep-quant `ForwardGguf` (vt::MatmulBT/GemmRowSlice) + the real-model greedy run vs the llama.cpp-laguna same-quant oracle remain the W5 close (73 GB single-GB10, host-orchestrated) | `MODEL-TEXT-laguna-laguna-for-causal-lm` | | 🚫 | `DeepseekV3ForCausalLM` / `DeepseekV32ForCausalLM` | DeepSeek-V3 / V3.2 | HW-blocked (671B, ~642 GiB fp8 vs 119 GiB unified memory); V3.2 additionally DEP-blocked (DSA indexer) | `MODEL-TEXT-deepseek-v2-deepseek-v3-for-causal-lm` | | 🚫 | `GlmMoeDsaForCausalLM` | GLM-5 (DSA) | HW-blocked (1404 GiB bf16) and DEP-blocked (GLM-5.x is DeepSeek-V3.2 verbatim) | `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` | diff --git a/.agents/parity-ledger.md b/.agents/parity-ledger.md index c669faf02..8fceb8e7f 100644 --- a/.agents/parity-ledger.md +++ b/.agents/parity-ledger.md @@ -919,4 +919,5 @@ Columns: | 2026-08-07 (`row/H3-BF16-SHARDED-DIT`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; CPU-only, no GPU and no download; lifecycle unchanged) | **MiniMax-H3 — the ORIGINAL bf16 release (13 safetensors shards, 66.3 GB) is now INDEXABLE.** Every H3 render so far used a QUANTIZED DiT and H3 is unusually quantization-sensitive (Q3_K_M -> Q4_K_M alone turned a murky lattice into a photoreal close-up; ComfyUI PR 15298 blames the partial split-half RoPE's channel-wise magnitude outliers), but the full-precision question was unaskable because every DiT loader took a SINGLE file. Adds (a) `MiniMaxH3ShardedCheckpoint::Open(dir)` (`src/vllm/model_executor/models/minimax_h3_sharded.cpp`), which resolves tensors through the checkpoint's own `model.safetensors.index.json` weight map (never by scanning) with one index over every shard, mirroring the in-tree multi-shard template `LoadMiniMaxH3EncoderWeights(const std::vector&, ...)`, and throws BY NAME when the index names a tensor its shard does not hold; (b) `EnumerateMiniMaxH3ShardedTensors`, the shapes-only manifest the geometry parser consumes; (c) `LoadMiniMaxH3DitFromShards`, the host-f32 reference loader; (d) `MiniMaxH3IsFp32IslandTensor`, single-sourcing the upstream fp32-ISLAND split the three existing streamers each hand-rolled; (e) `--dit ` in `examples/minimax_h3_gen` for both `--dump-params` and the run path, every existing `--dit` form unchanged. The DEVICE streamer is the stacked follow-up `row/H3-BF16-SHARDED-STREAM`, split out to stay inside the 900-line PR cap. | vLLM-Omni `vllm_omni/diffusion/models/minimax_h3/minimax_h3_transformer.py:85-101` (MINIMAX_H3_FP32_PARAM_NAMES / _BUFFER_NAMES, the island split) and `:906-922` (the parameter set); the shard-index container convention is HF safetensors' own `model.safetensors.index.json` weight_map, already consumed in-tree by `LoadSafetensorsIndex` and the multi-shard encoder/VAE loaders. No vLLM behavior changed; H3 remains BEYOND-PIN (vllm-omni, not the pinned vLLM repo). | **LANDED + CPU-GATED (loader brick; `benchmark_binding=false` — no throughput owed, and NO bf16-vs-quant render or speed number is claimed).** Re-gated AFTER the rebase onto `f34e0d17`: `test_minimax_h3` 72/72 cases / 54497 assertions, clean Release build of `libvllm.a`, `test_minimax_h3` and `minimax-h3-gen`. Two gates: (1) index+name mapping over a synthetic 4-shard set — every tensor resolves to the shard the index named AND to the bytes written there, a tensor missing from its shard throws WITH ITS NAME, and the derived geometry equals the single-file path field for field; (2) a SPARSE 13-shard release declaring the REAL 535 tensors at REAL shapes (66.3 GB declared, 144 KB on disk) derives the SHIPPED geometry 50/5376/56/128/14336/24/32/1x2x2/5120, and `minimax-h3-gen --dit --dump-params` prints all 20 fields on it. Also FIXES a real latent defect this row's sanitizer lane exposed: `MiniMaxH3ReadSafetensorF32` read 16-bit payloads through `reinterpret_cast`, which is UB on a safetensors file whose JSON header leaves the payload odd-aligned (the format does not require padding); now a byte-wise `memcpy`. RED-first proven: reverting it reproduces UBSan's `load of misaligned address` at the same line and exits 1. Honest residuals: no device load of the real 66.3 GB release, no measured peak RSS, and no bf16-vs-quantized render/speed comparison — the quality question is UNBLOCKED, not answered. | | 2026-08-07 (`row/H3-BF16-SHARDED-STREAM`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; stacked on `row/H3-BF16-SHARDED-DIT`; CPU-only, no GPU and no download; lifecycle unchanged) | **MiniMax-H3 — the ORIGINAL bf16 release (13 shards, 66.3 GB) now STREAMS to the device.** §8.13 made the checkpoint indexable but its only loader was host-f32 (~132 GB on the real release); on a 122 GiB UNIFIED pool that holds the model TWICE, and the non-streaming NVFP4 loader was already OOM-killed at anon-rss 125 GB on HALF this size, so the real release was not loadable at all. Adds `StreamMiniMaxH3ShardedToDeviceBf16` (`minimax_h3_device.cpp`, sharing `BindStreamedDitViews` with the GGUF and NVFP4 streamers): manifest first, then ONE tensor at a time, with a BF16-on-disk tensor bound for a bf16 device slot — essentially the whole 66.3 GB — uploaded DIRECTLY out of the read-only mmap with NO host buffer, each source range released via `MaybeReleaseSourcePages`, and `rope.inv_freq` kept HOST-resident. Adds `MiniMaxH3ShardStreamStats` (mirroring `Nvfp4W4A16Stats`) so the path is observable, and `--dit --device cuda` in `examples/minimax_h3_gen`. | vLLM-Omni `vllm_omni/diffusion/models/minimax_h3/minimax_h3_transformer.py:85-101` (MINIMAX_H3_FP32_PARAM_NAMES / _BUFFER_NAMES, the fp32-island split the stream honours). The streaming SHAPE is our own in-tree convention (`StreamMiniMaxH3Nvfp4ToDeviceBf16`), which exists because upstream never has to load this checkpoint on one unified-memory device; recorded as a deviation in porting-inventory §9 terms. No vLLM behavior changed; H3 remains BEYOND-PIN. | **LANDED + CPU-GATED (loader brick; `benchmark_binding=false` — no throughput owed, and NO bf16-vs-quant render or speed number is claimed).** `test_minimax_h3` 73/73 cases / 55203 assertions, clean Release build of `libvllm.a`, `test_minimax_h3` and `minimax-h3-gen`. Two gates: (1) streamed == non-streamed — all 46 weight views BIT-EXACT (`memcmp == 0`) vs `StageMiniMaxH3DitWeights(kBF16)`, dtypes included (12 fp32 islands), both device forwards IDENTICAL (video and audio max|diff| == 0.0), `rope.inv_freq` host-resident; (2) the loader RAN — counters ASSERTED, observed `shards=3 tensors=46 direct=37 converted=9 bytes=444504 host_peak=8192`, i.e. BOTH upload paths taken, every view owned by this loader, and `host_peak_bytes` bounded by one tensor (< 1/4 of bytes uploaded) so the peak cannot scale with the model. Honest residuals: the real 66.3 GB load, its measured peak RSS, and CUDA memcpy from a file-backed mmap are all UNVERIFIED (CPU-only row); the bf16-vs-quant A/B is now RUNNABLE and has not been run. | | 2026-08-06 (`row/H3-BF16-SHARDED-DIT`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; CPU-only, no GPU and no download; lifecycle unchanged) | **MiniMax-H3 — the ORIGINAL bf16 release (13 safetensors shards, 66.3 GB) now LOADS, and it STREAMS.** Every H3 render so far used a QUANTIZED DiT and H3 is unusually quantization-sensitive (Q3_K_M -> Q4_K_M alone turned a murky lattice into a photoreal close-up; ComfyUI PR 15298 blames the partial split-half RoPE's channel-wise magnitude outliers), but the full-precision question was unaskable because every DiT loader took a SINGLE file. Adds (a) `MiniMaxH3ShardedCheckpoint::Open(dir)` (`src/vllm/model_executor/models/minimax_h3_sharded.cpp`), which resolves tensors through the checkpoint's own `model.safetensors.index.json` weight map (never by scanning) with one index over every shard, mirroring the in-tree multi-shard template `LoadMiniMaxH3EncoderWeights(const std::vector&, ...)`, and throws BY NAME when the index names a tensor its shard does not hold; (b) `StreamMiniMaxH3ShardedToDeviceBf16` (`minimax_h3_device.cpp`, sharing `BindStreamedDitViews` with the GGUF and NVFP4 streamers), which converts+uploads ONE tensor at a time and uploads a BF16 tensor bound for a bf16 device slot DIRECTLY out of the read-only mmap with no host buffer at all, releasing each source range afterwards; (c) `LoadMiniMaxH3DitFromShards`, the host-f32 reference loader; (d) `MiniMaxH3IsFp32IslandTensor`, single-sourcing the upstream fp32-ISLAND split the three existing streamers each hand-rolled; (e) `--dit ` in `examples/minimax_h3_gen` for both `--dump-params` and the run path, every existing `--dit` form unchanged. It MUST stream: the pool is UNIFIED (122 GiB shared host+device) and the non-streaming NVFP4 loader was already OOM-killed at anon-rss 125 GB on half this size. | vLLM-Omni `vllm_omni/diffusion/models/minimax_h3/minimax_h3_transformer.py:85-101` (MINIMAX_H3_FP32_PARAM_NAMES / _BUFFER_NAMES, the island split) and `:906-922` (the parameter set); the shard-index container convention is HF safetensors' own `model.safetensors.index.json` weight_map, already consumed in-tree by `LoadSafetensorsIndex` and the multi-shard encoder/VAE loaders. No vLLM behavior changed; H3 remains BEYOND-PIN (vllm-omni, not the pinned vLLM repo). | **LANDED + CPU-GATED (loader brick; `benchmark_binding=false` — no throughput owed, and NO bf16-vs-quant render or speed number is claimed).** `test_minimax_h3` 68/68 cases / 49300 assertions, clean Release build of `libvllm.a`, `test_minimax_h3` and `minimax-h3-gen`. Four gates: (1) index+name mapping over a synthetic 4-shard set — every tensor resolves to the shard the index named AND to the bytes written there, a tensor missing from its shard throws WITH ITS NAME, and the derived geometry equals the single-file path field for field; (2) streamed == non-streamed — all 46 weight views BIT-EXACT (`memcmp == 0`) vs `StageMiniMaxH3DitWeights(kBF16)`, dtypes included (12 fp32 islands), both device forwards IDENTICAL (max|diff| == 0), `rope.inv_freq` HOST-resident; (3) the loader RAN — `MiniMaxH3ShardStreamStats` (mirroring `Nvfp4W4A16Stats`) proves shards opened, tensors streamed, BOTH upload paths taken, every view owned by this loader, and `host_peak_bytes` bounded by one tensor (< 1/4 of bytes uploaded), i.e. peak cannot scale with the model; (4) a SPARSE 13-shard release declaring the REAL 535 tensors at REAL shapes (66.3 GB declared, 144 KB on disk) derives the SHIPPED geometry 50/5376/56/128/14336/24/32/1x2x2/5120, and `minimax-h3-gen --dit --dump-params` prints all 20 fields on it. Honest residuals: the real 66.3 GB load and its measured peak RSS, CUDA memcpy from a file-backed mmap, and any bf16-vs-quantized render/speed comparison are all UNVERIFIED here (no GPU, no download, per the operator's instruction). Not pushed. | -| 2026-08-06 (`row/H3-ENC-BF16-COND-DIFF`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; lifecycle unchanged) | **MiniMax-H3 - the bf16 TEXT ENCODER (14 safetensors shards, 63 GB) now LOADS, it STREAMS, and `--encoder-only` runs the tower alone.** Every H3 render so far conditioned on a Q4_K_M Qwen3-VL-32B encoder and the encoder's contribution had never been measured, but `--encoder` accepted only a GGUF. Adds (a) `MiniMaxH3EncoderConfigFromShards`, deriving the geometry from the shard index's SHAPES alone with the SAME recovery rules AND the same non-shape defaults (`rope_theta`, `mrope_section`, `rms_norm_eps`, `selected_layer`) as the GGUF loader, so an A/B cannot be comparing two RoPEs; (b) `StreamMiniMaxH3EncoderShardsToDevice` (`src/vllm/model_executor/models/minimax_h3_encoder_sharded.cpp`), which fills the SAME `MiniMaxH3EncoderDeviceWeights::views` map the GGUF arm fills - over bf16 instead of ggml blocks - uploading projections DIRECTLY out of the read-only mmap and doing the `[q|k|v]` / `[gate|up]` row fusions ON THE DEVICE into offsets of one allocation, so even the transform costs no host copy; (c) `MiniMaxH3EncoderEmbedTokensFromShards`, a per-row gather out of the `[151936, 5120]` table; (d) a bf16-weight WIDEN step in `MiniMaxH3EncoderTextForwardDevice` (scratch reused across layers) because `vt::MatmulBT` needs one dtype for both operands and these activations are f32 - the 50 layers H3 runs are 48.8 GiB bf16 vs 97.5 GiB f32 on a 122 GiB UNIFIED pool; (e) `--encoder ` and `--encoder-only` in `examples/minimax_h3_gen`, which drops peak from ~96 GiB (DiT loaded first) to ~49 GiB. No new forward: the encoder graph is byte-for-byte the same code for both arms, which is what makes the quantization question measurable. | The name map is the one already gated in-tree for `LoadMiniMaxH3EncoderWeights(const std::vector&, ...)` (`model.language_model.layers.N.` -> `layers.N.`, q/k/v and gate/up FUSED, final `norm.weight` and `lm_head.weight` deliberately unbound because H3 reads the UNNORMALIZED truncated output); shard resolution reuses `MiniMaxH3ShardedCheckpoint` (§8.6, cherry-picked from `1a46ff17`), i.e. the checkpoint's own HF `model.safetensors.index.json` weight_map. Encoder truncation to `min(num_hidden_layers, 50)` is upstream vLLM-Omni's own. No vLLM behavior changed; H3 remains BEYOND-PIN (vllm-omni). | **LANDED + CPU-GATED (loader brick; `benchmark_binding=false`).** `test_minimax_h3` 70/70 cases / 49706 assertions (up from 68/68), clean Release build of `libvllm.a`, `test_minimax_h3`, `minimax-h3-gen`. Three gates: (1) a synthetic 4-shard encoder at the REAL name spellings resolves, and every fused view is `memcmp`-exact against `q ++ k ++ v` / `gate ++ up` for EVERY layer, unfused projections byte-exact, separate names gone, final norm + lm_head + vision tower NOT bound, truncation honoured, embedding gather exact and out-of-range throwing; (2) the loader RAN and is NOT the GGUF path - `MiniMaxH3EncoderShardStreamStats` asserted on shards/layers/views/fused groups/direct-vs-converted uploads with `host_peak_bytes` equal to ONE norm (peak cannot scale with the model), and the views are `kBF16`, a dtype the GGUF loader can never produce; (3) the WIDENING is exact - the same checkpoint written BF16 and F32 (bf16-rounded values) streams to `kBF16` and `kF32` views respectively and the two full encoder forwards are BIT-IDENTICAL (`memcmp == 0`), so the conditioning A/B cannot be confounded by the widening. Also repairs a PRE-EXISTING `check-public-doc-tables` ratchet red on `origin/main` (docs/STATUS.md 284114 > 284073). The real 63 GB load, its peak RSS, and the Q4_K_M-vs-bf16 conditioning numbers are the GPU follow-up in this row. | +| 2026-08-06 (`row/H3-ENC-BF16-COND-DIFF`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; lifecycle unchanged) | **MiniMax-H3 - the bf16 TEXT ENCODER (14 safetensors shards, 63 GB) now LOADS, it STREAMS, and `--encoder-only` runs the tower alone.** Every H3 render so far conditioned on a Q4_K_M Qwen3-VL-32B encoder and the encoder's contribution had never been measured, but `--encoder` accepted only a GGUF. Adds (a) `MiniMaxH3EncoderConfigFromShards`, deriving the geometry from the shard index's SHAPES alone with the SAME recovery rules AND the same non-shape defaults (`rope_theta`, `mrope_section`, `rms_norm_eps`, `selected_layer`) as the GGUF loader, so an A/B cannot be comparing two RoPEs; (b) `StreamMiniMaxH3EncoderShardsToDevice` (`src/vllm/model_executor/models/minimax_h3_encoder_sharded.cpp`), which fills the SAME `MiniMaxH3EncoderDeviceWeights::views` map the GGUF arm fills - over bf16 instead of ggml blocks - uploading projections DIRECTLY out of the read-only mmap and doing the `[q|k|v]` / `[gate|up]` row fusions ON THE DEVICE into offsets of one allocation, so even the transform costs no host copy; (c) `MiniMaxH3EncoderEmbedTokensFromShards`, a per-row gather out of the `[151936, 5120]` table; (d) a bf16-weight WIDEN step in `MiniMaxH3EncoderTextForwardDevice` (scratch reused across layers) because `vt::MatmulBT` needs one dtype for both operands and these activations are f32 - the 50 layers H3 runs are 48.8 GiB bf16 vs 97.5 GiB f32 on a 122 GiB UNIFIED pool; (e) `--encoder ` and `--encoder-only` in `examples/minimax_h3_gen`, which drops peak from ~96 GiB (DiT loaded first) to ~49 GiB. No new forward: the encoder graph is byte-for-byte the same code for both arms, which is what makes the quantization question measurable. | The name map is the one already gated in-tree for `LoadMiniMaxH3EncoderWeights(const std::vector&, ...)` (`model.language_model.layers.N.` -> `layers.N.`, q/k/v and gate/up FUSED, final `norm.weight` and `lm_head.weight` deliberately unbound because H3 reads the UNNORMALIZED truncated output); shard resolution reuses `MiniMaxH3ShardedCheckpoint` (§8.13, landed as its own row), i.e. the checkpoint's own HF `model.safetensors.index.json` weight_map. Encoder truncation to `min(num_hidden_layers, 50)` is upstream vLLM-Omni's own. No vLLM behavior changed; H3 remains BEYOND-PIN (vllm-omni). | **LANDED + CPU-GATED (loader brick; `benchmark_binding=false`).** Re-gated AFTER the rebase onto `row/H3-BF16-SHARDED-STREAM`: `test_minimax_h3` 75/75 cases / 55609 assertions, clean Release build of `libvllm.a`, `test_minimax_h3`, `minimax-h3-gen`. Three gates: (1) a synthetic 4-shard encoder at the REAL name spellings resolves, and every fused view is `memcmp`-exact against `q ++ k ++ v` / `gate ++ up` for EVERY layer, unfused projections byte-exact, separate names gone, final norm + lm_head + vision tower NOT bound, truncation honoured, embedding gather exact and out-of-range throwing; (2) the loader RAN and is NOT the GGUF path - `MiniMaxH3EncoderShardStreamStats` asserted on shards/layers/views/fused groups/direct-vs-converted uploads with `host_peak_bytes` equal to ONE norm (peak cannot scale with the model), and the views are `kBF16`, a dtype the GGUF loader can never produce; (3) the WIDENING is exact - the same checkpoint written BF16 and F32 (bf16-rounded values) streams to `kBF16` and `kF32` views respectively and the two full encoder forwards are BIT-IDENTICAL (`memcmp == 0`), so the conditioning A/B cannot be confounded by the widening. The real 63 GB load, its peak RSS, and the Q4_K_M-vs-bf16 conditioning numbers are the GPU follow-up in this row. | +| 2026-08-07 (`row/H3-ENC-BF16-COND-DIFF`; `ROAD-V1-H3`; model `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`; MEASUREMENT, lifecycle unchanged) | **MiniMax-H3 - THE NUMBER: what quantizing the TEXT ENCODER to Q4_K_M does to the conditioning.** Same prompt (wuxia, 233 tokens), same tokenizer, same 50-layer truncation, same `MiniMaxH3EncoderTextForwardDevice`, same f32 activations - only the weight bytes differ (Q4_K_M ggml blocks vs the original bf16 14-shard release); both arms self-report IDENTICAL geometry (50/5120/64/8/128/25600), which is what establishes they are the same model. Conditioning `[233, 5120]` f32 via `--encoder-only --save-embeds`. | Not a vLLM-parity change: H3 is BEYOND-PIN (vllm-omni), and vLLM-Omni serves NO quantized H3 at all (BF16-only), so there is no upstream arm to compare against - the bf16 release IS the reference here, and it is the one the loader added in `6d454b00` makes runnable. The quantization-sensitivity premise is ComfyUI PR 15298 (H3's partial split-half RoPE produces channel-wise magnitude outliers that corrupt even INT8), and the measurement CONFIRMS its mechanism concretely: token 0 is an attention sink at norm 15,522 vs a 366 mean (42x), carrying 68% of the total squared error with its DIRECTION intact (cos 0.99962). | **MEASURED on Thor sm_110, build `d1085374` (built and measured as `d1085374`, amended for the row-branch trailer; IDENTICAL tree `dd9283cf`, so the measurement binary IS this commit), GPU idle.** Q4_K_M vs bf16: max|diff| 154.0, RMS 0.5045, rel RMS **0.03403** (0.06849 excluding the sink token), per-token cosine min 0.90916 / mean **0.99745** / median 0.99810, rotation median 3.535 deg / max 24.61 deg, 232 of 233 tokens below cosine 0.999. NOT a scale change: norm ratio mean 0.99010 but the best global rescale only moves 0.03403 -> 0.03280, so it is DIRECTIONAL. CALIBRATION arm (bf16 encoder, ONE-WORD prompt edit `at night`->`at dawn`, also 233 tokens): rel RMS 0.01897 / 0.06666 excl. sink, cosine mean 0.99769 median 0.99963, 172 of 233 tokens above 0.999. So quantization moves the conditioning by the SAME total energy as rewriting a word of the prompt (6.85% vs 6.67%) but with the opposite SHAPE - diffuse over every token instead of concentrated on the words that changed. Cost: Q4 arm 40 s / 18.0 GiB peak; bf16 arm 40 s / 45.41 GiB uploaded / host conversion peak 0.0195 MiB / 51.95 GiB total peak, streamer counters `layers=50 tensors=400 direct=350 converted=200 fused=100` proving the shard path ran. `benchmark_binding=false` (no throughput claim). EXPLICITLY NOT ESTABLISHED: that the RENDER changes - nothing here measures the DiT's sensitivity to a 3.5-degree median rotation; the owed follow-up is a same-DiT/same-seed render A/B over the two saved embeds. | diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index dcef39ff4..3dcb855b9 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -78,7 +78,7 @@ models we already ship + benchmark. Full seam map + M0–M5 W-plan: | 13 | `ROAD-V1-D4` | **KV persistent state to disk, and external KV-cache provider interoperability with LMCache** (user-directed 2026-07-22: "let's do the KV persistent state to disk support, and LMCache support too", under the standing same-featureset-as-vLLM-and-better bar) | [`KV-OFFLOAD`](engine-matrix.md), [`KV-EXTERNAL-CACHE`](engine-matrix.md), [`KV-CONNECTORS`](engine-matrix.md), [coverage view §2](feature-matrix.md#2-kv-cache--memory), [LMCache quickstart](https://docs.lmcache.ai/getting_started/quickstart.html) | spike ACCEPTED [kv-persistence-lmcache.md](specs/kv-persistence-lmcache.md) — 60 enumerated features across `vllm/v1/kv_offload/`, the `KVConnectorBase_V1` ABI and the LMCache integration, each with a DONE/PARTIAL/MISSING verdict read out of our source. **The two halves of the user's ask are NOT the same kind of work.** Disk persistence is a faithful MIRROR job and is tractable: vLLM's `fs` tier is ~101 lines of `open`/`write`/`readv` with nothing Python-specific in the byte path, one raw file per block, temp-file + atomic rename under `O_DIRECT`. LMCache is NOT: the vLLM-facing glue is vendored in-tree (~2396 lines) but every file of it imports the EXTERNAL PyPI package at module scope, and the storage engine, wire protocol, config schema and CUDA-IPC handoff all live outside the tree with no upstream test that runs without it — so it is scoped as an interop STUDY with a go/no-go, never a from-scratch client. **REOPENED 2026-07-23 ([LMCache client wire analysis](specs/lmcache-cpp-client-connector.md)), and the "no specified wire protocol" half of that verdict is REFUTED by reading the LMCache package: vLLM connects to a RUNNING LMCache over TWO fully-specified portable wires — the `lm://` remote-store (plain TCP + fixed `struct` header + raw KV bytes, no ZMQ/msgpack/pickle/CUDA-IPC) and the MP server (ZMQ + `msgspec.msgpack` + CUDA-IPC, the user's "zmq" recollection). A from-scratch C++ client is FEASIBLE with ZERO `lmcache` in-process; both wires sidestep the hash blocker because LMCache keys on its own blake3 token hash. Recommend the `lm://` mode first. Residual risk is that LMCache is an unpinned moving target — an interop feature with a version-sync cost, not a mechanical core port. LMCACHE-CLIENT W1 LANDED 2026-07-23 (`CLAIM-LMCACHE-CPP-CLIENT`, `KV-EXTERNAL-CACHE` `SPIKE`→`ACTIVE`): the pure-CPU `lm://` wire codec — fixed-`struct` `ClientMetaMessage`/`ServerMetaMessage` framing, the `CacheEngineKey` string, the blake3 rolling token hash (vendored BLAKE3 1.5.5), and the `KV_2LTD` `[2,L,T,D]` repack — is BYTE/BIT-EXACT vs fixtures from the real Python codec (`test_lmcache_codec` 6/6, 2074 assertions), blake3 verified byte-identical on x86-64 + aarch64, and INERT (no call site; the connector is client-W3). LMCACHE-CLIENT W2 LANDED 2026-07-23 — the go/no-go PASSED: a blocking POSIX-socket `LMCacheRemoteClient` (PUT/GET/EXIST/HEALTH/LIST + `KV_2LTD` repack + `VT_LMCACHE_*` config) round-trips a REAL `lmcache.v1.server` (`8570aad`, run headless from source in a throwaway venv — torch imported before lmcache to dodge a torch circular import, the compiled `c_ops` ext stubbed as unused by the lm:// CPU store) byte-identical (`test_lmcache_client` 36/36), with BIDIRECTIONAL interop proven against LMCache's OWN Python protocol codec; the always-on CI gate is a same-binary C++ mock-server round-trip (45/45, no Python). STILL `ACTIVE`, not DONE. Resume at client-W3 (wire as a `KVConnector` over the parent W5 seam, then key-agreement + DGX every-axis gates).** **Blocking correction found in OUR source:** `NONE_HASH` is seeded from `std::random_device` with no escape hatch, so every block hash differs across processes and a content-addressed disk tier would score 0% hits on restart — we are WORSE than vLLM here, which at least exposes `PYTHONHASHSEED`. This also FALSIFIES the caching spike's §B2 claim that we are deterministic by construction. **Two upstream weaknesses recorded as beyond-parity targets:** the `fs` tier's `config.json` is written and never read (its only identity check is a path digest that omits checkpoint content, weight quantization, rope config and `sliding_window` — a silent-wrong-output hazard we will not copy), and the disk tier has no capacity accounting and no eviction. Three matrix rows `INVENTORIED` -> `SPIKE`; `SharedStorageConnector` found RENAMED to `ExampleConnector` and `P2pNcclConnector` found DELETED at the pin, both stale in the prior record **W1-W3 IMPLEMENTED 2026-07-22, CPU-only.** W1 deterministic block hashes: `init_none_hash` now resolves explicit arg > `$VLLM_PREFIX_CACHING_HASH_SEED` > `$PYTHONHASHSEED` > a fixed built-in default, so hashes are identical across processes with ZERO configuration — the blocking correction is CLOSED, and we now BEAT upstream on this axis rather than trailing it (upstream is random-by-default and documents `PYTHONHASHSEED` as the operator's problem). Proven by comparing hash chains emitted by SEPARATELY LAUNCHED processes, with a negative control confirming the opt-in `=random` mode genuinely disagrees. W2 CPU primary tier: `CachePolicy` (LRU + ARC) with the `ref_cnt == -1` tri-state and the ATOMIC evict, `CPUOffloadingManager` incl. the `prepare_store -> nullopt` skip control path, pinned backing store, and a side-queue event-polled device/host transfer worker. W3 disk `fs` tier: one raw file per block, temp-file + atomic rename publish, self-healing unlink, dual-queue read/write pool. **BOTH recorded upstream weaknesses are now EXCEEDED rather than merely noted** — the identity block is a VERIFIED header read on EVERY open that REFUSES on mismatch across 27 fields (tested per field, with a positive control), and the tier carries a byte budget with policy eviction honoured across restarts. `O_DIRECT` deliberately NOT ported (a header+payload file breaks its alignment requirement); the GIL-releasing batch-lookup C extension is unconditionally unnecessary without a GIL. `KV-OFFLOAD` `SPIKE` -> `PARTIAL`. **W4 IMPLEMENTED 2026-07-23:** the TIERING MANAGER (ONE manager over CPU primary + disk secondary — disk→CPU promotion RETRY→flush→HIT, cascade demotion, reset drains the secondary first and never resets it so a persisted cache survives) and the CONNECTOR/SCHEDULER HALF (`OffloadingConnector`, the semantics of `KVConnectorBase_V1`'s scheduler hooks — nullopt third state, `block_hashes` striding, load-before-compute — wired OPT-IN + DEFAULT-OFF into the scheduler). First measured offload speedup: a restarted-prefix workload through the REAL scheduler saved 32/48 prefill tokens (2/3 blocks HIT from disk), promoted bytes byte-identical to the cold store; identity refusal holds through a promotion. Ported the SEMANTICS not the Python plugin ABI (compile-time wiring); the full abstract ABI is W5. **W5 IMPLEMENTED 2026-07-23, CPU-only:** the connector seam is now a first-class C++ ABI — the abstract `KVConnector` base carrying the full scheduler + worker method set of `KVConnectorBase_V1` (the scheduler methods load-bearing, the worker hooks defaulted no-ops for our synchronous runner, documented), a compile-time `KVConnectorFactory` + `REGISTER_KV_CONNECTOR` (the C++ analogue of vLLM's `importlib` module path), and a `KVTransferConfig` selection surface (default `kv_connector` empty == no connector == zero behaviour change, `kv_role` validation, `fail`-default load policy). The W4 disk connector was refactored ONTO this base behaviour-identically — the restart-hit e2e reproduces byte-for-byte and a config-selected owning connector shortcuts prefill by the identical 32/48. `KV-CONNECTORS` `SPIKE`→`ACTIVE`. This closes the seam so LMCache client-W3 is 'implement the abstract `KVConnector` with the landed W2 `lm://` client'. **LMCACHE-CLIENT W3 LANDED 2026-07-23 — the `lm://` client wired as a `KVConnector` over the W5 seam (`LMCacheConnector`, `REGISTER_KV_CONNECTOR("LMCacheConnector", …)`, default OFF), the FIRST time the whole chain engine -> connector -> W2 client -> a running lm:// server -> back runs.** Scheduler side computes rolling-blake3 chunk hashes and `Exist`-probes the remote store for the longest cached prefix (synchronous `(n, false)`, mirroring `lmcache_connector.py:230-259`); worker `StoreChunk`/`LoadChunk` drive the W2 client with foreign-block REFUSAL. **Gate ACHIEVED = the connector-level round-trip: STORE a prefix -> a fresh "restarted" connector LOOKS UP + shortcuts prefill through the REAL scheduler (32/48 tokens saved) -> LOAD byte-identical; foreign-key REFUSAL; default-off INERT** (`test_lmcache_connector` 5 cases / 50 assertions vs an in-process mock; store->load ALSO GREEN vs a REAL `lmcache.v1.server` 8570aad, 16 assertions, `VT_LMCACHE_LIVE_*`). **LMCACHE-CLIENT W4 LANDED 2026-07-23 — REAL peer KEY-AGREEMENT + a peer->us interop LOAD, both PROVEN — the interop-correctness milestone is COMPLETE; `KV-EXTERNAL-CACHE` stays `ACTIVE` for the DGX full-model output-invariance + throughput arm.** The actual `lm://` key derivation is `ChunkedTokenDatabase` (NOT the blake3 MP hasher): chunk_size 256, a rolling prefix-hash over `(prefix_int, tuple(tokens), extra=())` keyed by vLLM's OWN hash (portable `sha256_cbor`), folded to uint64 each step, `NONE_HASH=fold8(sha256_cbor(str(PYTHONHASHSEED)))`. Mirrored byte-exact (`chunked_token_database.{h,cpp}`, reusing `CborValue`+`sha256_cbor`) and wired as connector `key_mode=kVllmSha256Cbor` (chunk 256) alongside W3's kept-green blake3 path. Key-agreement GREEN: `test_lmcache_key_agreement` 4/85 == the REAL lmcache `ChunkedTokenDatabase.process_tokens()` BYTE-FOR-BYTE (fixtures dumped from the unmodified real driver + vLLM's pinned `sha256_cbor`/`init_none_hash`), sample `meta-llama/Llama-3.1-8B@1@0@33d6862800fff40c@bfloat16`. Peer->us LOAD GREEN over the wire: a REAL lmcache `ChunkedTokenDatabase` derives a key + PUTs KV to a REAL `lmcache.v1.server`, our C++ re-derives the SAME key and GETs the 512 B byte-identical (`run_key_interop.sh`). ASan+UBSan clean. Text-only (mm-hash extra_keys deferred). **LMCACHE-CLIENT W5 LANDED 2026-07-24 — the LAST open arm, connector-ON full-model OUTPUT-INVARIANCE + throughput in a REAL generation loop, is CLOSED (spec gates 4/6 met).** The worker side is now wired into the engine: `GPUModelRunner::execute_model` calls `ConnectorLoadExternalKv` before the forward (writes the external-prefix KV into the allocated GPU blocks, load-before-compute) and `ConnectorStorePromptKv` after (stores each newly-complete prompt block), and `LoadedEngine` builds the connector from an `EngineParams` `KVTransferConfig` and wires it to BOTH the scheduler and the runner. **OUTPUT-INVARIANCE PROVEN on a real OPT-125m loop vs a live `lmcache.v1.server`: connector-ON generated tokens are BIT-IDENTICAL to connector-OFF (cold full prefill) — first-divergence index -1 — on BOTH (a) a store->restart->load cycle in one process AND (b) a genuinely cold second process that only hits the server; prefill saved on the hit = 48 tokens (3×16-token blocks).** `tests/vllm/models/test_lmcache_output_invariance.cpp` PASSES both modes via `scripts/lmcache/run_output_invariance.sh`. THROUGHPUT reported HONESTLY: on a 125M model the wall-clock delta is noise-dominated (fixed TCP/copy overhead ~ tiny compute saved), so NO binding speedup is claimed — a real speed number is owed by an every-axis grid on a larger model + long shared-prefix corpus (docs/BENCHMARKS.md). No-regression: OPT SACRED UNCHANGED default-off (6/6, 96/96, 63/63); connector unit tests green (codec 6/6, client 3/3, connector 5/5, key-agreement 4/4, kv_offload_connector 11/11); ASan+UBSan clean on the connector path; CUDA `-Werror` 0 warnings. Additive + default-off inert (all worker/loader changes are behind a null-connector guard)| `PARTIAL` | **W7 the one genuine beyond-parity item (imperative named per-sequence save/restore), which now has the verified header it depends on; and a binding every-axis LMCache throughput grid on a larger model vs vLLM's `--kv-transfer-config`.** W5 (the abstract ABI) is DONE; LMCache client W1 (codec) + W2 (client) + W3 (connector round-trip) + W4 (peer key-agreement + interop load) + W5 (full-model output-invariance) are DONE. **The benchmark blocker is CLEARED:** the caching spike's W1 prefix-cache counters landed earlier, so the W4 offload arm proved its hits | | 13a | `ROAD-V1-D4-APC` | **Prompt / prefix caching to full vLLM parity, then beyond (user-directed 2026-07-22: "same featureset of vLLM and better")** — the headline user-facing caching feature, previously mentioned only once in this roadmap despite being a shipped, default-ON behaviour for dense models | [`KV-PREFIX-CACHE`](engine-matrix.md), [`KV-BLOCK-POOL`](engine-matrix.md), [`KV-HYBRID-COORD`](engine-matrix.md), [`KV-MAMBA-ALIGN`](engine-matrix.md), [`KV-EVENTS`](engine-matrix.md), [`KV-PREFIX-MATCH-UNIT`](engine-matrix.md), [`ENG-CASCADE-ATTN`](engine-matrix.md), [coverage view §2](feature-matrix.md#2-kv-cache--memory) | umbrella spike ACCEPTED [prefix-prompt-caching-parity.md](specs/prefix-prompt-caching-parity.md) — enumerates the complete pinned-vLLM caching surface (38 features) with a per-feature DONE/PARTIAL/MISSING verdict grounded in our source. **The ported core is deeper than the record claimed** (chain hashing, block pool, all three coordinators, the full hybrid fixed-point intersection, four single-type managers); the real gaps are narrower and different: block-hash extra keys are a no-op stub, there are NO prefix-cache statistics at any level, KV events are inert, `cache_salt` and 3 of 4 hash algos are absent, and `reset_prefix_cache` is implemented but unreachable. Three matrix rows corrected, two of them in our favour. `ENG-CASCADE-ATTN` DISPOSITIONED as not owed (default-off, absent from the MRV2 runner we port, unreachable on Blackwell). llama.cpp comparison completed: its "prompt cache" is session/slot state serialization, strictly weaker than APC on every reuse axis, and vLLM already covers disk persistence via the `kv_offload` fs tier — the ONE genuine capability neither vLLM nor we have is an imperative named per-sequence save/restore **W1 IMPLEMENTED 2026-07-22: prefix-cache statistics exist for the first time.** `BaseCacheStats`/`PrefixCacheStats`/`CachingMetrics` ported 1:1 from `vllm/v1/metrics/stats.py:35-142`, recorded in `get_computed_blocks`, flagged by `reset_prefix_cache`, taken-and-swapped per step and folded into a 1000-request sliding window exposed on `Scheduler`/`EngineCore`/`LLMEngine`. Per the standing parity-enabler rule `log_stats` is DEFAULTED ON (mirroring upstream's `disable_log_stats=False`), so no benchmark arm is void for want of a counter. `Request::num_preemptions` un-deferred to feed the mutually-exclusive `preempted_*` triple. **FIRST MEASURED HIT RATE: 0.75** (1920 of 2560 queried tokens over 16 requests sharing a 128-token prefix), with a caching-OFF 0.0 negative control — the first demonstration in this project that APC actually serves cached tokens. The hard blocker on [`BACKEND-GATE-CUDA-SGLANG-PREFIX`](backend-matrix.md) is CLOSED **W2 DONE 2026-07-27 (`CLAIM-ROADMAP-D4APC`, CPU-gated on dgx GB10, NOT pushed):** `generate_block_hash_extra_keys` ported 1:1 (`kv_cache_utils.py:451-591`) — mm hash + LoRA name + `cache_salt`, fixed order lora->mm->salt (prompt_embeds deferred: no prompt-embeds path); `cache_salt`/`lora_name` carried on `Request`/`EngineCoreRequest`, set before the first hash in `FromEngineCoreRequest` (fixes a latent ordering bug). RED-first no-false-share PROVEN: with the stub a differently-salted request false-hits the prior tenant's 48 cached tokens (`n1==48`), with extra keys `n1==0`. Ported extra-key/ordering cases + hash- and manager-level no-false-share (`test_kv_cache_utils.cpp` 29/29, `test_kv_cache_manager.cpp` 10/10). **This unblocks the MM + LoRA cache consumers.** **W3 DONE 2026-07-27 (`CLAIM-ROADMAP-D4APC-W3`, dgx GB10, NOT pushed) — the FIRST-EVER cache-ON model gate:** on `Qwen/Qwen3-4B` (dense, full-attention, APC-default-ON — the vehicle the prior "vehicle-blocked" note missed) a shared-prefix workload runs APC-ON and APC-OFF through the full paged engine, gating token-identity + hits + prefill drop. **NO engine code changed** (`git diff --stat` = tests+scripts+goldens) ⇒ pure GATE over the already-shipped default-ON path; binary byte-identical ⇒ SACRED unaffected. RESULT (`test_qwen3_apc_e2e` 2/2, 84/84 asserts): APC-ON hits **2240/2777 (rate 0.807)**, APC-OFF 0; APC-ON == APC-OFF EXACT on 5/6 (the 1 diff a vLLM-confirmed 0.125-nat near-tie, RCA'd = attention-kernel-path near-tie flip, not a cache bug); **== vLLM-APC-ON** teacher-forced (APC-OFF 6/6 max gap 0.0 nats = exact argmax, APC-ON 6/6 max gap 0.125 nats, 0 outside top-20); **TTFT drop 70.1→39.9 ms = 1.76×** on a cache hit. Existing 4B SACRED gate 16/16 GREEN (no regression). Oracle vLLM 0.25.0 (0.26 venv broken — editable source disk-reclaimed; 4B byte-stable across the pin). | `DONE` (headline) | **Row DONE for the default dense APC path (block hashing incl. extra_keys, pool, coordinators, stats, scheduling, cache-ON e2e all gated).** Named NON-BLOCKING tails tracked in their own rows / future items: W4 KV events (`KV-EVENTS` — event GENERATION + `msgpack` PAYLOAD DONE 2026-07-27 `CLAIM-ROADMAP-D4-KV-EVENTS`, `SPIKE`→`ACTIVE`, byte-exact vs `msgspec`; live ZMQ transport + engine batch wiring DEFERRED), W5 partial-block primitive (upstream dead-code), W6 Mamba-`align` hybrid cache-on (`KV-MAMBA-ALIGN`, SPIKE — feeds `BACKEND-GATE-CUDA-SGLANG-PREFIX`), W7 `reset_prefix_cache` dev-endpoint + `--prefix-caching-hash-algo` + `skip_reading_prefix_cache`, W8/W9 the beyond-vLLM named session save/restore. The every-axis cache-on grid vs vLLM/SGLang is a separate perf follow-on under `ROAD-V1-A`. No `/metrics` route yet (`SERVE-METRICS`), so the hit rate is read from the engine API. **`--prefix-match-unit` (0.26-new fine-grained matching unit) W0 spike + W1 resolver LANDED 2026-07-28 (`CLAIM-PREFIX-MATCH-UNIT`, `KV-PREFIX-MATCH-UNIT` PARTIAL): `resolve_kv_cache_block_sizes` computes `hash_block_size = prefix_match_unit if set else gcd(group_block_sizes)`, RED-first unit-gated; config/CLI/ABI field (W2) + scheduler threading of `hash_block_size != block_size` (W3, needs the `KV-BLOCK-POOL` align path) + benchmark (W4) deferred.** | | 14 | `ROAD-V1-D5` | LoRA, local KV/weight offload, expert streaming, wider model zoo | [engine matrix](engine-matrix.md), [model matrix](model-matrix.md) | corrected expert-streaming spike accepted (`ENG-EXPERT-STREAM` READY): bank-only safetensors→Marlin bank, fixed contiguous cache slots matching Marlin dense strides, logical→slot remap after explicit router D2H, chunked C` working everywhere `--dit ` did; gated CPU-only (72/72, 54497 post-rebase) on index/name mapping and on the REAL 535-tensor geometry read from a sparse 13-shard release. The DEVICE streamer landed 2026-08-07 (`row/H3-BF16-SHARDED-STREAM`, spec §8.14): one tensor at a time, zero host buffer for the bulk, bit-exact vs the non-streamed reference (73/73, 55203). This unblocks the bf16-vs-quantized quality A/B; no render or speed number is claimed. Spec §8.13. **W-FP4a LANDED (CPU) 2026-08-06 (`row/H3-FP4-SPEED`)**: the NVFP4 DiT projections now keep FP4 PACKED and route through the shared `dense_nvfp4::MatmulNvfp4W4A16D` (Marlin W4A16 — vLLM's own forced-a16 selection; SAME kernel as Laguna/dense-Qwen3 NVFP4; no new quant code); fp4-vs-bf16 WIRING gate GREEN (62/62·30039, W4A16 dispatcher runs all 11 quantized GEMMs). **GB10 leg LANDED 2026-08-06 (`row/H3-FP4-GPU-E2E`, PR #64):** Marlin W4A16 RAN on sm_121a (`dense_gemms==11` default / `marlin_gemms==11` VT_MARLIN_DENSE=0, `fallback_gemms==0`), fp4-vs-bf16 BYTE-EXACT; fp4 is a MEMORY win (~16 vs ~66 GB), ~0.79–0.83× the bf16 arm per diffusion forward (compute-bound large M; 3.47× faster at small decode-like M). Real-checkpoint fp4-resident t2va e2e RUNS end-to-end (real 18.75 GB NVFP4 DiT + VAEs + GGUF Qwen3-VL-32B encoder → valid mp4/wav; DiT s/step 5.45/20.0/209 s @512/768/REF-768×1344-209f) but frames are a non-scene patch-grid at 12/20/50 steps → OPEN render-coherence bug (device VAE decode / denoise), separate from the fp4 speed work. vLLM-Omni serves NO quantized H3 (BF16-only) -> HW/loader-forced-indirect (4×B300 209f render 86.964 s vs 1×GB10 209 s/forward). | +| H3 | `ROAD-V1-H3` | **DIFFUSION generation — a new capability class.** MiniMax-H3 (`MiniMaxH3DiTModel`): omni-modal video+audio generation via a 50-step flow-matching denoise loop, ported from vLLM-Omni. Not autoregressive: no KV cache, sampler or logits. | [`MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit`](model-matrix.md) | [minimax-h3 spike](specs/minimax-h3.md) | `PARTIAL` | **W0-W2 landed 2026-08-03**: packed layout (fp64 grid bit-exact), latent packing, scheduler and the full DiT forward parity-gated vs the upstream vLLM-Omni modules at reduced dims (max abs diff 1.6e-7, 10/10 cases). **W2b device-resident forward LANDED (f32) and GPU-VERIFIED 2026-08-03** — the whole DiT graph runs with activations resident in device memory, gated vs the same upstream goldens on a Thor sm_110 GPU at video 1.49e-7 / audio 8.94e-8. Only 3 H3 kernels were needed; the port reuses the tuned shared ops. Next gate: bf16 stream + `vt::FusedChain` glue folds, then the FP4 path — which needs sm_121a, since sm_110 resolves every fp4/cutlass feature DISABLED. **HW verdict CORRECTED 2026-08-03: e2e is NOT blocked** — quantized H3 checkpoints fit (GGUF ~41 GB working set; NVFP4 likewise) and the ComfyUI-GGUF arm's 535-tensor manifest already resolves onto our contract, so e2e + a speed comparison are reachable. W7 `/v1/videos` still needs a NEW MP4/AV-encoder dependency decision. **bf16 13-SHARD RELEASE INDEXES 2026-08-07 (`row/H3-BF16-SHARDED-DIT`)**: the ORIGINAL 66.3 GB bf16 DiT (13 safetensors shards) is now resolvable through its own `model.safetensors.index.json`, with a host-f32 reference loader and `--dit ` working everywhere `--dit ` did; gated CPU-only (72/72, 54497 post-rebase) on index/name mapping and on the REAL 535-tensor geometry read from a sparse 13-shard release. The DEVICE streamer landed 2026-08-07 (`row/H3-BF16-SHARDED-STREAM`, spec §8.14): one tensor at a time, zero host buffer for the bulk, bit-exact vs the non-streamed reference (73/73, 55203). **ENCODER + THE NUMBER 2026-08-07 (`row/H3-ENC-BF16-COND-DIFF`, spec §8.15)**: the 14-shard bf16 text encoder streams too and `--encoder-only` runs it alone; measured over 233 tokens, Q4_K_M-vs-bf16 conditioning is cos 0.99745 mean / 6.85% rel RMS excl. sink / 3.5 deg median rotation — as much as a one-word prompt edit, but DIFFUSE. Whether the RENDER changes is NOT established (75/75, 55609). This unblocks the bf16-vs-quantized quality A/B; no render or speed number is claimed. Spec §8.13. **W-FP4a LANDED (CPU) 2026-08-06 (`row/H3-FP4-SPEED`)**: the NVFP4 DiT projections now keep FP4 PACKED and route through the shared `dense_nvfp4::MatmulNvfp4W4A16D` (Marlin W4A16 — vLLM's own forced-a16 selection; SAME kernel as Laguna/dense-Qwen3 NVFP4; no new quant code); fp4-vs-bf16 WIRING gate GREEN (62/62·30039, W4A16 dispatcher runs all 11 quantized GEMMs). **GB10 leg LANDED 2026-08-06 (`row/H3-FP4-GPU-E2E`, PR #64):** Marlin W4A16 RAN on sm_121a (`dense_gemms==11` default / `marlin_gemms==11` VT_MARLIN_DENSE=0, `fallback_gemms==0`), fp4-vs-bf16 BYTE-EXACT; fp4 is a MEMORY win (~16 vs ~66 GB), ~0.79–0.83× the bf16 arm per diffusion forward (compute-bound large M; 3.47× faster at small decode-like M). Real-checkpoint fp4-resident t2va e2e RUNS end-to-end (real 18.75 GB NVFP4 DiT + VAEs + GGUF Qwen3-VL-32B encoder → valid mp4/wav; DiT s/step 5.45/20.0/209 s @512/768/REF-768×1344-209f) but frames are a non-scene patch-grid at 12/20/50 steps → OPEN render-coherence bug (device VAE decode / denoise), separate from the fp4 speed work. vLLM-Omni serves NO quantized H3 (BF16-only) -> HW/loader-forced-indirect (4×B300 209f render 86.964 s vs 1×GB10 209 s/forward). | | 15 | `ROAD-V1-D6` | **llama.cpp device breadth folded into scope (user-directed 2026-08-05):** the 11 ggml backends vLLM has no platform for — cann, musa, opencl, openvino, rpc, webgpu, zdnn, zendnn, hexagon, blas, virtgpu — inventoried as `BACKEND-GGML-*`. **SPIKES FIRST:** no implementation before each row's `.agents/specs/.md` clears the spike contract, per the standing directive. vLLM stays the mirror source; llama.cpp is the breadth reference. | [backend matrix](backend-matrix.md) | ☐ per-row spike required | `INVENTORIED` | first spike accepted | An area row cannot enter `READY` without a real spike under `specs/`, and cannot diff --git a/.agents/specs/minimax-h3.md b/.agents/specs/minimax-h3.md index 0bf18181d..16c6e59c2 100644 --- a/.agents/specs/minimax-h3.md +++ b/.agents/specs/minimax-h3.md @@ -1003,7 +1003,7 @@ runnable; it has not been run. ## 8.15 The bf16 TEXT ENCODER — 14 shards, 63 GB, streamed; and `--encoder-only` (2026-08-06, `row/H3-ENC-BF16-COND-DIFF`) -**Why.** §8.6 made the full-precision *DiT* loadable, but every H3 render — including +**Why.** §8.13/§8.14 made the full-precision *DiT* loadable, but every H3 render — including the ones whose output looks competent-but-generic — conditioned on a **Q4_K_M** text encoder (`enc_q4km.gguf`, 14.6 GB, Qwen3-VL-32B). Nobody had ever measured the encoder's contribution. That matters because weak conditioning and a @@ -1051,7 +1051,7 @@ safetensors shards + `model.safetensors.index.json`, 63 GB**. Both encoder paths now go through ONE helper, so the conditioning a render consumes and the conditioning the A/B measures are produced by the same code. -**Gates (CPU, `test_minimax_h3` 70/70, 49706 assertions, up from 68/68).** +**Gates (CPU, re-run after the rebase onto `row/H3-BF16-SHARDED-STREAM`: `test_minimax_h3` 75/75, 55609 assertions).** 1. *Resolve, fuse and stream*: a synthetic 4-shard encoder at the REAL name spellings (`model.language_model.layers.N.*`, `model.visual.*`, `lm_head.weight`). Geometry from shapes matches; every fused view's bytes are `memcmp`-exact against @@ -1070,3 +1070,62 @@ safetensors shards + `model.safetensors.index.json`, 63 GB**. respectively (asserted), and the two full encoder forwards are **BIT-IDENTICAL** (`memcmp == 0`), not merely close. Without this, "we measured what quantizing the encoder costs" would be confounded by what the widening itself did. + +### 8.8 THE NUMBER — what Q4_K_M does to the conditioning (2026-08-06, `row/H3-ENC-BF16-COND-DIFF`, Thor sm_110, build `d1085374` (built and measured as `d1085374`, amended for the row-branch trailer; IDENTICAL tree `dd9283cf`, so the measurement binary IS this commit)) + +**Method.** Same prompt (`wuxia.txt`, **233 tokens**), same tokenizer, same 50-layer +truncation, same `MiniMaxH3EncoderTextForwardDevice`, same f32 activations — only +the weight bytes differ. Both arms self-report identical geometry +(`layers=50 hidden=5120 heads=64 kv_heads=8 head_dim=128 ffn=25600`), which is what +establishes they are the same model. Conditioning is `[233, 5120]` f32 via +`--encoder-only --save-embeds`. A CALIBRATION arm encodes a ONE-WORD edit of the +same prompt with the bf16 encoder (`bamboo forest at night` -> `at dawn`, also 233 +tokens), because a cosine has no meaning without a yardstick. + +| | max\|diff\| | RMS | rel RMS | rel RMS excl. sink | cos min | cos mean | cos median | angle mean | angle max | +|---|---:|---:|---:|---:|---:|---:|---:|---:|---:| +| **Q4_K_M vs bf16** | 154.0 | 0.5045 | **0.03403** | **0.06849** | 0.90916 | **0.99745** | 0.99810 | 3.793° | 24.61° | +| bf16, one-word edit | 31.1 | 0.2812 | 0.01897 | 0.06666 | 0.84736 | 0.99769 | 0.99963 | 2.228° | 32.07° | + +1. **Not a scale change.** Q4 conditioning is uniformly ~1% smaller (norm ratio mean + 0.99010) but the best single global rescale removes almost none of the difference + (0.03403 -> 0.03280). It is DIRECTIONAL, the kind that matters. +2. **Its energy is on par with a one-word prompt edit** (6.85% vs 6.67% excluding the + sink token). Quantizing the encoder moves the conditioning about as much as + rewriting a word of the prompt. +3. **The SHAPE is opposite.** The edit is SPARSE — 172/233 tokens stay above cosine + 0.999 (median rotation 0.16°), the change concentrating on ~6 tokens, max 32°. + Quantization is DIFFUSE — 232/233 tokens fall below 0.999, EVERY token rotates a + few degrees (median 3.5°), one token by 24.6°. A smear everywhere, not a different + prompt. +4. **`max|diff|` 154 is the attention SINK, not corruption.** Token 0 has norm 15,522 + against a 366 mean (42x) and carries 68% of the total squared error, yet its + direction is nearly untouched (cosine 0.99962) — the channel-wise magnitude-outlier + behaviour ComfyUI PR 15298 attributes to the partial split-half RoPE, showing up + concretely. This is why the sink-excluded column is the honest aggregate. + +**Verdict.** Q4_K_M does real, measurable, directional damage — comparable in +magnitude to editing the prompt — but DIFFUSELY. A uniform few-degree rotation of +every token is the signature that blunts fine-grained compositional instruction +(coverage, blocking, staging) toward a prompt's average semantics, which is exactly +the "competent but generic" symptom. The bf16 encoder is worth a render A/B. + +**What this does NOT establish.** It does not prove the RENDER changes; nothing here +measures the DiT's sensitivity to a 3.5°-median rotation. The owed follow-up is a +byte-identical-everything-else render A/B (same DiT, seed, steps; +`--prompt-embeds cond_q4km.bin` vs `cond_bf16.bin`), which is exactly what +`--save-embeds`/`--prompt-embeds` make controllable. + +**REPRODUCED.** Both arms were re-run from scratch (fresh process, fresh +load of the checkpoint) and each produced a BYTE-IDENTICAL conditioning file: +`cond_q4km.bin` md5 `a331232096ef1da2628f885950b2fc55` and `cond_bf16.bin` md5 +`9c096b63b9bd07f604daebb2fc090f46` on both runs. So every number above is +deterministic, not a sample — there is no noise band to argue about, and a +future change to either path shows up as an md5 change. + +**Cost.** Q4_K_M arm 40 s wall, peak **18.0 GiB**. bf16 arm 40 s wall (35 s +streaming), **45.41 GiB** uploaded, host conversion peak **0.0195 MiB** (one norm — +the projections never touch a host buffer), total peak **51.95 GiB** of the 122 GiB +UNIFIED pool. Streamer counters on the real checkpoint: +`layers=50 tensors=400 direct=350 converted=200 fused=100`, i.e. the shard path ran +and every projection took the no-host-copy upload. diff --git a/.agents/state.md b/.agents/state.md index 3220e2f58..d961007ff 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -41184,11 +41184,11 @@ quality question; nothing here claims it. `row/H3-ENC-BF16-COND-DIFF` (helper, branched from `origin/main` `ad231615`, with -§8.6's `MiniMaxH3ShardedCheckpoint` **cherry-picked** from `row/H3-BF16-SHARDED-DIT` +§8.13's `MiniMaxH3ShardedCheckpoint` (landed separately as `row/H3-BF16-SHARDED-DIT` + `row/H3-BF16-SHARDED-STREAM`) `1a46ff17` rather than mirrored - the shard index resolver is exactly the piece this row needs and duplicating 222 lines of it would drift). -**Why.** §8.6 made the full-precision *DiT* loadable, but every H3 render so far - +**Why.** §8.13/§8.14 made the full-precision *DiT* loadable, but every H3 render so far - including the ones that look competent-but-generic - conditioned on a **Q4_K_M** encoder (`enc_q4km.gguf`, 14.6 GB, Qwen3-VL-32B), and the encoder's contribution had NEVER been measured. Weak conditioning and quantization-damaged conditioning look @@ -41199,7 +41199,7 @@ produces channel-wise magnitude outliers that corrupt even INT8). The blocker wa mechanical - `--encoder` took only a GGUF, the unquantized tower ships as 14 safetensors shards. -**What landed** (full detail: [specs/minimax-h3.md](specs/minimax-h3.md) §8.7). +**What landed** (full detail: [specs/minimax-h3.md](specs/minimax-h3.md) §8.15). `MiniMaxH3EncoderConfigFromShards` (geometry from the index's SHAPES alone, SAME recovery rules and SAME non-shape defaults as the GGUF loader, so an A/B cannot be comparing two RoPEs), `StreamMiniMaxH3EncoderShardsToDevice` (fills the same @@ -41219,7 +41219,7 @@ is residency, not numerics - and it is GATED as such. `--encoder-only` matters f the same reason: the DiT was loaded FIRST in the normal path, so conditioning alone used to cost ~96 GiB peak instead of ~49 GiB. -**Gates (CPU, `test_minimax_h3` 70/70 cases / 49706 assertions, up from 68/68).** +**Gates (CPU, re-run POST-REBASE: `test_minimax_h3` 75/75 cases / 55609 assertions).** (1) resolve+fuse+stream over a synthetic 4-shard encoder at the REAL name spellings - every fused view `memcmp`-exact against `q ++ k ++ v` and `gate ++ up` for every layer, unfused projections byte-exact, separate names GONE, `norm.weight`/`lm_head`/ @@ -41232,15 +41232,86 @@ EXACT - the same checkpoint written BF16 and F32 (bf16-rounded values) streams t `kBF16` and `kF32` views respectively and the two full forwards are BIT-IDENTICAL (`memcmp == 0`), so the measurement below cannot be confounded by the widening. -**Record repair done in passing.** `docs/STATUS.md` was 41 chars OVER its -`check-public-doc-tables` ratchet on `origin/main` (284114 vs 284073) - a pre-existing -red. The H3 row's superseded narrative was collapsed to the binding result, which -brings the page back inside the ratchet. - -**Pre-existing red, NOT this row's:** `check-fusion-consistency` / -`test_check_fusion_consistency` (`minimax_h3_video_vae_device` gemm-merge drift), -verified RED on a clean `origin/main` tree before this work. +**Rebase note.** This row was written off `ad231615` and REBASED onto the landed +`row/H3-BF16-SHARDED-STREAM` before landing; the gate numbers here are the POST-REBASE +re-run. Two claims from the pre-rebase write-up no longer apply and are withdrawn: the +`docs/STATUS.md` ratchet red (284114 vs 284073) was repaired on main in the meantime, +and the earlier 70/70 count is superseded by the 75/75 recorded above. `origin/main` had ALSO grown +a vision-conditioning path in the encoder block (`--cond-image`, DeepStack taps) that +this row's text-only `EncodeH3Prompt` refactor would have silently dropped; the refactor +is now scoped to `--encoder-only` and main's inline block is kept for the run path, so +both capabilities survive. Next: the measurement itself - encode the wuxia prompt with both encoders on the Thor GPU and diff the `[tokens, 5120]` conditioning (max|diff|, RMS, relative RMS, per-token cosine). Numbers land in the same row. + +## 2026-08-07T13:40 - H3 THE NUMBER: Q4_K_M encoder moves the conditioning as much as a ONE-WORD prompt edit, but DIFFUSELY (row/H3-ENC-BF16-COND-DIFF, Thor) + + + +The measurement the loader existed for. Full tables: +[benchmark-record.md](benchmark-record.md) and +[specs/minimax-h3.md](specs/minimax-h3.md) §8.15. Build `d1085374` (built and measured as `d1085374`, amended for the row-branch trailer; IDENTICAL tree `dd9283cf`, so the measurement binary IS this commit), Thor sm_110, +CUDA 13.0.1 container, GPU idle (the LocalAI render had finished; it was never +touched). + +**Controlled the way it has to be.** Same prompt (`wuxia.txt`, 233 tokens), same +tokenizer, same 50-layer truncation, same `MiniMaxH3EncoderTextForwardDevice`, same +f32 activations - only the weight bytes differ. Both arms self-report IDENTICAL +geometry (50 / 5120 / 64 / 8 / 128 / 25600), which is what proves they are the same +model rather than two checkpoints that merely share a name. + +**A CALIBRATION arm, because a cosine means nothing without a yardstick.** The bf16 +encoder also encoded a ONE-WORD edit of the same prompt (`at night` -> `at dawn`, +also 233 tokens). + +| | rel RMS | rel RMS excl. sink | cos mean | cos median | cos min | angle median | angle max | +|---|---:|---:|---:|---:|---:|---:|---:| +| **Q4_K_M vs bf16** | **0.03403** | **0.06849** | **0.99745** | 0.99810 | 0.90916 | 3.535° | 24.61° | +| bf16 one-word edit | 0.01897 | 0.06666 | 0.99769 | 0.99963 | 0.84736 | 1.565° | 32.07° | + +max|diff| 154.0 / RMS 0.5045 (quant) vs 31.1 / 0.2812 (edit). + +**The read.** (1) NOT a scale change - the Q4 conditioning is uniformly ~1% smaller +(norm ratio 0.99010) but the best global rescale removes almost none of it +(0.03403 -> 0.03280), so it is DIRECTIONAL. (2) Its total energy is ON PAR with a +one-word prompt edit (6.85% vs 6.67% excluding the sink token). (3) The SHAPE is +opposite: the edit is SPARSE (172/233 tokens above cosine 0.999, median rotation +0.16°, concentrated on ~6 tokens), quantization is DIFFUSE (232/233 below 0.999, +median 3.5°, every token). (4) `max|diff|` 154 is the ATTENTION SINK, not +corruption: token 0 has norm 15,522 vs a 366 mean (42x) and carries 68% of the total +squared error while its DIRECTION is intact (cosine 0.99962) - ComfyUI PR 15298's +channel-wise magnitude outliers, concretely. + +**Verdict.** Q4_K_M does real, measurable, directional damage, comparable in +magnitude to editing the prompt, but diffusely - a uniform few-degree rotation of +every token, which is the signature that blunts fine-grained compositional +instruction toward a prompt's average semantics. That matches the +"competent but generic" symptom, so the bf16 encoder is worth a render A/B. + +**Explicitly NOT established: that the RENDER changes.** Nothing here measures the +DiT's sensitivity to a 3.5°-median rotation. The owed follow-up is a +byte-identical-everything-else render A/B - same DiT, seed and steps, with +`--prompt-embeds cond_q4km.bin` vs `cond_bf16.bin` (both saved on the box), which is +exactly what the save/replay seam makes controllable. + +**REPRODUCED.** Both arms were re-run from scratch (fresh process, fresh +checkpoint load) and each produced a BYTE-IDENTICAL file - `cond_q4km.bin` md5 +`a331232096ef1da2628f885950b2fc55`, `cond_bf16.bin` md5 +`9c096b63b9bd07f604daebb2fc090f46` on both runs. Every number above is +deterministic, not a sample; there is no noise band to argue about. + +**Cost, recorded for the next run.** Q4_K_M 40 s / 18.0 GiB peak; bf16 40 s (35 s +streaming) / 45.41 GiB uploaded / host conversion peak 0.0195 MiB (ONE norm - the +projections never touch a host buffer) / 51.95 GiB total peak of the 122 GiB unified +pool. Real-checkpoint streamer counters: `layers=50 tensors=400 direct=350 +converted=200 fused=100`. + +Also fixed this session: the streamer uploaded norms from a conversion buffer scoped +INSIDE its branch, i.e. freed before the Synchronize that guarantees the async +`cudaMemcpyAsync` landed. Norms are ~20 KB so CUDA would usually stage them and get +away with it, which is exactly why it could not ship (`d1085374`). + +Next: the render A/B, and the same treatment for the bf16 DiT (§8.13's loader is +already in this branch). diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index c827a51fd..a48909916 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -320,6 +320,7 @@ built on it rather than keeping the flattering one. | MiniMax-H3 render duration (audio halving) | **CLOSED**: a 124-frame render silently muxed as **61**. The decoded audio ran half the video's duration and the muxer passes `-shortest`. Every structural check passed: shapes were self-consistent, just halved | Gated on the duration invariant (latent steps / 40 Hz equals the video duration), not on shape self-consistency, which a halved pipeline satisfies | | MiniMax-H3 quantization floor | **Use Q4_K_M, not Q3_K_M.** H3's split-half RoPE produces channel-wise magnitude outliers 3-bit cannot hold; a controlled A/B (same prompt, seed, code) turned a murky lattice-covered silhouette into a photoreal close-up | Per-tensor mixed precision, if a smaller footprint is ever owed | | MiniMax-H3 image conditioning (`row/H3-CONDITIONED-E2E`, `row/H3-VISION-SCATTER`, `row/H3-REF2VA-ASSEMBLY`, `row/H3-NVFP4-LOADER-DIFF`, `row/H3-NVFP4-STREAM-DIFF`) | **fl2va COHERENT; ref2va NVFP4 nibble loader bug FIXED (byte-verified); grid residual DIAGNOSED §8.12** | Activation diff + fingerprints (PR #95): NO load-path defect; all weights/islands/RoPE quant-noise-close to the coherent GGUF; grid = community-NVFP4 quant fidelity, not a loader fix. See benchmark-record | +| MiniMax-H3 encoder quantization (`H3-ENC-BF16-COND-DIFF`) | **Measured Thor (`d1085374`).** Q4_K_M vs bf16 encoder, same 233-token prompt, same forward: rel RMS **0.0340** (0.0685 excl. sink), per-token cosine mean **0.99745** / min 0.909, median rotation **3.5°** | NOT a scale change (best rescale 0.0340->0.0328). Same energy as a ONE-WORD prompt edit but DIFFUSE: 232/233 tokens rotate vs 172/233 untouched. Render A/B owed. Detail: benchmark-record | | MXFP4 Qwen3-8B (W4A16 Marlin) | **`KERNEL-MARLIN-DENSE-EXEC` x3 (dense-ON default): c1 1.020, c2/c4/c8 0.962/0.966/0.969, GPU mem 2.63x less** (beats #51 1.005/0.925/0.939/0.953 EVERY axis); #44 3/3, 32B-NVFP4A16 6/6; -Werror test-guard fixes x2 | **VT_MARLIN_DENSE default-ON** (+951us). `FLASH-PTXAS` #82: cuModule A/B ties our+vLLM PTX across ptxas 13.0/13.2/driver-JIT (~144us); +10us is engine CONTEXT not codegen, no ptxas lever/flip (retires #75) | | Vulkan vs llama.cpp Vulkan (`BENCH-VK-LLAMA`) | **llama.cpp denominator MEASURED** on GB10 at full strength (`NV_coopmat2`): qwen3-0.6B F16 pp128 **11,730 t/s**, tg32 **161.4 t/s**. Our arm NOT RUN. [Detail](../.agents/benchmark-record.md) | BLOCKED: no model on dgx loads in BOTH engines (our GGUF path rejects arch `qwen3`; llama.cpp cannot load our Laguna file). Needs a shared arch or `VK-D` | | 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 dbf7c28e6..f53ea7b2b 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -133,7 +133,7 @@ they sit outside the gated list above. |---|---|---|---| | Voxtral audio (`VoxtralForConditionalGeneration`) | Voxtral-Mini-3B-2507 | near-tie-robust 16/16 vs vLLM 0.25.0 | decode 0.97x (beats vLLM); encoder TTFT ~17x, pending | | Whisper audio encoder | openai/whisper-small; whisper-large-v3 (Voxtral cfg) | encoder tower 77/77; large-v3 tower 203/203 | pending | -| MiniMax-H3 DiT (`MiniMaxH3DiTModel`, vllm-omni lane) | MiniMax-H3 (33.1B video+audio) | portable 72/72; t2va+fl2va COHERENT; ref2va NVFP4 grid = the community checkpoint's own quant fidelity, NO loader bug (§8.12); loads GGUF + NVFP4, STREAMS the bf16 13-shard DiT and the 14-shard bf16 text encoder | FP4/Marlin landed; ref2va NVFP4 render blocked on checkpoint quant (needs official modelopt NVFP4), speed pending; no bf16 render yet | +| MiniMax-H3 DiT (`MiniMaxH3DiTModel`, vllm-omni lane) | MiniMax-H3 (33.1B video+audio) | portable 75/75; t2va+fl2va COHERENT; ref2va NVFP4 grid = that checkpoint's own quant fidelity (§8.12); GGUF + NVFP4 + bf16 shards (DiT and encoder) all stream; Q4_K_M-vs-bf16 conditioning MEASURED, cos 0.99745 mean | FP4/Marlin landed; ref2va NVFP4 render blocked on checkpoint quant (needs official modelopt NVFP4), speed pending; no bf16 render yet | | MTP speculator | Qwen3.6-27B, Qwen3.6-35B-A3B | token-identical to vLLM `mtp` at c1 | ~4% faster c1; +16% output tput (MoE) | | DFlash block-diffusion | Qwen3 (DFlash draft) | near-tie e2e 27/27 vs vLLM | 2.9x over spec-off, 1.003x vs vLLM DFlash-on | | DeepSeek-V4 MTP | DeepSeek-V4-Flash (nextn head) | lossless 5/5; real-model weight-blocked | pending | diff --git a/docs/STATUS.md b/docs/STATUS.md index 308b82229..78bf980c9 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