mudler · localai-bot · Aug 6, 2026 · Aug 6, 2026
diff --git a/.agents/NOW.md b/.agents/NOW.md
@@ -22,7 +22,7 @@ checkpoint on `upstream/main` at `59674cf1d`.
 | Kimi-Linear-48B (KDA+NoPE-MLA+MoE) | **e2e RUNS** (bf16-resident §13): 13/13·656. Token gate **NEAR-TIE 106/128** | device GDN/MLA islands; 1.59 tok/s; default OFF |
 | 35B fresh grid | **BOUND** @`1ea26427`: 0.93-1.03x, c16 0.93x. INTAKE + Option A both NEGATIVE | Lever left: prefill glue (#61) |
 | Qwen3.5-4B revalidation | 0.9971x @`59674cf1` (#35); TTFT/PSS pass, TPOT/ITL open | `docs/bench-evidence/` |
-| MXFP4 parity | **c1 1.020, c2-c8 0.962-0.969, mem 2.63x**. `FLASH-AUDIT` #68: c8 flash +12.5us/call = occupancy/L2-bound; `-use_fast_math` TRIED→REGRESSES, rejected | real lever: flash occupancy/L2; no flip |
+| MXFP4 parity | c1 1.020, c2-c8 0.962-0.969 (BELOW-FLOOR). `FLASH-OCCUPANCY` #75 CLOSED: matched-c8 ncu, occupancy IDENTICAL 8.33%; matched vLLM's exact reg+instr, STILL +10us = ptxas gap, no lever/flip (#75) | c8: flash ptxas ~40% + glue ~18% |
 | ROW-SERVE-ASYNC-DENSE-MIRROR | **LANDED+dgx-VERIFIED** (`f9c969ae`): async mirror on classic dense Qwen3; SACRED 184/184 | Residual: sibling scope one-liner |
 | CPU levers (`QUANT-GGUF-CIQ-GEMM`) | Op-dispatch profile DONE: decode **47% threadpool sync**, prefill **~39% paged attn**. **G5 not next** | Parakeet encoder; attn dtype hoist |
 | Supported-models list (`row/DOCS-SUPPORTED-MODELS-MATRIX`) | **DRAFT PR**: FEATURES per-arch table CI-bound to registry (30 archs) | Reviewer merge |

diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md
@@ -19,6 +19,46 @@ from relative link targets repointed for this file's location.
 
 # Benchmarks
 
+## QUANT-CT-MXFP4-FLASH-OCCUPANCY — the owed ours-vs-vLLM flash decode ncu diff: occupancy is IDENTICAL (8.33%, smem-limited); the flash gap is a COMPILER-CODEGEN difference (vLLM's wheel SASS runs 13% fewer instructions on the byte-identical kernel), NOT occupancy/L2/arch/fast-math — no lever exists on our stack (2026-08-06, `row/QUANT-CT-MXFP4-FLASH-OCCUPANCY`, base `f7a1e322`, GB10 sm_121a, PR #75)
+
+#69 closed the compile lens (flash source byte-identical to vLLM's `2c839c33`; `-use_fast_math` REJECTED, +21us/call) and OWED the ours-vs-vLLM ncu diff (vLLM side lost to a box OOM-reboot). This row runs that diff to a MEASURED verdict on an idle box and re-frames the #69 premise. **Vehicle** `Yi30/Qwen3-8B-MXFP4` (dense W4A16 Marlin), oracle arm `VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel`.
+
+**W0 (free, existing c8 nsys CSVs) — the #69 "8.3% occupancy, register-limited" figure was the WRONG kernel.** The prior ncu pair was context-MISMATCHED: ours ran `vllm-cli` with a ~5-token prompt (context ~5 → `num_splits=1`, Split=false — a different kernel instance), vLLM ran `lens=[1024]` (Split=true). At the REAL c8 decode kernel (`gpu_trace_c8_dflt` vs `vllm_offline_trace_c8`) both run the IDENTICAL grid `1x3x64` (num_splits=3, gridZ=64=batch×kv_heads) and ours uses FEWER registers (216 vs 241) → more occupancy headroom, not less.
+
+**W1 — matched-c8 ncu diff (both engines, grid `1x3x64`, full section set; ours = `vllm-bench --input-len 1024 --output-len 128 -c 8`; vLLM = offline NPROMPTS=8, `enforce_eager` to dodge the FULL-cudagraph node-profiling driver-resource conflict; vLLM `gpu_memory_utilization` dropped 0.5→0.15 after full-section replay + the 0.5-util reservation OOM-rebooted the 119 GiB unified pool once — the SAME #69 failure mode, now root-caused: ncu kernel-replay on the vLLM process is memory-heavy):**
+
+| metric | OURS (native sm_121a) | vLLM (sm_80-PTX driver-JIT wheel) |
+|---|---|---|
+| grid (X×Y×Z) | **192 (1×3×64)** | **192 (1×3×64)** — identical |
+| block / dyn smem per block | 128 / **81.92 KB** | 128 / **81.92 KB** — identical |
+| Block Limit Shared Mem | **1** | **1** (both smem-limited to 1 CTA/SM) |
+| **Achieved Occupancy** | **8.33%** | **8.33%** — IDENTICAL |
+| Registers / thread | 216 | 241 |
+| **Executed Instructions** | **3,695,180** (19,245/sched) | **3,267,840** (17,020/sched) — ours **+13.1%** |
+| Duration (ncu, median) | ~165 us | ~157 us |
+| L2 hit / L1 hit | 1.03% / 0.39% | 1.23% / 1.07% |
+| SM / Memory throughput | 13.5% / 10.4% | 14.2% / 11.0% |
+| Active warps/sched (eligible) | ~1.0 (0.06) | ~1.0 (0.06) |
+| stall: short-scoreboard (LDS→MMA) | ~63% of ~16.6 cyc | ~58% of ~16.9 cyc — same KIND |
+
+**Answer to the owed W1 question: YES — vLLM's instance runs at the SAME 8.33% occupancy.** Both are SMEM-limited to 1 CTA/SM by the byte-identical 81.92 KB `kSmemSize` (the 216-vs-241 reg delta is MOOT — smem binds first). Occupancy is refuted as the lever. L2-warmth/adjacency is refuted: both stream KV from DRAM at ~1% L2 hit (c8 KV working set ≫ L2 → nothing a preceding kernel leaves stays hot), so the flash kernel's own memory profile is context-independent. num_splits/grid/GQA-pack are identical. The ONLY measured difference is **executed instructions: ours +13.1%** (the same +11.6% scalar bloat #69's cuobjdump saw on the non-split kernel, now on the actual c8 decode kernel with EVERYTHING ELSE held identical). That made instruction-count the leading CODEGEN hypothesis — but W2 tests it directly and **REFUTES it as the cause**: matching vLLM's exact instruction count does NOT close the gap.
+
+**W2 — mirror-first lever exploration (controlled ncu on the SAME c8 kernel; flash TUs recompiled per cell into an isolated build):**
+
+| flash-TU build variant | reg | executed instr | ncu duration | vs vLLM 3.27M/157us |
+|---|---|---|---|---|
+| A native sm_121a (main default) | 216 | 3.70M (19,245/sched) | ~165 us | baseline |
+| B compute_80 PTX driver-JIT (mirror vLLM's arch) | 236 | 3.69M (19,212/sched) | ~165 us | **NEGATIVE — no change** |
+| **C** compute_80 + `-use_fast_math` (vLLM's EXACT recipe) | **241** | **3,267,072 (17,008/sched)** | **~167 us (161-170)** | **matches vLLM reg+instr EXACTLY, STILL ~10us slower** |
+| #68 native sm_121a + `-use_fast_math` | 255 | (fewer) | 189.8 us | REGRESSED (prior row) |
+| vLLM `_vllm_fa2_C` (sm_80+PTX+fast-math wheel) | 241 | 3,267,840 (17,020/sched) | ~157 us | the target |
+
+**The decisive cell is C.** Arch-mirror alone (B) is NEUTRAL (compute_80 PTX driver-JIT = native instr/duration). Adding fast-math (C) reproduces vLLM's SASS profile EXACTLY — 241 reg and 17,008 instr vs vLLM's 241/17,020, a −13% instruction drop from A's 19,245 — and YET C is ~167 us, essentially the SAME as A/B's ~165 us and still ~10 us SLOWER than vLLM's ~157 us. So the −13% instruction reduction does NOT translate to speed (the kernel is NOT instruction-bound — our toolchain sits at ~165 us regardless of 19,245 or 17,008 instr), and matching vLLM's exact register+instruction profile still leaves the gap. The residual is the specific SASS instruction-SCHEDULING/selection quality of vLLM's wheel `ptxas`, applied to the byte-identical source — which our nvcc-13.0 `ptxas` does not reproduce even when told to emit the same arch, fast-math, register count and instruction count. (Sub-finding: native+fast-math REGRESSES to 189.8 us (#68) but compute_80+fast-math is neutral at ~167 us — the arch target changes how fast-math's SASS lands, but neither is a win.) Build variants were made by swapping `--generate-code` and adding `-use_fast_math` in the exact ninja compile command, then ar+link'ing `libvllm.a`/binaries manually (CMake has no per-source `CUDA_ARCHITECTURES` property — `set_source_files_properties(... CUDA_ARCHITECTURES ...)` silently no-ops; each variant verified via `cuobjdump` and an ncu regs/instr read).
+
+**W3 — THE PARITY VERDICT: MXFP4 stays BELOW-FLOOR at c2-c8 (binding unchanged: c1 1.020 / c2 0.962 / c4 0.966 / c8 0.969 — nothing landed to move it).** Every candidate the diff could name is refuted by measurement: occupancy (identical 8.33%, smem-bound), L2-warmth (~1% both, KV ≫ L2), num_splits/grid (identical), register count (matched vLLM's 241 in C), instruction count (matched vLLM's 17k in C — and it did NOT help: the kernel is not instruction-bound), `__launch_bounds__`/reg-pressure (moot — smem-bound), arch-mirror (neutral), fast-math (regresses on native, neutral on compute_80). The flash decode term (+12.5 us/call ≈ +450 us/step, ~40% of the c8 residual after the dense-direct marlin banked the rest) is an IRREDUCIBLE (for us) **ptxas SASS-scheduling-quality gap**: on the byte-identical source, at identical grid/occupancy/smem/L2/stall-structure AND matched register+instruction counts, vLLM's wheel-`ptxas` schedules the SAME work ~10 us/call tighter than our nvcc-13.0 `ptxas`. Closing it requires vLLM's exact build `ptxas`, not something reachable by source, arch, flag, register or instruction control on our stack — so it is NOT a code, kernel, occupancy, or scheduling lever available to us. Honest c8 residual map: ~0.969 = flash ptxas-quality ~40% (this row, un-closable-by-us) + glue ~18% (portable-fusion, numerics-delicate, sub-parity) + residual-marlin/host the remainder. NO default flip owed (no throughput win). The CMakeLists NOTE records why -use_fast_math and the arch-mirror are NOT used, so they are not re-tried.
+
+Evidence: `dgx:~/mxfp4-nsys/{ours_flash_c8_ncu,vllm_flash_c8_ncu,buildB_flash_c8_ncu,buildC_flash_c8_ncu}.ncu-rep`; drivers `{ncu_flash_c8,ncu_vllm_eager2,perf_ab,buildc_ncu}.sh` + `vllm_offline_decode8_eager015.py`. Box OOM-rebooted once (util-0.5 vLLM ncu-replay), recovered, left clean.
+
 ## QUANT-CT-MXFP4-FLASH-AUDIT — `--use_fast_math` on the FA2 TUs REJECTED (measured flash REGRESSION); the flash decode gap vs vLLM is a runtime memory/occupancy effect, not the SASS instruction count (2026-08-06, `row/QUANT-CT-MXFP4-FLASH-AUDIT`, base `origin/main` `4ce9fb74`, GB10 sm_121a, PR #69)
 
 The #67 refutation re-attributed the c2-c8 MXFP4 residual FLASH-dominant and OWED a

diff --git a/.agents/specs/qwen3-mxfp4-flash-occupancy-2026-08-06.md b/.agents/specs/qwen3-mxfp4-flash-occupancy-2026-08-06.md
@@ -0,0 +1,77 @@
+# QUANT-CT-MXFP4-FLASH-OCCUPANCY — the owed ours-vs-vLLM flash decode ncu diff + the occupancy/L2/codegen lever
+
+<!-- spec-status: CLOSED -->
+Row: `QUANT-CT-MXFP4-FLASH-OCCUPANCY` (helper, `row/QUANT-CT-MXFP4-FLASH-OCCUPANCY`).
+Base: `origin/main` `f7a1e322`. Vehicle: `Yi30/Qwen3-8B-MXFP4` (dense
+`Qwen3ForCausalLM`, W4A16 Marlin keep-quant); oracle arm
+`VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel`. GB10 sm_121a.
+
+## Why this row exists
+
+`#68/#69` (FLASH-AUDIT) closed the compile lens: our flash source is byte-identical
+to vLLM's pinned `2c839c33`; `-use_fast_math` was TRIED and REJECTED (measured
++21 us/call regression — the 246→255 reg bump cuts occupancy on a latency-bound
+kernel). Fresh same-tool nsys: ours 168.8 vs vLLM 156.3 us/call at c8 (+12.5,
++450 us/step). `#69` characterized OURS (occ 8.3%, L2 53%, ~38% smem-scoreboard +
+~37% CTA-barrier stalls) but the vLLM-side ncu was LOST to a box OOM-reboot — the
+precise ours-vs-vLLM ncu diff is OWED. This row runs that diff on an idle box and
+takes whichever lever it names.
+
+## W0 free finding (from existing c8 nsys CSVs, no GPU) — the premise is REFRAMED
+
+`gpu_trace_c8_dflt` (ours) and `vllm_offline_trace_c8` (vLLM), flash_fwd_splitkv
+non-combine, c8 dominant decode:
+
+| engine | grid (GrdX×Y×Z) | blk | Reg/Trd | mean us/call | combine |
+|---|---|---|---|---|---|
+| OURS  | 1×3×64 | 128 | **216** | **174.2** | 64×1×1 reg46 3.7us |
+| vLLM  | 1×3×64 | 128 | **241** | **155.8** | 64×1×1 reg64 3.9us |
+
+The grid is **IDENTICAL** (num_splits=3, gridZ=64=batch×kv_heads — same GQA-pack,
+same split heuristic), and **ours uses FEWER registers (216 < 241)** → ours has
+MORE occupancy headroom, not less. The `#69` "8.3% occupancy, register-limited"
+number was the batch-1 short-context num_splits=1 kernel (a DIFFERENT regime); the
+prior ncu pair was context-MISMATCHED (ours ~5-token cli prompt vs vLLM lens=1024).
+So at the real c8 decode kernel, occupancy is NOT the vLLM advantage. The residual
+is codegen (our native sm_121a SASS vs vLLM's sm_80-PTX driver-JIT SASS) or L2.
+
+## Work items
+
+- **W1** — matched-workload ncu diff on BOTH engines at c8 (grid 1×3×64):
+  OURS `vllm-bench --input-len 1024 --output-len 128 --concurrency 8`; vLLM
+  `vllm_offline_decode.py NPROMPTS=8`. Full section set (LaunchStats, Occupancy,
+  SOL, MemoryWorkload/L2/DRAM, WarpStateStats, SchedulerStats). Table the diff;
+  disambiguate codegen (SM throughput / instruction stats differ) vs L2 (memory
+  section differs).
+- **W2** — the lever the diff names, mirror-first. Leading candidate: compile our
+  flash TUs for `80-virtual` (compute_80 PTX → driver JIT to sm_121), MIRRORING
+  vLLM's exact build path, if the diff shows codegen. A build-flag change like
+  fast-math — runtime arbitrates; keep the `#69` lesson (measure, never assume).
+  Byte-exact-first; near-tie razor + full battery if the reduction order shifts.
+- **W3** — binding c1..c8 ×3 production defaults vs `1.020/0.962/0.966/0.969`;
+  THE PARITY VERDICT (≥1.0 every axis ⇒ MXFP4 DONE; short ⇒ honest residual map).
+
+## Gates
+
+Byte-exact razor: #44 MXFP4-8B smoke 3/3 token-exact + coherent. If reduction
+order shifts: SACRED 0.6B/4B distributional + 32B strict, async, memcheck, eager
++ graphed. Box safety: BOTH flock locks, free -g ≥ 90, worker STOPPED, tmux +
+done-markers, sequential arms, single-load steady-state.
+
+## CLOSED — verdict (2026-08-06)
+
+W1 DONE (matched-c8 ncu, both engines, grid 1x3x64): **vLLM runs at the SAME
+8.33% occupancy** (both smem-limited to 1 CTA/SM by the byte-identical 81.92 KB
+smem), L2 ~1% both, identical stall structure. The ONLY measured diff was +13%
+executed instructions — the CODEGEN hypothesis. W2 tested every mirror-first
+lever and REFUTED all of them: arch-mirror (compute_80 PTX driver-JIT) is
+neutral; and building vLLM's EXACT recipe (compute_80 + `-use_fast_math`)
+reproduces vLLM's SASS profile EXACTLY (241 reg, 17,008 instr vs vLLM 241/17,020)
+and is STILL ~167 us, ~10 us slower than vLLM's ~157 us. So matching vLLM's arch,
+fast-math, register AND instruction count does NOT close the gap: the kernel is
+not instruction-bound, and the residual is vLLM's wheel-`ptxas` SASS-scheduling
+quality, un-reachable from our nvcc 13.0. W3 VERDICT: MXFP4 stays BELOW-FLOOR
+(binding unchanged 1.020/0.962/0.966/0.969), NO lever exists on our stack, NO
+default flip owed. No functional code shipped; a CMakeLists NOTE + benchmark
+record (#75) capture the closed levers so they are not re-tried. NO byte-exact
+razor / SACRED battery owed (nothing shipped to gate).
diff --git a/.agents/state.md b/.agents/state.md
@@ -39537,3 +39537,45 @@ are committed on the row for the next session. fp4 speed path stays CLOSED.
 Box left clean (GPU idle, both locks free, worker down, ~39 G ckpt cached at
 `dgx:~/h3fp4/ckpt`, diagnostic latents at `dgx:~/h3fp4/diag`). Evidence:
 `dgx:~/h3fp4/{diag,rt_out,out_small,out_cond}`; PR #70.
+
+## 2026-08-06T21:45 - QUANT-CT-MXFP4-FLASH-OCCUPANCY: the owed matched-c8 flash ncu diff - occupancy is IDENTICAL (8.33%), the gap is an irreducible-for-us ptxas SASS-scheduling quality difference (matched vLLM's exact reg+instr, still +10us), NO lever on our stack
+
+<!-- state: 2026-08-06T21:45 -->
+
+Ran the #69-owed ours-vs-vLLM flash decode ncu diff to a MEASURED verdict on an
+idle box (PR #75, `row/QUANT-CT-MXFP4-FLASH-OCCUPANCY`, base `f7a1e322`).
+
+**The #69 premise is REFRAMED (and its "occupancy/L2-bound" line CORRECTED).** The
+prior ncu pair was context-MISMATCHED: ours ran a ~5-token cli prompt
+(`num_splits=1`, Split=false), vLLM ran lens=1024 (Split=true) - different kernel
+instances. At the REAL c8 decode kernel both engines run the IDENTICAL grid
+`1x3x64`. The matched-c8 ncu (both engines, full section set; vLLM `enforce_eager`
+to dodge a FULL-cudagraph node-profiling driver-resource conflict; vLLM util
+0.5->0.15 after ncu-replay + the 0.5 reservation OOM-rebooted the 119 GiB pool
+once - the same #69 failure mode, now root-caused) shows: **occupancy IDENTICAL
+8.33%** (both smem-limited to 1 CTA/SM by the byte-identical 81.92 KB smem - the
+216-vs-241 reg delta is moot), **L2 ~1% on both** (KV streamed from DRAM, no
+adjacency/warmth to win), identical short-scoreboard stall structure. The only
+measured diff: ours +13% executed instructions.
+
+**W2 refuted every mirror-first lever, DECISIVELY via the exact vLLM recipe.**
+Arch-mirror (compute_80 PTX driver-JIT) is neutral. Building vLLM's EXACT recipe
+(compute_80 + `-use_fast_math`) reproduces vLLM's SASS profile EXACTLY - 241 reg,
+17,008 instr vs vLLM 241/17,020 - and is STILL ~167 us, the same as our native
+~165 us and ~10 us slower than vLLM's ~157 us. So the -13% instruction reduction
+does NOT translate to speed (the kernel is not instruction-bound) and matching
+vLLM's arch/fast-math/reg/instr does not close the gap. The residual is vLLM's
+wheel-`ptxas` SASS-scheduling quality, un-reachable from nvcc 13.0. (native+fast-
+math REGRESSES to 189.8 us per #68; compute_80+fast-math is neutral - the arch
+changes how fast-math lands but neither wins.) CMake has no per-source
+`CUDA_ARCHITECTURES` property, so build variants were made by swapping
+`--generate-code` in the exact ninja command + manual ar/link, cuobjdump-verified.
+
+**Verdict.** MXFP4 stays BELOW-FLOOR at c2-c8 (binding unchanged
+1.020/0.962/0.966/0.969). NO lever exists on our stack: occupancy, L2, num_splits,
+reg count, instruction count, `__launch_bounds__`, arch-mirror and fast-math are
+ALL refuted by measurement. The flash term (+450 us/step, ~40% of the c8 residual)
+is an irreducible-for-us ptxas quality gap. NO default flip owed; no functional
+code shipped (CMakeLists NOTE + benchmark-record #75 record the closed levers).
+Box left clean (GPU idle, both locks free, worker down). Evidence:
+`dgx:~/mxfp4-nsys/{ours,vllm,buildB,buildC}_flash_c8_ncu.ncu-rep`; PR #75.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -1412,15 +1412,27 @@ if(VLLM_CPP_CUDA)
       src/vt/cuda/cuda_flash_attn_fa2.cu
       ${_FA2_KERNEL_SRCS})
     target_compile_definitions(vllm PUBLIC VLLM_CPP_FLASH_ATTN)
-    # NOTE (QUANT-CT-MXFP4-FLASH-AUDIT #68): -use_fast_math was TRIED here to mirror
-    # vLLM's own flash-attn build (its _vllm_fa2_C is fast-math; it makes our decode
-    # kernel's SASS match vLLM's 4832/REG255 vs 4880/REG255). But a controlled
-    # same-build nsys A/B REFUTED it as a speed lever: this decode kernel is
-    # memory-latency-bound (ncu occ ~10.7%), so the +11.6% instruction reduction does
-    # not help, while fast-math's higher register count (246->255) LOWERS occupancy
-    # and REGRESSED flash 168.8 -> 189.8 us/call. So we do NOT pass -use_fast_math.
-    # The residual flash gap to vLLM (168.8 vs 156.3 us/call) is a runtime memory/L2
-    # effect, not the SASS; see .agents/benchmark-record.md.
+    # NOTE (QUANT-CT-MXFP4-FLASH-AUDIT #68 + FLASH-OCCUPANCY #75): the c8 decode
+    # flash gap vs vLLM (ours ~168-174 vs vLLM ~156 us/call) is a COMPILER-CODEGEN
+    # gap, MEASURED to a verdict — NOT occupancy, NOT L2, NOT num_splits, and NOT a
+    # flag we can set. The owed matched-c8 ncu (both engines, IDENTICAL grid
+    # 1x3x64) shows occupancy is IDENTICAL 8.33% (both smem-limited to 1 CTA/SM by
+    # the byte-identical 81.92 KB kSmemSize — the 216-vs-241 reg delta is moot),
+    # L2 hit ~1% on BOTH (KV streamed from DRAM: no adjacency/warmth to win), and
+    # the same ~60% short-scoreboard stall structure. The ONLY difference: vLLM's
+    # SASS executes ~13% FEWER instructions (3.27M vs 3.70M) in this latency-bound
+    # ~1-warp/scheduler regime — but that instruction count is NOT the cause. #75
+    # built the exact vLLM recipe: compiling the kernel TUs for compute_80 PTX +
+    # -use_fast_math REPRODUCES vLLM's SASS profile EXACTLY (241 reg, 17,008 instr
+    # vs vLLM 241/17,020) and is STILL ~167 us — the SAME as our native ~165 us and
+    # still ~10 us slower than vLLM's ~157 us. So matching vLLM's arch, fast-math,
+    # register AND instruction counts does NOT close the gap: the kernel is not
+    # instruction-bound, and the residual is vLLM's wheel-`ptxas` SASS-SCHEDULING
+    # quality, not reachable from our nvcc 13.0. (native+fast-math REGRESSES to
+    # 189.8 us, #68; compute_80 PTX driver-JIT alone is neutral, #75 Build B.) So
+    # we keep the plain native sm_121a compile and pass no -use_fast_math. The
+    # residual is an irreducible-for-us ptxas SASS-quality gap on a byte-identical
+    # kernel; see .agents/benchmark-record.md (#75).
     set_source_files_properties(
       src/vt/cuda/cuda_flash_attn_fa2.cu
       ${_FA2_KERNEL_SRCS} PROPERTIES