D1/B1: bare-matmul op-kind + NVIDIA GEMM arbiter candidates (sm_120)

docs/audit/compiler/COMPILER_REFACTOR_PLAN.md

@@ -308,13 +308,21 @@ chains, small attention). Crown-jewel GEMM stays Tier 2/3.
   documented *instruction-encoding skeleton* (`emit_wgmma_matmul_ptx` — its own header
   says "NOT a complete assemblable kernel: needs smem matrix descriptors + TMA/cp.async
   + the full accumulator operand list"), so completing it is a real Hopper WGMMA-kernel
-  build (assemble-only here; execution needs Hopper) — **not** a bug fix; sm_100
-  tcgen05; and wiring the emit lane as a **Tier-2 EMITTED** arbiter candidate (blocked
-  on a bare-matmul op-kind — the candidate registry has no matmul op today, only
-  fused_region/attention/gated/pointwise). Unlike C3, NVIDIA has no *fused* shipped
-  kernel to register as a Tier-3 `FusedRegion` candidate yet (the shipped kernel is a
-  pure GEMM served by the
-  jit `nvidia_mma` executor).
+  build (assemble-only here; execution needs Hopper) — **not** a bug fix; and sm_100
+  tcgen05.
+  **B1 landed 2026-07-07 — the emit lane is now a first-class arbiter candidate.**
+  A bare-matmul op-kind (`candidate.OP_MATMUL` + `fusion_core.MatmulRegion` +
+  `verify_synthesized_matmul`) unblocks the D1 arbiter for plain GEMM. Two NVIDIA
+  matmul candidates register under `(nvidia, matmul)`: the **shipped** mma.sync GEMM
+  (`NvidiaMmaGemmShippedCandidate`, Tier-3 via `runtime._nvidia_mma_gemm_2d`) and the
+  **compiler-emitted** `ptx_emit` GEMM (`NvidiaMmaGemmEmittedCandidate`, Tier-2 via the
+  launch bridge, `runtime._nvidia_ptx_gemm_2d`). Both F4-gated by the universal oracle;
+  tier-priority picks the shipped lane by default (lead-safe, Decision #28), the E3
+  `force` hatch selects the emitted lane. Live-proven on sm_120 (bf16/f16 ×
+  16x8x16/32x16x32/64x64x64). So NVIDIA now has its Tier-3 *hand-tuned* GEMM candidate
+  (the pure-GEMM shipped kernel, previously only reachable via the jit `nvidia_mma`
+  executor) alongside the Tier-2 emitted lane — D2's measured loop is the follow-on
+  that lets Tier-2 win where faster.
 - **C3 · ROCm generic synth → HIP** — **generic lane LANDED 2026-07-06**, `[MAC]`
   author → `[AMD]` proof. `emit/rocm_hip.py` is now a **full three-seam plugin**
   (parallel to x86): `RocmHipEmitter` turns a `FusedRegion` into HIP source (a

python/tessera/compiler/emit/candidate.py

@@ -53,6 +53,9 @@ class Tier(IntEnum):
 OP_ATTENTION = "attention"
 OP_GATED_MATMUL = "gated_matmul"
 OP_POINTWISE = "pointwise"
+#: A bare matmul ``D = A @ B`` (no fusion) — the op-kind for plain-GEMM candidates
+#: (NVIDIA emitted mma.sync vs shipped GEMM). Keyed on ``fusion_core.MatmulRegion``.
+OP_MATMUL = "matmul"
 
 #: op-kind → (verify function name, KernelRunner method) — the seam that lets the
 #: arbiter reuse the exact universal F4 oracle for each candidate.
@@ -61,6 +64,7 @@ class Tier(IntEnum):
     OP_ATTENTION: ("verify_synthesized_attention", "run_fused_attention"),
     OP_GATED_MATMUL: ("verify_synthesized_gated", "run_gated_matmul_region"),
     OP_POINTWISE: ("verify_synthesized_pointwise", "run_pointwise_graph"),
+    OP_MATMUL: ("verify_synthesized_matmul", "run_matmul"),
 }
 
 
@@ -180,6 +184,9 @@ def run_gated_matmul_region(self, region, *a, **k):
         def run_pointwise_graph(self, region, *a, **k):
             return self._route("run_pointwise_graph", region, a, k)
 
+        def run_matmul(self, region, *a, **k):
+            return self._route("run_matmul", region, a, k)
+
         def _route(self, called, region, a, k):
             if called != method:
                 raise NotImplementedError(

python/tessera/compiler/emit/nvidia_cuda.py

@@ -46,6 +46,7 @@
 from tessera.compiler.emit._fused_scalar_body import row_compute_body
 from tessera.compiler.emit.candidate import (
     OP_FUSED_REGION,
+    OP_MATMUL,
     Candidate,
     Tier,
     register_candidate,
@@ -61,7 +62,7 @@
     register_emitter,
     register_runner,
 )
-from tessera.compiler.fusion_core import FusedRegion
+from tessera.compiler.fusion_core import FusedRegion, MatmulRegion
 
 _TARGET = "nvidia"
 _LANG = "cuda"
@@ -275,9 +276,98 @@ def run(self, region: Any, A: Any, B: Any, bias: Any = None,
                                                residual=residual)
 
 
+# ── D1 matmul candidates (B1) — bare GEMM, Tier-2 emitted vs Tier-3 shipped ────
+#
+# The arbiter enumerates these per (target="nvidia", op=matmul) and F4-gates each.
+# Tier-priority (Decision #28) prefers the hand-tuned shipped lane by default; D2's
+# measured loop lets the emitted lane win where it is faster + in accuracy budget.
+# Both are 16-bit storage (bf16/f16) → f32 accumulate, so they declare the f16
+# budget the oracle honors. Off an NVIDIA GPU / without the built libs they decline
+# to the reference and drop out of the enumeration.
+
+_GEMM_F16_ATOL = 5e-3          # 16-bit storage vs the f32 reference (Decision #28)
+_GEMM_DTYPES = ("bfloat16", "float16")
+
+
+def _aligned_2d(A: Any, B: Any) -> bool:
+    """A (M,K) @ B (K,N) with the emitted kernel's tile alignment (M%16,N%8,K%16)."""
+    import numpy as np
+    Aa, Ba = np.asarray(A), np.asarray(B)
+    if Aa.ndim != 2 or Ba.ndim != 2 or Aa.shape[1] != Ba.shape[0]:
+        return False
+    M, K = Aa.shape
+    _, N = Ba.shape
+    return M % 16 == 0 and N % 8 == 0 and K % 16 == 0
+
+
+class NvidiaMmaGemmShippedCandidate(Candidate):
+    """Tier-3 (hand-tuned): the shipped ``libtessera_nvidia_gemm`` mma.sync GEMM —
+    the crown-jewel lane, arbiter default until D2 measures otherwise. Serves any
+    (unaligned OK) bf16/f16 matmul; declines off an NVIDIA GPU."""
+
+    name = "nvidia_mma_gemm_shipped"
+    tier = Tier.HAND_TUNED
+    target = _TARGET
+    op = OP_MATMUL
+    accuracy_atol = _GEMM_F16_ATOL
+
+    def available(self) -> bool:
+        try:
+            from tessera import runtime as rt
+            return rt._nvidia_mma_runtime_available()
+        except Exception:
+            return False
+
+    def applies_to(self, region: Any) -> bool:
+        return isinstance(region, MatmulRegion) and region.dtype in _GEMM_DTYPES
+
+    def run(self, region: Any, A: Any, B: Any, *a: Any, **k: Any) -> tuple[Any, str]:
+        try:
+            from tessera import runtime as rt
+            return rt._nvidia_mma_gemm_2d(A, B, region.dtype), "nvidia_mma_shipped"
+        except Exception:
+            return region.reference(A, B), "reference"
+
+
+class NvidiaMmaGemmEmittedCandidate(Candidate):
+    """Tier-2 (emitted): the compiler-EMITTED ``ptx_emit`` mma.sync GEMM driven
+    through the launch bridge — the C2 emit lane as a first-class arbiter candidate.
+    Serves ALIGNED (M%16/N%8/K%16) bf16/f16 matmuls; declines (to the reference) for
+    ragged shapes or off an NVIDIA GPU / without the built bridge."""
+
+    name = "nvidia_mma_gemm_emitted"
+    tier = Tier.EMITTED
+    target = _TARGET
+    op = OP_MATMUL
+    accuracy_atol = _GEMM_F16_ATOL
+
+    def available(self) -> bool:
+        try:
+            from tessera import runtime as rt
+            return (rt._load_nvidia_ptx_launch() is not None
+                    and rt._nvidia_mma_runtime_available())
+        except Exception:
+            return False
+
+    def applies_to(self, region: Any) -> bool:
+        return isinstance(region, MatmulRegion) and region.dtype in _GEMM_DTYPES
+
+    def run(self, region: Any, A: Any, B: Any, *a: Any, **k: Any) -> tuple[Any, str]:
+        if not _aligned_2d(A, B):              # emitter is aligned-only (for now)
+            return region.reference(A, B), "reference"
+        try:
+            from tessera import runtime as rt
+            return rt._nvidia_ptx_gemm_2d(A, B, region.dtype), "nvidia_ptx_gemm"
+        except Exception:
+            return region.reference(A, B), "reference"
+
+
 # ── registration (import side effect, exactly like rocm_hip / x86_llvm) ────────
 register_emitter(NvidiaCudaEmitter())
 register_compiler(_TARGET, _nvidia_cuda_compile_fn)
 register_runner(NvidiaCudaRunner(), default=False)
 
 register_candidate(NvidiaGenericCudaCandidate())
+# Bare-GEMM lanes: hand-tuned shipped (Tier 3) + compiler-emitted (Tier 2).
+register_candidate(NvidiaMmaGemmShippedCandidate())
+register_candidate(NvidiaMmaGemmEmittedCandidate())

python/tessera/compiler/fusion.py

@@ -17,6 +17,7 @@
     FusedRegion as FusedRegion,
     FusionCost as FusionCost,
     GatedMatmulRegion as GatedMatmulRegion,
+    MatmulRegion as MatmulRegion,
     NormChainRegion as NormChainRegion,
     POINTWISE_OPS as POINTWISE_OPS,
     PointwiseGraphRegion as PointwiseGraphRegion,
@@ -68,6 +69,7 @@
     should_fuse_region as should_fuse_region,
     verify_synthesized_attention as verify_synthesized_attention,
     verify_synthesized_gated as verify_synthesized_gated,
+    verify_synthesized_matmul as verify_synthesized_matmul,
     verify_synthesized_pointwise as verify_synthesized_pointwise,
     verify_synthesized_region as verify_synthesized_region,
 )
@@ -196,6 +198,8 @@
     "discover_fusable_regions",
     "discover_attention_regions",
     "GatedMatmulRegion",
+    "MatmulRegion",
+    "verify_synthesized_matmul",
     "synthesize_gated_matmul_msl",
     "run_gated_matmul_region",
     "should_fuse_gated",

python/tessera/compiler/fusion_core.py

@@ -302,6 +302,41 @@ def reference(self, A: np.ndarray, B: np.ndarray,
         return out.astype(np.float32)
 
 
+def _round_to_storage(x: "np.ndarray", dtype: str) -> "np.ndarray":
+    """Round f32 ``x`` to the 16-bit storage ``dtype`` and back to f32, so a
+    reference sees the same rounded operands the 16-bit GEMM kernels do."""
+    a = np.ascontiguousarray(x, np.float32)
+    if dtype in ("float16", "f16"):
+        return a.astype(np.float16).astype(np.float32)
+    if dtype in ("bfloat16", "bf16"):
+        try:
+            import ml_dtypes
+            return a.astype(ml_dtypes.bfloat16).astype(np.float32)
+        except Exception:                      # round-to-nearest-even fallback
+            u = a.view(np.uint32).astype(np.uint64)
+            bits = (((u + 0x7FFF + ((u >> 16) & 1)) >> 16) << 16).astype(np.uint32)
+            return bits.view(np.float32)
+    raise ValueError(f"unsupported matmul storage dtype {dtype!r}")
+
+
+@dataclass(frozen=True)
+class MatmulRegion:
+    """A bare matmul ``D = A @ B`` (no fusion) — the region kind the D1 arbiter
+    keys a plain-GEMM candidate on (there is no matmul feature in ``FusedRegion``,
+    which always carries at least one fused op). ``dtype`` is the 16-bit storage
+    (``bfloat16``/``float16``); the accumulate is f32, matching the emitted
+    ``mma.sync`` and shipped GEMM kernels."""
+
+    dtype: str = "bfloat16"
+
+    def reference(self, A: np.ndarray, B: np.ndarray) -> np.ndarray:
+        """The f32 result of ``A @ B`` with both operands rounded to ``dtype``
+        first — the horizontal-oracle ground truth the GEMM candidate matches."""
+        Aq = _round_to_storage(A, self.dtype)
+        Bq = _round_to_storage(B, self.dtype)
+        return (Aq @ Bq).astype(np.float32)
+
+
 # ─────────────────────────────────────────────────────────────────────────────
 # F2b-tiled — threadgroup-tiled synthesis for large N (the stack kernel caps at
 # SYNTH_MAX_N; this lifts it to SYNTH_MAX_N_TILED via dynamic threadgroup memory)
@@ -934,6 +969,36 @@ def verify_synthesized_region(region: FusedRegion, *, seed: int = 0,
     return verdict
 
 
+def verify_synthesized_matmul(region: "MatmulRegion", *, seed: int = 0,
+                              atol: float = 1e-3, force: bool = False,
+                              runner: KernelRunner | None = None) -> bool:
+    """Codegen-gated oracle for a bare-matmul candidate (see
+    ``verify_synthesized_region``): run the candidate's GEMM on an aligned probe
+    and compare to the dtype-rounded ``A @ B`` reference. Arbiter-only — the
+    ``runner`` is a candidate adapter exposing ``run_matmul``; a runner without it
+    (a real backend runner) trusts the reference, since nothing device-emitted is
+    in play for this op."""
+    r = runner or _runner()
+    run = getattr(r, "run_matmul", None)
+    if run is None:
+        return True
+    key = (r.target, "M", region.dtype)
+    if not force and key in _VERIFY_CACHE:
+        return _VERIFY_CACHE[key]
+    rng = np.random.default_rng(seed)
+    M, N, K = 32, 16, 32                        # aligned probe (M%16, N%8, K%16)
+    A = (rng.standard_normal((M, K)) * 0.4).astype(np.float32)
+    B = (rng.standard_normal((K, N)) * 0.4).astype(np.float32)
+    out, execution = run(region, A, B)
+    if execution in REFERENCE_EXECUTIONS:
+        verdict = True
+    else:
+        verdict = bool(np.allclose(out, region.reference(A, B),
+                                   atol=_effective_atol(r, atol)))
+    _VERIFY_CACHE[key] = verdict
+    return verdict
+
+
 def verify_synthesized_attention(region: AttentionRegion, *, seed: int = 0,
                                  atol: float = 1e-3, force: bool = False,
                                  runner: KernelRunner | None = None) -> bool: