README: state the compiler direction; correct five stale claims

README.md

@@ -102,7 +102,7 @@ Use these status words consistently:
 | scaffolded | Directory, API shape, or design skeleton exists, but behavior is incomplete or artifact-only. |
 | planned | Design direction only. |
 
-Current status snapshot (reviewed 2026-08-24). Generated dashboards are the
+Current status snapshot (reviewed 2026-08-30). Generated dashboards are the
 source of truth for exact counts and executable lanes:
 [`runtime_execution_matrix.md`](docs/audit/generated/runtime_execution_matrix.md),
 [`runtime_abi.md`](docs/audit/generated/runtime_abi.md), and
@@ -120,8 +120,11 @@ source of truth for exact counts and executable lanes:
 | Mathematical and model IR surfaces | implemented / lit-testable | GA/EBM, reasoning-attention families, DFlash, DiffusionGemma, and frontier MoE model-class contracts are compiler-visible. Native execution is claimed only where a backend row below or a generated audit proves it. |
 | Runtime ABI and audits | implemented | Runtime C ABI surfaces and generated audit dashboards are drift-gated; exact counts are listed in the support snapshot below. |
 
-The fast unit suite passes under `-m "not slow"`; the full Python
-suite collects ~14,400 tests including slow/heavy benchmark contracts.
+The fast unit suite passes under `-m "not slow"`; the full Python suite
+additionally collects the slow/heavy benchmark contracts. Per the note above
+and Decision #26, the collected count is deliberately not quoted — it drifts
+with every landing, and the figure previously written here had drifted by
+several thousand.
 
 ### Current Support Snapshot
 
