Add compiler north-star plan pair + fix sm_120 capability drift - #268
Merged
Merged
Conversation
Compiler direction (docs): - New paired plan + theory: COMPILER_THEORY_OF_OPERATION.md (three-tier kernel model, accuracy-budgeted measured arbiter, three-system fleet, W1-W8 scope register) and COMPILER_REFACTOR_PLAN.md (workstreams A-E spine + F-K world-class, coordination + §9 source-verified seam verdicts). - Reassess OPTIMIZING_COMPILER_PLAN.md: F0-F5 landed on Apple; rewrite F6 (the backend-build seam) since its "Mac can't run CUDA/ROCm" premise is dead (Strix Halo gfx1151 + NR2 Pro sm_120 now execute) and scope the anti-goal. - Dynamic-shapes decision pulled into the spine: symbolic-dim-aware KernelEmitter/TargetPlugin API + shape-bucket arbiter key, bucket-specialize first. - Wire the north star into MASTER_AUDIT.md, docs/audit/README.md, README.md, docs/README.md, and CLAUDE.md (new Decision #28 + reference row). sm_120 capability fix (code): - gpu_target.py: route all coarse capability properties (supports_wgmma / tcgen05 / tmem / cta_pairs / mbarrier / tma / block_scaled_mma) through the authoritative _CUDA_13_3_FEATURES matrix via cuda_feature_status instead of isa >= SM_x. Fixes consumer Blackwell sm_120 wrongly reporting Hopper wgmma + datacenter tcgen05/TMEM/CTA-pairs as supported (it is NOT a superset of sm_100; its matrix path is mma.sync). Also fixes sm_120 wrongly inheriting the SM_90 FA-4 attn default in jit.py. - test_gpu_target.py: bug-encoding test_sm120_runtime_arch becomes a test_sm120_consumer_blackwell_capabilities regression guard; drop stale "rubin_placeholder" naming. Doc drift cleanups: - CLAUDE.md: fix stale "Execution reality" (gfx1151 + sm_120 now execute). - docs/README.md: ROCm row artifact-only -> gfx1151 hardware-runtime. - CANONICAL_API.md: sm_120 "Rubin placeholder" -> Blackwell consumer; correct the WGMMA column + footnote to match the fixed properties. Gates: mypy clean, generated-doc drift in sync, doc lint passed, test_gpu_target + test_audit_docs green. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
![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: 832ff9cef0
ℹ️ 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".
Address PR review: GPUTargetProfile(SM_120) was fixed, but the string-alias capability path (get_target_capability / backend_capabilities) still advertised invalid consumer-Blackwell features. - capabilities.py: drop wgmma / wgmma_sparse / tcgen05 / tcgen05_pair / tmem from nvidia_sm120.features (consumer Blackwell is NOT a superset of datacenter sm_100; FP4 goes through mma.sync.block_scale). Now mirrors the cuda_feature_set(SM_120) "ready" flags. - test_compiler_capabilities.py: add test_nvidia_features_match_cuda_matrix — a single-source-of-truth guard asserting no NVIDIA capability entry advertises a feature the CUDA matrix marks not_supported, plus a positive lock that sm_120 excludes the datacenter/Hopper flags. Prevents this drift from recurring. - Regenerate test_coverage dashboards (deterministic negative_refs count shift from the added guard test). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Fold the AOCL-DLP reference (amd/aocl-dlp — AMD's BLIS-family DL primitives: low-precision GEMM/batch GEMM, pre/post-ops, INT4/FP16, symmetric quant, OpenMP) into the north-star plan: - Theory Tier-3 list: add it to the x86 line (CPU analog of cuBLAS/rocWMMA). - Refactor Plan C1: the x86 TargetPlugin registers AOCL-DLP as a Zen-family Tier-3 candidate — AVX512-based (fits the Zen 5 fleet box, no AMX), fills the x86 OpenMP + INT4/FP16 gaps, opt-in behind a build flag (BLAS-family lib like Accelerate, Decision #23-clean), arbiter-selected only where measured faster; license check before a shipped lane. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
gstoner
added a commit
that referenced
this pull request
Jul 3, 2026
…271) Consumer Blackwell sm_120 (RTX 50-series) was lumped with datacenter sm_100 in the Target-IR lowering and emitted `tcgen05_mma` / `tmem_alloc`. sm_120 is NOT a superset of sm_100: it has no tcgen05/TMEM (datacenter sm_100a only) and no Hopper wgmma. Its matrix path is warp-level `mma.sync` (FP4 via `mma.sync.aligned...block_scale`), with tma + mbarrier available. This is the remaining lowering site of the same sm_120!=sm_100-superset bug fixed for the capability queries in #268. - target_ir.py `_lower_nvidia_op`: split the branch. sm_100 keeps tmem_alloc + tcgen05_mma; new sm_120 branch emits tessera_nvidia.mma_sync (m16n8k16, bf16 -> f32) + tma_async_copy + mbarrier. Grounded in gpu_target._CUDA_13_3_FEATURES[ISA.SM_120] (wgmma/tcgen05/tmem not_supported; tma/mbarrier/block_scaled_mma ready). - target_ir.py `_verify_nvidia_op`: added the mma_sync contract branch (arch, shape, dtype_ab, dtype_c; no warpgroup -- mma.sync is warp-scoped). - matmul_pipeline.py `_render_nvidia_target_ir`: same split (currently an unused renderer, hardened so a future wire-up can't resurrect the bug). - TARGET_IR_SPEC.md: corrected tcgen05/tmem rows to datacenter SM100 only; added a tessera_nvidia.mma_sync row for consumer SM120. - Tests: re-added the nvidia_sm120 golden fixture (regenerated golden shows mma_sync, not tcgen05_mma/tmem_alloc) and added a sm_120 -> mma_sync lowering assertion. All host-free; hardware execution stays gated. Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
gstoner
added a commit
that referenced
this pull request
Jul 3, 2026
…MEM (#273) C++ half of the consumer-Blackwell sm_120 fix — complements #271 (the Python emitter half). NVIDIALowering.cpp lumped sm_120 with datacenter sm_100 (`smVersion >= kBlackwellSM`) and emitted tcgen05_mma/tmem_alloc, but consumer Blackwell (RTX 50-series) has NO tcgen05/TMEM and no Hopper wgmma — its matrix path is warp-level `mma.sync`. Same sm_120≠sm_100-superset bug already fixed in the capability queries (#268) and the Python emitter (#271). - NVIDIALowering.cpp: add kConsumerBlackwellSM=120; route sm_120 tile.mma to a `tessera_nvidia.mma_sync` op (m16n8k16, bf16/f32) checked BEFORE the datacenter tcgen05 branch; add the mma_sync -> llvm.nvvm.mma.sync.contract NVVM mapping; tighten the TMEM guard to reject consumer sm_120 (not just <SM100). Op name + shape match #271 for cross-lane consistency. - New lit fixture sm120_tile_to_nvidia.mlir (mma_sync, not tcgen05/tmem); tmem_requires_blackwell.mlir updated + a sm_120 rejection RUN line. Verified host-free on the dev Mac: built tessera-nvidia-opt with TESSERA_ENABLE_CUDA=OFF (no GPU/ptxas needed); all 7 nvidia lit fixtures pass (sm_120 new, sm_100/sm_90 unchanged, tmem rejects sm_90 + sm_120). 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.
Overview
Two related threads: (1) establish the go-forward compiler direction as a paired plan + theory set (the "north star" for building out backends), and (2) fix a real sm_120 capability-matrix bug surfaced while grounding those docs. Plus doc-drift cleanups from the last two days of hardware bring-up.
Everything here is docs + a self-contained code fix; no runtime behavior changes beyond the corrected capability queries.
What changed
Compiler direction (docs)
docs/audit/compiler/:COMPILER_THEORY_OF_OPERATION.md— the durable model: three-tier kernel model (generic synthesizer / per-arch codegen plugin / hand-tuned library), an accuracy-budgeted measured arbiter keyed on(op, shape-bucket, dtype, target), the three-system fleet (Apple dev Mac / Strix Halo ROCm / NR2 Pro CUDA), and the W1–W8 world-class scope register.COMPILER_REFACTOR_PLAN.md— execution: Workstreams A–E (kernel spine) + F–K (world-class), coordination/routing, and §9 source-verified seam verdicts.OPTIMIZING_COMPILER_PLAN.md— F0–F5 landed on Apple; rewrote F6 (the backend-build seam) because its "the Mac can't run CUDA/ROCm, silicon deferred" premise is dead (gfx1151 + sm_120 now execute), and scoped the "no new hand-written kernel" anti-goal to Apple/default-path.KernelEmitter/TargetPlugininterface is symbolic-dim-aware from day one (static | bucket | dynamicpolicy), the arbiter keys on shape-bucket, first impls bucket-specialize.MASTER_AUDIT.md,docs/audit/README.md,README.md,docs/README.md, andCLAUDE.md(new Decision Apple GPU MLA: same-length B>1 batching in decode_batch (throughput) #28 + a Key Reference Files row).sm_120 capability fix (code)
gpu_target.py: route all coarse capability properties (supports_wgmma/tcgen05/tmem/cta_pairs/mbarrier/tma/block_scaled_mma) through the authoritative_CUDA_13_3_FEATURESmatrix viacuda_feature_status, instead ofisa >= SM_x.wgmmaand datacentertcgen05/TMEM/CTA-pairs as supported — it is not a superset of datacenter sm_100; its low-precision matrix path ismma.sync(incl. FP4block_scale).jit.pywas applying the SM_90 FA-4 attn default to sm_120 (which has no wgmma).test_gpu_target.py: the bug-encodingtest_sm120_runtime_arch(assertedsupports_tcgen05 is True) becomes atest_sm120_consumer_blackwell_capabilitiesregression guard; dropped the stale "rubin_placeholder" test name.Doc-drift cleanups (last-2-days hardware bring-up)
CLAUDE.md: fix stale "Execution reality" — gfx1151 + sm_120 now execute; no longer "purely hardware-gated."docs/README.md: ROCm status rowartifact-only→hardware-runtime on gfx1151.CANONICAL_API.md: sm_120 "Rubin placeholder" → "Blackwell consumer"; corrected the WGMMA column + footnote to match the fixed properties.Reviewer notes
.supports_*values for sm_120 (and semantically-precisembarrier→mbarrier_arrive_tx). SM_80/86/89/90/100 results are unchanged. No generated dashboard consumes these properties (drift gate stays green).MASTER_AUDIT.md+ generated dashboards remain status truth (Decision Apple GPU MLA: paged-cache decoder (production-serving wiring) #26).ptx_emit.pyhas no serialize→ptxas→launch bridge yet (today's sm_120 matmul runs via a shipped.so), so that bridge — not shape breadth — is the long pole of the NVIDIA compiled lane.Gates
mypyclean · generated-doc drift in sync (19) · doc lint passed ·test_gpu_target.py+test_audit_docs.pygreen · 152-test NVIDIA/target sweep passed.🤖 Generated with Claude Code