From ed0324ae0c73b9fec077c721db0fc62ca4a850d7 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 14:52:25 +0000 Subject: [PATCH 1/3] =?UTF-8?q?test(minimax=5Fh3):=20W-FP4a=20GB10=20leg?= =?UTF-8?q?=20=E2=80=94=20CUDA=20Marlin-W4A16=20speed=20case=20+=20fp4-res?= =?UTF-8?q?ident=20driver=20flag?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The CPU wiring gate cannot reach the sm_121a FP4 tensor-core route: the existing "an NVFP4 checkpoint loads into a runnable DiT" case runs the fp4/bf16 forwards on a CPU queue, so on a CUDA build the W4A16 dispatcher still lands in `fallback_gemms`, never `marlin_gemms`. This adds the GB10 leg (spec 8.4a): - Factor the synthetic compressed-tensors NVFP4 file writer out of the CPU case into `WriteMiniMaxH3Nvfp4File(params, path)` (behavior-identical; CPU gate still 63/63 · 30039 assertions), generalized to loop num_layers/token_refiner_layers. - New CUDA case "the NVFP4 fp4 forward runs Marlin W4A16 on CUDA (speed)": builds the file at REAL H3 geometry (one AdaLN block + one refiner, so exactly 11 quantized GEMMs), streams both the bf16 and the fp4-resident arms to the device off the SAME file, runs both device forwards, and asserts `marlin_gemms == 11` (the this-path-ran counter — on CPU these are fallback_gemms). Records the real fp4-vs-bf16 numeric delta (Marlin bf16-TC accumulate vs the bf16 arm's dequant+MatmulBT — same fp4 bytes, so the delta is reduction slack, not quant), steady per-forward times for both arms (median over reps, cold leg discarded), and a per-GEMM microbench of qkv/out/fc1/fc2 (Marlin fp4 vs the bf16 twin). Env-tunable (H3_FP4_{LT,LH,LW,AT,AC,TEXT,REPS}) so the one binary sweeps sequence length and reps on dgx. SKIPs cleanly with no CUDA backend. - examples/minimax-h3-gen: `--fp4-resident` routes the NVFP4 safetensors CUDA path through StreamMiniMaxH3Nvfp4ToDeviceFp4 (packed FP4 on device, ~16 GB vs the bf16 stream's ~66 GB) so the real-checkpoint t2va e2e fits the 119 GiB pool. GB10 build+run of the CUDA case, and the real-checkpoint render, land in follow-up commits with the measured numbers. Row row/H3-FP4-GPU-E2E, ROAD-V1-H3. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] --- docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- examples/minimax_h3_gen/main.cpp | 13 +- tests/vllm/models/test_minimax_h3.cpp | 407 +++++++++++++++++++------- 4 files changed, 321 insertions(+), 103 deletions(-) diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 5a7e8121..d4b4442e 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -296,7 +296,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) | Full-model GB10 e2e RUNS (bf16-resident §13), NEAR-TIE 106/128, pool math CLOSES; default OFF | Full model RUNS on GB10 (bf16-resident, RSS peak 1.7 GiB, min-avail 21 GiB, no OOM). Token NEAR-TIE 106/128 (6/8 prompts exact, numerics vs deterministic oracle). 1.59 tok/s. Detail: spec §13 | | vLLM 0.26 re-benchmark | Pending | Re-run the binding grids on the advanced pin | -| MiniMax-H3 FP4 speed (W-FP4a) | Pending. fp4-resident Marlin-W4A16 routing CPU-landed (62/62); GB10 delta + per-step unmeasured (disk window); real e2e disk-blocked; vLLM-Omni has no quantized H3 (BF16-only) | Build CUDA `test_minimax_h3` on dgx, run the NVFP4 case (Marlin via `marlin_gemms`), capture delta + s/step. Detail: benchmark-record + spec §8 | +| MiniMax-H3 FP4 speed (W-FP4a) | Pending (measurement in flight, `row/H3-FP4-GPU-E2E`). fp4-resident Marlin-W4A16 routing CPU-landed (63/63); the GB10 CUDA speed case (asserts `marlin_gemms==11`, records fp4-vs-bf16 delta + per-forward/per-GEMM times) and `--fp4-resident` driver flag landed; dgx build+run + real e2e in flight; vLLM-Omni has no quantized H3 (BF16-only) | Build+run CUDA `test_minimax_h3` on dgx (Marlin via `marlin_gemms`), capture delta + s/step; fp4-resident real-checkpoint t2va e2e. Detail: benchmark-record + spec §8 | | 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**. `QUANT-CT-MXFP4-FINAL-STACK` TERMINAL: 2 last levers exhausted (num_splits cap gated-OFF c1-only; glue folds via FusedChain, residual out-of-catalog). c2-c8 GPU-intrinsic; see record | | SGLang floor arms | Never ran | Both arms of the SGLang comparison | | cuBLAS invocation-parity guard | CI guard landed (CPU); `kGemvHeuristicAlgos` refactor build-verify owed | `nvcc` rebuild + SACRED gate on dgx | diff --git a/docs/STATUS.md b/docs/STATUS.md index 1daac77a..830b5026 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -73,7 +73,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 device (bf16) in " + std::cerr << " streamed NVFP4 DiT -> device (" << (fp4_resident ? "fp4-resident" : "bf16") + << ") in " << std::chrono::duration(std::chrono::steady_clock::now() - t0).count() << " s\n"; } else { diff --git a/tests/vllm/models/test_minimax_h3.cpp b/tests/vllm/models/test_minimax_h3.cpp index 7c297187..ec80d9f1 100644 --- a/tests/vllm/models/test_minimax_h3.cpp +++ b/tests/vllm/models/test_minimax_h3.cpp @@ -17,8 +17,11 @@ #include +#include +#include #include #include +#include #include #include #include @@ -358,6 +361,114 @@ std::unique_ptr BuildDitForwardCase(const MiniMaxH3DitParams& p) return c; } +// Serialize a synthetic compressed-tensors NVFP4 (W4A16) MiniMax-H3 DiT file at +// the geometry `want`, exactly as the real `lilcheaty/MiniMax-H3-NVFP4` file +// stores it: quantized projections as U8 packed [out, in/2] + E4M3 group-16 +// weight_scale + F32 weight_scale_2; islands (patch/time/output/norms) plain F32. +// Factored out of the CPU "NVFP4 checkpoint loads" case so the CUDA speed case +// can build the SAME file at real geometry without duplicating 100 lines. Both +// callers set num_layers == token_refiner_num_layers == 1, which is what makes +// the quantized-GEMM count exactly 11 (refiner 4 + block 5 + condition + final). +void WriteMiniMaxH3Nvfp4File(const MiniMaxH3DitParams& want, const std::string& path) { + struct Entry { + std::string name; + std::string dtype; + std::vector shape; + std::string bytes; + }; + std::vector entries; + + auto add_plain = [&](const std::string& name, const std::vector& shape) { + int64_t numel = 1; + for (int64_t d : shape) numel *= d; + const std::vector values = MakeParam("nvfp4." + name, numel, 0.1); + entries.push_back({name, "F32", shape, + std::string(reinterpret_cast(values.data()), + values.size() * sizeof(float))}); + }; + auto add_quant = [&](const std::string& name, int64_t out_dim, int64_t in_dim) { + REQUIRE(in_dim % 16 == 0); + std::string packed(static_cast(out_dim * (in_dim / 2)), '\0'); + for (size_t i = 0; i < packed.size(); ++i) { + packed[i] = static_cast((i * 37 + 11) & 0xFF); // deterministic nibbles + } + std::string scales(static_cast(out_dim * (in_dim / 16)), '\0'); + for (size_t i = 0; i < scales.size(); ++i) { + scales[i] = static_cast(0x38); // e4m3 ~ 1.0 + } + const float global = 0.5f; + entries.push_back({name, "U8", {out_dim, in_dim / 2}, packed}); + entries.push_back({name + "_scale", "F8_E4M3", {out_dim, in_dim / 16}, scales}); + entries.push_back({name + "_scale_2", "F32", {}, + std::string(reinterpret_cast(&global), sizeof(float))}); + }; + + const int64_t inner = want.num_attention_heads * want.attention_head_dim; + const int64_t video_width = want.video_row_width(); + // Islands stay unquantized, exactly as the real checkpoint has them. + add_plain("video_patch_proj.weight", {want.hidden_size, video_width}); + add_plain("video_patch_proj.bias", {want.hidden_size}); + add_plain("audio_patch_proj.weight", {want.hidden_size, want.audio_latents_dim}); + add_plain("audio_patch_proj.bias", {want.hidden_size}); + add_plain("condition_proj.bias", {want.hidden_size}); + add_plain("time_embedder.proj_in.weight", {want.time_embed_hidden_size, want.timestep_input_dim}); + add_plain("time_embedder.proj_in.bias", {want.time_embed_hidden_size}); + add_plain("time_embedder.proj_out.weight", {want.time_embed_dim, want.time_embed_hidden_size}); + add_plain("time_embedder.proj_out.bias", {want.time_embed_dim}); + add_plain("rope.inv_freq", {want.rope_inv_freq_len}); + add_quant("condition_proj.weight", want.hidden_size, want.text_dim); + auto add_block = [&](const std::string& prefix, bool with_adaln) { + add_plain(prefix + ".norm1.weight", {want.hidden_size}); + add_plain(prefix + ".norm2.weight", {want.hidden_size}); + add_plain(prefix + ".attn.q_norm.weight", {want.attention_head_dim}); + add_plain(prefix + ".attn.k_norm.weight", {want.attention_head_dim}); + add_quant(prefix + ".attn.qkv_proj.weight", 3 * inner, want.hidden_size); + add_quant(prefix + ".attn.out_proj.weight", want.hidden_size, inner); + add_quant(prefix + ".mlp.fc1.weight", 2 * want.ffn_hidden_size, want.hidden_size); + add_quant(prefix + ".mlp.fc2.weight", want.hidden_size, want.ffn_hidden_size); + if (with_adaln) { + add_quant(prefix + ".adaln_proj.linear.weight", want.adaln_out_features, want.time_embed_dim); + add_plain(prefix + ".adaln_proj.linear.bias", {want.adaln_out_features}); + } + }; + for (int64_t i = 0; i < want.token_refiner_num_layers; ++i) + add_block("token_refiner.blocks." + std::to_string(i), false); + add_plain("token_refiner.final_norm.weight", {want.hidden_size}); + for (int64_t i = 0; i < want.num_layers; ++i) + add_block("blocks." + std::to_string(i), true); + add_plain("final_layer.norm.weight", {want.hidden_size}); + add_quant("final_layer.adaln_proj.linear.weight", want.final_adaln_out_features, want.time_embed_dim); + add_plain("final_layer.adaln_proj.linear.bias", {want.final_adaln_out_features}); + add_plain("final_layer.video_out.weight", {video_width, want.hidden_size}); + add_plain("final_layer.video_out.bias", {video_width}); + add_plain("final_layer.audio_out.weight", {want.audio_latents_dim, want.hidden_size}); + add_plain("final_layer.audio_out.bias", {want.audio_latents_dim}); + + std::string header = "{"; + size_t offset = 0; + bool first = true; + for (const Entry& e : entries) { + if (!first) header += ","; + first = false; + header += "\"" + e.name + "\":{\"dtype\":\"" + e.dtype + "\",\"shape\":["; + for (size_t i = 0; i < e.shape.size(); ++i) { + if (i) header += ","; + header += std::to_string(e.shape[i]); + } + header += "],\"data_offsets\":[" + std::to_string(offset) + "," + + std::to_string(offset + e.bytes.size()) + "]}"; + offset += e.bytes.size(); + } + header += "}"; + FILE* fh = std::fopen(path.c_str(), "wb"); + REQUIRE(fh != nullptr); + const uint64_t n = header.size(); + std::fwrite(&n, sizeof(n), 1, fh); + std::fwrite(header.data(), 1, header.size(), fh); + for (const Entry& e : entries) std::fwrite(e.bytes.data(), 1, e.bytes.size(), fh); + std::fclose(fh); +} + } // namespace TEST_CASE("minimax_h3: the deterministic weight stream matches the generator") { @@ -2884,107 +2995,10 @@ TEST_CASE("minimax_h3: an NVFP4 checkpoint loads into a runnable DiT") { want.final_adaln_out_features = 2 * want.hidden_size; want.rope_inv_freq_len = 2; - struct Entry { - std::string name; - std::string dtype; - std::vector shape; - std::string bytes; - }; - std::vector entries; - - auto add_plain = [&](const std::string& name, const std::vector& shape) { - int64_t numel = 1; - for (int64_t d : shape) numel *= d; - const std::vector values = MakeParam("nvfp4." + name, numel, 0.1); - entries.push_back({name, "F32", shape, - std::string(reinterpret_cast(values.data()), - values.size() * sizeof(float))}); - }; - // A quantized projection: packed U8 [out, in/2] + E4M3 [out, in/16] + F32 scalar. - auto add_quant = [&](const std::string& name, int64_t out_dim, int64_t in_dim) { - REQUIRE(in_dim % 16 == 0); - std::string packed(static_cast(out_dim * (in_dim / 2)), '\0'); - for (size_t i = 0; i < packed.size(); ++i) { - packed[i] = static_cast((i * 37 + 11) & 0xFF); // deterministic nibbles - } - std::string scales(static_cast(out_dim * (in_dim / 16)), '\0'); - for (size_t i = 0; i < scales.size(); ++i) { - scales[i] = static_cast(0x38); // e4m3 ~ 1.0 - } - const float global = 0.5f; - entries.push_back({name, "U8", {out_dim, in_dim / 2}, packed}); - entries.push_back({name + "_scale", "F8_E4M3", {out_dim, in_dim / 16}, scales}); - entries.push_back({name + "_scale_2", "F32", {}, - std::string(reinterpret_cast(&global), sizeof(float))}); - }; - const int64_t inner = want.num_attention_heads * want.attention_head_dim; const int64_t video_width = want.video_row_width(); - // Islands stay unquantized, exactly as the real checkpoint has them. - add_plain("video_patch_proj.weight", {want.hidden_size, video_width}); - add_plain("video_patch_proj.bias", {want.hidden_size}); - add_plain("audio_patch_proj.weight", {want.hidden_size, want.audio_latents_dim}); - add_plain("audio_patch_proj.bias", {want.hidden_size}); - add_plain("condition_proj.bias", {want.hidden_size}); - add_plain("time_embedder.proj_in.weight", {want.time_embed_hidden_size, want.timestep_input_dim}); - add_plain("time_embedder.proj_in.bias", {want.time_embed_hidden_size}); - add_plain("time_embedder.proj_out.weight", {want.time_embed_dim, want.time_embed_hidden_size}); - add_plain("time_embedder.proj_out.bias", {want.time_embed_dim}); - add_plain("rope.inv_freq", {want.rope_inv_freq_len}); - // condition_proj is quantized in the real file. - add_quant("condition_proj.weight", want.hidden_size, want.text_dim); - auto add_block = [&](const std::string& prefix, bool with_adaln) { - add_plain(prefix + ".norm1.weight", {want.hidden_size}); - add_plain(prefix + ".norm2.weight", {want.hidden_size}); - add_plain(prefix + ".attn.q_norm.weight", {want.attention_head_dim}); - add_plain(prefix + ".attn.k_norm.weight", {want.attention_head_dim}); - add_quant(prefix + ".attn.qkv_proj.weight", 3 * inner, want.hidden_size); - add_quant(prefix + ".attn.out_proj.weight", want.hidden_size, inner); - add_quant(prefix + ".mlp.fc1.weight", 2 * want.ffn_hidden_size, want.hidden_size); - add_quant(prefix + ".mlp.fc2.weight", want.hidden_size, want.ffn_hidden_size); - if (with_adaln) { - add_quant(prefix + ".adaln_proj.linear.weight", want.adaln_out_features, want.time_embed_dim); - add_plain(prefix + ".adaln_proj.linear.bias", {want.adaln_out_features}); - } - }; - add_block("token_refiner.blocks.0", false); - add_plain("token_refiner.final_norm.weight", {want.hidden_size}); - add_block("blocks.0", true); - add_plain("final_layer.norm.weight", {want.hidden_size}); - add_quant("final_layer.adaln_proj.linear.weight", want.final_adaln_out_features, want.time_embed_dim); - add_plain("final_layer.adaln_proj.linear.bias", {want.final_adaln_out_features}); - add_plain("final_layer.video_out.weight", {video_width, want.hidden_size}); - add_plain("final_layer.video_out.bias", {video_width}); - add_plain("final_layer.audio_out.weight", {want.audio_latents_dim, want.hidden_size}); - add_plain("final_layer.audio_out.bias", {want.audio_latents_dim}); - - // Serialize a safetensors file: 8-byte header length, JSON header, then data. - std::string header = "{"; - size_t offset = 0; - bool first = true; - for (const Entry& e : entries) { - if (!first) header += ","; - first = false; - header += "\"" + e.name + "\":{\"dtype\":\"" + e.dtype + "\",\"shape\":["; - for (size_t i = 0; i < e.shape.size(); ++i) { - if (i) header += ","; - header += std::to_string(e.shape[i]); - } - header += "],\"data_offsets\":[" + std::to_string(offset) + "," + - std::to_string(offset + e.bytes.size()) + "]}"; - offset += e.bytes.size(); - } - header += "}"; const std::string path = "/tmp/minimax_h3_nvfp4_test.safetensors"; - { - FILE* fh = std::fopen(path.c_str(), "wb"); - REQUIRE(fh != nullptr); - const uint64_t n = header.size(); - std::fwrite(&n, sizeof(n), 1, fh); - std::fwrite(header.data(), 1, header.size(), fh); - for (const Entry& e : entries) std::fwrite(e.bytes.data(), 1, e.bytes.size(), fh); - std::fclose(fh); - } + WriteMiniMaxH3Nvfp4File(want, path); const vllm::SafetensorsFile st = vllm::SafetensorsFile::Open(path); const vllm::MiniMaxH3GgufDit loaded = vllm::LoadMiniMaxH3DitFromNvfp4(st); @@ -3152,6 +3166,203 @@ TEST_CASE("minimax_h3: an NVFP4 checkpoint loads into a runnable DiT") { std::remove(path.c_str()); } +TEST_CASE("minimax_h3: the NVFP4 fp4 forward runs Marlin W4A16 on CUDA (speed)") { + // The GB10 leg the CPU wiring gate cannot reach (spec 8.4a): on the CUDA backend + // the W4A16 dispatcher hits dense_nvfp4::MatmulNvfp4MarlinD, so the fp4 path RAN + // is provable by `marlin_gemms == 11` (not the CPU `fallback_gemms`). Built at the + // REAL H3 geometry (one AdaLN block + one refiner) so the per-forward and + // per-GEMM times, and the fp4-vs-bf16 numeric delta, are the production shapes, + // not the reduced-dim wiring model. Env-tunable so the same binary sweeps + // sequence length and reps on dgx: H3_FP4_{LT,LH,LW,AT,AC,TEXT,REPS}. + namespace dnv = vllm::dense_nvfp4; + vt::Backend* cuda = nullptr; + try { + cuda = &vt::GetBackend(vt::DeviceType::kCUDA); + } catch (...) { + MESSAGE("SKIP: no CUDA backend registered"); + return; + } + vt::Queue q = cuda->CreateQueue(); + + auto env_i = [](const char* k, int64_t dflt) -> int64_t { + const char* v = std::getenv(k); + return (v && *v) ? std::atoll(v) : dflt; + }; + const int64_t lt = env_i("H3_FP4_LT", 16), lh = env_i("H3_FP4_LH", 32), + lw = env_i("H3_FP4_LW", 32), at = env_i("H3_FP4_AT", 8), + ac = env_i("H3_FP4_AC", 2), text_len = env_i("H3_FP4_TEXT", 64); + const int reps = static_cast(env_i("H3_FP4_REPS", 8)); + + // Real geometry, single-layer: this is what makes the quantized-GEMM count 11. + MiniMaxH3DitParams want; // defaults ARE the shipped H3 geometry + want.num_layers = 1; + want.token_refiner_num_layers = 1; + const int64_t video_width = want.video_row_width(); + const std::string path = "/tmp/minimax_h3_nvfp4_cuda_speed.safetensors"; + WriteMiniMaxH3Nvfp4File(want, path); + const vllm::SafetensorsFile st = vllm::SafetensorsFile::Open(path); + + // Stream BOTH arms to the DEVICE off the SAME file. bf16 dequantizes to bf16 and + // runs vt::MatmulBT; fp4 keeps the packed FP4 resident and routes each quantized + // projection through the Marlin W4A16 grouped GEMM. + vllm::MiniMaxH3DitParams bf16_params, fp4_params; + const vllm::MiniMaxH3DitDeviceWeights bf16_staged = + vllm::StreamMiniMaxH3Nvfp4ToDeviceBf16(q, st, &bf16_params); + const vllm::MiniMaxH3DitDeviceWeights fp4_staged = + vllm::StreamMiniMaxH3Nvfp4ToDeviceFp4(q, st, &fp4_params); + CHECK(!fp4_staged.weights.blocks[0].qkv_fp4.Empty()); + CHECK(fp4_staged.weights.blocks[0].qkv_proj.data == nullptr); + CHECK(bf16_staged.weights.blocks[0].qkv_fp4.Empty()); + CHECK(bf16_staged.weights.blocks[0].qkv_proj.data != nullptr); + + // Production-shaped packed sequence. + const MiniMaxH3PackedSequence packed = BuildMiniMaxH3PackedSequence( + text_len, lt, lh, lw, at, ac, /*include_keyframe_cond=*/false, {}, 0); + const int64_t seq = packed.seq_len; + const int64_t num_img = static_cast(packed.img_pos.size()); + const int64_t num_audio = static_cast(packed.audio_pos.size()); + const int64_t num_text = static_cast(packed.text_pos.size()); + const std::vector x(static_cast(seq * video_width), 0.25f); + const std::vector audio_x(static_cast(seq * want.audio_latents_dim), 0.1f); + const std::vector prompt(static_cast(num_text * want.text_dim), 0.2f); + const std::vector unique_ts = {0.4f}; + const std::vector inverse(static_cast(seq), 0); + const std::vector refiner_cu = {0, static_cast(num_text), + static_cast(num_text)}; + MiniMaxH3DitInputs in; + in.seq_len = seq; + in.x = x.data(); + in.audio_x = audio_x.data(); + in.img_position_ids = packed.img_position_ids.data(); + in.unique_timesteps = unique_ts.data(); + in.num_unique_timesteps = 1; + in.inverse_indices = inverse.data(); + in.token_tags = packed.token_tags.data(); + in.prompt_embeds = prompt.data(); + in.img_pos = packed.img_pos.data(); + in.num_img_pos = num_img; + in.audio_pos = packed.audio_pos.data(); + in.num_audio_pos = num_audio; + in.text_pos = packed.text_pos.data(); + in.num_text_pos = num_text; + in.infer_out_pos = packed.img_pos.data(); + in.num_infer_out_pos = num_img; + in.update_mask = packed.update_mask.data(); + in.cu_seqlens = packed.cu_seqlens.data(); + in.num_cu_seqlens = static_cast(packed.cu_seqlens.size()); + in.refiner_cu_seqlens = refiner_cu.data(); + in.num_refiner_cu_seqlens = static_cast(refiner_cu.size()); + + auto now = [] { return std::chrono::steady_clock::now(); }; + auto us = [](std::chrono::steady_clock::time_point a, + std::chrono::steady_clock::time_point b) { + return std::chrono::duration(b - a).count(); + }; + auto median = [](std::vector v) { + std::sort(v.begin(), v.end()); + return v.empty() ? 0.0 : v[v.size() / 2]; + }; + + // ── Per-FORWARD timing: the whole DiT step, both arms (the forward returns host + // vectors, so each call syncs the device — no extra Synchronize needed). ── + MiniMaxH3DitOutputs bf16_out = + MiniMaxH3DitForwardDevice(q, bf16_params, bf16_staged.weights, in, vt::DType::kBF16); + std::vector tb; + for (int r = 0; r < reps; ++r) { + const auto t0 = now(); + bf16_out = MiniMaxH3DitForwardDevice(q, bf16_params, bf16_staged.weights, in, vt::DType::kBF16); + tb.push_back(us(t0, now())); + } + + // Cold fp4 forward first (builds the Marlin resident/repack), then the counter + // is asserted on ONE clean forward, then timed. + MiniMaxH3DitOutputs fp4_out = + MiniMaxH3DitForwardDevice(q, fp4_params, fp4_staged.weights, in, vt::DType::kBF16); + dnv::ResetW4A16Stats(); + fp4_out = MiniMaxH3DitForwardDevice(q, fp4_params, fp4_staged.weights, in, vt::DType::kBF16); + const dnv::Nvfp4W4A16Stats stats = dnv::GetW4A16Stats(); + MESSAGE("H3FP4FWD counters marlin_gemms=" << stats.marlin_gemms + << " dense_gemms=" << stats.dense_gemms + << " fallback_gemms=" << stats.fallback_gemms + << " fused_gate_up=" << stats.fused_gate_up); + // The POSITIVE proof the Marlin W4A16 path RAN on the GPU for all 11 quantized + // projections. On CPU these would be fallback_gemms; here they are marlin_gemms. + CHECK(stats.marlin_gemms == 11); + std::vector tf; + for (int r = 0; r < reps; ++r) { + const auto t0 = now(); + fp4_out = MiniMaxH3DitForwardDevice(q, fp4_params, fp4_staged.weights, in, vt::DType::kBF16); + tf.push_back(us(t0, now())); + } + + REQUIRE(fp4_out.video_logits.size() == bf16_out.video_logits.size()); + for (float v : fp4_out.video_logits) REQUIRE(std::isfinite(v)); + for (float v : fp4_out.audio_logits) REQUIRE(std::isfinite(v)); + const double video_delta = + MaxAbsDiff(fp4_out.video_logits, bf16_out.video_logits.data(), fp4_out.video_logits.size()); + const double audio_delta = + MaxAbsDiff(fp4_out.audio_logits, bf16_out.audio_logits.data(), fp4_out.audio_logits.size()); + double vmax = 0.0; + for (float v : bf16_out.video_logits) vmax = std::max(vmax, std::abs(static_cast(v))); + const double mb = median(tb), mf = median(tf); + MESSAGE("H3FP4FWD seq=" << seq << " reps=" << reps + << " per_forward_bf16_ms=" << (mb / 1000.0) + << " per_forward_fp4_ms=" << (mf / 1000.0) + << " ratio(bf16/fp4)=" << (mf > 0 ? mb / mf : 0.0) + << " avg_per_marlin_gemm_us_fp4=" << (mf / 11.0)); + MESSAGE("H3FP4FWD fp4-vs-bf16(CUDA Marlin) video max|diff|=" << video_delta + << " audio max|diff|=" << audio_delta + << " (bf16 video max|val|=" << vmax << ", rel=" + << (vmax > 0 ? video_delta / vmax : 0.0) << ")"); + // Both arms consume the SAME fp4 bytes; the delta is the GEMM path (Marlin's + // bf16 tensor-core accumulate vs the bf16 arm's dequant+MatmulBT), so it is + // matmul-reduction slack amplified through one block, not a quantization error. + CHECK(std::isfinite(video_delta)); + CHECK(vmax > 0.0); + CHECK(video_delta <= 0.10 * vmax + 1e-3); // <=10% relative: sanity, not a tie + + // ── Per-GEMM microbench: the dominant per-token projections at M = seq. Marlin + // W4A16 (fp4-resident) vs the bf16 arm's own dequant+GEMM, SAME numbers. ── + dnv::Dev d{*cuda, q}; + auto bench = [&](const char* nm, const vllm::Nvfp4Weight& w) { + const int64_t N = w.n, K = w.k; + std::vector xh(static_cast(seq) * K, 0x3DCCu); // bf16 ~0.1 + dnv::DBuf xb(d, vt::DType::kBF16, {seq, K}, xh.data()); + for (int i = 0; i < 3; ++i) { auto o = dnv::MatmulNvfp4W4A16D(d, xb.t(), w, vt::DType::kBF16); (void)o; } + cuda->Synchronize(q); + std::vector f; + for (int r = 0; r < reps; ++r) { + const auto t0 = now(); + auto o = dnv::MatmulNvfp4W4A16D(d, xb.t(), w, vt::DType::kBF16); + cuda->Synchronize(q); + f.push_back(us(t0, now())); + (void)o; + } + const std::vector wb = dnv::DequantNvfp4ToBLayout(w); // [K, N] bf16 + dnv::DBuf wbd(d, vt::DType::kBF16, {K, N}, wb.data()); + for (int i = 0; i < 3; ++i) { dnv::DBuf o(d, vt::DType::kBF16, {seq, N}); vt::Matmul(q, o.t(), xb.t(), wbd.t()); } + cuda->Synchronize(q); + std::vector b; + for (int r = 0; r < reps; ++r) { + dnv::DBuf o(d, vt::DType::kBF16, {seq, N}); + const auto t0 = now(); + vt::Matmul(q, o.t(), xb.t(), wbd.t()); + cuda->Synchronize(q); + b.push_back(us(t0, now())); + } + const double gf = median(f), gb = median(b); + MESSAGE("H3FP4GEMM " << nm << " M=" << seq << " N=" << N << " K=" << K + << " fp4_marlin_us=" << gf << " bf16_us=" << gb + << " ratio(bf16/fp4)=" << (gf > 0 ? gb / gf : 0.0)); + }; + bench("qkv", fp4_staged.weights.blocks[0].qkv_fp4); + bench("out", fp4_staged.weights.blocks[0].out_fp4); + bench("fc1", fp4_staged.weights.blocks[0].fc1_fp4); + bench("fc2", fp4_staged.weights.blocks[0].fc2_fp4); + + std::remove(path.c_str()); +} + TEST_CASE("minimax_h3: the encoder VISION block matches upstream") { // The repeated unit of the H3-Encoder's Qwen3-VL vision tower. It differs from // the text tower in ways that all matter numerically: LayerNorm (with bias) not From f3066ae38d860ef0dd06f16aa2dbc97879b1c1bc Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 15:03:04 +0000 Subject: [PATCH 2/3] test(minimax_h3): GB10 leg asserts the Marlin W4A16 FAMILY + fallback==0 MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit First dgx run of the CUDA speed case: on the production default the counter is `dense_gemms==11`, not `marlin_gemms` — VT_MARLIN_DENSE is default-ON, so each W4A16 projection takes vLLM's OWN dense Marlin GEMM (`vt::MarlinDenseGemm`), not the single-expert MoE-grouped Marlin. Both are the Marlin W4A16 path; the "this-path-ran" proof is `marlin_gemms + dense_gemms == 11` AND `fallback_gemms == 0` (no CPU/redundant-dequant fallback). VT_MARLIN_DENSE=0 opts into the grouped route and yields `marlin_gemms==11`. Assertion updated accordingly; the spec's "grouped MoE marlin" routing note (8.1) is corrected in the record. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] --- tests/vllm/models/test_minimax_h3.cpp | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/tests/vllm/models/test_minimax_h3.cpp b/tests/vllm/models/test_minimax_h3.cpp index ec80d9f1..2cfa1615 100644 --- a/tests/vllm/models/test_minimax_h3.cpp +++ b/tests/vllm/models/test_minimax_h3.cpp @@ -3286,8 +3286,14 @@ TEST_CASE("minimax_h3: the NVFP4 fp4 forward runs Marlin W4A16 on CUDA (speed)") << " fallback_gemms=" << stats.fallback_gemms << " fused_gate_up=" << stats.fused_gate_up); // The POSITIVE proof the Marlin W4A16 path RAN on the GPU for all 11 quantized - // projections. On CPU these would be fallback_gemms; here they are marlin_gemms. - CHECK(stats.marlin_gemms == 11); + // projections, and NOT the CPU/redundant-dequant fallback. On CPU these would be + // fallback_gemms. On CUDA the production default (VT_MARLIN_DENSE on, sm_121a) + // routes each projection through vLLM's OWN dense Marlin GEMM -> `dense_gemms`; + // VT_MARLIN_DENSE=0 opts into the single-expert MoE-grouped Marlin -> `marlin_gemms`. + // Either Marlin sub-route is the "this-path-ran" signal; fallback_gemms MUST be 0. + const uint64_t marlin_family = stats.marlin_gemms + stats.dense_gemms; + CHECK(marlin_family == 11); + CHECK(stats.fallback_gemms == 0); std::vector tf; for (int r = 0; r < reps; ++r) { const auto t0 = now(); From 70cf16ebfce7672b8ab7c5b71b027ed02af90041 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 15:59:16 +0000 Subject: [PATCH 3/3] =?UTF-8?q?docs(minimax-h3):=20W-FP4a=20GB10=20leg=20M?= =?UTF-8?q?EASURED=20=E2=80=94=20Marlin=20byte-exact,=20fp4=3Dmemory-win,?= =?UTF-8?q?=20e2e=20RUNS=20but=20frame=20is=20a=20non-scene?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit GB10 sm_121a results for row/H3-FP4-GPU-E2E (PR #64): Leg 1 (Marlin W4A16 + timing): the fp4 path RAN on GPU — default dense_gemms==11 (VT_MARLIN_DENSE default-ON => vLLM's OWN dense Marlin GEMM, not the grouped route the spec assumed), marlin_gemms==11 under VT_MARLIN_DENSE=0, fallback_gemms==0 both. fp4-vs-bf16 BYTE-EXACT (max|diff|=0). Timing crossover per-forward bf16/fp4: 3.47x @seq64 (fp4 faster, memory-bound) -> 0.79-0.83x @seq4224-7040 (fp4 slower, compute-bound). fp4 W4A16 is a WEIGHT-BANDWIDTH win; for H3's large-M diffusion forward it is a MEMORY win (~16 vs ~66 GB), not speed. Leg 2 (real-checkpoint fp4-resident t2va e2e): the real 18.75 GB NVFP4 DiT + both VAEs + the GGUF Qwen3-VL-32B encoder run the WHOLE chain end-to-end -> valid mp4/wav. BUT the decoded frame is a structured non-scene patch-grid at 12/20/50 steps, conditioned or not -> OPEN render-coherence bug (device VAE decode / denoise), separate from the fp4 speed work. DiT s/step 5.45/20.0/209 s @512/768/REF-209f. Leg 3 (indirect): HW/loader-forced-indirect — 4xB300 renders 209f in 86.964 s vs 1xGB10 209 s/forward; vLLM-Omni has no quantized H3 arm. Updates: benchmark-record entry, spec 8.1 routing correction + 8.4 status, STATUS/BENCHMARKS/FEATURES rows, model-matrix + roadmap H3 rows, state + NOW. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] --- .agents/NOW.md | 2 +- .agents/benchmark-record.md | 123 ++++++++++++++++++++++++++++++++++++ .agents/model-matrix.md | 2 +- .agents/roadmap_v1.md | 2 +- .agents/specs/minimax-h3.md | 52 +++++++++++---- .agents/state.md | 48 ++++++++++++++ docs/BENCHMARKS.md | 2 +- docs/FEATURES.md | 2 +- docs/STATUS.md | 2 +- 9 files changed, 217 insertions(+), 18 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index fbee0e88..c6eafd73 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -18,7 +18,7 @@ checkpoint on `upstream/main` at `59674cf1d`. | DeepSeek-V4-Flash decode | **Closed: BEATS ds4 1.144x** (`VT_V4_RESIDENT_W`, byte-exact). Phase-2 routed-expert residency NEGATIVE (−3.4%), default-OFF | — | | f32-out GEMV audit | Only laguna + deepseek_v4 bf16 tower affected; gate/on-framework dense unaffected | Re-verify deepseek_v4 tower same-tool | | Invocation-parity prevention | CI guard (`check-gemv-invocation-consistency.py`) + AGENTS.md checklist landing | Review + merge; CUDA build-verify `kGemvHeuristicAlgos` on dgx | -| MiniMax-H3 lane | **W-FP4a fp4 routing CPU-landed** (`row/H3-FP4-SPEED`): NVFP4 kept packed -> Marlin W4A16, no new quant code; gate 62/62 | GB10 CUDA gate; real e2e disk-blocked | +| MiniMax-H3 lane | **W-FP4a GB10 leg LANDED** (`row/H3-FP4-GPU-E2E`, PR #64): Marlin W4A16 byte-exact, fp4=MEMORY win; e2e RUNS but frame=non-scene patch-grid | **OPEN: render bug**; fp4 CLOSED | | Kimi-Linear-48B (KDA+NoPE-MLA+MoE) | **Full-model GB10 e2e RUNS** (bf16-resident §13): CPU+CUDA 13/13·656, no OOM. **Token gate NEAR-TIE 106/128** (6/8 token-exact) | device GDN/MLA islands + bf16 stream; 1.59 tok/s; default OFF | | 35B fresh grid | **BOUND** @`1ea26427`: tput 0.93-1.03x, c16 0.93x. INTAKE + Option A both **RESOLVED NEGATIVE** (H2D-out-of-capture tput WASH) | Real lever left: prefill glue (task #61) | | Qwen3.5-4B revalidation | 0.9971x @`59674cf1` (#35); TTFT/PSS pass, TPOT/ITL open | `docs/bench-evidence/` | diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 1298a0ba..8692bc12 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -13142,6 +13142,129 @@ vllm-omni source; the documented anchor is the recipe's 4×B300 BF16 evidence (F 209f 1248×768 = 86.964 s; two-video Ref2VA = 784.394 s, `recipe:298-311`). +## MiniMax-H3 W-FP4a GB10 leg — Marlin W4A16 RAN on sm_121a (byte-exact); fp4 is a MEMORY win not a diffusion-forward SPEED win; e2e RENDERS on real weights but the frame is a non-scene patch-grid (2026-08-06, `row/H3-FP4-GPU-E2E`, `ROAD-V1-H3`) + +**Setup.** dgx.casa GB10 sm_121a, CUDA 13, cutlass 4.5.0, Release. git-archive of +`row/H3-FP4-GPU-E2E` (goldens md5 byte-identical both sides). Targeted build of +`test_minimax_h3` only. New CUDA case `minimax_h3: the NVFP4 fp4 forward runs +Marlin W4A16 on CUDA (speed)` builds a synthetic compressed-tensors NVFP4 file at +REAL H3 geometry (1 AdaLN block + 1 refiner = 11 quantized projections), streams +BOTH the bf16 arm and the fp4-resident arm off the SAME file, runs both device +forwards, and times them (median over 12 reps, cold leg discarded) on an idle box +(free=117 G after drop_caches). Env-tunable seq via `H3_FP4_{LT,LH,LW,AT,AC,TEXT}`. + +**The Marlin path RAN (the mission's `marlin_gemms==11`, corrected for the real route).** +On CUDA the W4A16 dispatcher counter is: +- **default (VT_MARLIN_DENSE ON): `dense_gemms=11, marlin_gemms=0, fallback_gemms=0`** — + each of the 11 quantized projections took vLLM's OWN dense Marlin GEMM + (`vt::MarlinDenseGemm`), NOT the single-expert MoE-grouped Marlin the spec §8.1 + assumed. +- **VT_MARLIN_DENSE=0: `marlin_gemms=11, dense_gemms=0, fallback_gemms=0`** — the + grouped MoE-Marlin route, the mission's literal counter. +`fallback_gemms==0` in BOTH ⇒ no CPU/redundant-dequant fallback; the sm_121a +Marlin W4A16 tensor-core path genuinely executed for every projection. (Spec §8.1 +routing note corrected: the production default is the DENSE marlin, not grouped.) + +**fp4-vs-bf16 numeric delta = 0 (byte-exact), all sizes.** Both arms consume the +SAME packed NVFP4 bytes; the dense Marlin GEMM (fp32-C reduce) is byte-preserving +vs the bf16 arm's dequant+MatmulBT, so the two device forwards produce BIT-IDENTICAL +video+audio logits (bf16 video max|val|=1243.18, max|diff|=0). The fp4-resident arm +costs nothing in accuracy. + +**Timing — the crossover (median/12, cold discarded).** Per-forward (whole 2-block +DiT step; attention + fp32 islands + the 11 W4A16 GEMMs), dense route: + +| seq (M) | bf16 ms | fp4 ms | ratio bf16/fp4 | +|---|---|---|---| +| 64 | 12.57 | 3.62 | **3.47** (fp4 faster) | +| 4224 | 76.61 | 92.83 | 0.825 (fp4 slower) | +| 7040 | 129.36 | 164.23 | 0.788 (fp4 slower) | + +(grouped route @seq4224: bf16 76.07 / fp4 94.28 = 0.807 — same story, dense marginally faster.) + +Per-GEMM microbench, `ratio(bf16/fp4)` (fp4 faster when >1): + +| GEMM [N,K] | M=64 | M=4224 | M=7040 | +|---|---|---|---| +| qkv [21504,5376] | 3.16 | 0.582 | 0.529 | +| out [5376,7168] | 3.90 | 1.02 | 1.06 | +| fc1 [28672,5376] | 3.32 | 0.596 | 0.533 | +| fc2 [5376,14336] | 2.86 | 1.22 | 0.90 | + +**Verdict.** Marlin W4A16 is a WEIGHT-BANDWIDTH optimization: it reads ¼ the weight +bytes but does the SAME bf16 tensor-core FLOPs plus an on-the-fly dequant. At small +M (decode-like, memory-bound) it is ~3× faster; in H3's diffusion-forward regime +(one forward over the WHOLE packed sequence, M = thousands, compute-bound) the big +projections (qkv/fc1) are ~1.7–1.9× SLOWER and out/fc2 ~tied, netting ~0.8× per +forward. **So for H3 the fp4-resident arm's value is MEMORY** (~16 GB packed vs +~66 GB bf16 = 4×, which is what lets the DiT + encoder + VAEs fit the 119 GiB +pool), not diffusion-forward throughput. A bf16 arm would be ~1.2× faster per step +if it fit — it does not (66 GB DiT + 13 GB encoder + VAEs). + +**Leg 2 — real-checkpoint t2va e2e (fp4-resident), it RENDERS on real weights.** +Files (byte-complete, sizes verified): DiT `minimax_h3_ref2va_nvfp4_full.safetensors` +18 745 492 256 B (the UNPRUNED variant our forward implements — the pruned files use +a timestep-LUT AdaLN we do not) + `vae/minimax_h3_video_vae_fp16.safetensors` +5 207 808 496 B + `vae/minimax_h3_audio_vae_fp32.safetensors` 605 254 808 B + +encoder `qwen3vl-32B-MiniMax-H3-Q4_K_M.gguf` 14 576 977 888 B, downloaded 8-way +parallel (HF xet CDN throttles a single connection to ~5 MB/s). VAE configs +(latents_mean/std) taken from the project's own baked `minimax_h3_vae_configs.inc`; +tokenizer from public `Qwen/Qwen3-VL-32B-Instruct/tokenizer.json` (the GGUF is +weights-only). `examples/minimax-h3-gen --fp4-resident --device cuda` (the new flag +routes NVFP4 safetensors through `StreamMiniMaxH3Nvfp4ToDeviceFp4`): stages the real +DiT in **6–18 s**, geometry recovered = 50 layers / 5376 / 56 heads. + +Small render 256×256 / 22 frames (1 VAE tile), UNCONDITIONED (synthetic embeds, +12 steps) AND CONDITIONED (GGUF Qwen3-VL-32B encoder keep-quant 13.24 GiB → +[16, 5120] conditioning; encoder+DiT+VAEs coresident ~35 GB, free stayed ≥80 G): +both run the full path — 22 PPM frames + `audio.wav` (32 kHz stereo) + a VALID +`h264/yuv420p 256×256 + AAC` mp4 (ffprobe-confirmed). **FRAME SANITY:** a decoded +frame is a regular grid of small multicolour blocks (16 px = the vae_ratio-16 latent +cell) with smooth internal gradients on a dark ground — structured image data (the +VAE IS decoding, NOT white noise), **but NOT a coherent scene, IDENTICALLY at 12, 20 +AND 50 steps (the reference count) whether conditioned or not.** Step count and +conditioning are ruled out; the grid is at exactly the latent-cell scale, i.e. each +latent decodes to an independent block with no cross-patch coherence. **Honest +verdict: the composed fp4-resident pipeline RUNS end to end on the real checkpoint +and every stage executes, but it does NOT yet render a coherent scene — a real OPEN +render bug (device video-VAE decode and/or denoise convergence at real geometry), +separate from and NOT blocking the fp4 speed work.** This is precisely what the +frame-sanity gate catches: every numeric unit gate green + valid mp4, yet the frame +is a structured non-scene. First eyeball of an H3 frame (the prior Thor run reached +`succeeded`/valid ffprobe but its content was never viewed). + +DiT s/step, FULL 50-layer fp4-resident, `--denoise-only` (per FORWARD, drop_caches, +locks held): + +| config | seq | s/forward | +|---|---|---| +| 512×512 / 22f | 1 898 | **5.45** | +| 768×768 / 61f | 12 948 | **20.03** | +| **768×1344 / 209f (vllm-omni REF canvas)** | 63 224 | **209.09** | + +The reference canvas FITS in the pool (a forward runs) but a full 50-step render is +49 × 209 ≈ **2.85 h** of DiT + the 209-frame tiled VAE decode — memory-feasible but +hours-long, so the full reference render was NOT run (the honest "largest fitting +config" per the mission). The composed path is proven on the real checkpoint at the +small canvas. + +**Indirect speed statement (mission #3), HW/loader-forced-INDIRECT.** vLLM-Omni has +NO quantized H3 arm (BF16-only, single-GPU only via `--enable-cpu-offload`; +source-audited a4ea67a2/v0.26.0), so a quant-matched one-GB10 comparison is +impossible (DeepSeek-GGUF precedent). The documented vLLM-Omni anchor is 4× B300 +BF16: FL2VA 209f 1248×768 = **86.964 s** mean client latency for a WHOLE 50-step +render (recipe:298-311). Ours, single GB10 fp4-resident, is **209 s for ONE forward** +at the comparable 209-frame reference canvas (768×1344) — so 4× B300 completes an +entire ~49-forward render + VAE in 87 s (~1.8 s/forward-equivalent) while one GB10 +takes 209 s for a single forward, i.e. **~116× per-forward**. That is the expected +HW/loader-forced-indirect gap: 4 datacenter B300s + BF16 + USP-4 sequence-parallel + +torch.compile + block-caching vs one GB10 + fp4 + none of those — and vLLM-Omni +cannot run a quantized H3 on one GPU at all. Our honest SAME-box number is the +fp4-vs-bf16-arm ratio (leg 1, DiT-only real geometry): fp4 is **0.79–0.83×** the +bf16 arm per diffusion forward — the fp4 arm trades ~20% forward speed for 4× less +weight memory, which is what makes the DiT + encoder + VAEs fit one GB10 at all. + + ## Laguna-S-2.1-NVFP4 decode — router top-k KERNEL-EFFICIENCY (`VT_LAGUNA_TOPK_SHFL`), BYTE-EXACT, SigmoidTopK 1.67×, −0.57% decode-step GPU (2026-08-03, `CLAIM-LAGUNA-TOPK-SHFL`) Measured results for vllm.cpp, against the reference engine each workload actually competes with. Every number here was produced on real hardware, greedy, diff --git a/.agents/model-matrix.md b/.agents/model-matrix.md index 83470a46..9ef0052e 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 | **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. **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); GB10 speed delta + real e2e pending (disk-blocked); 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. **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/roadmap_v1.md b/.agents/roadmap_v1.md index 64ed24d4..9e6894ef 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -76,7 +76,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 HW/loader-forced-indirect. | +| 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. **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 0c578861..4a7d4b29 100644 --- a/.agents/specs/minimax-h3.md +++ b/.agents/specs/minimax-h3.md @@ -335,7 +335,13 @@ alpha==0) is **forced to the Marlin W4A16 grouped GEMM**, bypassing the capability-based kernel registry. So on sm_121a **every quantized H3 projection takes the SAME kernel** — `dense_nvfp4::MatmulNvfp4MarlinD` (single-expert `vt::MoeGroupedGemmNvfp4Marlin`), the exact path the Laguna routed-experts -(`laguna.cpp`) and the dense Qwen3-32B NVFP4 arm (`qwen3_5.cpp`) use. The +(`laguna.cpp`) and the dense Qwen3-32B NVFP4 arm (`qwen3_5.cpp`) use. +**MEASURED CORRECTION (GB10, 2026-08-06):** the PRODUCTION default is one level up — +`VT_MARLIN_DENSE` is default-ON, so `MatmulNvfp4W4A16D` routes each projection through +vLLM's OWN **dense** Marlin GEMM (`vt::MarlinDenseGemm`, counter `dense_gemms`), and +the grouped `MoeGroupedGemmNvfp4Marlin` (`marlin_gemms`) is taken only under +`VT_MARLIN_DENSE=0`. Both are Marlin W4A16 and both are byte-exact to the bf16 arm; +dense is marginally faster. See §8.4 + the benchmark record. The cutlass-FP4 / true-W4A4 route (`MatmulNvfp4Fp4D`) is NOT taken here: it needs fp4 ACTIVATIONS this checkpoint does not carry, and is deliberately private to `qwen3_5.cpp` (`dense_nvfp4_gemm.h:12-18`). @@ -436,14 +442,36 @@ video flow_shift 12 / audio 3, no CFG; default canvas **768×1344**, default fra - **W-FP4a: CPU-LANDED + gated** — fp4-resident loader + Marlin-W4A16 routing + the fp4-vs-bf16 wiring gate. No new quant code. -- **W-FP4a GB10 leg: PENDING** — build the CUDA `test_minimax_h3` on dgx.casa, - run the CUDA `an NVFP4 checkpoint loads into a runnable DiT` case (the same gate - exercises the Marlin path via `marlin_gemms`), and capture the fp4-vs-bf16 delta + - steady per-step time. Gated on a safe disk/build window (dgx: 33 G free vs 15 G - floor; the 570 G shared `.cache` is not ours to prune). -- **W-FP4b real-checkpoint t2va e2e: DISK-BLOCKED.** The NVFP4 arm working set is - DiT + Qwen3-VL-32B NVFP4 encoder + both VAEs (~41 GB+ on disk to download); with - ~18 GB usable above the disk floor it does not fit, and the encoder is required - for a real conditioned render. Recorded HW/disk-forced-indirect. A DiT-only - steady-per-step at real geometry (fp4 vs bf16) is the reachable speed number once - a single NVFP4 DiT variant download fits. +- **W-FP4a GB10 leg: LANDED (2026-08-06, `row/H3-FP4-GPU-E2E`).** A dedicated CUDA + case `minimax_h3: the NVFP4 fp4 forward runs Marlin W4A16 on CUDA (speed)` (the + existing "loads into a runnable DiT" case runs the forwards on a CPU queue, so it + could never bump the GPU counter) builds the synthetic NVFP4 file at REAL geometry + and runs both arms on a CUDA queue. **Marlin RAN:** default `dense_gemms==11` + (VT_MARLIN_DENSE is default-ON → vLLM's OWN dense Marlin GEMM, NOT the grouped + route §8.1 assumed), `marlin_gemms==11` under VT_MARLIN_DENSE=0, `fallback_gemms==0` + in both. **fp4-vs-bf16 delta = 0 (byte-exact).** **Timing crossover** (median/12, + cold discarded): per-forward ratio bf16/fp4 = 3.47× @seq64 (fp4 faster, + memory-bound), 0.825× @seq4224, 0.788× @seq7040 (fp4 slower, compute-bound). So + fp4 W4A16 is a WEIGHT-BANDWIDTH win (decode-like small M) and a ~1.2× LOSS in H3's + large-M diffusion forward; its H3 value is MEMORY (~16 GB vs ~66 GB bf16). Benchmark + record has the full per-GEMM tables. +- **W-FP4b real-checkpoint t2va e2e: RUNS on real weights; frame COHERENCE is an open + bug.** dgx now has room; the real NVFP4 DiT (`minimax_h3_ref2va_nvfp4_full`, + 18.75 GB, unpruned) + both VAEs + the GGUF Qwen3-VL-32B encoder were downloaded and + the WHOLE t2va chain runs with `--fp4-resident` (new driver flag → the fp4-resident + streamer, ~16 GB device vs ~66 GB bf16): encoder → [16,5120] text conditioning → + fp4-resident DiT → both VAEs → ffmpeg, producing a valid `h264 256×256 + AAC 32 kHz` + mp4 + wav. **But the decoded frame is a structured multicolour patch-grid at the + latent-cell scale, NOT a coherent scene — identically at 12/20/50 steps, conditioned + or not.** So the composed path is proven to RUN e2e on the real checkpoint, but a + coherent render is an OPEN bug (device video-VAE decode and/or denoise convergence + at real geometry), independent of the fp4 speed work. **DiT s/step (full 50-layer + fp4-resident, per forward):** 5.45 s @512×512/22f, 20.03 s @768×768/61f, 209.09 s + @768×1344/209f (the vllm-omni REF canvas). The REF canvas fits in the pool but a full + 50-step render is ~2.85 h, so it was not run (largest-fitting-config honesty). +- **Comparability (mission #3):** HW/loader-forced-INDIRECT. 4× B300 BF16 renders a + whole 50-step FL2VA 209f in 86.964 s (~1.8 s/forward-equiv); one GB10 fp4-resident + is 209 s for ONE forward at the comparable canvas (~116× per-forward) — 4 datacenter + GPUs + BF16 + USP-4 + torch.compile + block-caching vs one GB10 + fp4 + none, and + vLLM-Omni cannot serve a quantized H3 on one GPU at all. The honest same-box number + is the fp4-vs-bf16 ratio (0.79–0.83× per forward, 4× less weight memory). diff --git a/.agents/state.md b/.agents/state.md index 7eb2ebf2..9443dcbf 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -39275,3 +39275,51 @@ DISK-BLOCKED: NVFP4 working set (DiT + Qwen3-VL-32B NVFP4 encoder + 2 VAEs) is render. Comparability: vLLM-Omni CANNOT serve a quantized H3 on one GPU (BF16-only in practice; source-audited `a4ea67a2`) -> HW/loader-forced-indirect, DeepSeek-GGUF precedent. Draft PR is the claim. + +## 2026-08-06T16:05 - MiniMax-H3 W-FP4a GB10 leg: Marlin W4A16 GB10-verified byte-exact; fp4 = memory win; real-checkpoint e2e RUNS but frame is a non-scene patch-grid + + +Row `row/H3-FP4-GPU-E2E` (PR #64, off `ROAD-V1-H3`). The two GPU legs the disk +previously blocked, run on dgx.casa GB10 sm_121a (CUDA 13, cutlass 4.5.0, Release); +git-archive transfer (goldens md5 byte-identical both sides), targeted builds only. + +**Leg 1 — Marlin W4A16 + timing (spec §8.4a).** The existing NVFP4 case runs the +fp4/bf16 forwards on a CPU queue, so on CUDA it could never bump the GPU counter. A +new CUDA case `minimax_h3: the NVFP4 fp4 forward runs Marlin W4A16 on CUDA (speed)` +builds the synthetic NVFP4 file at REAL geometry (1 block + 1 refiner = 11 quantized +projections) and runs both arms on a CUDA queue. Marlin RAN: default +`dense_gemms==11` (VT_MARLIN_DENSE is DEFAULT-ON → vLLM's OWN dense Marlin GEMM, NOT +the grouped MoE-marlin the spec §8.1 assumed), `marlin_gemms==11` under +VT_MARLIN_DENSE=0, `fallback_gemms==0` in both. fp4-vs-bf16 delta = 0 (BYTE-EXACT). +Timing crossover (median/12, cold discarded, idle box): per-forward bf16/fp4 = 3.47× +@seq64 (fp4 FASTER, memory-bound decode-like) → 0.825× @seq4224 → 0.788× @seq7040 +(fp4 SLOWER, compute-bound). Per-GEMM: qkv/fc1 ~1.7–1.9× slower at large M, out/fc2 +~tied. **fp4 W4A16 is a WEIGHT-BANDWIDTH optimization; for H3's large-M diffusion +forward it is a MEMORY win (~16 GB packed vs ~66 GB bf16), not a speed win.** + +**Leg 2 — real-checkpoint fp4-resident t2va e2e.** Downloaded the real NVFP4 arm +(DiT `ref2va_nvfp4_full` 18.75 GB unpruned + video/audio VAEs + GGUF Qwen3-VL-32B +encoder) 8-way parallel; VAE configs from our baked `minimax_h3_vae_configs.inc`, +tokenizer from public Qwen3-VL-32B. New `examples/minimax-h3-gen --fp4-resident` +routes NVFP4 safetensors through `StreamMiniMaxH3Nvfp4ToDeviceFp4` (~16 GB device). +The WHOLE chain runs — encoder → [16,5120] conditioning → fp4-resident DiT → both +VAEs → ffmpeg → valid h264/AAC mp4 + wav. **But the decoded frame is a structured +multicolour patch-grid at the latent-cell scale, NOT a coherent scene, IDENTICALLY +at 12/20/50 steps, conditioned or not.** So the composed path is PROVEN to run e2e +on the real checkpoint, but coherent rendering is an OPEN bug (device video-VAE +decode / denoise convergence at real geometry), independent of the fp4 speed work. +DiT s/step (full 50-layer fp4-resident): 5.45/20.03/209.09 s @512×512-22f / +768×768-61f / 768×1344-209f. The REF canvas fits but a full 50-step render is ~2.85 h, +so it was not run (largest-fitting-config honesty). + +**Leg 3 — indirect speed (HW/loader-forced-indirect).** vLLM-Omni has no quantized +H3 arm (BF16-only, single-GPU only via cpu-offload). 4×B300 BF16 renders 209f +1248×768 in 86.964 s (whole 50-step render, ~1.8 s/forward-equiv); one GB10 +fp4-resident is 209 s for ONE forward at the comparable REF canvas (~116×). Honest +same-box number: fp4 0.79–0.83× the bf16 arm per forward, 4× less weight memory. + +**Next.** Root-cause the H3 render-coherence bug (device VAE decode and/or denoise +convergence at real geometry) — the frame-sanity gate caught it (all unit gates +green + valid mp4, yet a non-scene). fp4 speed path is CLOSED. Box left clean (GPU +idle, locks free, worker down, disk ≥15 G; checkpoint cached for reruns). Benchmark +record + spec §8 + STATUS/BENCHMARKS/FEATURES + model-matrix/roadmap updated. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index d4b4442e..1240855e 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -296,7 +296,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) | Full-model GB10 e2e RUNS (bf16-resident §13), NEAR-TIE 106/128, pool math CLOSES; default OFF | Full model RUNS on GB10 (bf16-resident, RSS peak 1.7 GiB, min-avail 21 GiB, no OOM). Token NEAR-TIE 106/128 (6/8 prompts exact, numerics vs deterministic oracle). 1.59 tok/s. Detail: spec §13 | | vLLM 0.26 re-benchmark | Pending | Re-run the binding grids on the advanced pin | -| MiniMax-H3 FP4 speed (W-FP4a) | Pending (measurement in flight, `row/H3-FP4-GPU-E2E`). fp4-resident Marlin-W4A16 routing CPU-landed (63/63); the GB10 CUDA speed case (asserts `marlin_gemms==11`, records fp4-vs-bf16 delta + per-forward/per-GEMM times) and `--fp4-resident` driver flag landed; dgx build+run + real e2e in flight; vLLM-Omni has no quantized H3 (BF16-only) | Build+run CUDA `test_minimax_h3` on dgx (Marlin via `marlin_gemms`), capture delta + s/step; fp4-resident real-checkpoint t2va e2e. 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) but frame is a non-scene patch-grid (OPEN bug) | Root-cause H3 render coherence (VAE/denoise); fp4 speed CLOSED. Detail: benchmark-record + spec §8 | | 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**. `QUANT-CT-MXFP4-FINAL-STACK` TERMINAL: 2 last levers exhausted (num_splits cap gated-OFF c1-only; glue folds via FusedChain, residual out-of-catalog). c2-c8 GPU-intrinsic; see record | | SGLang floor arms | Never ran | Both arms of the SGLang comparison | | cuBLAS invocation-parity guard | CI guard landed (CPU); `kGemvHeuristicAlgos` refactor build-verify owed | `nvcc` rebuild + SACRED gate on dgx | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index b1a30412..5c96d595 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -115,7 +115,7 @@ oversight. | Image | ✅ correctness-gated | ✅ | ✅ | ◐ | | Video | ✅ correctness-gated | ✅ | ✅ | ☐ | | Audio | ✅ correctness-gated | ✅ | ◐ | ◐ | -| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ portable path complete; e2e on real weights (Thor); FP4-resident Marlin-W4A16 routing CPU-landed, GB10 speed pending | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | +| Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ portable path complete; fp4-resident e2e RUNS on GB10 (real 18.75 GB NVFP4 DiT + VAEs + GGUF encoder → valid mp4/wav); Marlin W4A16 GB10-verified byte-exact; render COHERENCE open (non-scene patch-grid) | ✅ (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 830b5026..4ebe4f41 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -73,7 +73,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