diff --git a/README.md b/README.md index 579b313f2..a8e6f6b74 100644 --- a/README.md +++ b/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 diff --git a/docs/audit/generated/bootstrap_prune_gap.csv b/docs/audit/generated/bootstrap_prune_gap.csv new file mode 100644 index 000000000..bc876be1c --- /dev/null +++ b/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 diff --git a/docs/audit/generated/bootstrap_prune_gap.md b/docs/audit/generated/bootstrap_prune_gap.md new file mode 100644 index 000000000..a9e96d91f --- /dev/null +++ b/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_` 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. diff --git a/python/tessera/compiler/bootstrap_prune_audit.py b/python/tessera/compiler/bootstrap_prune_audit.py new file mode 100644 index 000000000..e87d15653 --- /dev/null +++ b/python/tessera/compiler/bootstrap_prune_audit.py @@ -0,0 +1,374 @@ +"""Gap analysis for pruning the bootstrap (Python) backend compiler. + +Direction (2026-08-30): the original Python->backend path was a **bootstrap +compiler**. The architecture is core MLIR/LLVM -- Graph -> Schedule -> Tile -> +Target driven by ``tessera-opt`` -- and the per-backend ``package_*`` families +are the prune target. This dashboard answers the question that has to be +settled *before* any of that is deleted: **which families does the mainline +compiler already cover, and which would lose their only lowering?** + +It is deliberately a generated dashboard rather than a hand-written table. +The whole point is to watch a gap close, and a hand table would be stale by +the second landing (Decision #26). + +What is derived vs declared +--------------------------- +Everything countable is **derived by AST** from the live sources: the +``package_*`` inventory per backend, the family names ``native_package_kind`` +returns, and the ``supports_scheduled_*`` predicates the driver consults. +Nothing here is a transcribed number. + +The one thing that cannot be derived is *which compiled predicate serves which +family*, because that correspondence lives in the driver's control flow rather +than in any table. It is therefore **declared** in ``_FAMILY_TO_COMPILED`` +below and **verified**: if a named module or predicate stops existing, the +generator raises rather than silently reporting a family as covered. A +mis-declared mapping would produce exactly the false "already covered" that +would make a prune lossy, so it fails closed. +""" + +from __future__ import annotations + +import ast +from dataclasses import dataclass +from pathlib import Path + +_COMPILER = Path(__file__).resolve().parent + +#: Backend bootstrap packager modules, in fleet-lead order. +_BACKEND_MODULES: tuple[tuple[str, str], ...] = ( + ("nvidia_sm120", "nvidia_native.py"), + ("rocm_gfx1151", "rocm_native.py"), + ("x86", "x86_native.py"), + ("apple_cpu", "apple_cpu_native.py"), +) + +#: family name -> (scheduled module stem, admission predicate). +#: DECLARED, not derived -- the correspondence lives in driver.py control flow. +#: Verified for existence below; a rename fails the generator rather than +#: silently marking a family covered. +_FAMILY_TO_COMPILED: dict[str, tuple[str, str]] = { + "matmul": ("scheduled_matmul", "supports_scheduled_matmul"), + "attention": ("scheduled_attention", "supports_scheduled_attention"), + "attention_backward": ( + "scheduled_attention_backward", + "supports_scheduled_attention_backward", + ), + "depth_attention": ( + "scheduled_depth_attention", + "supports_scheduled_depth_attention", + ), +} + +#: The generic compiled fallback the driver tries last. It admits by op +#: structure rather than by family name, so it is reported separately: a +#: family it happens to accept is covered, but not *by that family's name*. +_GENERIC_COMPILED = ("scheduled_kernel", "supports_scheduled_kernel") + + +@dataclass(frozen=True) +class BackendInventory: + """What one backend's bootstrap module contains.""" + + target: str + module: str + #: (name, first-parameter type) for every package_* in the module. + packagers: tuple[tuple[str, str], ...] + families: tuple[str, ...] + lines: int + + @property + def bootstrap(self) -> tuple[str, ...]: + """Packagers that re-enter Graph IR — the prune target.""" + return tuple(n for n, t in self.packagers if _is_bootstrap(t)) + + @property + def compiled_packagers(self) -> tuple[str, ...]: + """Packagers that consume an already-lowered artifact — not a target.""" + return tuple(n for n, t in self.packagers if not _is_bootstrap(t)) + + +def _parse(path: Path) -> ast.Module | None: + try: + return ast.parse(path.read_text(encoding="utf-8")) + except (OSError, SyntaxError): + return None + + +def _first_param_type(node: ast.FunctionDef) -> str: + """Annotation of the first positional parameter, or '' if unannotated.""" + args = node.args.args or node.args.posonlyargs + if not args or args[0].annotation is None: + return "" + try: + return ast.unparse(args[0].annotation) + except Exception: # pragma: no cover - defensive on exotic annotations + return "" + + +def _packagers(tree: ast.Module) -> tuple[tuple[str, str], ...]: + """(name, first-parameter type) for every ``package_*`` function. + + The first parameter is what separates the two populations, and it is a + real data-flow fact rather than a naming convention: + + * ``GraphIRModule`` -- the function reads Graph IR and emits target code + itself, bypassing Schedule and Tile. That is the bootstrap compiler. + * ``Scheduled*Artifact`` -- the function packages an artifact the compiled + route already lowered ("without Graph re-entry"). That is the mainline + compiler's packaging step and is NOT a prune target. + + An earlier version of this audit classified by name suffix and wrongly + counted six compiled-route packagers as bootstrap surface. + """ + return tuple( + sorted( + (node.name, _first_param_type(node)) + for node in tree.body + if isinstance(node, ast.FunctionDef) and node.name.startswith("package_") + ) + ) + + +def _is_bootstrap(param_type: str) -> bool: + """Whether a packager re-enters Graph IR rather than consuming an artifact.""" + return "GraphIRModule" in param_type + + +def _classified_families(tree: ast.Module) -> tuple[str, ...]: + """String literals returned by ``native_package_kind``. + + These are the families the backend's own classifier recognises, which is + the set the driver dispatches on. + """ + for node in tree.body: + if not isinstance(node, ast.FunctionDef) or node.name != "native_package_kind": + continue + names = [ + sub.value.value + for sub in ast.walk(node) + if isinstance(sub, ast.Return) + and isinstance(sub.value, ast.Constant) + and isinstance(sub.value.value, str) + ] + return tuple(dict.fromkeys(names)) + return () + + +def _predicate_exists(module_stem: str, predicate: str) -> bool: + tree = _parse(_COMPILER / f"{module_stem}.py") + if tree is None: + return False + return any( + isinstance(node, ast.FunctionDef) and node.name == predicate + for node in tree.body + ) + + +def collect_inventories() -> tuple[BackendInventory, ...]: + """Derive each backend's bootstrap surface from its live source.""" + out: list[BackendInventory] = [] + for target, filename in _BACKEND_MODULES: + path = _COMPILER / filename + tree = _parse(path) + if tree is None: + continue + text = path.read_text(encoding="utf-8") + out.append( + BackendInventory( + target=target, + module=filename, + packagers=_packagers(tree), + families=_classified_families(tree), + lines=text.count("\n") + 1, + ) + ) + return tuple(out) + + +def verify_declared_mapping() -> None: + """Fail closed if a declared compiled route no longer exists. + + A stale entry here would report a family as already covered by the + mainline compiler when it is not -- the single error that would make a + prune silently lossy. + """ + missing = [ + f"{family} -> {mod}.{pred}" + for family, (mod, pred) in _FAMILY_TO_COMPILED.items() + if not _predicate_exists(mod, pred) + ] + generic_mod, generic_pred = _GENERIC_COMPILED + if not _predicate_exists(generic_mod, generic_pred): + missing.append(f" -> {generic_mod}.{generic_pred}") + if missing: + raise RuntimeError( + "bootstrap_prune_audit: declared compiled routes no longer exist: " + + "; ".join(sorted(missing)) + + ". Update _FAMILY_TO_COMPILED rather than letting a family be " + "reported as covered when it is not." + ) + + +def family_rows() -> tuple[tuple[str, str, str, str], ...]: + """(target, family, compiled_route, status) for every classified family.""" + verify_declared_mapping() + rows: list[tuple[str, str, str, str]] = [] + for inv in collect_inventories(): + for family in inv.families: + route = _FAMILY_TO_COMPILED.get(family) + if route is None: + rows.append((inv.target, family, "—", "gap")) + else: + rows.append( + (inv.target, family, f"{route[0]}.{route[1]}", "compiled") + ) + return tuple(rows) + + +def orphan_packagers() -> tuple[tuple[str, str], ...]: + """Packagers whose name matches no classified family. + + ``package_`` is the naming convention, so a packager with no + matching family is reached some other way -- a sibling entry point, a + dtype specialisation, or dead code. Each needs its own disposition before + the prune; none can be assumed covered. + """ + out: list[tuple[str, str]] = [] + for inv in collect_inventories(): + families = set(inv.families) + for packager in inv.bootstrap: + suffix = packager.removeprefix("package_") + if suffix in families or suffix == "native": + continue + out.append((inv.target, packager)) + return tuple(out) + + +def summary() -> dict[str, int]: + rows = family_rows() + inventories = collect_inventories() + return { + "backends": len(inventories), + "packagers": sum(len(i.packagers) for i in inventories), + "bootstrap": sum(len(i.bootstrap) for i in inventories), + "compiled_packagers": sum(len(i.compiled_packagers) for i in inventories), + "lines": sum(i.lines for i in inventories), + "families": len(rows), + "compiled": sum(1 for r in rows if r[3] == "compiled"), + "gap": sum(1 for r in rows if r[3] == "gap"), + "orphan_packagers": len(orphan_packagers()), + } + + +def render_markdown() -> str: + inventories = collect_inventories() + rows = family_rows() + orphans = orphan_packagers() + s = summary() + + out: list[str] = [ + "# 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 |", + "|---|---|", + f"| Backends with a bootstrap module | {s['backends']} |", + f"| `package_*` functions total | {s['packagers']} |", + f"| — **bootstrap** (re-enter Graph IR; prune target) | {s['bootstrap']} |", + f"| — compiled-route packagers (consume a lowered artifact) | {s['compiled_packagers']} |", + f"| Lines in those modules | {s['lines']} |", + f"| Classified families | {s['families']} |", + f"| — covered by a compiled route | {s['compiled']} |", + f"| — **gap (no compiled route)** | {s['gap']} |", + f"| Packagers matching no family | {s['orphan_packagers']} |", + "", + "## Per-backend bootstrap surface", + "", + "| Target | Module | bootstrap | compiled-route | Families | Lines |", + "|---|---|---|---|---|---|", + ] + for inv in inventories: + out.append( + f"| `{inv.target}` | `{inv.module}` | {len(inv.bootstrap)} " + f"| {len(inv.compiled_packagers)} | {len(inv.families)} | {inv.lines} |" + ) + + out += [ + "", + "## 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 |", + "|---|---|---|---|", + ] + for target, family, route, status in rows: + mark = "✅ compiled" if status == "compiled" else "🔴 **gap**" + route_cell = f"`{route}`" if route != "—" else "—" + out.append(f"| `{target}` | `{family}` | {route_cell} | {mark} |") + + out += [ + "", + "## Packagers matching no classified family", + "", + "`package_` 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 |", + "|---|---|", + ] + for target, packager in orphans: + out.append(f"| `{target}` | `{packager}` |") + + out += [ + "", + "## 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.", + "", + ] + return "\n".join(out) + + +def render_csv() -> str: + lines = ["target,family,compiled_route,status"] + lines += [f"{t},{f},{r},{s}" for t, f, r, s in family_rows()] + return "\n".join(lines) + "\n" + + +if __name__ == "__main__": # pragma: no cover - manual inspection aid + print(render_markdown()) diff --git a/python/tessera/compiler/generated_docs.py b/python/tessera/compiler/generated_docs.py index c63430420..c9798db9d 100644 --- a/python/tessera/compiler/generated_docs.py +++ b/python/tessera/compiler/generated_docs.py @@ -68,6 +68,16 @@ def _r_contract_consumers_csv() -> str: return contract_consumers.render_csv() +def _r_bootstrap_prune() -> str: + from . import bootstrap_prune_audit + return bootstrap_prune_audit.render_markdown() + + +def _r_bootstrap_prune_csv() -> str: + from . import bootstrap_prune_audit + return bootstrap_prune_audit.render_csv() + + def _r_support_table() -> str: from . import audit return audit.render_markdown() @@ -472,6 +482,13 @@ def render_canonical(self) -> str: render_csv=_r_compiler_progress_csv, also_gate_md=True, ), + # ── Bootstrap-prune gap analysis (E2E-REAL-6 backend half) ── + GeneratedDoc( + "bootstrap_prune", "compiler_progress", + _GEN / "bootstrap_prune_gap.md", _r_bootstrap_prune, + csv_path=_GEN / "bootstrap_prune_gap.csv", + render_csv=_r_bootstrap_prune_csv, + ), # ── Op / primitive coverage ── GeneratedDoc( "support_table", "op_coverage", _GEN / "support_table.md", _r_support_table, diff --git a/tests/unit/test_bootstrap_prune_audit.py b/tests/unit/test_bootstrap_prune_audit.py new file mode 100644 index 000000000..4383668be --- /dev/null +++ b/tests/unit/test_bootstrap_prune_audit.py @@ -0,0 +1,87 @@ +"""The prune gap analysis must fail closed, not report a false 'covered'. + +This dashboard exists to decide what may safely be deleted. Its dangerous +failure is not being wrong loudly — it is reporting a family as already +served by the mainline compiler when it is not, because acting on that +deletes the family's only lowering. Every test here guards that direction. +""" + +from __future__ import annotations + +import pytest + +from tessera.compiler import bootstrap_prune_audit as audit + + +def test_declared_compiled_routes_all_exist(): + """A renamed predicate must raise, never silently mark a family covered.""" + audit.verify_declared_mapping() # raises on drift + + +def test_a_stale_declared_route_fails_closed(monkeypatch): + """The guard above is only worth having if it actually fires.""" + monkeypatch.setitem( + audit._FAMILY_TO_COMPILED, "matmul", ("scheduled_matmul", "no_such_predicate") + ) + with pytest.raises(RuntimeError, match="no longer exist"): + audit.verify_declared_mapping() + + +def test_bootstrap_is_classified_by_graph_ir_re_entry_not_by_name(): + """`package_scheduled_*` consumes a lowered artifact and is NOT a target. + + Classifying by name suffix counted six compiled-route packagers as + bootstrap surface — an overstatement that would have made the prune look + larger than it is. The real discriminator is whether the function takes a + `GraphIRModule` (re-enters Graph IR, bypassing Schedule/Tile) or an + already-lowered `Scheduled*Artifact`. + """ + assert audit._is_bootstrap("GraphIRModule") is True + assert audit._is_bootstrap("ScheduledAttentionArtifact") is False + assert audit._is_bootstrap("") is False + + by_target = {inv.target: inv for inv in audit.collect_inventories()} + rocm = by_target["rocm_gfx1151"] + assert "package_scheduled_attention" in rocm.compiled_packagers + assert "package_scheduled_attention" not in rocm.bootstrap + assert "package_softmax" in rocm.bootstrap + + +def test_every_packager_lands_in_exactly_one_population(): + """No packager may be silently dropped from the accounting.""" + for inv in audit.collect_inventories(): + assert len(inv.bootstrap) + len(inv.compiled_packagers) == len(inv.packagers) + assert not set(inv.bootstrap) & set(inv.compiled_packagers) + + +def test_gap_rows_are_only_ever_families_with_no_declared_route(): + """`gap` must mean exactly one thing, or the table cannot be acted on.""" + for _target, family, route, status in audit.family_rows(): + if status == "gap": + assert family not in audit._FAMILY_TO_COMPILED + assert route == "—" + else: + assert status == "compiled" + assert family in audit._FAMILY_TO_COMPILED + + +def test_the_analysis_is_not_vacuous(): + """A gap table that found nothing would pass every check above. + + The prune has not started, so both populations must be non-empty; if this + ever fails because `gap` reached zero, the guard should be replaced by the + evidence that closed it rather than deleted. + """ + s = audit.summary() + assert s["backends"] >= 3 + assert s["bootstrap"] > 0, "no bootstrap packagers found — parsing likely broke" + assert s["families"] > 0 + assert s["compiled"] > 0, "no family resolves to a compiled route" + + +def test_markdown_states_the_two_legitimate_exits(): + """The doc must not read as a delete-list.""" + text = audit.render_markdown() + assert "Absorbed" in text and "Re-expressed" in text + assert "abi_call" in text + assert "Decision #31" in text