feat(nvidia): the shipped GEMM becomes the first declared delegate

docs/audit/backend/nvidia/todo.md

@@ -183,6 +183,58 @@ that box probes `avx512=False` and its x86 lanes skip honestly. Princess-Luna
 `hardware_amx` must never be used to mean "x86 hardware" — see the standing
 section in `docs/audit/backend/x86/todo.md`.
 
+## `NVIDIA-TIER-PRIORITY-IS-WRONG-AT-SCALE-2026-08-30` — measured, not argued
+
+**The first thing the delegate's device timer produced, and it contradicts the
+arbiter's default.** Decision #28 displaces a hand-tuned kernel when a compiled
+one measures **faster and in accuracy budget**. On sm_120 (RTX 5070), f16,
+square, device-resident CUDA-event timing, spreads of 0.000–0.008 ms across
+repeats:
+
+| shape | `nvidia_mma_gemm_shipped` (T3) | `nvidia_tile_matmul_shared` (T2) | faster | max\|err\| |
+|---|---|---|---|---|
+| 512³ | **0.043 ms** | 0.059 ms | delegate, by 37% | both 2.48e-05 |
+| 1024³ | 0.320 ms | **0.312 ms** | compiled, by 2.3% | both 6.10e-05 |
+| 2048³ | 2.448 ms | **2.051 ms** | compiled, by 16.2% | both 1.54e-04 |
+
+The error columns are **equal at every shape**, so the in-budget half is
+satisfied outright. The displacement condition therefore holds at 1024³ and
+above — and `arbitrate()` still returns the delegate, because tier priority is
+the default and D2's measured loop is not wired into this path.
+
+Two things follow, and neither was visible before:
+
+* **The compiled Tessera kernel beats the hand-tuned one at scale.** That is a
+  result about the compiler, not about the arbiter.
+* **The crossover is shape-dependent**, which is the concrete argument for
+  shape-bucketed measured selection rather than a single global winner. A
+  flat "measurement beats tier" switch would regress 512³ by 37%.
+
+**Do not read this as "delete the delegate."** It wins by 37% at 512³, and
+Decision #28's lead-safety exists precisely so a crown-jewel lane is displaced
+per shape by evidence rather than wholesale by policy.
+
+**Why it was invisible until now.** End-to-end wall time ranks the two the
+*other* way — 9.4 ms vs 33.1 ms at 2048³ — because it is host-dominated: the
+Tile lane spends 2.99 ms on device inside 34.0 ms of wall time, and the two
+lanes do not share a host path, so e2e compares numpy conversions. The Tier-3
+lane had no device timer at all, so the honest comparison could not be made.
+Pinned by `tests/device/nvidia/test_shipped_gemm_delegate.py`.
+
+**Open follow-ups.**
+1. Wire shape-bucketed measured selection into the `OP_MATMUL` NVIDIA path so
+   the 1024³+ crossover is acted on. The `measure` hook and the autotune
+   corpus already exist; nothing calls them for this bucket.
+2. `nvidia_mma_gemm_emitted` still has no device timer. The NVIDIA backend
+   carries **two block-index conventions**: `ptx_emit` and the shipped AOT
+   kernel map x→M, y→N, while `NVIDIALowering.cpp` and the launch bridge's
+   `benchmarkTileGemm16` map x→N, y→M. Driving the emitted kernel through the
+   harness returns rc=5, and registering its geometry would launch a
+   transposed grid (at 512×512: rows to 1024, columns only to 256 — half the
+   output unwritten, with a plausible-looking latency). Unify the convention,
+   or give the harness an explicit axis-order field. A unit test pins the
+   current mapping so "fixing" one side fails loudly.
+
 ## `SM120-BUILD-CONFIG-RESOLVED-2026-08-30` — there was no trade; use CUDA=ON
 
 **Superseded the "fleet-config decision" framing below: configuring the NVIDIA

python/tessera/compiler/emit/candidate.py

