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

[gstoner]

Section B, item 1 — wire the emitted GEMM as a first-class arbiter candidate. The blocker you flagged (no bare-matmul op-kind in the D1 registry) is cleared. Verified on the live sm_120 box.

The blocker

The D1 candidate registry only had fused_region/attention/gated/pointwise op-kinds. A plain D = A @ B isn't a FusedRegion (that always carries ≥1 fused op), so the compiler-emitted mma.sync GEMM lane had nowhere to plug in.

What landed

• fusion_core: MatmulRegion (bare GEMM, 16-bit storage / f32 accumulate) + verify_synthesized_matmul (arbiter-only F4 oracle, dtype-rounded reference) + a _round_to_storage helper; re-exported through the fusion facade.
• candidate.py: OP_MATMUL op-kind + its (verify, run_matmul) map entry + a run_matmul method on the _as_runner adapter.
• runtime.py: two 2D GEMM execution helpers — _nvidia_mma_gemm_2d (shipped libtessera_nvidia_gemm, row-major B) and _nvidia_ptx_gemm_2d (compiler-emitted ptx_emit via the launch bridge, col-major B), keyed by 16-bit dtype; + a bridge loader mirroring the shipped-GEMM one.
• emit/nvidia_cuda.py: NvidiaMmaGemmShippedCandidate (Tier-3 hand-tuned) + NvidiaMmaGemmEmittedCandidate (Tier-2 emitted, aligned-only) under (nvidia, matmul); both F4-gated, f16 accuracy budget.

Behavior

Tier-priority picks the shipped lane by default (lead-safe, Decision #28); the E3 force hatch selects the emitted lane. NVIDIA now has its Tier-3 hand-tuned GEMM candidate (previously only reachable via the jit nvidia_mma executor) alongside the Tier-2 emitted lane.

Verified on sm_120 (RTX 5070 Ti)

Both lanes F4-verify + execute + match the dtype-rounded reference across bf16/f16 × 16×8×16 / 32×16×32 / 64×64×64; the arbiter selects shipped by default and force selects emitted. 152 passed / 135 skipped across the plugin+fusion suite; ruff + mypy clean; drift green.

Next (this section)

• D2: measured autotune loop on sm_120 — replace tier-priority with real on-device latency per shape-bucket (lets the Tier-2 emitted lane win where faster). Now unblocked by this op-kind.
• D3: fallback logging for NVIDIA.

🤖 Generated with Claude Code

[chatgpt-codex-connector]

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 4569a9920a

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

[chatgpt-codex-connector]

Validate B's K dimension before launching GEMM

When run_arbitrated selects the shipped NVIDIA matmul candidate, mismatched rank-2 inputs such as A.shape == (M, K) and B.shape == (K2, N) with K2 != K reach this helper. This line discards K2 and passes K to the native C ABI, which then copies K*N elements from B even if the NumPy buffer only contains K2*N; that can overread host memory or silently compute against the wrong slice instead of raising like MatmulRegion.reference/the existing JIT path does. Check Bc.shape[0] == K before calling the C function.

Useful? React with 👍 / 👎.