fix(cuda-arch): arch-gate the WMMA selectors — pre-Ampere guards were a live trap

.agents/specs/cuda-arch-breadth-fp16.md

@@ -428,7 +428,8 @@ until a card exists.
 | W1 ✅ **DONE** | guard the bf16-WMMA TU | `#if __CUDA_ARCH__ >= 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<T>` 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 `<sm_80` behind `__trap()`, but every predicate that SELECTS them was host-side only (shape/dtype/env), so a pre-Ampere board would still pick a trapping kernel. Three chokepoints, each now requiring `DeviceCaps::sm_major >= 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 `<sm_80`. So pre-Ampere CORRECTNESS is covered by paths that exist today, and this port buys prefill throughput only. ~2,000 lines adapted from llama.cpp's contiguous KV to our paged block-table layout, against a floor (llama.cpp on-card) that needs hardware we do not have. **Do this only for a genuine speed claim on T4/V100, and only once a card is reachable** | W1b; hardware for any payoff |
 | W3 | FEATURE-TABLE + tactic registration | `fattn-fp16` feature row + `sm_75/70/61/60` cells; register `fattn-fp16-tile/vec` tactics; selector arch term at `:2562`; `CudaArchFeaturesTest.cmake` + registry-selection tests | W2 |
 | W4 | **Turing derive-and-ship** | `sm_75` `-Werror` build + `cuobjdump` SASS proof; row → `DERIVED+BUILD-VERIFIED (testing-welcome)`; labeled untested | W3, nvcc 13 (**doable now**) |
 | W5 | wire a `<13` toolkit | provision CUDA 12.x; Volta/Pascal build-verify (`70`/`61`/`60` SASS); those rows → `DERIVED+BUILD-VERIFIED` | a 12.x toolkit |
@@ -471,7 +472,16 @@ until a card exists.
    list that honours it. Worth a standing check: any TU whose kernels are
    arch-specific should sit behind its cell, and `CMakeLists.txt:903` currently
    does not.
-9. **Guarding a body orphans its helpers, and `-Werror=all-warnings` turns that
+9. **A COMPILE guard without a SELECTOR guard is a trap, not a fix.** This is the
+   single most important lesson of W1a-W1c. `#if __CUDA_ARCH__ >= 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<T>` 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 `<sm_80` guarding.
-10. **A per-TU compile sweep is NOT a link.** W1a compiled 20 TUs to throwaway
+11. **A per-TU compile sweep is NOT a link.** W1a compiled 20 TUs to throwaway
    objects; it proves no TU has an `<sm_80` *compile* blocker beyond the two
    named, and nothing about undefined symbols, `__trap()` stubs reachable at
    link, or fatbin assembly. Do not quote 18/20 as "the library builds".

docs/BENCHMARKS.md

@@ -312,7 +312,7 @@ built on it rather than keeping the flattering one.
 | MXFP4 Qwen3-8B (W4A16 Marlin) | #51 x3: c1 1.005, c2/c4/c8 0.925/0.939/0.953, mem 2.18x. `MARLIN-STRUCT`: decode-graph + gate_up FUSION default-ON, marlin 180->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), 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
 

docs/STATUS.md

@@ -1538,36 +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). **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
-established that bf16 needs no fp16 model path on these arches: there are zero
-bf16 *arithmetic* intrinsics in the CUDA tree (the pattern is convert-on-load,
-compute in fp32), so models stay bf16 and only WMMA fragment instantiation is
-Ampere-gated. **Volta (`sm_70`, V100) and Pascal (`sm_60`/`sm_61`, P100/P40)
-are not-yet-buildable** because CUDA 13 dropped their code generation and no
-CUDA 12.x toolkit is provisioned here; both failing TUs fail for capability
-reasons that hold on Volta by construction, so the fix list transfers but the
-SASS proof does not. There is no vLLM oracle on these cards
-(vLLM will not run there), so a real correctness test uses llama.cpp on the same
-card plus a newer-card/CPU cross-check; nothing is runtime-verified yet.
+20 unconditionally-built CUDA TUs at `sm_75` (nvcc 13.0.88) first measured 18 PASS
+/ 2 FAIL and **now measures 20/20 PASS, 0 errors and 0 warnings**. The two
+residuals (`cuda_gdn.cu`, `cuda_matmul_nvfp4.cu`) carried bf16 WMMA fragments plus
+a harder second class, `wmma::precision::tf32` — also Ampere+, but failing at the
+type-alias definition rather than at a use site, so a body-only guard does not
+compile. Both are now wrapped `#if __CUDA_ARCH__ >= 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
 

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,

src/vt/cuda/cuda_gdn.cu

@@ -32,6 +32,7 @@
 #include <vector>
 
 #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
+  // <sm_80 to GdnScanCuda, the portable sequential scan that already serves the
+  // arbitrary-dim corners. Fails safe: caps invalid -> 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;
   }

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 <typename Tout>
 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 <sm_80 arch term lives inside WmmaEnabled(), so this branch (and the
+  // other five call sites) route to the WMMA-free naive / split-K path on a
+  // pre-Ampere board without a local check.
   if (p < kTileMinRows || !WmmaEnabled()) {
     const int64_t y = p < 65535 ? p : 65535;
     constexpr int kBlock = 256;

src/vt/cuda/cuda_paged_attn.cu

@@ -2608,7 +2608,17 @@ void LaunchPaged(cudaStream_t s, Tensor& out, const Tensor& query, const Tensor&
   // so they fall through to the correctness-grade f32 CUDA-core flash below.
   // (Latent bug the first additive dense model, MODEL-TEXT-qwen3, forced out; the
   // gate models keep the WMMA path unchanged — they run d=256, TQ/TKV=bf16.)
-  const bool wmma = is_prefill && d == 256 && PrefillWmmaEnabled() &&
+  // ARCH TERM (required, not an optimisation). The five WMMA prefill kernels are
+  // compiled `#if __CUDA_ARCH__ >= 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 <sm_80 falls through to the portable CUDA-core flash below,
+  // which is correctness-grade and needs no tensor cores. On sm_80+ this term is
+  // always true, so the gate models' selection is unchanged.
+  // See .agents/specs/cuda-arch-breadth-fp16.md §V0-a / W1c.
+  const bool arch_has_bf16_mma = GetDeviceCaps().sm_major >= 8;
+  const bool wmma = is_prefill && d == 256 && arch_has_bf16_mma && PrefillWmmaEnabled() &&
                     std::is_same<TKV, __nv_bfloat16>::value &&
                     std::is_same<TQ, __nv_bfloat16>::value;
   // GQA K/V reuse: eligible when qpk = hq/num_kv_heads is a multiple of the reuse