@@ -166,6 +166,19 @@ def applies_to(self, region: Any) -> bool:
         region with a reduction epilogue it cannot fuse)."""
         return True
 
+    def accuracy_budget(self, region: Any) -> "tuple[float | None, float | None]":
+        """``(atol, rtol)`` the F4 oracle must hold this candidate to for
+        ``region``. Defaults to the flat class attributes.
+
+        A region hook rather than a constant because a candidate serving more
+        than one storage dtype does not have one budget: bf16 carries 8
+        significand bits to f16's 11, so holding both to f16's number either
+        rejects a correct bf16 kernel or accepts a bad f16 one. Backends that
+        register one candidate per dtype are unaffected -- they simply keep
+        answering with their class attributes.
+        """
+        return self.accuracy_atol, self.accuracy_rtol
+
     def measure_device_latency(self, region: Any, *inputs: Any, reps: int = 100,
                                warmup: int = 10) -> float | None:
         """Optional device-resident latency in milliseconds.
@@ -248,7 +261,7 @@ def candidates_for(target: str, op: str) -> list[Candidate]:
 _PROBE_NS = "candidate"
 
 
-def _as_runner(candidate: Candidate) -> Any:
+def _as_runner(candidate: Candidate, region: Any = None) -> Any:
     """Wrap a :class:`Candidate` as a :class:`KernelRunner` so the existing
     ``fusion_core.verify_synthesized_*`` oracle gates it unchanged — the whole
     point of D1's F4 reuse. Only the candidate's own op method is wired; the other
@@ -263,11 +276,14 @@ def _as_runner(candidate: Candidate) -> Any:
     from tessera.compiler.emit.kernel_emitter import KernelRunner
 
     _, method = _OP_VERIFY[candidate.op]
+    # Resolved per region so a multi-dtype candidate is gated on the budget it
+    # declared for THIS dtype rather than on a single class-level number.
+    _atol, _rtol = candidate.accuracy_budget(region)
 
     class _CandidateRunner(KernelRunner):
         target = f"{_PROBE_NS}::{candidate.target}::{candidate.name}"
-        accuracy_atol = candidate.accuracy_atol
-        accuracy_rtol = candidate.accuracy_rtol
+        accuracy_atol = _atol
+        accuracy_rtol = _rtol
         last_execution: str | None = None
 
         def run_fused_region(self, region, *a, **k):
@@ -323,7 +339,7 @@ def verify_candidate(candidate: Candidate, region: Any, *, atol: float = 1e-3,
 
     verify_name, _ = _OP_VERIFY[candidate.op]
     verify = getattr(fusion_core, verify_name)
-    adapter = _as_runner(candidate)
+    adapter = _as_runner(candidate, region)
     matched = bool(verify(region, runner=adapter, force=True, atol=atol, seed=seed))
     if not matched:
         return False

python/tessera/compiler/emit/delegate_contract.py

@@ -292,12 +292,58 @@ class DelegatedCandidate(Candidate):
     output rather than trusting it above.
     """
 
-    def __init__(self, contract: DelegateContract, *, target: str, op: str) -> None:
+    #: Contract fields that describe the delegate itself rather than one dtype
+    #: route. Every member of a `variants` family must agree on these, or the
+    #: representative contract misdescribes the family.
+    _FAMILY_INVARIANT = ("arch", "accuracy", "determinism", "covers",
+                         "binding", "provenance")
+
+    def __init__(self, contract: DelegateContract, *, target: str, op: str,
+                 name: str | None = None,
+                 variants: "dict[str, DelegateContract] | None" = None) -> None:
         self.delegate_contract = contract
