From d76121d5cde89cc490f3d60f066f6b10d226c49b Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 07:32:48 +0000 Subject: [PATCH 1/2] =?UTF-8?q?fix(cuda-arch):=20arch-gate=20the=20WMMA=20?= =?UTF-8?q?prefill=20selector=20=E2=80=94=20the=20guards=20were=20a=20live?= =?UTF-8?q?=20trap?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit W1/W1b compiled the five bf16-WMMA prefill kernels under `#if __CUDA_ARCH__ >= 800` with an `#else __trap()`. The predicate that SELECTS them (cuda_paged_attn.cu:2611) is entirely host-side — shape, dtype and env only — and never consulted the device capability. So on a Turing/Volta/Pascal board a d=256 bf16 prefill would still select a kernel whose body is a trap: the guard made the TU compile and left a runtime crash behind it. Gate the predicate on the cached DeviceCaps::sm_major >= 8. The whole WMMA ladder (gqa / flash2 / flash2vec) derives from `wmma`, so one term covers all four launchers, and 144 GEMM/step (vLLM-structural); #44 3/3, 32B-NVFP4A16 142/142 | nsys c8: marlin +1,177us (CTA 144 vs 48, dominant); `VT_MARLIN_E1_PAR1` E=1->48 CTAs near-parity but flips a strict 32B token (OFF). Byte-preserving `KERNEL-MARLIN-DENSE-PORT` landed gated-OFF; GPU binding pending | | 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 | -| Pre-Ampere breadth (Turing `sm_75` / Volta `sm_70` / Pascal) | **NO NUMBER OWED, nothing executes on these arches.** 2026-08-06 sm_75 compile audit (nvcc 13.0.88): 20 unconditional sm_80+ constructs enumerated; detail in .agents/benchmark-record.md | Port the llama.cpp `fattn-tile`/`fattn-vec` fp16 body. Perf floor when a card exists is **llama.cpp on the same card** (vLLM does not run there) | +| Pre-Ampere breadth (Turing `sm_75` / Volta `sm_70` / Pascal) | **No number owed; nothing runs on these arches.** 2026-08-06 `sm_75`: 20/20 TUs PASS (0 err/warn) after WMMA body guards; prefill selector arch-gated; GB10 SASS byte-identical. [Detail](../.agents/benchmark-record.md) | Port the llama.cpp `fattn-tile`/`fattn-vec` fp16 body. Perf floor when a card exists is **llama.cpp on the same card** (vLLM does not run there) | ## Reproduce diff --git a/docs/STATUS.md b/docs/STATUS.md index 9ff46e26..8f10c01e 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -1555,9 +1555,21 @@ gate build is unaffected, by measurement:** at `sm_121a` both TUs compile `-Werror=all-warnings` 0-warn, `cuda_gdn.cu` SASS is bit-identical across 824,704 lines, and `cuda_matmul_nvfp4.cu` shows zero instruction-level differences (the only 148 differing lines are `Function :` headers carrying the anonymous-namespace -hash, which shifts on any edit to a file). **No `sm_75` library link exists yet** -— a per-TU compile sweep is not a link, and no Turing, Volta or Pascal board has -run any of this. Separately, the audit +hash, which shifts on any edit to a file). + +**The prefill selector is arch-gated as well**, which the guards alone did not +cover: the predicate choosing the bf16-WMMA prefill kernels was entirely +host-side (shape, dtype, env) and never consulted the device, so a pre-Ampere +board would have selected a kernel whose body is a `__trap()`. It now also +requires `sm_major >= 8`, and `= 800` with an `#else __trap()` (W1/W1b) because + // bf16 WMMA fragments are Ampere+. This predicate is otherwise entirely + // host-side (shape + dtype + env), so on a Turing/Volta/Pascal board it would + // still SELECT a kernel whose body is a trap. Gate on the cached device + // capability so = 8; + const bool wmma = is_prefill && d == 256 && arch_has_bf16_mma && PrefillWmmaEnabled() && std::is_same::value && std::is_same::value; // GQA K/V reuse: eligible when qpk = hq/num_kv_heads is a multiple of the reuse From 89aa9e7b6b8d51b50f4683601f847f025db74670 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 08:09:31 +0000 Subject: [PATCH 2/2] =?UTF-8?q?fix(cuda-arch):=20W1c=20=E2=80=94=20arch-ga?= =?UTF-8?q?te=20the=20GDN=20and=20MoE=20selectors;=20W2=20rescoped?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Completes the class the prefill-selector fix opened. W1/W1b made the WMMA bodies COMPILE on = 8 and each falling through to a portable path that already exists: - cuda_gdn.cu:5359 (GdnPrefillKernelCuda) -> GdnScanCuda, the sequential scan that already serves the arbitrary-dim corners. This is the single chokepoint: all 7 guarded GDN launches live in LaunchChunkedPrefill, itself reached only from here, and the TSc=float instantiations use the TF32 WmmaCfg, so f32 is not a way around the guard. - cuda_matmul_nvfp4.cu WmmaEnabled():77 -> the naive / tiled / split-K CUDA-core kernels. Folded into the predicate rather than its six call sites because all six mean the same thing and each already has a CUDA-core fallthrough. Queried LIVE rather than latched in the static env cache: the device context need not exist at static-init time, and a wrong value cached there is unrecoverable. Both fail safe — caps invalid selects the portable path. Verified on dgx (nvcc 13.0.88). All three TUs -Werror=all-warnings rc=0, 0 warnings, at BOTH sm_75 and sm_121a. GB10 inert by measurement: sm_121a SASS IDENTICAL for all three against the pre-change build (933,178 + 825,294 + 137,542 lines, anon-namespace hash normalized) — the change is host-side and no device code moved. W2 is rescoped in the spec from a correctness brick to a SPEED brick. The bf16-WMMA path is only `is_prefill && d == 256 && bf16 q+KV`; all decode and every other prefill shape already run portable kernels, and these gates route d=256 to the CUDA-core flash on 256224 in the same change per the gate's own instruction. Detail moved to .agents/benchmark-record.md. No number is claimed or owed: nothing executes on these arches, no library link exists, and sm_70 stays uncompilable until a CUDA 12.x toolkit is wired. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-5 [ClaudeCode] --- .agents/benchmark-record.md | 65 +++++++++++++++++++++++ .agents/specs/cuda-arch-breadth-fp16.md | 16 ++++-- docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 70 ++++++++++--------------- scripts/check-public-doc-tables.py | 2 +- src/vt/cuda/cuda_gdn.cu | 16 +++++- src/vt/cuda/cuda_matmul_nvfp4.cu | 18 ++++++- 7 files changed, 140 insertions(+), 49 deletions(-) diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 269a16bf..0aaba56c 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -12680,6 +12680,71 @@ vLLM side reused from #52 (same pin `55596792`). ours-default = graph+fuse (par1 | SPAN | 30,732 | 29,800 | 28,491 | +2,241 | +1,309 | | TPOT ms (client, nsys-inflated) | 37.22 | 36.23 | 34.58 | +2.64 | +1.65 | +**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 **DECODE-SPEED ATTRIBUTED (W7 profile-only, 2026-07-31, `CLAIM-LAGUNA-W7-SPEED`): `nsys` of the W6 decode on the real UD-Q4_K_XL GGUF (GB10) shows the 0.66 s/tok (~1.5 tok/s, vs llama.cpp 27.8 tok/s on the identical bytes, ~15-18x) is 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` sync in `LagunaForwardGgufCached` + scalar host glue; 39.4% of the GPU time is `QuantizeQ8K` activation-quant (re-quantized per GEMM), the weight GEMVs run un-grouped at ~22% of the 240 GB/s peak (llama.cpp ~76%). Ranked levers (all in-tree from ds4): device-resident decode (kill the syncs) 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 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= 8`), so `` alias definition rather than as an incomplete type at a use +site, so a body-only guard does not compile it. `cuda_matmul_nvfp4.cu` needed 5 +body guards and stays compiled for every arch: despite its name it also carries +the generic bf16 MoE grouped GEMMs, so gating the TU on its `fp4-mma` cell (the +spec's earlier suggestion, now RETRACTED) would have stripped those from +`sm_80/90a/100a/110`. Guarding bodies then orphaned their helpers, and +`-Werror=all-warnings` promotes `#177-D "declared but never referenced"` to an +error, so the residual walked 110 → 13 → 2 → 0 across `WmmaCfg::WK`, the +`V128` staging helpers and `WyMerge` (which needed a WHOLE-function guard, +since a body-only guard leaves an emitted-but-uncalled definition). + +**The compile guards alone were a live trap, which is the transferable lesson.** +Every predicate SELECTING a guarded kernel was host-side only (shape, dtype, +env), so a pre-Ampere board would still have picked a `__trap()` body — the +"fix" converted a build error into a runtime crash on exactly the boards it was +meant to enable. Three chokepoints now require `DeviceCaps::sm_major >= 8`, each +falling through to a portable path that already existed: +`cuda_paged_attn.cu:2611` → `LaunchPrefillFlash`; `cuda_gdn.cu:5359` → +`GdnScanCuda` (the single chokepoint — all 7 guarded GDN launches live in +`LaunchChunkedPrefill`, reached only from there, and the `TSc=float` +instantiations use the TF32 config so f32 is no way around it); and +`cuda_matmul_nvfp4.cu` `WmmaEnabled():77`, folded into the predicate rather than +its six call sites because all six mean the same thing and each already has a +CUDA-core fallthrough. All three fail safe (caps invalid → portable), and the +nvfp4 one is queried LIVE rather than latched in the static env cache, since the +device context need not exist at static-init time. + +**GB10 inert, by measurement, not by argument:** all three TUs compile 0-warn at +`sm_121a` and their SASS is IDENTICAL to the pre-change build — 933,178 + +825,294 + 137,542 lines, anon-namespace hash normalized (that hash shifts on any +edit to a file, which is why the raw diff shows `Function :` headers only). The +change is host-side; no device code moved. + +**W2 is RESCOPED to a speed brick.** The bf16-WMMA path is only +`is_prefill && d == 256 && bf16 q+KV`; all decode and every other prefill shape +already run portable kernels, and the selector gates now route d=256 to the +CUDA-core flash on `= 800` wraps the bodies of all 5 bf16-WMMA prefill kernels (`cuda_paged_attn.cu:732,958,1197,1472,1716`; `#else __trap()`), so `__CUDA_ARCH__ < 800` compiles the TU selecting the existing scalar path. **Build-verified (dgx nvcc 13.0.88 + cutlass 4.5.0, base `034be66e`):** single-arch `75` `-Werror=all-warnings` 0-warn EXIT=0, `cuobjdump -lelf` → real `cuda_paged_attn.cu.1.sm_75.cubin`; the `:1797 __nv_bfloat16 fragment` error GONE (RED: unguarded HEAD FAILS 21 errors). GB10 sm_121a byte-identical — same TU `-Werror` 0-warn AND 0 SASS instruction diffs vs unguarded. NO Turing board ran it | DONE — (mechanical, nvcc 13) | | W1a ✅ **DONE** | **V0 full-library compile audit** — every unconditionally-built CUDA TU compiled at `sm_75`, failures enumerated and classified. **MEASURED 2026-08-06 (base `249697b7`, dgx nvcc 13.0.88): 20 TUs, 18 PASS (0 err / 0 warn), 2 FAIL** (`cuda_gdn.cu` 110 errors, `cuda_matmul_nvfp4.cu` 10). All 8 fast-path FEATURE-TABLE cells confirmed DISABLED at `75`, bounding the surface at the `CMakeLists.txt:896-916` list. Compile-only, no GPU. Full detail + classification in §V0 | DONE — nvcc 13, no card | | W1b ✅ **DONE** | **finish the guard set.** `cuda_gdn.cu`: both `WmmaCfg` specializations' members guarded (bf16 fragments AND the tf32 alias block — a lookup failure at the alias, so a body-only guard does NOT compile), 8 device bodies guarded `#if __CUDA_ARCH__ >= 800` / `#else __trap()`, plus the `V128` staging helpers and `WyMerge`. `cuda_matmul_nvfp4.cu`: 5 WMMA bodies guarded, TU left compiled for every arch (see the §V0-b correction — gating it on `fp4-mma` would have stripped the generic bf16 MoE GEMMs from `sm_80/90a/100a/110`). **VERIFIED (dgx nvcc 13.0.88): `sm_75` 20/20 TUs PASS, 0 errors 0 warnings** (was 18/20). **`sm_121a` byte-identity HELD:** both TUs `-Werror=all-warnings` 0-warn, `cuda_gdn.cu` SASS bit-identical across 824,704 lines, `cuda_matmul_nvfp4.cu` **zero instruction-level diffs** (all 148 differing lines are `Function :` headers carrying the anon-namespace hash, which shifts on any edit — same artifact W1 recorded). NO board ran any of it | DONE | -| W2 | port `fattn-tile`+`fattn-vec` fp16 body | new `cuda_paged_attn_fp16.cu`, 1:1 from `fattn-tile.cuh`/`fattn-vec.cuh:21`; fp16 accum + `sm_61` fp32 variant; C1 numerics vs CPU oracle | W1b | +| W1c ✅ **DONE** | **arch-gate the SELECTORS — the guards alone were a live trap.** W1/W1b made the WMMA bodies compile on `= 8` and each falling through to an EXISTING portable path: (a) `cuda_paged_attn.cu:2611` → `LaunchPrefillFlash` (CUDA-core register-tiled flash); (b) `cuda_gdn.cu:5359` `GdnPrefillKernelCuda` → `GdnScanCuda` (sequential scan) — the single chokepoint, since all 7 guarded GDN launches live in `LaunchChunkedPrefill`, itself reached only from there, and the `TSc=float` instantiations use the TF32 `WmmaCfg` so f32 is not a way around it; (c) `cuda_matmul_nvfp4.cu` `WmmaEnabled():77` → naive/tiled/split-K. (c) is folded into the predicate itself rather than its six call sites because all six mean the same thing and each already has a CUDA-core fallthrough; it is queried LIVE, not latched in the static env cache, since the device context need not exist at static-init time. All three fail safe (caps invalid → portable). **VERIFIED (dgx nvcc 13.0.88): all 3 TUs `-Werror=all-warnings` rc=0 0-warn at BOTH `sm_75` and `sm_121a`; `sm_121a` SASS IDENTICAL for all three** (933,178 + 825,294 + 137,542 lines, anon-namespace hash normalized) — host-side change, device code untouched. NO board ran it | W1b | +| W2 | port `fattn-tile`+`fattn-vec` fp16 body — **RESCOPED to a SPEED brick by W1c.** The bf16-WMMA path is only `is_prefill && d == 256 && bf16 q+KV` (`cuda_paged_attn.cu:2611`); all decode and every other prefill shape already run portable kernels, and W1c routes the d=256 case to `LaunchPrefillFlash` on `= 800` / + `#else __trap()` makes a TU *compile* on an old arch; it does nothing about the + host code that decides to launch it. Every predicate selecting a guarded kernel + here was pure host-side shape/dtype/env, so the "fix" would have turned a + compile error into a runtime crash on exactly the boards it was meant to + enable. Any future arch-guarding must pair each `#if` with an arch term on its + selector and name the portable path it falls through to — and if no portable + path exists, that is a design problem, not a mechanical edit. +10. **Guarding a body orphans its helpers, and `-Werror=all-warnings` turns that into a build failure.** W1b needed three iterations for exactly this: after the 8 GDN bodies were guarded, `nvcc` reported `#177-D "declared but never referenced"` for `WmmaCfg::WK`, every `V128` member, and finally `WyMerge` @@ -480,7 +490,7 @@ until a card exists. `__device__` function whose callers are all guarded, wrap the WHOLE function (a body-only guard leaves an emitted-but-uncalled definition that still trips `#177-D`). Expect this cascade on any future `144 GEMM/step (vLLM-structural); #44 3/3, 32B-NVFP4A16 142/142 | nsys c8: marlin +1,177us (CTA 144 vs 48, dominant); `VT_MARLIN_E1_PAR1` E=1->48 CTAs near-parity but flips a strict 32B token (OFF). Byte-preserving `KERNEL-MARLIN-DENSE-PORT` landed gated-OFF; GPU binding pending | | 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 | -| Pre-Ampere breadth (Turing `sm_75` / Volta `sm_70` / Pascal) | **No number owed; nothing runs on these arches.** 2026-08-06 `sm_75`: 20/20 TUs PASS (0 err/warn) after WMMA body guards; prefill selector arch-gated; GB10 SASS byte-identical. [Detail](../.agents/benchmark-record.md) | Port the llama.cpp `fattn-tile`/`fattn-vec` fp16 body. Perf floor when a card exists is **llama.cpp on the same card** (vLLM does not run there) | +| Pre-Ampere breadth (Turing `sm_75` / Volta `sm_70` / Pascal) | **No number owed; nothing runs on these arches.** 2026-08-06 `sm_75`: 20/20 TUs PASS (0 err/warn), WMMA bodies + all 3 selectors arch-gated; GB10 SASS byte-identical. [Detail](../.agents/benchmark-record.md) | Port the llama.cpp `fattn-tile`/`fattn-vec` fp16 body. Perf floor when a card exists is **llama.cpp on the same card** (vLLM does not run there) | ## Reproduce diff --git a/docs/STATUS.md b/docs/STATUS.md index 8f10c01e..962d19d7 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -1538,48 +1538,34 @@ no Turing board ran it; a green compile + SASS is not execution evidence. **Scope correction (2026-08-06): that was ONE translation unit, not a library build**, and this page previously read as the latter. A full compile audit of all -20 unconditionally-built CUDA TUs at `sm_75` (base `249697b7`, nvcc 13.0.88) now -measures **18 PASS (0 errors, 0 warnings) and 2 FAIL**: `cuda_gdn.cu` (110 error -lines — bf16 WMMA fragments plus `wmma::precision::tf32`, which is also Ampere+ -and fails at the type-alias definition rather than at a use site) and -`cuda_matmul_nvfp4.cu` (10 errors — bf16 fragments; the TU is compiled -unconditionally even though its own `fp4-mma` feature cell resolves DISABLED). -The audit also confirms all eight fast-path feature cells resolve DISABLED at -`75`, so the remaining surface is those two files. - -**Both are now guarded, and the audit is green: `sm_75` compiles 20/20 TUs, 0 -errors and 0 warnings** (2026-08-06). The bf16 and TF32 WMMA bodies in both files -are wrapped `#if __CUDA_ARCH__ >= 800` with a `__trap()` fallback, together with -the helper structs and device functions that only those bodies use. **The GB10 -gate build is unaffected, by measurement:** at `sm_121a` both TUs compile -`-Werror=all-warnings` 0-warn, `cuda_gdn.cu` SASS is bit-identical across 824,704 -lines, and `cuda_matmul_nvfp4.cu` shows zero instruction-level differences (the -only 148 differing lines are `Function :` headers carrying the anonymous-namespace -hash, which shifts on any edit to a file). - -**The prefill selector is arch-gated as well**, which the guards alone did not -cover: the predicate choosing the bf16-WMMA prefill kernels was entirely -host-side (shape, dtype, env) and never consulted the device, so a pre-Ampere -board would have selected a kernel whose body is a `__trap()`. It now also -requires `sm_major >= 8`, and `= 800` with a `__trap()` +fallback, along with the helpers only those bodies use. All eight fast-path +feature cells resolve DISABLED at `75`. + +**The three WMMA selectors are arch-gated too, which the compile guards alone did +not cover:** each predicate was host-side only (shape, dtype, env), so a +pre-Ampere board would still have selected a `__trap()` body — a build error +turned into a runtime crash. All three now require `sm_major >= 8` and fall +through to paths that already exist: the portable CUDA-core flash (attention), the +sequential scan (GDN), and the naive/tiled/split-K kernels (NVFP4/MoE). Each fails +safe if the device capability is unreadable. **GB10 is unaffected by measurement** +— the three TUs compile 0-warn at `sm_121a` with byte-identical SASS. + +**No `sm_75` library link exists yet** — a per-TU compile sweep is not a link, and +no Turing, Volta or Pascal board has run any of this. bf16 needs no fp16 model +path here: there are zero bf16 *arithmetic* intrinsics in the CUDA tree +(convert-on-load, compute in fp32), so models stay bf16 and only WMMA fragment +instantiation is Ampere-gated. **Volta (`sm_70`, V100) and Pascal are +not-yet-buildable** — CUDA 13 dropped their code generation and no 12.x toolkit is +provisioned here; the fix list transfers to Volta by construction, the SASS proof +does not. There is no vLLM oracle on these cards, so real correctness testing +means llama.cpp on the same card plus a newer-card/CPU cross-check; nothing is +runtime-verified yet. ## Serving and API notes diff --git a/scripts/check-public-doc-tables.py b/scripts/check-public-doc-tables.py index 29dd3b53..01587c75 100755 --- a/scripts/check-public-doc-tables.py +++ b/scripts/check-public-doc-tables.py @@ -321,7 +321,7 @@ def features_errors(text: str) -> list[str]: # Lowering these numbers as the page is compacted is the gate closing. STATUS = ROOT / "docs/STATUS.md" STATUS_RATCHET = { - "chars": 284_329, + "chars": 284081, "h2_sections": 11, "long_paragraphs": 89, "oversized_cells": 47, diff --git a/src/vt/cuda/cuda_gdn.cu b/src/vt/cuda/cuda_gdn.cu index 8e9d7b52..b0264ca7 100644 --- a/src/vt/cuda/cuda_gdn.cu +++ b/src/vt/cuda/cuda_gdn.cu @@ -32,6 +32,7 @@ #include #include "vt/cuda/conv_update_fast.h" +#include "vt/cuda/cuda_device_caps.h" #include "vt/cuda/cuda_gdn_internal.h" #include "vt/cuda/gdn_packed_decode_triton.h" #include "vt/cuda/gdn_prefill_conv.h" @@ -5356,9 +5357,22 @@ void GdnPrefillKernelCuda(Queue& q, Tensor& out, const Tensor& q_in, const Tenso // fallback (VT_GDN_CHUNKED=0). The bf16 chunked path is WMMA (tensor-core), // which tiles at 16 and 32 — bf16 dims that are not WMMA-friendly fall back // to the sequential scan (real gate dims Dk=Dv=128 satisfy both). + // ARCH TERM (required, not an optimisation). Every kernel LaunchChunkedPrefill + // reaches — the WU, delta_h (wmma/reg/regring/tma) and chunk_o bodies — is + // compiled `#if __CUDA_ARCH__ >= 800` with an `#else __trap()`, because bf16 + // fragments AND wmma::precision::tf32 are Ampere+. That holds for BOTH scratch + // dtypes: the TSc=float instantiations use the TF32 WmmaCfg, so f32 is not a + // way around it. This routing is the single chokepoint (all 7 launches live in + // LaunchChunkedPrefill, itself reached only from here), so gating it sends + // scan. On sm_80+ the term + // is always true, so the gate models' path is unchanged. + // See .agents/specs/cuda-arch-breadth-fp16.md §V0-a / W1c. + const DeviceCaps& gdn_caps = GetDeviceCaps(); + const bool arch_has_mma = gdn_caps.valid && gdn_caps.sm_major >= 8; const bool wmma_ok = q_in.dtype != DType::kBF16 || (dk % kWM == 0 && dv % kNB == 0); if (ChunkedPrefillEnabled() && dk <= kChunkMaxDim && dv <= kChunkMaxDim && args.scale != 0.0f && - wmma_ok) { + wmma_ok && arch_has_mma) { GdnPrefillChunkedCuda(q, out, q_in, k, v, g, beta, state, qsl, args); return; } diff --git a/src/vt/cuda/cuda_matmul_nvfp4.cu b/src/vt/cuda/cuda_matmul_nvfp4.cu index 837f647d..d3fb0074 100644 --- a/src/vt/cuda/cuda_matmul_nvfp4.cu +++ b/src/vt/cuda/cuda_matmul_nvfp4.cu @@ -79,7 +79,20 @@ bool WmmaEnabled() { const char* e = std::getenv("VT_NVFP4_WMMA"); return e == nullptr || (e[0] != '0'); }(); - return on; + // ARCH TERM (required, not an optimisation). Every WMMA body this predicate + // selects is compiled `#if __CUDA_ARCH__ >= 800` with an `#else __trap()` + // (bf16 WMMA fragments are Ampere+), while all six call sites are otherwise + // host-side shape/env tests — so on a pre-Ampere board they would select a + // trap. Folded in HERE rather than at each call site because every one of the + // six means the same thing ("take the bf16 tensor-core path"), and each already + // has a CUDA-core fallthrough (naive / tiled / split-K) that is + // correctness-grade. Queried live rather than cached in the static above: the + // device context need not exist at static-init time, and a wrong value latched + // there would be unrecoverable. Fails safe: caps invalid -> CUDA-core path. On + // sm_80+ this is always true, so gate-model selection is unchanged. + // See .agents/specs/cuda-arch-breadth-fp16.md §V0-b / W1c. + const DeviceCaps& caps = GetDeviceCaps(); + return on && caps.valid && caps.sm_major >= 8; } // M=1/decode-path 128-bit vectorized fp4 weight loads (A/B; default ON). Set @@ -1413,6 +1426,9 @@ template void LaunchGroupedBf16(cudaStream_t s, Tensor& out, const Tensor& act, const Tensor& expert_ids, const Tensor* row_map, const Tensor& weight_ptrs, int64_t p, int64_t n, int64_t k, int64_t e_count) { + // The