Section E: E3 escape-hatch test + sm_120 kernel inventory + NVIDIA_AUDIT reconcile - #298
Merged
Merged
Conversation
The generic NVIDIA CUDA lane now covers the two remaining region kinds that declined to the reference (attention landed in C4), so all four fusion_core region kinds execute on-device: * _synthesize_gated_cuda(): the SwiGLU gate O = gate_act(A@Wg) ⊙ (A@Wu), one thread per output row sharing the A load across the two K-contractions; the gate activation reuses the CUDA-clean _fused_scalar_body.pointwise_snippet. * _synthesize_pointwise_cuda(): a same-shape pointwise DAG, one thread per element, emitted straight from the POINTWISE_OPS C-expression table in topo order (+ a `sign` device shim, the one op with no CUDA builtin). One kernel per region (DAG + input count baked in); the runner binds argtypes to the input count. * run_gated_matmul_region / run_pointwise_graph build + launch (were declining); NvidiaGatedCandidate (op=gated_matmul) + NvidiaPointwiseCandidate (op=pointwise) registered so the arbiter serves them. Live-proven on sm_120 (RTX 5070 Ti): gated matches across silu/gelu/relu, the pointwise DAG relu(a+b)*sigmoid(c) matches, both pass the universal F4 oracle. Still open (C5 tail): dtypes beyond f32 for these lanes — the generic lane is the f32 correctness-first floor; 16-bit storage is served by the B1 matmul lane, and bf16/f16 fused/attention/gated kernels are a follow-on. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…eview) Two P2s in the C5 pointwise lane: * gelu's POINTWISE_OPS template contains an unqualified clamp(...) (an MSL builtin), which CUDA doesn't define — so any DAG with gelu (e.g. add->gelu) failed to compile and silently fell back. Add a clamp device shim. * the sign shim returned 0 for NaN, but POINTWISE_OPS["sign"] == np.sign preserves NaN — so a sign chain on NaN data verified on finite F4 probes yet diverged on real NaN inputs. Make sign (and clamp) NaN-preserving (isnan(x) ? x : ...), matching np.sign / np.clip. Live-proven on sm_120: a gelu DAG now compiles + runs (~6e-8), and a sign chain on NaN data preserves NaN in place; + a host-free emit assertion.
…DIT reconcile E3 (escape hatch, tests/unit/test_nvidia_e3_escape_hatch.py): assert the arbiter never orphans a hand-tuned kernel — the crown-jewel Tier-3 lane wins by default (tier priority) and is forceable; a lower tier is forceable too (tier priority is overridable, the D2/E3 seam); forcing an unknown candidate raises ArbiterError. Host-free with fake candidates + live on the NVIDIA shipped mma.sync GEMM. sm_120 kernel inventory (docs/nvidia_sm120_mma_sync_kernel_inventory.md): a hardware-verified companion to nvidia_cuda13_kernel_inventory.md enumerating the three kernel families — compiler-emitted CUDA (fused/attn/gated/pointwise), emit-path PTX (mma.sync GEMM general/single + NVFP4 assemble-only + wgmma skeleton), shipped C-ABI GEMM — with entry symbols, shapes, dtypes, execution status, and the D1/D2/D3 arbiter tier mapping. NVIDIA_AUDIT reconcile: the "compiler-generated NVIDIA lane" Still-Open item is now the headline LANDED work (#290-#297); flash_attn proven (C4); NVFP4 emit+assemble landed (#291). Next Work rewritten to the real remainder (NVFP4 numerics, tensor-core perf lanes + fleet autotune corpus, dtypes, wgmma/sm_100). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
gstoner
enabled auto-merge (squash)
July 7, 2026 20:54
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 7d0caab470
ℹ️ 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".
gstoner
added a commit
that referenced
this pull request
Jul 7, 2026
) * docs: reconcile status surfaces with the merged C/D/E backend wave (PRs #286–#302) The plan doc carried "landed" annotations for Workstreams C/D/E, but the all-up status surfaces had drifted behind. Refresh them (prose only; counts stay in the generated dashboards per Decision #26): - MASTER_AUDIT.md: frontmatter + body date → 2026-07-07; added a 2026-07-07 reconcile banner (A–E spine wave); refreshed the NVIDIA (Phase G) row — no longer "one op × one arch": now a generic CUDA lane over all four fusable region kinds + hand-emitted tensor-core mma.sync GEMM/flash-attn Tier-2 lanes the arbiter selects, via a shipped PTX launch bridge, all F4-gated and hardware-proven; updated the Runtime/backend + ROCm rows to note the generic plugin lanes + arbiter + x86 Zen 5 execution. - backend/nvidia/NVIDIA_AUDIT.md: frontmatter date → 2026-07-07 (body was already reconciled in #298). - COMPILER_REFACTOR_PLAN.md: the stale Phase-0 E3 row (⬜⬜⬜) → landed (E3 escape hatch lives in the D1 `force` path + PR #298 test); added the E3-landed annotation to the bullet. - README.md: refreshed the NVIDIA support-snapshot row; reframed "Where the compiler is going" from future-direction to largely-built (the three-tier framework + arbiter now live in python/tessera/compiler/emit/, NVIDIA/ROCm/x86 executing). - docs/README.md: same north-star reframe + a new NVIDIA sm_120 executing-lane status row. Doc-only. Frontmatter contract 8/8, generated-doc drift gate in sync (19), docs lint passed. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * docs: scope NVIDIA prose to the execution matrix (Codex P2) The README's Current Support Snapshot says generated audits are the authority when prose disagrees, but runtime_execution_matrix.md records only sm_120's shipped `nvidia_mma` GEMM. My reconcile prose claimed the new generic CUDA + tensor-core lanes as executable, which disagrees with the cited dashboard. Root cause: those lanes run through the arbiter/emit subsystem (emit/nvidia_cuda.py + candidate.py + in-process nvcc), a different path from the execution-matrix executor registry — so they're hardware-proven by the plugin/perf/conformance test gates (test_nvidia_plugin.py, test_nvidia_perf_ratchet.py, test_conformance_execute_compare_nvidia.py), not recorded in the matrix, and a regen won't add them. Fix (README, MASTER_AUDIT, docs/README): scope the execution-matrix-authoritative rows to `nvidia_mma`, attribute the arbiter/emit + tensor-core lanes to their real proof surface (the test gates), and add "promote the arbiter/emit lanes into the execution matrix" as an explicit still-open item — which is the genuine gap. Doc-only. Frontmatter 8/8, drift gate in sync (19), docs lint passed. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Section E — the guardrail test + doc reconcile after the B/C/D landings.
E3 — escape-hatch test (
tests/unit/test_nvidia_e3_escape_hatch.py)Asserts the arbiter never orphans a hand-tuned kernel (the lead-safety contract, Decision #28):
ArbiterError(never silently picks another).Host-free with fake candidates + live on the NVIDIA shipped mma.sync GEMM (the Tier-3 matmul lane): forced + wins by default; the emitted Tier-2 lane is forceable too.
sm_120 kernel inventory (
docs/nvidia_sm120_mma_sync_kernel_inventory.md)A hardware-verified companion to
nvidia_cuda13_kernel_inventory.md(which is the SM_90+ planning inventory). Enumerates the three kernel families with entry symbols / shapes / dtypes / execution status + the D1/D2/D3 arbiter tier mapping:tessera_nvidia_{fused,attn,gated,pointwise}(Tier-1);tessera_mma_gemm_{bf16,f16}(Tier-2) + single-tile + NVFP4 (assemble-only) + wgmma (skeleton);tessera_nvidia_mma_gemm_{f16,bf16,tf32}(Tier-3).NVIDIA_AUDIT reconcile
Its "Still Open #1: compiler-generated NVIDIA lane" is now the headline landed work (#290–#297) — the doc predated the merges. Flipped to landed; flash_attn proven (C4); NVFP4 emit+assemble landed (#291). "Next Work" rewritten to the real remainder (NVFP4 numerics, tensor-core perf lanes + fleet autotune corpus, dtypes beyond f32, wgmma/sm_100).
Verified
E3 host-free + live pass on sm_120; ruff clean; drift gate green (docs are hand-authored, not generated).
This completes the Section E items. (Disjoint from the in-flight C5 #297 — different files.)
🤖 Generated with Claude Code