gstoner · gstoner · Jul 2, 2026 · Jul 2, 2026 · Jul 2, 2026 · Jul 2, 2026
diff --git a/CLAUDE.md b/CLAUDE.md
@@ -22,11 +22,16 @@ heuristics. "Standalone" means runtime-independent of PyTorch / JAX / Flax
 Target hardware: NVIDIA (SM90 Hopper, SM100 Blackwell), AMD ROCm,
 x86 AMX/AVX512, Apple M-series CPU/GPU.
 
-**Execution reality:** the **x86 AMX/AVX512** backend and **Apple CPU
-(Accelerate) + GPU (MPS/MSL/MPSGraph)** backends execute natively today. NVIDIA
-and ROCm are toolchain-pinned Target IR + lit fixtures with native execution
-**hardware-gated** (Phase G/H). Other backends produce IR/artifacts until a
-hardware-gated proof row says otherwise. See
+**Execution reality (reconciled 2026-06-24 hardware bring-up):** the **x86
+AMX/AVX512** backend and **Apple CPU (Accelerate) + GPU (MPS/MSL/MPSGraph)**
+backends execute natively. Non-Apple hardware is **no longer purely gated**:
+**ROCm gfx1151** (Strix Halo, RDNA 3.5) executes a compiler-generated matmul +
+flash-attention family via `runtime.launch()`, and **NVIDIA sm_120** (RTX 5070 Ti,
+consumer Blackwell) has a hardware-verified `mma.sync` matmul. Broader op coverage
+and datacenter archs (ROCm CDNA/MI300; NVIDIA Hopper sm_90 / datacenter sm_100)
+stay hardware-gated (Phase G/H). Everything else produces IR/artifacts until a
+hardware-gated proof row says otherwise — read the generated dashboards for what
+is actually proven, never counts copied into prose (Decision #26). See
 [`docs/audit/backend/BACKEND_AUDIT.md`](docs/audit/backend/BACKEND_AUDIT.md).
 
 ---
@@ -190,6 +195,8 @@ Per-phase deliverables and the open-work priority queue live in
 
 27. **Ground every Metal / Apple GPU API claim in a real source before declaring it possible or "blocked."** Authoritative sources, in reliability order: **(1) on-machine SDK headers** — `xcrun --show-sdk-path` → `…/System/Library/Frameworks/{Metal,MetalPerformanceShaders,MetalPerformanceShadersGraph,MetalPerformancePrimitives}.framework/Headers/`; **(2) user-provided doc dumps**; **(3) the `apple-metal-docs-urls` memory file**. **WebFetch caveat:** developer.apple.com is a JS-rendered SPA — `WebFetch` returns only the page title, not the API body — so it is NOT a reliable Metal-doc source. Anti-pattern: writing a "blocked / no API path" conclusion from absence of evidence in one source.
 
+28. **The forward compiler direction is the three-tier / measured-arbiter model — leads set the ceiling, the generic framework raises the floor.** (North star, 2026-07-02.) Kernels come from three tiers: **(1)** a generic synthesizer (arch-agnostic region-IR + F4 oracle + synth→compile→cache→launch loop), **(2)** a per-arch codegen plugin (`KernelEmitter`/`TargetPlugin`: MSL / PTX / AMDGCN / C-LLVM), **(3)** hand-tuned kernels. A **measured, accuracy-budgeted arbiter** picks the fastest *in-budget* candidate per `(op, shape-bucket, dtype, target)`. **ROCm and CUDA are the lead performance targets: shared infra must never cap their ceiling** — hand-emitted `wgmma`/`mma.sync`/MFMA/WMMA stay first-class arbiter candidates, displaced only when a compiled kernel is both faster and in accuracy budget. The synthesizer/plugin interface is **symbolic-dim-aware from day one** (`static | bucket | dynamic` policy; first impls bucket-specialize) so dynamic shapes never force an API break. Full model: [`docs/audit/compiler/COMPILER_THEORY_OF_OPERATION.md`](docs/audit/compiler/COMPILER_THEORY_OF_OPERATION.md); execution: [`docs/audit/compiler/COMPILER_REFACTOR_PLAN.md`](docs/audit/compiler/COMPILER_REFACTOR_PLAN.md). These are *direction*; MASTER_AUDIT + generated dashboards stay status truth.
+
 ---
 
 ## Key Design Contracts
@@ -317,6 +324,7 @@ python3 benchmarks/run_all.py --backends x86 --output tessera_benchmarks.json
 | What you need | Where |
 |---------------|-------|
 | **START HERE — status + open-work queue** | `docs/audit/MASTER_AUDIT.md` (+ theme audits; `docs/audit/README.md` for the map) |
+| **Forward compiler direction (north star)** — three-tier/arbiter model + coordination (Decision #28) | `docs/audit/compiler/COMPILER_THEORY_OF_OPERATION.md` (read first) + `COMPILER_REFACTOR_PLAN.md` + reassessed `OPTIMIZING_COMPILER_PLAN.md` |
 | **Generated dashboards** (count/status truth — never hand-edit) | `docs/audit/generated/` |
 | Authoritative API naming | `docs/CANONICAL_API.md` |
 | Canonical tensor attributes & dtypes | `docs/reference/tessera_tensor_attributes.md` |

diff --git a/README.md b/README.md
@@ -256,6 +256,31 @@ kernels → one synthesized symbol set, `synth_matmul_epilogue{,_tiled,_f16}`),
 retirement proven bit-close to the kernel it replaced on Metal. The catalog
 *shrinks* as the general path absorbs it.
 
+### Where the compiler is going (north star)
+
+Apple proved the middle-end; the forward direction generalizes it across all four
+backends. The go-forward plan is a paired plan + theory set under
+[`docs/audit/compiler/`](docs/audit/compiler/) (direction, not status — status
+stays in the generated dashboards):
+
+- [`COMPILER_THEORY_OF_OPERATION.md`](docs/audit/compiler/COMPILER_THEORY_OF_OPERATION.md)
+  — the **three-tier kernel model** (generic synthesizer / per-arch codegen plugin
+  / hand-tuned library) with a **measured, accuracy-budgeted arbiter** that picks
+  the fastest in-budget candidate per `(op, shape-bucket, dtype, target)`.
+- [`COMPILER_REFACTOR_PLAN.md`](docs/audit/compiler/COMPILER_REFACTOR_PLAN.md)
+  — workstreams to share the lowering spine, lift the synthesizer behind a
+  per-arch `KernelEmitter` plugin (MSL / PTX / AMDGCN / C-LLVM), and coordinate
+  development across the three build systems (Apple dev Mac, Strix Halo / ROCm,
+  NR2 Pro / CUDA).
+- [`OPTIMIZING_COMPILER_PLAN.md`](docs/audit/compiler/OPTIMIZING_COMPILER_PLAN.md)
+  — the F0–F6 middle-end plan (F0–F5 landed on Apple; F6 = the backend-build seam).
+
+Governing rule: **ROCm and CUDA are the lead performance targets; the generic
+framework raises the floor and must never cap their ceiling.** Hand-tuned
+`wgmma` / `mma.sync` / MFMA / WMMA kernels stay first-class candidates the arbiter
+measures — a compiled kernel wins only when it is both faster and in accuracy
+budget.
+
 ---
 
 ## Mathematical IR Surfaces

diff --git a/docs/CANONICAL_API.md b/docs/CANONICAL_API.md
@@ -1,12 +1,12 @@
 ---
 status: Normative
 classification: Normative
-last_updated: 2026-06-16
+last_updated: 2026-07-02
 ---
 
 # Tessera Canonical API Quick Reference
 **Status:** Normative - grounded in `python/tessera/` active implementation  
-**Last updated:** June 16, 2026  
+**Last updated:** July 2, 2026  
 **Use this document** to resolve any naming disagreement in other docs. If something here conflicts with another doc, this file wins.
 
 > **Tensor attribute vocabulary lives in [`docs/reference/tessera_tensor_attributes.md`](reference/tessera_tensor_attributes.md).** That document is the canonical reference for the six tensor attributes (`shape`, `dtype`, `layout`, `device`/`target`, `distribution`, `numeric_policy`), the canonical dtype names + accepted aliases, the planned/gated dtype set (`uint*`, `complex*`, packed `int4`, `mxfp*`), and the JAX-like promotion direction. Every dtype string used in this file (`"fp32"`, `"bf16"`, `"fp16"`, `"fp8_e4m3"`, etc.) follows that vocabulary; aliases like `"f32"` are accepted at API boundaries but should normalize before storage.
@@ -235,12 +235,25 @@ def flash_attn_fwd(Q, K, V):
 | `ISA.SM_86` | RTX 30xx | ❌ | ❌ |
 | `ISA.SM_89` | RTX 40xx | ❌ | ❌ |
 | `ISA.SM_90` | H100 / GH200 | ✅ | ✅ |
-| `ISA.SM_100` | B100 / GB200 | ✅ | ✅ |
-| `ISA.SM_120` | Rubin placeholder | ✅ | ✅ |
+| `ISA.SM_100` | B100 / GB200 (Blackwell datacenter) | ✅ | ✅ |
+| `ISA.SM_120` | RTX 50-series — Blackwell consumer (GB20x, CC 12.0) | ❌¹ | ✅ |
 
 `GPUTargetProfile` key parameters: `isa`, `warps_per_cta` (default 4, must be power of 2), `shared_mem_bytes` (None = SM default), `pipeline_stages` (default 2).
 
-Key properties: `.supports_wgmma` / `.supports_tma` / `.supports_mbarrier` → `isa >= SM_90`; `.supports_tcgen05` / `.supports_tmem` / `.supports_cta_pairs` / `.supports_block_scaled_mma` → `isa >= SM_100`; `.runtime_arch` emits CUDA architecture strings such as `sm_90a`, `sm_100a`, and `sm_120`.
+Key properties (**all sourced from the per-SM CUDA 13.3 feature matrix in
+`gpu_target.py`, not a coarse `isa`-ordering**): `.supports_wgmma` / `.supports_tma`
+/ `.supports_mbarrier` (SM_90 / SM_100; sm_120 has `tma`+`mbarrier` but **not**
+`wgmma`); `.supports_tcgen05` / `.supports_tmem` / `.supports_cta_pairs` (datacenter
+SM_100 only — **not** consumer sm_120); `.supports_block_scaled_mma` (SM_100 via
+`tcgen05`, sm_120 via `mma.sync.block_scale`); `.runtime_arch` emits CUDA
+architecture strings such as `sm_90a`, `sm_100a`, and `sm_120`.
+
+> **¹ Consumer Blackwell sm_120 does not implement Hopper `wgmma`** (so
+> `.supports_wgmma` is `False`). It is **not** a superset of datacenter sm_100:
+> no `wgmma`, no `tcgen05`/TMEM. Its low-precision matrix path is 5th-gen Tensor
+> Core `mma.sync` (incl. FP4 `block_scale`, compile target `sm_120a`). The
+> `.supports_*` queries read the authoritative per-arch feature matrix, so they
+> return the hardware-correct answer per ISA.
 
 ### String target aliases
 

diff --git a/docs/README.md b/docs/README.md
@@ -2,7 +2,7 @@
 status: Normative
 classification: Normative
 authority: Documentation authority tree
-last_updated: 2026-06-16
+last_updated: 2026-07-02
 ---
 
 # Tessera Documentation Map
@@ -58,6 +58,14 @@ Compiler planning readers should use
 source-base review, next compiler milestones, and the Visual Complex Analysis
 assessment. Older audit files remain useful historical context.
 
+The **forward compiler direction (north star)** is the paired plan + theory set:
+`docs/audit/compiler/COMPILER_THEORY_OF_OPERATION.md` (read first — the three-tier
+kernel model + accuracy-budgeted measured arbiter + the three-system fleet),
+`docs/audit/compiler/COMPILER_REFACTOR_PLAN.md` (workstreams + coordination), and
+the reassessed `docs/audit/compiler/OPTIMIZING_COMPILER_PLAN.md` (F6 = the
+backend-build seam). These are *direction*; `docs/audit/MASTER_AUDIT.md` and the
+generated dashboards stay status truth.
+
 Reliability and validation readers should start with
 `docs/guides/Tessera_QA_Reliability_Guide.md`. It is the hands-on guide for
 correctness, numerical stability, determinism, expected failures, performance
@@ -136,4 +144,4 @@ Use these labels consistently in active docs:
 | Runtime C ABI and Python wrapper | mock-runtime / hardware-runtime where C backend is built and device-present | `docs/spec/RUNTIME_ABI_SPEC.md`, `python/tessera/runtime.py`, `src/runtime/` |
 | Apple CPU backend | implemented / lit-testable / hardware-runtime via Accelerate + BNNS | `python/tessera/compiler/target_ir.py`, `src/compiler/codegen/Tessera_Apple_Backend/`, `python/tessera/runtime.py`, Apple target-contract tests |
 | Apple GPU backend | implemented / lit-testable / hardware-runtime on Darwin via MPS, MPSGraph, custom MSL, Metal 4 lanes, and packaged `.mtlpackage` ABI validation | `python/tessera/compiler/target_ir.py`, `python/tessera/apple_mlpkg.py`, `python/tessera/compiler/apple_packaged_manifest.py`, `src/compiler/codegen/Tessera_Apple_Backend/`, Apple target-contract tests |
-| ROCm backend | implemented / lit-testable / artifact-only until Phase H hardware execution | `python/tessera/compiler/target_ir.py`, `src/compiler/codegen/Tessera_ROCM_Backend/`, ROCm target-contract tests |
+| ROCm backend | implemented / lit-testable / hardware-runtime on gfx1151 (Strix Halo, RDNA 3.5) via compiler-generated HIP `runtime.launch()` lanes; CDNA/MI300-class remains hardware-gated | `python/tessera/compiler/target_ir.py`, `src/compiler/codegen/Tessera_ROCM_Backend/`, `docs/audit/backend/rocm/ROCM_AUDIT.md`, ROCm target-contract tests |
diff --git a/docs/audit/MASTER_AUDIT.md b/docs/audit/MASTER_AUDIT.md
@@ -32,6 +32,20 @@ work, and remaining work across the compiler, runtime/backend, platform
 backends, coverage, and domain tracks. Generated dashboards remain the source of
 truth for counts; theme audit documents carry the reasoning and work plan.
 
+> **North star for forward development (2026-07-02).** The go-forward compiler
+> direction is the paired plan + theory set under `compiler/`:
+> [`COMPILER_THEORY_OF_OPERATION.md`](compiler/COMPILER_THEORY_OF_OPERATION.md)
+> (read first — the three-tier kernel model, the **accuracy-budgeted measured
+> arbiter**, the Mac/Strix-Halo/NR2-Pro fleet, and the W1–W8 world-class scope
+> register), [`COMPILER_REFACTOR_PLAN.md`](compiler/COMPILER_REFACTOR_PLAN.md)
+> (Workstreams A–E spine + F–K world-class), and the reassessed
+> [`OPTIMIZING_COMPILER_PLAN.md`](compiler/OPTIMIZING_COMPILER_PLAN.md) (F6 = the
+> backend-build seam). Governing rule: **ROCm/CUDA are the lead performance
+> targets; the generic framework raises the floor and must never cap their
+> ceiling.** These docs are *direction*; this page + the generated dashboards
+> stay *status truth* (Decision #26). Full map in
+> [`README.md`](README.md#forward-plans--the-compiler-north-star).
+
 ## Current Truth Snapshot
 
 | Area | Current state | Still open |

diff --git a/docs/audit/README.md b/docs/audit/README.md
@@ -79,6 +79,26 @@ and gates and are excluded from this contract.
 | [generated/compiler_progress.md](generated/compiler_progress.md) | Generated all-up compiler-progress dashboard: phase/IR state, primitives, integration, backend/codegen pathways, and open work. |
 | [generated/](generated/) | Script/test-owned generated dashboards. |
 
+## Forward Plans — the compiler north star
+
+The go-forward direction for compiler development (the **new north star**) is a
+paired plan + theory set under `compiler/`. Read these before starting backend or
+middle-end work; they supersede the pre-2026-07 "op-library" framing.
+
+| Doc | Role |
+|---|---|
+| [compiler/COMPILER_THEORY_OF_OPERATION.md](compiler/COMPILER_THEORY_OF_OPERATION.md) | **Read first.** Durable conceptual model — three-tier kernel model (generic framework / per-arch plugin / hand-tuned library), the accuracy-budgeted measured arbiter, the three-system fleet, and the W1–W8 world-class scope register. |
+| [compiler/COMPILER_REFACTOR_PLAN.md](compiler/COMPILER_REFACTOR_PLAN.md) | Execution plan — Workstreams A–E (kernel spine) + F–K (world-class dimensions), sequencing, and the Mac/Strix-Halo/NR2-Pro coordination + routing matrix. |
+| [compiler/OPTIMIZING_COMPILER_PLAN.md](compiler/OPTIMIZING_COMPILER_PLAN.md) | Middle-end synthesis (F0–F5 landed on Apple); **F6 = the backend-build seam** (reassessed 2026-07-02). |
+| [compiler/EVALUATOR_PLAN.md](compiler/EVALUATOR_PLAN.md) | The scoring engine that gates every promotion in the plans above. |
+| [compiler/STAGE_A_EMIT_PLAN.md](compiler/STAGE_A_EMIT_PLAN.md) | Cross-vendor emit-ladder grounding. |
+
+Governing rule across all of them: **ROCm/CUDA are the lead performance targets;
+the generic framework raises the floor and must never cap their ceiling** (Theory
+§1). Status of what has landed vs. open still flows through
+[MASTER_AUDIT.md](MASTER_AUDIT.md) and the generated dashboards — the plans are
+*direction*, not status truth.
+
 ## Themes
 
 | Theme | Start here | Scope |