gstoner · gstoner · Jul 7, 2026 · Jul 7, 2026 · Jul 7, 2026 · Jul 7, 2026
diff --git a/docs/audit/backend/nvidia/NVIDIA_AUDIT.md b/docs/audit/backend/nvidia/NVIDIA_AUDIT.md
@@ -63,14 +63,24 @@ This document consolidates NVIDIA-specific audit material.
 ## Still Open
 
 The original "no execution row / not hardware-proven" gaps are **closed**
-(above). What remains is breadth beyond the first proven kernel:
+(above), and the **compiler-generated lane is now landed + hardware-proven** on
+sm_120 (RTX 5070 Ti, PRs #290–#297):
 
-- **Compiler-GENERATED NVIDIA lane** — the analog of ROCm's `rocm_compiled`
-  (a `tessera-opt` NVIDIA pipeline that generates + serializes the kernel
-  in-process, rather than dispatching the hand-shipped `libtessera_nvidia_gemm`
-  symbol). Today's `nvidia_mma` lane mirrors `rocm_wmma` (shipped symbol), not
-  the compiler-generated path.
-- **NVFP4 block-scale matmul (#9)** — the warp `mma.sync…block_scale`
+- **Compiler-GENERATED NVIDIA lane — LANDED.** `emit/nvidia_cuda.py` is a full
+  three-seam plugin (emitter + `nvcc` compile + ctypes runner) that synthesizes,
+  compiles, and launches kernels in-process for **all four `fusion_core` region
+  kinds** — fused matmul-epilogue, flash-attention (C4), SwiGLU gate + pointwise
+  DAG (C5) — each F4-gated on-device. The emit-path `mma.sync` GEMM
+  (`ptx_emit.py` → the shipped `tessera_nvidia_ptx_launch` bridge: driver-JIT +
+  `cuLaunchKernel`) executes the *emitted* PTX, distinct from the hand-shipped
+  `libtessera_nvidia_gemm` symbol. Both are first-class D1 arbiter candidates:
+  the shipped GEMM is **Tier-3 hand-tuned**, the emitted GEMM **Tier-2 emitted**
+  (B1), with D2 measured autotune + D3 fallback logging choosing/observing between
+  them. So NVIDIA now has the `rocm_compiled` analog it lacked, plus the arbiter
+  surface.
+- **NVFP4 block-scale matmul (#9)** — **emit + ptxas-assemble landed** (#291,
+  `emit_nvfp4_block_scale_mma_ptx`); on-device execution + non-unit-scale numerics
+  stay gated on the PTX-ISA scale-distribution spec. The warp `mma.sync…block_scale`
   instruction already assembles + executes on `sm_120a` (see
   `spikes/sm120_mma_sync/`); productization is pending the PTX ISA
   scale-distribution spec for numerics.
@@ -87,14 +97,20 @@ The original "no execution row / not hardware-proven" gaps are **closed**
 
 ## Next Work
 
-1. Build the compiler-generated NVIDIA lane (the `rocm_compiled` analog) via a
-   `tessera-opt` NVIDIA pipeline; add its executable matrix row.
-2. Land NVFP4 block-scale matmul once the scale-distribution numerics are
-   grounded; flip its manifest row when execute-and-compare passes on `sm_120a`.
-3. Bring the sm_120 `mma.sync` flash-attention forward to the same
-   execute-and-compare bar (attention analog of the GEMM proof).
-4. Author an sm_120 `mma.sync` kernel inventory (sibling to
-   `docs/nvidia_cuda13_kernel_inventory.md`).
+Done (2026-07-07): the compiler-generated lane (#290–#297), the sm_120 `mma.sync`
+flash-attention execute-compare (C4), and the sm_120 kernel-inventory doc
+(`docs/nvidia_sm120_mma_sync_kernel_inventory.md`). Remaining:
+
+1. **NVFP4 block-scale execution + numerics** — bind the fp4 fragment packing and
+   flip the manifest row once execute-and-compare passes on `sm_120a` and the
+   scale-distribution numerics are grounded (emit + ptxas-assemble already land).
+2. **mma.sync tensor-core versions** of the flash-attention + fused lanes (the
+   correctness-first synth lanes are proven; these are the perf follow-on) + the
+   **D2 fleet-shared autotune corpus** (persist `MeasureCache`, Theory §7.5).
+3. **Dtypes beyond f32** for the fused / attention / gated lanes (16-bit storage
+   is served by the B1 matmul lane today).
+4. **`wgmma` sm_90a** — complete the instruction-encoding skeleton into a real
+   Hopper WGMMA kernel (assemble-only until a Hopper box) — and **sm_100 tcgen05**.
 5. Promote sm_80/90/100 manifest rows only when their own silicon is available
    and the generated dashboards agree.
 

diff --git a/docs/nvidia_sm120_mma_sync_kernel_inventory.md b/docs/nvidia_sm120_mma_sync_kernel_inventory.md
@@ -0,0 +1,102 @@
+---
+status: Informative
+classification: Reference / Kernel Inventory
+authority: Companion to docs/nvidia_cuda13_kernel_inventory.md
+last_updated: 2026-07-07
+---
+
+# NVIDIA sm_120 (consumer Blackwell) mma.sync Kernel Inventory
+
+> Hardware-verified companion to
+> [`docs/nvidia_cuda13_kernel_inventory.md`](nvidia_cuda13_kernel_inventory.md)
+> (which is the SM_90+ WGMMA *planning* inventory). This doc enumerates the
+> kernels Tessera actually **synthesizes, compiles, and runs on sm_120** — proven
+> on an RTX 5070 Ti (compute cap 12.0, CUDA 13.3, driver ≥610.43.02). Statuses
+> here are execution truth, not roofline targets; the generated dashboards
+> (`docs/audit/generated/`) remain the drift-gated count/status surface.
+
+Three kernel families feed the NVIDIA lane, each a distinct compile path:
+
+1. **Compiler-emitted CUDA** — `emit/nvidia_cuda.py` synthesizes CUDA C, compiles
+   it with `nvcc -arch=sm_120a`, and launches via `ctypes` (the three-seam generic
+   plugin). Correctness-first, arch-agnostic; F4-gated on-device.
+2. **Emit-path PTX** — `ptx_emit.py` emits PTX text, the shipped
+   `tessera_nvidia_ptx_launch` bridge driver-JITs it (`cuModuleLoadDataEx`, cached
+   by kernel name) and launches it (`cuLaunchKernel`).
+3. **Shipped C-ABI** — `libtessera_nvidia_gemm.so` NVRTC-compiles a hand-written
+   CUDA-C `mma.sync` GEMM at first call (the hand-tuned lane).
+
+Execution status legend: **✅ proven** (execute-and-compare on sm_120) ·
+**🟡 assemble-only** (ptxas accepts; not launched/numerically gated) ·
+**⬜ skeleton** (instruction-encoding only, not assemblable).
+
+---
+
+## 1. Compiler-emitted CUDA lanes (`emit/nvidia_cuda.py`)
+
+| Entry symbol | Op / region | Shape model | dtype | Kernel shape | Status |
+|---|---|---|---|---|---|
+| `tessera_nvidia_fused` | `FusedRegion` — matmul + prologue/epilogue/residual/reduction | runtime M/N/K | f32 | one thread per output row | ✅ |
+| `tessera_nvidia_attn` | `AttentionRegion` — `O = softmax(scale·Q·Kᵀ)·V` | runtime M/Nk/D/Dv (Dv ≤ 256) | f32 | flash: one query/thread, online softmax, streaming KV, causal + transpose flags | ✅ (C4) |
+| `tessera_nvidia_gated` | `GatedMatmulRegion` — SwiGLU gate `gate_act(A·Wg) ⊙ (A·Wu)` | runtime M/K/H | f32 | one output row/thread, shared A load | ✅ (C5) |
+| `tessera_nvidia_pointwise` | `PointwiseGraphRegion` — same-shape pointwise DAG | runtime numel | f32 | one element/thread; DAG from `POINTWISE_OPS` C-expr table (+ NaN-safe `sign`/`clamp` shims) | ✅ (C5) |
+
+All four are Tier-1 **synthesized** D1 arbiter candidates
+(`Nvidia{Generic,FlashAttn,Gated,Pointwise}…Candidate`). f32 is the
+correctness-first floor; 16-bit storage is served by the GEMM lanes below.
+
+## 2. Emit-path PTX lanes (`ptx_emit.py` → `tessera_nvidia_ptx_launch`)
+
+| Entry symbol | Op | Shape | dtype | Instruction | Status |
+|---|---|---|---|---|---|
+| `tessera_mma_m16n8k16_bf16` | matmul (single tile) | 16×8×16 | bf16→f32 | `mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32` | ✅ |
+| `tessera_mma_gemm_bf16` | matmul (general) | aligned M%16/N%8/K%16, index < 2³¹ | bf16→f32 | m16n8k16 tile, K-loop + grid-tiled | ✅ (C2 breadth) |
+| `tessera_mma_gemm_f16` | matmul (general) | aligned M%16/N%8/K%16, index < 2³¹ | f16→f32 | m16n8k16 tile, K-loop + grid-tiled | ✅ |
+| `tessera_nvfp4_mma_m16n8k64` | matmul (block-scale) | 16×8×64 | fp4 e2m1 + ue4m3 scales → f32 | `mma.sync…m16n8k64…kind::mxf4nvf4.block_scale.scale_vec::4X` | 🟡 assemble-only (numerics gated on PTX-ISA scale spec) |
+| `tessera_wgmma_matmul_bf16` | matmul (Hopper) | m64n{64,128,256}k16 | bf16→f32 | `wgmma.mma_async…` | ⬜ skeleton (needs smem descriptors + TMA; sm_90a, no Hopper box) |
+
+The emitted general GEMM (`tessera_mma_gemm_{bf16,f16}`) is the **Tier-2 emitted**
+D1 candidate (`NvidiaMmaGemmEmittedCandidate`); the launch bridge rejects
+i32-index-overflow shapes (element count > 2³¹) honestly.
+
+## 3. Shipped C-ABI lane (`libtessera_nvidia_gemm.so`)
+
+| Entry symbol | Op | Shape | dtype | Status |
+|---|---|---|---|---|
+| `tessera_nvidia_mma_gemm_f16` | matmul (general tiled) | any M/N/K | f16→f32 | ✅ |
+| `tessera_nvidia_mma_gemm_bf16` | matmul (general tiled) | any M/N/K | bf16→f32 | ✅ |
+| `tessera_nvidia_mma_gemm_tf32` | matmul (general tiled) | any M/N/K | fp32/tf32-math→f32 | ✅ |
+
+The shipped GEMM is the **Tier-3 hand-tuned** D1 candidate
+(`NvidiaMmaGemmShippedCandidate`) — the arbiter default (lead-safe, Decision #28),
+displaced only when D2's measured loop proves the emitted lane faster + in budget.
+
+---
+
+## Arbiter mapping (D1/D2/D3)
+
+| op | Tier-3 hand-tuned | Tier-2 emitted | Tier-1 synthesized |
+|---|---|---|---|
+| `matmul` | `tessera_nvidia_mma_gemm_*` (shipped) | `tessera_mma_gemm_*` (PTX bridge) | — |
+| `fused_region` | — | — | `tessera_nvidia_fused` |
+| `attention` | — | — | `tessera_nvidia_attn` |
+| `gated_matmul` | — | — | `tessera_nvidia_gated` |
+| `pointwise` | — | — | `tessera_nvidia_pointwise` |
+
+Selection: tier-priority by default; `emit/autotune.py` measures on-device and
+caches the fastest per `(device, op, shape-bucket, dtype)` (D2); every dispatch is
+recorded in the arbiter fallback log (D3, `arbiter_dispatch_histogram`).
+
+## Toolchain pin
+
+CUDA Toolkit **13.3** (PTX ISA 9.3); target `sm_120a` (FP4
+`mma.sync.block_scale`); driver ≥610.43.02; smem 100 KB/SM. The emit-path and
+shipped lanes need only the host compiler + CUDA driver (libcuda) + NVRTC at load
+time; the compiler-emitted CUDA lane needs `nvcc`.
+
+## Still open
+
+NVFP4 execution + non-unit-scale numerics; mma.sync tensor-core versions of the
+attention + fused lanes (perf); dtypes beyond f32 for the fused/attention/gated
+lanes; the Hopper `wgmma` completion (sm_90a) and sm_100 `tcgen05` (their own
+silicon). See [`docs/audit/backend/nvidia/NVIDIA_AUDIT.md`](audit/backend/nvidia/NVIDIA_AUDIT.md).