-        self.name = contract.identity()
+        # A delegate that binds a DIFFERENT symbol per dtype has more than one
+        # identity, and `callee` is identity. The shipped NVIDIA GEMM is the
+        # first real case: one candidate reaching
+        # `tessera_nvidia_mma_gemm_f16` or `..._bf16` by dtype. Declaring one
+        # of those callees and sometimes calling the other is exactly the
+        # Python-vs-IR drift this contract exists to stop, so the family is
+        # declared instead.
+        #
+        # Callee identity is the WHOLE justification. A per-dtype *tolerance*
+        # looked like a second one -- bf16 carries 8 significand bits to f16's
+        # 11 -- and measurement on sm_120 refuted it: the two agree to within
+        # 25% at every K from 32 to 8192. The reason is that the oracle's
+        # reference rounds its operands to the storage dtype first, so input
+        # rounding cancels on both sides and the residual is f32
+        # accumulation-order error, which does not depend on the storage
+        # dtype. The hook is still per-dtype because a delegate family whose
+        # members really do differ numerically can say so; this one does not.
+        self.contract_variants = dict(variants or {})
+        for dtype, variant in self.contract_variants.items():
+            differing = [f for f in self._FAMILY_INVARIANT
+                         if getattr(variant, f) != getattr(contract, f)]
+            if differing:
+                raise DelegateContractError(
+                    f"delegate variant {dtype!r} disagrees with the family on "
+                    f"{', '.join(differing)}; those fields describe the "
+                    "delegate, not one dtype route, so a representative "
+                    "contract carrying different values would misdescribe it"
+                )
+        # `name` is a dispatch/cache key, NOT a claim -- so unlike tier and
+        # budget it is the registrant's to choose. Deriving it from `callee`
+        # is a good default and a bad requirement: the autotune corpus and the
+        # E3 `force` escape hatch key on this string, so binding it to a C
+        # symbol means renaming that symbol silently invalidates every
+        # persisted verdict and breaks `force` with no error. First use found
+        # this -- the shipped NVIDIA GEMM already had a stable name predating
+        # its contract.
+        self.name = name or contract.identity()
         self.target = target
         self.op = op
-        # Derived, never assigned by the subclass.
+        # Derived, never assigned by the subclass. These ARE claims: a delegate
+        # must not be able to assert a tier or a budget in Python that it did
+        # not declare to the verifier.
         self.tier = contract.arbiter_tier()
         self.accuracy_atol = contract.arbiter_accuracy_atol()
         self.accuracy_rtol = contract.arbiter_accuracy_rtol()
@@ -364,6 +410,26 @@ def applies_to(self, region: Any) -> bool:
             op_list_field=self._OP_LIST_FIELD,
         )
 
+    def contract_for(self, region: Any) -> DelegateContract:
+        """The contract governing this delegate for `region`.
+
+        Falls back to the representative when the region names no dtype or the
+        family declares no variant for it -- the representative is a real
+        declared contract, so the fallback still carries a bound rather than
+        defaulting to none (Decision #21a).
+        """
+        dtype = getattr(region, "dtype", None)
+        if isinstance(dtype, str):
+            variant = self.contract_variants.get(dtype)
+            if variant is not None:
+                return variant
+        return self.delegate_contract
+
+    def accuracy_budget(self, region: Any) -> "tuple[float | None, float | None]":
+        """`(atol, rtol)` for `region`, from that region's declared contract."""
+        c = self.contract_for(region)
+        return c.arbiter_accuracy_atol(), c.arbiter_accuracy_rtol()
+
     def render_target_ir(self, operands: str = "",
                          signature: str = "() -> ()") -> str:
         """The Target IR op declaring this candidate's delegation."""

python/tessera/compiler/emit/nvidia_cuda.py

@@ -55,6 +55,10 @@
     Tier,
     register_candidate,
 )