@@ -183,18 +186,59 @@ Graph IR    (tessera dialect: math ops, shape/dtype/layout metadata, diagnostics
 Schedule IR (schedule.* dialect: mesh.define/region, pipeline.region, stage, yield)
      |
      v
-Tile IR     (tile.* ops, tessera.attn.* FA-4 ops, tessera.queue.* barriers)
+Tile IR     (tile.* ops, tessera.attn.* FA-4 ops)
      |
      v
 Target IR   (backend-specific artifacts: x86, NVIDIA, ROCm, Apple, ...)
 ```
 
+New backends expose a **hardware-free Target IR dialect** before any
+hardware-specific lowering (Decision #19) — Tile IR is never lowered straight
+to PTX/HIP/Metal source.
+
+### Direction: MLIR/LLVM is the architecture; the Python backend path is bootstrap
+
+The original Python→backend compiler was a **bootstrap compiler**. The
+direction is to **prune it and build the core out through MLIR and LLVM IR**,
+so that the four layers above are the compiler rather than a description of
+one. Concretely, that means a program's physical code comes from
+Graph → Schedule → Tile → Target lowering driven by `tessera-opt`, not from
+Python emitting target source directly.
+
+This is the backend half of **E2E-REAL-6 ("one compiler authority")**, whose
+frontend half is already retiring the decoration-time AST extractor. It
+inherits that program's discipline: **duplicate lowering authorities are
+removed only after the surviving path is proven to carry what they carried**
+(Decision #31's ordering caveat), never by deletion first.
+
+Fast paths remain legitimate and expected — inline PTX or GCN assembly, and
+Tessera Standard Library native libraries. The requirement is that they arrive
+through a **standard interface** rather than a silent Python-side branch: a
+declared Target IR op that names the boundary, so Decision #28's arbiter can
+tell compiler-generated work from delegated work and score them against each
+other. `tessera_x86.abi_call` is the shipped precedent for that boundary;
+giving NVIDIA and ROCm the equivalent is the enabling step that has to land
+*before* the bootstrap path is pruned, since today several legitimate fast
+paths live inside it and have nowhere else to go.
+
 The Python compiler carries object models and verifier checks for Graph IR,
 Schedule IR, Tile IR, and CPU/x86, NVIDIA/CUDA, Apple, and ROCm Target IR.
 The JIT artifact spine emits textual MLIR-like inspection strings from those
 objects; native hardware execution remains target-specific and is claimed only
 where backend docs say so.
 
+Read that Python surface in two parts, because they are on opposite sides of
+the direction above. The **object models, verifiers and contract registries**
+are load-bearing and stay. The **per-backend packagers** — the
+`package_*` families in `{nvidia,rocm,x86,apple_cpu}_native.py` that emit
+target artifacts directly from Graph IR — are the bootstrap compiler, and are
+the prune target. They are entered today on a *tool-presence check*: when
+`tessera-opt` is findable the compiled route runs, and when it is not the
+packager does, which means a second lowering authority can be selected by
+whether a binary happens to be installed. Supplying a packager-only key on the
+compiled route now raises `SCHEDULE_KEY_NOT_HONORED_ON_COMPILED_ROUTE` rather
+than being dropped in silence (Decision #21a).
+
 For compiler-readiness audits, keep three lanes separate:
 
 - **Reference / contract lane:** Python reference behavior, public APIs,
@@ -215,9 +259,9 @@ authoritative status source). The canonical lowering pipelines registered in
 
 | Pipeline | Status |
 |------|--------|
-| `tessera-lower-to-x86` | implemented / lit-testable; hardware-runtime via the CPU JIT + native CPU ABI + AVX-512 compiled lanes (the AMX lane emits but is artifact-only — no AMX hardware in the fleet) |
+| `tessera-lower-to-x86` | implemented / lit-testable; hardware-runtime via the CPU JIT + native CPU ABI + AVX-512 compiled lanes. The AMX lane emits and stays artifact-only **by direction, not by hardware availability**: AMX is a retired target, superseded by ACE (AI Compute Extensions, agreed jointly by AMD and Intel), so the AMX ops remain an IR-level contract with no `amx.*` lowering. x86 native execution proof means AVX-512 |
 | `tessera-lower-to-gpu` (NVIDIA SM90 WGMMA/TMA) | implemented / lit-testable (SM90 WGMMA has no hardware-execution proof yet) |
-| `tessera-nvidia-pipeline-{sm90,sm100,sm120}` (per-SM aliases) | implemented / lit-testable; the sm_120 `mma.sync` GEMM additionally **executes** on consumer Blackwell hardware via a separate emit/runtime lane (`ptx_emit.py` + `libtessera_nvidia_gemm.so`), not through this IR pipeline |
+| `tessera-nvidia-pipeline-{sm90,sm100,sm120}` (per-SM aliases) | implemented / lit-testable. Two sm_120 lanes execute on consumer Blackwell, and they are **not** the same thing: the canonical matmul compiles Graph→Schedule→Tile through `tessera-opt` and is packaged to PTX via `tessera-nvidia-opt` + `mlir-opt` (this is the compiled route, and the one the direction above builds out), while the older `ptx_emit.py` + `libtessera_nvidia_gemm.so` GEMM is a separate bootstrap emit/runtime lane that bypasses the IR pipeline. **Requires `-DTESSERA_ENABLE_CUDA=ON` with `-DTESSERA_BUILD_NVIDIA_BACKEND=ON`**; a CUDA-less NVIDIA build produces a lean driver that never registers the NVIDIA Target IR dialect, and every scheduled lane then fails at `--tessera-schedule-to-tile` having touched no GPU |
 | `tessera-lower-to-rocm` | implemented / lit-testable / hardware-runtime on capable gfx1151 (RDNA3.5) hosts via HIP (WMMA matmul + attention family) |
 | `tessera-lower-to-apple_cpu` (artifact) / `tessera-lower-to-apple_cpu-runtime` (Accelerate) | implemented / lit-testable / hardware-runtime |
 | `tessera-lower-to-apple_gpu` (artifact) / `tessera-lower-to-apple_gpu-runtime` (MPS + custom MSL) | implemented / lit-testable / hardware-runtime |
@@ -322,6 +366,14 @@ framework raises the floor and must never cap their ceiling.** Hand-tuned
 measures — a compiled kernel wins only when it is both faster and in accuracy
 budget.
 
+This is what makes the bootstrap prune above safe rather than lossy, and the
+two must be read together. A hand-tuned or library kernel keeps its place as a
+**Tier-3 candidate the arbiter measures**; what it loses is the ability to be
+reached as an unmeasured fallback because a tool was missing. The distinction
+is *chosen* versus *defaulted into* — which is why the enabling work is the
+declared Target IR boundary (`abi_call` and its NVIDIA/ROCm equivalents), not
+the deletion.
+
 ---
 
 ## Mathematical IR Surfaces
@@ -452,7 +504,7 @@ pip install -e ".[dev]"
 # Daily edit-loop sanity check (fast tests only, < 512 MB RAM)
 pytest tests/unit/ -m "not slow" -q
 
-# Full Python suite including heavy benchmarks (~14,400 collected)
+# Full Python suite including heavy benchmarks
 pytest tests/unit/ -q
 
 # GA + EBM native Apple GPU health check; skip-recording on non-Darwin
@@ -490,11 +542,12 @@ cmake -S . -B build \
 
 cmake --build build --parallel
 
-# On Ubuntu 26.04 LTS (x86 + TheRock ROCm 7.14): bootstrap the toolchain once with
+# On Ubuntu 26.04 LTS (x86 + ROCm 10 series, HIP 7.15): bootstrap the toolchain once with
 #   bash scripts/setup_ubuntu.sh           # LLVM/MLIR 23.1 from apt.llvm.org + venv
 #   source .venv/bin/activate
 #   source scripts/_rocm_env.sh
-# then configure against upstream LLVM/MLIR and TheRock at /opt/rocm/core:
+# then configure against upstream LLVM/MLIR and ROCm. Both /opt/rocm and
+# /opt/rocm/core exist on that box and either works as CMAKE_PREFIX_PATH:
 cmake -S . -B build -G Ninja \
   -DLLVM_DIR=/usr/lib/llvm-23/lib/cmake/llvm \
   -DMLIR_DIR=/usr/lib/llvm-23/lib/cmake/mlir \
@@ -527,7 +580,7 @@ NVIDIA / ROCm toolchain checks (skip cleanly when toolchains absent):
 # Validate CUDA 13.3 PTX patterns against installed nvcc
 python scripts/validate_nvcc_compile.py
 
-# Validate ROCm 7.2.4 AMDGCN intrinsics against installed hipcc
+# Validate AMDGCN intrinsics against installed hipcc (ROCm 10 / HIP 7.15)
 python scripts/validate_hipcc_compile.py
 
 # Probe NCCL/RCCL ≥ 2.22 symbols at runtime

docs/audit/generated/bootstrap_prune_gap.csv

@@ -0,0 +1,25 @@
+target,family,compiled_route,status
+nvidia_sm120,attention_backward_lse,—,gap
+nvidia_sm120,attention_lse,—,gap
+nvidia_sm120,attention_backward,scheduled_attention_backward.supports_scheduled_attention_backward,compiled
+nvidia_sm120,paged_kv,—,gap
+nvidia_sm120,attention,scheduled_attention.supports_scheduled_attention,compiled
+nvidia_sm120,softmax,—,gap
+nvidia_sm120,norm,—,gap
+nvidia_sm120,reduction,—,gap
+nvidia_sm120,nvfp4_matmul,—,gap
+nvidia_sm120,int4_matmul,—,gap
+nvidia_sm120,mx_matmul,—,gap
+nvidia_sm120,matmul,scheduled_matmul.supports_scheduled_matmul,compiled
+rocm_gfx1151,softmax,—,gap
+rocm_gfx1151,reduction,—,gap
+rocm_gfx1151,paged_kv,—,gap
+rocm_gfx1151,attention,scheduled_attention.supports_scheduled_attention,compiled
+rocm_gfx1151,moe_dispatch,—,gap
+x86,softmax,—,gap
+x86,reduction,—,gap
+x86,matmul,scheduled_matmul.supports_scheduled_matmul,compiled
+x86,attention,scheduled_attention.supports_scheduled_attention,compiled
+x86,cohort2,—,gap
+x86,breadth,—,gap
+x86,elementwise,—,gap

docs/audit/generated/bootstrap_prune_gap.md

@@ -0,0 +1,108 @@
+# Bootstrap Prune — Mainline Coverage Gap
+
+**Generated. Do not hand-edit.** Regenerate with
+`python -m tessera.compiler.generated_docs --write`.
+
+The Python per-backend `package_*` families are the **bootstrap
+compiler**; the architecture is core MLIR/LLVM (Graph → Schedule →
+Tile → Target via `tessera-opt`). This dashboard answers what must be
+settled before any of it is deleted: **which families does the
+mainline compiler already cover, and which would lose their only
+lowering?** Decision #31's ordering caveat is the rule — a duplicate
+authority is removed only after the survivor is proven to carry what
+it carried.
+
+A `gap` row is *not* a defect. It is scope: work the mainline
+compiler must absorb, or a fast path that must be re-expressed
+through a declared Target IR boundary (Decision #28 Tier 3) before
+the bootstrap row can go.
+
+## Summary
+
+| Metric | Count |
+|---|---|
+| Backends with a bootstrap module | 4 |
+| `package_*` functions total | 49 |
+| — **bootstrap** (re-enter Graph IR; prune target) | 34 |
+| — compiled-route packagers (consume a lowered artifact) | 15 |
+| Lines in those modules | 8738 |
+| Classified families | 24 |
+| — covered by a compiled route | 6 |
+| — **gap (no compiled route)** | 18 |
+| Packagers matching no family | 9 |
+
+## Per-backend bootstrap surface
+
+| Target | Module | bootstrap | compiled-route | Families | Lines |
+|---|---|---|---|---|---|
+| `nvidia_sm120` | `nvidia_native.py` | 19 | 5 | 12 | 3783 |
+| `rocm_gfx1151` | `rocm_native.py` | 7 | 5 | 5 | 2894 |
+| `x86` | `x86_native.py` | 7 | 5 | 7 | 1846 |
+| `apple_cpu` | `apple_cpu_native.py` | 1 | 0 | 0 | 215 |
+
+## Family coverage
+
+`compiled` means a compiled-route admission predicate serves that
+family. It does **not** assert the compiled route reaches parity on
+every shape and dtype — that is per-family evidence the backend
+queues own.
+
+| Target | Family | Compiled route | Status |
+|---|---|---|---|
+| `nvidia_sm120` | `attention_backward_lse` | — | 🔴 **gap** |
+| `nvidia_sm120` | `attention_lse` | — | 🔴 **gap** |
+| `nvidia_sm120` | `attention_backward` | `scheduled_attention_backward.supports_scheduled_attention_backward` | ✅ compiled |
+| `nvidia_sm120` | `paged_kv` | — | 🔴 **gap** |
+| `nvidia_sm120` | `attention` | `scheduled_attention.supports_scheduled_attention` | ✅ compiled |
+| `nvidia_sm120` | `softmax` | — | 🔴 **gap** |
+| `nvidia_sm120` | `norm` | — | 🔴 **gap** |
+| `nvidia_sm120` | `reduction` | — | 🔴 **gap** |
+| `nvidia_sm120` | `nvfp4_matmul` | — | 🔴 **gap** |
+| `nvidia_sm120` | `int4_matmul` | — | 🔴 **gap** |
+| `nvidia_sm120` | `mx_matmul` | — | 🔴 **gap** |
+| `nvidia_sm120` | `matmul` | `scheduled_matmul.supports_scheduled_matmul` | ✅ compiled |
+| `rocm_gfx1151` | `softmax` | — | 🔴 **gap** |
+| `rocm_gfx1151` | `reduction` | — | 🔴 **gap** |
+| `rocm_gfx1151` | `paged_kv` | — | 🔴 **gap** |
+| `rocm_gfx1151` | `attention` | `scheduled_attention.supports_scheduled_attention` | ✅ compiled |
+| `rocm_gfx1151` | `moe_dispatch` | — | 🔴 **gap** |
+| `x86` | `softmax` | — | 🔴 **gap** |
+| `x86` | `reduction` | — | 🔴 **gap** |
+| `x86` | `matmul` | `scheduled_matmul.supports_scheduled_matmul` | ✅ compiled |
+| `x86` | `attention` | `scheduled_attention.supports_scheduled_attention` | ✅ compiled |
+| `x86` | `cohort2` | — | 🔴 **gap** |
+| `x86` | `breadth` | — | 🔴 **gap** |
+| `x86` | `elementwise` | — | 🔴 **gap** |
+
+## Packagers matching no classified family
+
+`package_<family>` is the convention, so these are reached by some
+other entry point — a sibling call site, a dtype specialisation, or
+dead code. Each needs its own disposition; none may be assumed
+covered because a same-named family is compiled.
+
+| Target | Packager |
+|---|---|
+| `nvidia_sm120` | `package_bf16_matmul` |
+| `nvidia_sm120` | `package_bf16_softmax` |
+| `nvidia_sm120` | `package_f16_matmul` |
+| `nvidia_sm120` | `package_f16_softmax` |
+| `nvidia_sm120` | `package_f32_softmax` |
+| `nvidia_sm120` | `package_paged_kv_read` |
+| `nvidia_sm120` | `package_scheduled_matmul` |
+| `rocm_gfx1151` | `package_attention_backward` |
+| `rocm_gfx1151` | `package_paged_kv_read` |
+
+## How to read a closing gap
+
+A family leaves this table one of two ways, and only these two:
+
+1. **Absorbed** — the mainline compiler grows an admission predicate
+   and lowering for it, proven against the bootstrap row it replaces.
+2. **Re-expressed** — it stays hand-written or library-backed, but is
+   reached through a declared Target IR boundary
+   (`tessera_x86.abi_call` and its per-backend equivalents) so the
+   Decision #28 arbiter can score it. Chosen, never defaulted into.
+
+Deleting a `gap` row without one of those is capability loss, which
+is the failure mode Decision #31's ordering caveat exists to prevent.