+from tessera.compiler.emit.delegate_contract import (
+    DelegateContract,
+    DelegatedCandidate,
+)
 from tessera.compiler.emit.kernel_cache import build, register_compiler
 from tessera.compiler.emit.kernel_emitter import (
     EmitError,
@@ -125,6 +129,18 @@
 _GEMM_BF16_ATOL = 5e-2
 _GEMM_DTYPES = ("bfloat16", "float16")
 
+#: The C symbol the shipped GEMM delegate binds per storage dtype — `callee` in
+#: its Target IR contract. Declared here rather than imported so the contract
+#: can be built without loading `tessera.runtime`, and drift-gated against
+#: `runtime._NVIDIA_GEMM_SYMBOLS` by `test_nvidia_delegate_contract.py`: a
+#: rename on one side and not the other would make the delegate declare a
+#: callee it does not call, which is the exact drift the contract exists to
+#: catch.
+_SHIPPED_GEMM_CALLEES = {
+    "float16": "tessera_nvidia_mma_gemm_f16",
+    "bfloat16": "tessera_nvidia_mma_gemm_bf16",
+}
+
 
 # ── CUDA source synthesis (generic FusedRegion lane) ──────────────────────────
 
@@ -5132,19 +5148,90 @@ def _aligned_2d(A: Any, B: Any) -> bool:
     return M % 16 == 0 and N % 8 == 0 and K % 16 == 0
 
 
-class NvidiaMmaGemmShippedCandidate(Candidate):
+#: Measured on the fleet's sm_120 (RTX 5070) against the oracle's
+#: dtype-rounded reference, M=N=256 with the probe's 0.4 operand scaling:
+#:
+#:     K      f16 max|err|   f16 rel      bf16 max|err|  bf16 rel
+#:     32     9.54e-07       2.44e-07     4.77e-07       1.22e-07
+#:     1024   5.15e-05       2.47e-06     3.62e-05       1.74e-06
+#:     8192   8.20e-04       1.23e-05     6.68e-04       9.98e-06
+#:
+#: The absolute error grows about K^1.2 while the relative error grows near
+#: sqrt(K), so a fixed `tolerance` is the wrong shape for this claim: 5e-3 has
+#: 6x headroom at K=8192 and would be breached somewhere past K~65536, on a
+#: kernel that is not wrong. The relative bound is the one that holds, and
+#: `tolerance_rel` is declared with ~8x headroom over the largest measured.
+#: Both are declared because the oracle combines them as
+#: `|a-b| <= atol + rtol*|ref|`, so the relative term is what carries large K.
+_SHIPPED_GEMM_TOLERANCE_REL = 1e-4
+
+#: The delegate is not sm_120-only. `tessera_nvidia_gemm.cpp` NVRTC-compiles
+#: `--gpu-architecture=compute_%d%d` for the LIVE device, and the kernel needs
+#: only `mma.sync.aligned.m16n8k16` (f16/bf16), which is sm_80 and later. The
+#: sm_120 cubin is an AOT fast path, not the envelope. `mma_arch` below stays
+#: "sm_120" for a different reason: it keys the analytical footprint model,
+#: whose `_STATIC_ISAS` currently holds exactly one NVIDIA record.
+_SHIPPED_GEMM_ARCH = "sm_80+"
+
+
+def _shipped_gemm_contract(callee: str) -> DelegateContract:
+    """The contract for one dtype route of the shipped GEMM.
+
+    `determinism` is a proof obligation, not a preference, so it is grounded
+    rather than assumed: the kernel assigns one 16x8 output tile to one warp
+    and reduces K serially into four accumulator registers
+    (`aot/tessera_nvidia_mma_f16_sm120_v1.cu`). No atomics, no split-K, no
+    cross-block reduction -- so the result is reproducible run to run and the
+    delegate may be selected inside `@jit(deterministic=True)`.
+    """
+    return DelegateContract(
+        callee=callee,
+        binding="c_abi",
+        provenance="handwritten_kernel",
+        arch=_SHIPPED_GEMM_ARCH,
+        accuracy="tolerance_bounded",
+        tolerance=_GEMM_F16_ATOL,
+        tolerance_rel=_SHIPPED_GEMM_TOLERANCE_REL,
+        determinism="deterministic",
+        # The kernel is a bare GEMM: offered a matmul+epilogue region it would
+        # implement only the root. It is registered for OP_MATMUL and checks
+        # `isinstance(region, MatmulRegion)`, so declaring this costs it
+        # nothing here -- but the contract describes the KERNEL, and claiming
+        # `whole_region` would assert an epilogue-absorbing ability it does
+        # not have.
+        covers="root_only",
+    )
+
+
+class NvidiaMmaGemmShippedCandidate(DelegatedCandidate):
     """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."""
+    (unaligned OK) bf16/f16 matmul; declines off an NVIDIA GPU.
+
+    **The first declared delegate.** Its tier and accuracy budget are derived
+    from the `DelegateContract` above rather than hand-set, so it cannot claim
+    in Python a budget it did not declare to the Target IR verifier. It binds a
+    different C symbol per dtype, which is why it declares a contract *family*:
+    `callee` is identity, and one candidate silently reaching two symbols under
+    one declared callee is the drift the contract exists to prevent.
+    """
 
-    name = "nvidia_mma_gemm_shipped"
-    tier = Tier.HAND_TUNED
-    target = _TARGET
-    op = OP_MATMUL
-    accuracy_atol = _GEMM_F16_ATOL
     mma_target = "nvidia"
+    #: Footprint-model ISA key, NOT an architecture claim — see
+    #: `_SHIPPED_GEMM_ARCH`. The contract carries the real envelope.
     mma_arch = "sm_120"
 
+    def __init__(self) -> None:
+        variants = {dtype: _shipped_gemm_contract(_SHIPPED_GEMM_CALLEES[dtype])
+                    for dtype in _GEMM_DTYPES}
+        super().__init__(
+            variants["float16"], target=_TARGET, op=OP_MATMUL,
+            # Kept rather than derived from `callee`: the autotune corpus and
+            # the E3 `force` hatch key on this string and predate the contract.
+            name="nvidia_mma_gemm_shipped",
+            variants=variants,
+        )
+
     def available(self) -> bool:
         try:
             from tessera import runtime as rt
@@ -5166,6 +5253,26 @@ def run(self, region: Any, A: Any, B: Any, *a: Any, **k: Any) -> tuple[Any, str]
         except Exception:
             return region.reference(A, B), "reference"
 
+    def measure_device_latency(self, region: Any, *inputs: Any, reps: int = 100,
+                               warmup: int = 10) -> float | None:
+        """CUDA-event kernel time, operands resident.
+
+        Without this the delegate could only be compared end-to-end, and
+        end-to-end is host-dominated: the compiled Tile lane measures 2.99 ms
+        on device inside 34.0 ms of wall time at 2048^3 on sm_120. A Tier-3
+        delegate that cannot be measured on device can never be displaced by a
+        faster compiled kernel, which is the whole of Decision #28.
+        """
+        if len(inputs) != 2:
+            return None
+        try:
+            from tessera import runtime as rt
+            An, Bn = region._natural(inputs[0], inputs[1], cast=False)
+            return rt._nvidia_mma_gemm_device_latency(
+                An, Bn, region.dtype, reps=reps, warmup=warmup)
+        except Exception:
+            return None
+
 
 class NvidiaMmaGemmEmittedCandidate(Candidate):
     """Tier-2 (emitted): the compiler-EMITTED ``ptx_emit`` mma.sync GEMM driven
@@ -5205,6 +5312,24 @@ def run(self, region: Any, A: Any, B: Any, *a: Any, **k: Any) -> tuple[Any, str]
         except Exception:
             return region.reference(A, B), "reference"
 
+    # No `measure_device_latency`: this lane cannot use the launch bridge's
+    # benchmark harness, and the reason is a real divergence rather than a
+    # missing table entry.
+    #
+    # `benchmarkTileGemm16` launches `gx = ceil(N/tileN)`, `gy = ceil(M/tileM)`
+    # -- x maps to N. The NVIDIA Tile lowering agrees (`NVIDIALowering.cpp`:
+    # `mt = blockY*16`, `nt = blockX*8`), which is why the two Tile candidates
+    # time correctly. The `ptx_emit` kernel uses the OPPOSITE convention
+    # (`ptx_emit.py`: `mt = ctaid.x*16`, `nt = ctaid.y*8`), as does the shipped
+    # AOT kernel, so the harness would launch it transposed: at 512x512 it
+    # would cover rows to 1024 and columns only to 256, leaving half the output
+    # unwritten while reporting a plausible latency. Registering swapped tile
+    # dims does not fix it either -- that happens to line up only when M == N.
+    #
+    # Decision #28's displacement test is still satisfiable: the delegate is
+    # measured against the two compiled Tile candidates. Unifying the grid
+    # convention is tracked in the NVIDIA backend queue.
+
 
 class NvidiaTileMatmulCandidate(Candidate):
     """Compiler-generated Tile→fragment GEMM schedule.