From 4764a426d8578a49eb77390fef55d822d277259a Mon Sep 17 00:00:00 2001 From: gstoner Date: Mon, 6 Jul 2026 19:15:02 -0600 Subject: [PATCH 1/2] =?UTF-8?q?C1:=20x86=20codegen=20plugin=20+=20ROCm=20r?= =?UTF-8?q?unner=E2=86=92F4=20gate=20+=20oracle=20accuracy=20budget?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Doc — reconcile COMPILER_REFACTOR_PLAN §C1 with what B2/B4a+C0 actually shipped: the plugin is THREE registered seams (emitter / compile_fn / runner), not one `TargetPlugin` struct. Map the 7 sketched fields onto reality (emit_kernel / compile_fn / spec_policy exist; shape_table+cost_model → A4/D1; intrinsic_set → a compile_fn build-flag; async_model → GPU-only, deferred). Split AOCL-DLP out to C1b. Add a phase-status table + soften the "Phase 0 gates everything" wording + fix a typo/date. C1 — x86 (Zen 5) plugin `emit/x86_llvm.py` mirroring `apple_msl.py`: - X86CEmitter: FusedRegion → portable f32 C (matmul + prologue/epilogue/residual/ reduction); its own op→C-snippet tables matching each fusion_core numpy ref (EpilogueOp.emit(target) raises for non-Metal by design). Rejects non-Fused regions / DYNAMIC / non-f32 via EmitError (Decision #21). - _x86_compile_fn: cc/clang -O3 -march=native -shared → .so (real AOT, not Apple's compile-on-launch); shape-anonymous (M/N/K are runtime args → one artifact serves all shapes). - X86CRunner: ctypes dlopen + launch → (out, "x86_native"), else numpy reference "reference". F4-verified on this box across relu/gelu/silu/sigmoid/tanh/bias + softmax/rmsnorm/layer_norm + prologue + residual. ROCm — `emit/rocm_hip.py`: wire the SHIPPED gfx1151 kernels into the universal F4 oracle (cross-backend differential equivalence), runner-only (no emitter — ROCm's kernels are shipped, not synthesized; the generic emit lane is C3). run_fused_ attention runs the compiled FA-2 lane on-device ("rocm_hip"); other kinds decline to reference. Oracle accuracy budget (D2 seed) — the ROCm lanes are f16 storage, so a fixed 1e-3 f32 tolerance misreads f16 rounding (~2.5e-3) as a miscompile. Add `KernelRunner.accuracy_atol`; the four verify_synthesized_* widen tolerance to `max(atol, runner.accuracy_atol)`. Non-breaking: Apple/x86 declare no budget → unchanged. ROCm declares 5e-3 — loose enough for f16, tight enough that an O(1) miscompile is still caught (tested). Verify: x86 18/18, rocm 9/9 (incl. 3 live gfx1151 attention gates), kernel emitter/cache 36/36; fusion/synthesis sweep 144 pass; mypy python/tessera 0 (348 files); ruff clean; audit-docs 8/8; generated-doc drift gate in sync. Co-Authored-By: Claude Opus 4.8 (1M context) --- docs/audit/compiler/COMPILER_REFACTOR_PLAN.md | 112 ++++++-- .../tessera/compiler/emit/kernel_emitter.py | 10 + python/tessera/compiler/emit/rocm_hip.py | 82 ++++++ python/tessera/compiler/emit/x86_llvm.py | 271 ++++++++++++++++++ python/tessera/compiler/fusion_core.py | 23 +- tests/unit/test_rocm_plugin.py | 129 +++++++++ tests/unit/test_x86_plugin.py | 132 +++++++++ 7 files changed, 737 insertions(+), 22 deletions(-) create mode 100644 python/tessera/compiler/emit/rocm_hip.py create mode 100644 python/tessera/compiler/emit/x86_llvm.py create mode 100644 tests/unit/test_rocm_plugin.py create mode 100644 tests/unit/test_x86_plugin.py diff --git a/docs/audit/compiler/COMPILER_REFACTOR_PLAN.md b/docs/audit/compiler/COMPILER_REFACTOR_PLAN.md index 06dd765a1..3cda7f5d1 100644 --- a/docs/audit/compiler/COMPILER_REFACTOR_PLAN.md +++ b/docs/audit/compiler/COMPILER_REFACTOR_PLAN.md @@ -1,5 +1,5 @@ --- -last_updated: 2026-07-02 +last_updated: 2026-07-06 audit_role: plan plan_state: open --- @@ -52,6 +52,41 @@ regression-gated on real silicon. --- +## 2a. Status at a glance (updated 2026-07-06) + +Landed state per §4 phase. Inline **landed** notes in §3 carry the detail; this +table is the single skim surface. `✅` done · `🟡` partial · `⬜` not started. + +| Phase | Task | Mac (`[MAC]`) | AMD (`[AMD]`) | NV (`[NV]`) | +|---|---|:--:|:--:|:--:| +| **0** | E1 golden-IR harness + determinism roundtrip | ✅ | — | — | +| **0** | E2 real-hardware perf ratchet | ✅ (shape gate) | ✅ gfx1151 (matmul+flash, PR #284) | ⬜ sm_120 (needs box) | +| **0** | E3 escape-hatch test | ⬜ | ⬜ | ⬜ | +| **1** | A1 shared `extractPtr`/`ensureExternalDecl` | ✅ | — | — | +| **1** | A2–A4 fusion matcher / verifiers / MMA selector | ⬜ | — | — | +| **2** | B1 split `fusion.py` | ✅ | — | — | +| **2** | B2a–c `KernelEmitter`/`Runner`/`SpecPolicy` | ✅ | — | — | +| **2** | B3 F4 oracle universal (backend-agnostic, C0) | ✅ | — | — | +| **2** | B4a `kernel_cache` synth→compile→cache loop | ✅ | — | — | +| **2** | B4 real AOT `compile_fn`s (`clang`/`ptxas`/`hipcc`) | ⬜ (per-arch, → C) | ⬜ | ⬜ | +| **3** | C0 backend-plugin handoff + non-Apple F4 gate | ✅ (PR #285) | — | — | +| **3** | C1 x86 plugin (`emit/x86_llvm.py`: emitter + `cc` compile + ctypes runner) | ✅ emit host-free | ✅ execute on Zen 5 | — | +| **3** | Oracle accuracy budget (`KernelRunner.accuracy_atol`, D2 seed) | ✅ | — | — | +| **3** | C1b x86 AOCL-DLP Tier-3 candidate (opt-in) | — | ⬜ | — | +| **4** | C2 NVIDIA emit pipeline + launch bridge | ⬜ | — | ⬜ | +| **5** | C3 ROCm emit pipeline (generic synth → HIP) | ⬜ | ⬜ | — | +| **3.5** | ROCm runner → F4 gate (`emit/rocm_hip.py`, shipped kernels, f16 budget) | ✅ | ✅ gfx1151 attn | — | +| **6** | D1–D3 arbitration + measured autotune | ⬜ | ⬜ | ⬜ | + +**Gate reality (softens §4/§9.2):** "Phase 0 gates everything" holds only for the +*lead-execution* proofs. The Mac-side E1 gate is green and gfx1151 E2 is recorded, +so `[MAC]` + `[AMD]` work (A1, B1–B4a, C0, C1 authoring) has correctly proceeded; +only the **`[NV]` sm_120 proofs** still gate on the NR2 Pro box (its E2 baseline +can't be recorded yet). Do not read §4's hard-gate phrasing as blocking Mac/AMD +authoring once E1 is green. + +--- + ## 3. Workstreams Each workstream lists tasks with an owning system tag `[MAC] [AMD] [NV]` (see §7 @@ -169,22 +204,47 @@ chains, small attention). Crown-jewel GEMM stays Tier 2/3. > framework calls into, with a copy-paste skeleton, the F4-verification recipe, > and the per-backend task cards (C1 x86 · C2 NVIDIA · C3 ROCm). -- **C1 · `TargetPlugin` interface** `[MAC]` — `{emit_kernel, shape_table, - cost_model, intrinsic_set, async_model, compile_fn, spec_policy}`. Apple + x86 - are the first two reference impls (simplest to validate host-free / on Zen 5). - `spec_policy` declares which specialization modes the plugin supports (`static | - bucket | dynamic`) + its bucketing strategy — so the static-shape gate now in - the lowering (`"requires static shapes"` in `TileToX86Pass` / `MatmulToAppleCPU`) - is replaced by a *policy*, not re-hardcoded per backend. The **x86 plugin's - Tier-3 candidate set** should register **AOCL-DLP** ([amd/aocl-dlp](https://github.com/amd/aocl-dlp)) - for the Zen family — AMD's BLIS-family DL primitives (low-precision GEMM/batch - GEMM incl. INT4/FP16, pre/post-ops matching `fused_epilogue`, symmetric quant, - OpenMP). It's AVX512-based (fits the Zen 5 fleet box, which has no AMX), fills - the x86 backend's OpenMP-threading + INT4/FP16 gaps, and is opt-in behind a - build flag (a BLAS-family library like Accelerate — Decision #23-clean, kept - behind the hardware-free Target IR). The arbiter (D1) selects it only where it - measures faster than the generic kernels on Zen; check its license before it - becomes a shipped/linked lane. +- **C1 · Per-arch plugin = three registered seams (NOT one `TargetPlugin` + struct)** `[MAC]` author → `[AMD]` execute on Zen 5. **Interface reconciled + 2026-07-06:** the plan originally sketched a single `TargetPlugin` object with + seven fields; what B2/B4a + the C0 handoff actually shipped is **three separate + registries** a backend self-registers into on import (mirroring + `emit/apple_msl.py`) — this is the real, tested seam, and there is no bundled + struct. A backend adds one module `emit/.py` implementing: + 1. **`KernelEmitter`** (`register_emitter`) — `emit(region, spec, dtype, dims) + → KernelSource`. This is the original `emit_kernel` field. + 2. **`compile_fn`** (`register_compiler`) — `source → artifact` (x86: `clang + -O3 -mavx512f -mavx512bf16 -shared` → `.so`). The original `compile_fn` field. + 3. **`KernelRunner`** (`register_runner`, `default=False`) — `run_*(region, + *inputs) → (out, execution_tag)`; a real tag (`"x86_native"`) gets F4-gated, + a `REFERENCE_EXECUTIONS` tag declines. + + The original seven fields map onto shipped reality as: **`emit_kernel`** = + `KernelEmitter`; **`compile_fn`** = `register_compiler`; **`spec_policy`** = + the `SpecPolicy(static|bucket|dynamic)` a `KernelEmitter` accepts + + `bucket_key`'s strategy (this is what replaces the hard `"requires static + shapes"` gate in `TileToX86Pass` / `MatmulToAppleCPU` — a policy, not a + per-backend hardcode). The remaining four were speculative and are **not** + first-class plugin fields: **`shape_table`** + **`cost_model`** live in the + shared MMA selector (A4 `MmaDescriptor`) and the arbiter (D1), keyed per + `(op, shape-bucket, dtype, target)` — not on the emitter; **`intrinsic_set`** + is a `compile_fn` build-flag detail (x86 = `-mavx512f -mavx512bf16`, **never + `-mavx*` AMX** on this AVX-512-only fleet), not a declared field; + **`async_model`** is a no-op for the synchronous CPU/x86 lane and is deferred + to the GPU emit lanes (C2/C3) that actually need an async-token model — do not + add it to the x86 plugin. **DoD splits by system:** `emit` is pure/host-free + (`[MAC]`: mypy + ruff + emitter unit tests); the clang compile + `ctypes` + launch + F4 execute-compare require the Zen 5 box (`[AMD]`). +- **C1b · x86 Tier-3 candidate: AOCL-DLP** `[AMD]`, opt-in, **separated from + C1** — register **AOCL-DLP** ([amd/aocl-dlp](https://github.com/amd/aocl-dlp)) + as a hand-tuned candidate the D1 arbiter measures, NOT part of the core plugin. + AMD's BLIS-family DL primitives (low-precision GEMM/batch GEMM incl. INT4/FP16, + pre/post-ops matching `fused_epilogue`, symmetric quant, OpenMP); AVX512-based + (fits the Zen 5 box, no AMX), fills the x86 backend's OpenMP-threading + + INT4/FP16 gaps, opt-in behind a build flag (a BLAS-family library like + Accelerate — Decision #23-clean, behind the hardware-free Target IR). The + arbiter selects it only where it measures faster than the generic kernels on + Zen; **check its license before it becomes a shipped/linked lane.** - **C2 · NVIDIA in-process emit pipeline** `[MAC]` authoring → `[NV]` proof — `tessera-opt --tessera-emit-nvidia`: Tile IR → `ptx_emit.py` (keep sm_120 `mma.sync`; extend `wgmma` for sm_90a; stub sm_100 tcgen05) → serialize → @@ -200,6 +260,22 @@ chains, small attention). Crown-jewel GEMM stays Tier 2/3. `--tessera-emit-rocm`: drives the existing gfx1151 WMMA + CDNA MFMA `Generate*` passes through the shared loop into the launch bridge; reuses the async-token SSA model in `ROCMWaveLdsPipeline`. +- **C3-precursor (landed 2026-07-06): ROCm runner → F4 gate + oracle accuracy + budget** `[MAC]` author → `[AMD]` proof. Ahead of the full emit pipeline, the + *shipped* gfx1151 kernels are now wired into the universal F4 oracle: + `emit/rocm_hip.py` registers a **runner-only** plugin (no emitter/`compile_fn` + — ROCm's kernels are shipped, not synthesized) whose `run_fused_attention` + runs the compiled FA-2 lane on-device (tag `"rocm_hip"`) and is gated against + the numpy reference; other region kinds decline to the reference (honest — the + fused-epilogue GPU kernel is C3 proper). Because those kernels are **f16 + storage**, this required the **accuracy-budget** seed (plan D2): a + `KernelRunner.accuracy_atol` the oracle widens its tolerance to, so f16 + rounding (~2.5e-3 on the probes) is not misread as a miscompile while an O(1) + bug still is. Apple/x86 (f32/exact) declare no budget → unchanged. This is the + cross-backend differential-equivalence superpower (Theory §7.5) applied to the + lead's shipped kernels, and the first concrete slice of the accuracy-budgeted + arbiter. Proven live: `tests/unit/test_rocm_plugin.py` gates gfx1151 attention + across scale/causal on-device. ### Workstream D — Candidate arbitration + measured autotune @@ -270,7 +346,7 @@ truth in `primitive_coverage.py`). | Risk | Mitigation | |---|---| -| Shared abstraction crips a lead | Theory rule #1 + E1/E2: lead opts out per op; IR/perf gated | +| Shared abstraction cripples a lead | Theory rule #1 + E1/E2: lead opts out per op; IR/perf gated | | Synthesizer split regresses Apple | B1–B3 pure relocation, oracle-gated, no new codegen; existing differential harness proves it | | NVIDIA/ROCm emit pipelines are the big new build | Additive lanes; shipped-symbol path stays until compiled lane ≥ parity (arbiter decides) | | Silicon boxes become a bottleneck | Mac-first routing (§7): only execute-compare + perf ratchet require a box | diff --git a/python/tessera/compiler/emit/kernel_emitter.py b/python/tessera/compiler/emit/kernel_emitter.py index 8f57f351d..e29bb2d4e 100644 --- a/python/tessera/compiler/emit/kernel_emitter.py +++ b/python/tessera/compiler/emit/kernel_emitter.py @@ -252,6 +252,16 @@ class KernelRunner(ABC): #: Backend identity, e.g. ``"apple_gpu"``. Subclasses set this. target: str = "" + #: Precision budget: the largest absolute error this backend's *correct* + #: kernel may show vs the f32 numpy reference, for the F4 oracle to treat as + #: "right" rather than "buggy". ``None`` = use the oracle's default (an f32 / + #: exact backend, e.g. Apple, x86). A half-precision lead backend (ROCm f16 + #: WMMA / flash-attn) sets a looser value so f16 rounding is not misread as a + #: miscompile — while an O(1) miscompile is still caught. This is the simplest + #: slice of the accuracy-budgeted arbiter (Decision #28 / plan D2); the oracle + #: takes ``max(caller_atol, accuracy_atol)``. + accuracy_atol: float | None = None + @abstractmethod def run_fused_region(self, region: Any, *args: Any, **kwargs: Any) -> tuple[Any, str]: """Run a matmul-epilogue region on ``(A, B, bias=None, ...)``.""" diff --git a/python/tessera/compiler/emit/rocm_hip.py b/python/tessera/compiler/emit/rocm_hip.py new file mode 100644 index 000000000..c9784b716 --- /dev/null +++ b/python/tessera/compiler/emit/rocm_hip.py @@ -0,0 +1,82 @@ +"""Workstream C — ROCm gfx1151 plugin: wire the shipped hardware-verified +gfx1151 kernels into the target-agnostic F4 oracle. + +Unlike the x86 plugin (which *emits* + compiles C), ROCm's lead kernels are the +already-shipped, hardware-verified gfx1151 lanes (WMMA GEMM, compiled FA-2 +flash-attention). So this module registers a :class:`KernelRunner` **only** — no +emitter, no ``compile_fn`` — which lets ROCm's real kernels be gated by the same +universal F4 correctness oracle as the synthesized backends (the cross-backend +differential-equivalence superpower, Theory §7.5) *without* claiming a generic +ROCm emit lane (that is C3). For region kinds ROCm has no single fused GPU kernel +for yet (matmul+epilogue / gated / pointwise), it declines to the numpy reference +— honest, never a mislabeled kernel (Decision #21). + +Precision: the flash-attn lane is f16 storage / f32 accumulate, so the runner +declares an f16 accuracy budget (:attr:`accuracy_atol`); the oracle widens its +tolerance to it so f16 rounding is not misread as a miscompile, while an O(1) +miscompile is still caught (Decision #28, the accuracy-budgeted arbiter). + +Runs only where a live gfx1151 + the compiled flash lane are present (probed via +``runtime._rocm_compiled_flash_attn_available``); on any other host it declines +to the reference, so authoring/tests stay host-free. +""" +from __future__ import annotations + +from typing import Any + +from tessera.compiler.emit.kernel_emitter import KernelRunner, register_runner + +_TARGET = "rocm" +_REAL_TAG = "rocm_hip" +#: f16 storage budget for the WMMA / flash lanes vs the f32 reference. Loose +#: enough for f16 rounding (measured max ~2.5e-3 on the oracle probes), tight +#: enough that an O(1) miscompile (transpose / wrong softmax / wrong scale) is +#: still caught. +_F16_ATOL = 5e-3 + + +class RocmHipRunner(KernelRunner): + target = _TARGET + accuracy_atol = _F16_ATOL + + def run_fused_attention(self, region: Any, Q: Any, K: Any, V: Any, + *a: Any, **k: Any) -> tuple[Any, str]: + import numpy as np + try: + from tessera import runtime as rt + if not rt._rocm_compiled_flash_attn_available(): + return region.reference(Q, K, V), "reference" + Qn, Kn = region._natural(Q, K) # orient per transpose flags (f32) + Vn = np.asarray(V, np.float32) + M, D = Qn.shape + Nk, Dk = Kn.shape + if D % 16 != 0 or Dk != D: # WMMA needs head_dim % 16 == 0 + return region.reference(Q, K, V), "reference" + q = np.ascontiguousarray(Qn.reshape(1, 1, M, D), np.float16) + kk = np.ascontiguousarray(Kn.reshape(1, 1, Nk, D), np.float16) + v = np.ascontiguousarray(Vn.reshape(1, 1, Nk, D), np.float16) + out = np.asarray(rt._rocm_flash_attn(q, kk, v, scale=region.scale, + causal=region.causal)) + return out.reshape(M, D).astype(np.float32), _REAL_TAG + except Exception: + return region.reference(Q, K, V), "reference" + + # No single fused GPU kernel for these yet (that is the C3 emit lane) — + # decline honestly to the numpy reference. + def run_fused_region(self, region: Any, A: Any, B: Any, bias: Any = None, + *a: Any, residual: Any = None, + **k: Any) -> tuple[Any, str]: + return region.reference(A, B, bias, residual), "reference" + + def run_gated_matmul_region(self, region: Any, A: Any, Wg: Any, Wu: Any, + *a: Any, **k: Any) -> tuple[Any, str]: + return region.reference(A, Wg, Wu), "reference" + + def run_pointwise_graph(self, region: Any, arrays: Any, + *a: Any, **k: Any) -> tuple[Any, str]: + return region.reference(*arrays), "reference" + + +# Runner only (no emitter/compile_fn — ROCm's kernels are shipped, not +# synthesized here). default=False so Apple stays the active default runner. +register_runner(RocmHipRunner(), default=False) diff --git a/python/tessera/compiler/emit/x86_llvm.py b/python/tessera/compiler/emit/x86_llvm.py new file mode 100644 index 000000000..b0a93a2cb --- /dev/null +++ b/python/tessera/compiler/emit/x86_llvm.py @@ -0,0 +1,271 @@ +"""Workstream C1 — x86 (Zen 5) codegen plugin. Mirrors ``emit/apple_msl.py``. + +Three registered seams against the target-agnostic synthesizer (``fusion_core``), +exactly the shape ``WORKSTREAM_C_HANDOFF.md`` prescribes: + +* :class:`X86CEmitter` (``register_emitter``) — a ``FusedRegion`` → portable C + source (matmul + prologue/epilogue/residual/reduction), f32. +* :func:`_x86_compile_fn` (``register_compiler``) — ``cc``/``clang -O3 + -march=native -shared`` → a ``.so`` path (real ahead-of-time compile, not the + Apple compile-on-launch deferral). +* :class:`X86CRunner` (``register_runner``, ``default=False``) — ``ctypes`` + dlopen + launch → ``(out, "x86_native")`` when the kernel ran, else the numpy + reference tagged ``"reference"`` (Decision #21: never mislabel a fallback). + +The op-name → C-snippet tables are maintained HERE because the shared +``EpilogueOp.emit(target)`` deliberately raises for non-Metal targets (no silent +wrong-language emit). Each C snippet matches its ``EPILOGUE_OPS``/``REDUCTION_OPS`` +numpy reference so the F4 oracle (``verify_synthesized_region``) gates this +backend for real on the Zen 5 box; on a host without a C compiler the runner +skip-cleans to the reference (honest, host-free-safe). + +Scope: the f32 ``FusedRegion`` hot path (the fusable-DAG middle ground). Other +region kinds / dtypes decline via :class:`EmitError` (emit) or a ``"reference"`` +tag (run) — never a mislabeled kernel. Widen ``can_emit`` as more kinds land. +""" +from __future__ import annotations + +import ctypes +import os +import shutil +import subprocess +import tempfile +from typing import Any + +from tessera.compiler.emit.kernel_cache import build, register_compiler +from tessera.compiler.emit.kernel_emitter import ( + EmitError, + KernelEmitter, + KernelRunner, + KernelSource, + SpecPolicy, + bucket_key, + register_emitter, + register_runner, +) +from tessera.compiler.fusion_core import FusedRegion + +_TARGET = "x86" +_LANG = "c" +_ENTRY = "tessera_x86_fused" +_REAL_TAG = "x86_native" + + +# ── op-name → C snippet tables (match the fusion_core numpy references) ──────── + +def _pointwise_c(op: str, var: str) -> str: + """C statement applying pointwise epilogue op ``op`` to lvalue ``var`` (f32). + Mirrors ``EPILOGUE_OPS[op].ref`` numerically. ``bias`` is handled separately + (it reads ``bias[n]``); the rest are pure activations valid in a prologue too.""" + if op == "relu": + return f"{var} = {var} > 0.0f ? {var} : 0.0f;" + if op == "gelu": # tanh approximation, clamped — identical to fusion_core._gelu + return ( + f"{{ float _t = 0.7978845608028654f*({var}+0.044715f*{var}*{var}*{var});" + f" _t = _t < -30.0f ? -30.0f : (_t > 30.0f ? 30.0f : _t);" + f" {var} = 0.5f*{var}*(1.0f+tanhf(_t)); }}" + ) + if op == "silu": + return f"{var} = {var} / (1.0f + expf(-{var}));" + if op == "sigmoid": + return f"{var} = 1.0f / (1.0f + expf(-{var}));" + if op == "tanh": + return f"{var} = tanhf({var});" + raise EmitError(f"x86: no C snippet for pointwise op {op!r}") + + +def _reduction_c(name: str, eps: float) -> str: + """C block reducing over the length-``N`` row ``row`` in place (f32). Mirrors + ``REDUCTION_OPS[name].ref``.""" + if name == "rmsnorm": + return ( + " { float _ss = 0.0f;\n" + " for (int n = 0; n < N; ++n) _ss += row[n]*row[n];\n" + f" float _inv = 1.0f/sqrtf(_ss/(float)N + {eps!r}f);\n" + " for (int n = 0; n < N; ++n) row[n] = row[n]*_inv; }\n" + ) + if name == "softmax": + return ( + " { float _mx = -INFINITY;\n" + " for (int n = 0; n < N; ++n) if (row[n] > _mx) _mx = row[n];\n" + " float _sm = 0.0f;\n" + " for (int n = 0; n < N; ++n) { row[n] = expf(row[n]-_mx); _sm += row[n]; }\n" + " for (int n = 0; n < N; ++n) row[n] = row[n]/_sm; }\n" + ) + if name == "layer_norm": + return ( + " { float _mean = 0.0f;\n" + " for (int n = 0; n < N; ++n) _mean += row[n];\n" + " _mean /= (float)N;\n" + " float _var = 0.0f;\n" + " for (int n = 0; n < N; ++n) { float _d = row[n]-_mean; _var += _d*_d; }\n" + f" float _inv = 1.0f/sqrtf(_var/(float)N + {eps!r}f);\n" + " for (int n = 0; n < N; ++n) row[n] = (row[n]-_mean)*_inv; }\n" + ) + raise EmitError(f"x86: no C snippet for reduction op {name!r}") + + +def _synthesize_fused_c(region: FusedRegion) -> str: + """Emit the C source for a ``FusedRegion`` (f32). Signature is dims-invariant + (M/N/K are runtime args), so one kernel serves every shape — the arbiter/cache + key it shape-anonymously.""" + prologue = "".join( + f" {_pointwise_c(op, 'a')}\n" for op in region.prologue + ) + epi_lines = [] + for op in region.epilogue: + if op == "bias": + epi_lines.append(" v = v + bias[n];") + else: + epi_lines.append(f" {_pointwise_c(op, 'v')}") + epilogue = "\n".join(epi_lines) + residual = (" v = v + residual[(long)m*N + n];\n" + if region.has_residual else "") + reduction = _reduction_c(region.reduction, region.eps) if region.reduction else "" + return ( + "#include \n" + f"int {_ENTRY}(const float* A, const float* B, const float* bias,\n" + " const float* residual, float* out,\n" + " int M, int N, int K) {\n" + " for (int m = 0; m < M; ++m) {\n" + " float* row = out + (long)m * N;\n" + " for (int n = 0; n < N; ++n) {\n" + " float v = 0.0f;\n" + " for (int k = 0; k < K; ++k) {\n" + " float a = A[(long)m*K + k];\n" + f"{prologue}" + " v += a * B[(long)k*N + n];\n" + " }\n" + f"{epilogue}\n" + f"{residual}" + " row[n] = v;\n" + " }\n" + f"{reduction}" + " }\n" + " return 1;\n" + "}\n" + ) + + +# ── Seam 1: emitter ─────────────────────────────────────────────────────────── + +class X86CEmitter(KernelEmitter): + target = _TARGET + lang = _LANG + + def can_emit(self, region: Any) -> bool: + return isinstance(region, FusedRegion) + + def emit(self, region: Any, *, spec: SpecPolicy = SpecPolicy.BUCKET, + dtype: str = "f32", dims: tuple[int, ...] | None = None) -> KernelSource: + if not isinstance(region, FusedRegion): + raise EmitError( + f"X86CEmitter cannot emit a region of type {type(region).__name__} " + "(only FusedRegion so far)") + if spec is SpecPolicy.DYNAMIC: + raise EmitError("X86CEmitter does not yet support SpecPolicy.DYNAMIC " + "(bucket/static only)") + if dtype != "f32": + raise EmitError(f"X86CEmitter only supports f32 so far, got {dtype!r}") + source = _synthesize_fused_c(region) + key = bucket_key(dims, spec, dim_names=getattr(region, "dim_names", None)) + return KernelSource(source=source, entry=_ENTRY, lang=self.lang, + spec=spec, shape_key=key) + + +# ── Seam 2: compile_fn (source → .so) ───────────────────────────────────────── + +def _cc() -> str: + """The C compiler to use: ``$TESSERA_X86_CC`` override, else clang, else cc. + (Zen 5 has no AMX — ``-march=native`` enables AVX-512 without hardcoding a + flag that could fail on the NR2 Pro's non-AVX-512 host.)""" + return (os.environ.get("TESSERA_X86_CC") + or shutil.which("clang") or shutil.which("cc") + or shutil.which("gcc") or "cc") + + +def _x86_compile_fn(source: KernelSource) -> str: + """Compile the emitted C to a shared object and return its path. Raises on a + toolchain/compile failure; ``build`` wraps it in ``CompileError`` (never a + silent no-op).""" + d = tempfile.mkdtemp(prefix="tessera_x86_") + src = os.path.join(d, "kernel.c") + so = os.path.join(d, "kernel.so") + with open(src, "w") as f: + f.write(source.source) + subprocess.run( + [_cc(), "-O3", "-march=native", "-fPIC", "-shared", src, "-o", so, "-lm"], + check=True, capture_output=True, text=True) + return so + + +# ── Seam 3: runner (execute → (out, tag)) ───────────────────────────────────── + +_LIB_CACHE: dict[str, Any] = {} + + +def _load_entry(artifact: str): + """dlopen ``artifact`` (cached) and return its bound entry symbol with the + fixed C ABI: ``int(A, B, bias, residual, out, M, N, K)``.""" + lib = _LIB_CACHE.get(artifact) + if lib is None: + lib = ctypes.CDLL(artifact) + _LIB_CACHE[artifact] = lib + fn = getattr(lib, _ENTRY) + fn.restype = ctypes.c_int + fn.argtypes = [ctypes.c_void_p] * 5 + [ctypes.c_int] * 3 + return fn + + +def _ptr(arr): + return arr.ctypes.data_as(ctypes.c_void_p) if arr is not None else None + + +class X86CRunner(KernelRunner): + target = _TARGET + + def run_fused_region(self, region: Any, A: Any, B: Any, bias: Any = None, + *args: Any, residual: Any = None, + **kwargs: Any) -> tuple[Any, str]: + import numpy as np + try: + Af = np.ascontiguousarray(A, np.float32) + Bf = np.ascontiguousarray(B, np.float32) + M, K = Af.shape + _, N = Bf.shape + # Shape-anonymous build: the C kernel takes M/N/K at runtime, so one + # compiled artifact serves every shape (dims=None → no shape key). + compiled = build(region, _TARGET, dtype="f32", dims=None) + fn = _load_entry(compiled.artifact) + bias_arr = (np.ascontiguousarray(bias, np.float32) + if bias is not None else None) + res_arr = (np.ascontiguousarray(residual, np.float32) + if residual is not None else None) + out = np.zeros((M, N), np.float32) + rc = fn(_ptr(Af), _ptr(Bf), _ptr(bias_arr), _ptr(res_arr), + _ptr(out), M, N, K) + if rc == 1: + return out, _REAL_TAG + except Exception: + pass + return region.reference(A, B, bias, residual), "reference" + + # x86 has no fused GPU-style kernel for these yet — decline honestly (the + # numpy reference, tagged so the oracle trusts rather than gates it). + def run_fused_attention(self, region: Any, Q: Any, K: Any, V: Any, + *a: Any, **k: Any) -> tuple[Any, str]: + return region.reference(Q, K, V), "reference" + + def run_gated_matmul_region(self, region: Any, A: Any, Wg: Any, Wu: Any, + *a: Any, **k: Any) -> tuple[Any, str]: + return region.reference(A, Wg, Wu), "reference" + + def run_pointwise_graph(self, region: Any, arrays: Any, + *a: Any, **k: Any) -> tuple[Any, str]: + return region.reference(*arrays), "reference" + + +# ── registration (import side effect, exactly like apple_msl) ───────────────── +register_emitter(X86CEmitter()) +register_compiler(_TARGET, _x86_compile_fn) +register_runner(X86CRunner(), default=False) diff --git a/python/tessera/compiler/fusion_core.py b/python/tessera/compiler/fusion_core.py index a6428b92c..4e1cede56 100644 --- a/python/tessera/compiler/fusion_core.py +++ b/python/tessera/compiler/fusion_core.py @@ -49,6 +49,17 @@ def _runner() -> KernelRunner: return r +def _effective_atol(runner: KernelRunner, atol: float) -> float: + """Oracle tolerance for ``runner``: the caller's ``atol`` widened to the + backend's declared precision budget (``runner.accuracy_atol``), so an f16 lead + kernel's rounding is not misread as a miscompile while an O(1) miscompile is + still caught. A ``None`` budget (f32/exact backends — Apple, x86) leaves + ``atol`` unchanged. Simplest slice of the accuracy-budgeted arbiter + (Decision #28 / plan D2).""" + budget = getattr(runner, "accuracy_atol", None) + return atol if budget is None else max(atol, budget) + + SYNTH_MAX_N = 1024 #: Cap on head_dim for the ONLINE-softmax attention kernel (M2): it streams keys @@ -678,7 +689,8 @@ def verify_synthesized_gated(region: GatedMatmulRegion, *, seed: int = 0, if execution in REFERENCE_EXECUTIONS: verdict = True else: - verdict = bool(np.allclose(out, region.reference(A, Wg, Wu), atol=atol)) + verdict = bool(np.allclose(out, region.reference(A, Wg, Wu), + atol=_effective_atol(r, atol))) _GATED_VERIFY_CACHE[key] = verdict return verdict @@ -911,7 +923,8 @@ def verify_synthesized_region(region: FusedRegion, *, seed: int = 0, if execution in REFERENCE_EXECUTIONS: verdict = True # no synthesized kernel to distrust else: - verdict = bool(np.allclose(out, region.reference(A, B, bias), atol=atol)) + verdict = bool(np.allclose(out, region.reference(A, B, bias), + atol=_effective_atol(r, atol))) _VERIFY_CACHE[key] = verdict return verdict @@ -933,7 +946,8 @@ def verify_synthesized_attention(region: AttentionRegion, *, seed: int = 0, if execution in REFERENCE_EXECUTIONS: verdict = True else: - verdict = bool(np.allclose(out, region.reference(Q, K, V), atol=atol)) + verdict = bool(np.allclose(out, region.reference(Q, K, V), + atol=_effective_atol(r, atol))) _VERIFY_CACHE[key] = verdict return verdict @@ -968,7 +982,8 @@ def verify_synthesized_pointwise(region: PointwiseGraphRegion, *, seed: int = 0, verdict = True # no synthesized kernel to distrust else: verdict = bool(np.allclose(out, region.reference(*probes), - atol=atol, equal_nan=True)) + atol=_effective_atol(r, atol), + equal_nan=True)) _VERIFY_CACHE[key] = verdict return verdict diff --git a/tests/unit/test_rocm_plugin.py b/tests/unit/test_rocm_plugin.py new file mode 100644 index 000000000..228488199 --- /dev/null +++ b/tests/unit/test_rocm_plugin.py @@ -0,0 +1,129 @@ +"""Workstream C — ROCm gfx1151 plugin: wire shipped kernels into the F4 oracle. + +Three layers: + +1. **Registration + decline paths (host-free)** — a runner-only plugin registers + for target "rocm" (no emitter/compiler); it declines matmul-epilogue / gated / + pointwise regions to the numpy reference (ROCm has no single fused GPU kernel + for those yet — that is C3). +2. **Accuracy-budget wiring (host-free)** — the F4 oracle widens its tolerance to + a runner's declared ``accuracy_atol`` so an f16 lead kernel's rounding is not + misread as a miscompile, while an O(1) miscompile still is. +3. **Live attention gate (needs a live gfx1151 + compiled flash lane)** — the + shipped compiled flash-attn kernel is gated by the same universal oracle: + `run_fused_attention` runs on-device ("rocm_hip") and matches the numpy + reference within the f16 budget. +""" +from __future__ import annotations + +import numpy as np +import pytest + +import tessera.compiler.fusion as F +import tessera.compiler.emit.rocm_hip as rocm # noqa: F401 — self-registers +from tessera.compiler.emit.kernel_emitter import KernelRunner, get_runner + + +def _rocm_flash_live() -> bool: + try: + from tessera import runtime as rt + return rt._rocm_compiled_flash_attn_available() + except Exception: + return False + + +# ── 1. Registration + decline paths (host-free) ─────────────────────────────── + +def test_rocm_runner_registered_with_f16_budget(): + r = get_runner("rocm") + assert r.target == "rocm" + assert r.accuracy_atol == 5e-3 # f16 storage budget + + +def test_rocm_registers_no_emitter(): + # ROCm's kernels are shipped, not synthesized here — emit("rocm") still raises. + from tessera.compiler.emit.kernel_emitter import EmitError, emit_kernel + with pytest.raises(EmitError, match="no KernelEmitter registered"): + emit_kernel(F.FusedRegion(epilogue=("relu",)), "rocm") + + +def test_rocm_declines_non_attention_regions(): + r = get_runner("rocm") + A = np.zeros((8, 12), np.float32) + B = np.zeros((12, 16), np.float32) + _, ex = r.run_fused_region(F.FusedRegion(epilogue=("relu",)), A, B, None) + assert ex == "reference" + _, ex = r.run_gated_matmul_region(F.GatedMatmulRegion(), + A, np.zeros((12, 16), np.float32), + np.zeros((12, 16), np.float32)) + assert ex == "reference" + + +# ── 2. Accuracy-budget wiring (host-free) ───────────────────────────────────── + +class _FakeF16Attn(KernelRunner): + """Returns the exact reference plus a fixed 3e-3 perturbation under a real + tag — models an f16 kernel whose rounding is within budget but past 1e-3.""" + target = "rocm_budget_test" + accuracy_atol: float | None = 5e-3 + + def run_fused_region(self, region, *a, **k): raise NotImplementedError + def run_fused_attention(self, region, Q, K, V, *a, **k): + return region.reference(Q, K, V) + np.float32(3e-3), "rocm_hip" + def run_gated_matmul_region(self, region, *a, **k): raise NotImplementedError + def run_pointwise_graph(self, region, *a, **k): raise NotImplementedError + + +def test_oracle_honors_accuracy_budget(): + F.clear_verification_cache() + region = F.AttentionRegion(scale=0.25) + within = _FakeF16Attn() + assert F.verify_synthesized_attention(region, runner=within, force=True) is True + + +def test_oracle_without_budget_rejects_same_error(): + # Same 3e-3 error but NO declared budget (accuracy_atol=None) → the default + # 1e-3 oracle rejects it. Proves the budget is what admits the f16 kernel. + F.clear_verification_cache() + + class _NoBudget(_FakeF16Attn): + target = "rocm_nobudget_test" + accuracy_atol = None + + assert F.verify_synthesized_attention( + F.AttentionRegion(scale=0.25), runner=_NoBudget(), force=True) is False + + +def test_oracle_budget_still_catches_gross_miscompile(): + # A budget must not blind the oracle to a real bug: an O(1) wrong result is + # still rejected even under the f16 budget. + F.clear_verification_cache() + + class _Wrong(_FakeF16Attn): + target = "rocm_wrong_test" + def run_fused_attention(self, region, Q, K, V, *a, **k): + return np.full_like(region.reference(Q, K, V), 9.0), "rocm_hip" + + assert F.verify_synthesized_attention( + F.AttentionRegion(scale=0.25), runner=_Wrong(), force=True) is False + + +# ── 3. Live attention gate (needs a live gfx1151 + compiled flash lane) ─────── + +@pytest.mark.slow +@pytest.mark.skipif(not _rocm_flash_live(), + reason="live gfx1151 + compiled flash-attn lane required") +@pytest.mark.parametrize("scale,causal", [(1.0, False), (0.25, False), (0.125, True)]) +def test_live_rocm_attention_gated(scale, causal): + F.clear_verification_cache() + region = F.AttentionRegion(scale=scale, causal=causal) + runner = get_runner("rocm") + # It actually runs on the GPU... + rng = np.random.default_rng(0) + Q = rng.standard_normal((8, 16)).astype(np.float32) + K = rng.standard_normal((8, 16)).astype(np.float32) + V = rng.standard_normal((8, 16)).astype(np.float32) + out, execution = runner.run_fused_attention(region, Q, K, V) + assert execution == "rocm_hip" + # ...and the universal F4 oracle gates it within the f16 budget. + assert F.verify_synthesized_attention(region, runner=runner, force=True) is True diff --git a/tests/unit/test_x86_plugin.py b/tests/unit/test_x86_plugin.py new file mode 100644 index 000000000..315ebd078 --- /dev/null +++ b/tests/unit/test_x86_plugin.py @@ -0,0 +1,132 @@ +"""Workstream C1 — x86 (Zen 5) codegen plugin contract + F4 gating. + +Mirrors the Apple emitter/runner contract tests for the new x86 backend +(`emit/x86_llvm.py`). Three layers, matching the handoff's definition of done: + +1. **Registration + emit (host-free)** — the three seams register for target + "x86"; `emit` produces C source and rejects unsupported regions/policies/dtypes + (Decision #21). +2. **F4 gating (host-free-safe)** — the universal oracle gates the x86 runner: + a wrong kernel is rejected, a correct one trusted. On a host without a C + compiler the runner skip-cleans to the numpy reference (tag "reference"), + which the oracle trusts — so the layer stays green everywhere and becomes a + real silicon check on the Zen 5 box. +3. **Real execution (needs a C compiler)** — compile + `ctypes` launch on this + box; assert the kernel ran ("x86_native") and matches numpy across the + epilogue / reduction / prologue chains. +""" +from __future__ import annotations + +import numpy as np +import pytest + +import tessera.compiler.fusion as F +import tessera.compiler.emit.x86_llvm as x86 # noqa: F401 — self-registers +from tessera.compiler.emit.kernel_emitter import ( + EmitError, SpecPolicy, get_emitter, get_runner, +) + +_HAVE_CC = x86._cc() is not None and __import__("shutil").which(x86._cc()) is not None + + +# ── 1. Registration + emit (host-free) ──────────────────────────────────────── + +def test_x86_seams_registered(): + from tessera.compiler.emit.kernel_cache import get_compiler + assert get_emitter("x86").target == "x86" + assert get_runner("x86").target == "x86" + assert callable(get_compiler("x86")) + + +def test_x86_does_not_hijack_active_runner(): + # Registered default=False, so Apple stays the active default runner. + from tessera.compiler.emit.kernel_emitter import active_runner + ar = active_runner() + assert ar is None or ar.target != "x86" + + +def test_emit_produces_c_source(): + src = get_emitter("x86").emit(F.FusedRegion(epilogue=("bias", "gelu")), dtype="f32") + assert src.lang == "c" + assert src.entry == "tessera_x86_fused" + assert "int tessera_x86_fused(" in src.source + assert "bias[n]" in src.source and "tanhf" in src.source # bias + gelu emitted + + +def test_emit_rejects_non_fused_region(): + with pytest.raises(EmitError, match="cannot emit"): + get_emitter("x86").emit(F.AttentionRegion()) + + +def test_emit_rejects_dynamic_spec(): + with pytest.raises(EmitError, match="DYNAMIC"): + get_emitter("x86").emit(F.FusedRegion(epilogue=("relu",)), spec=SpecPolicy.DYNAMIC) + + +def test_emit_rejects_non_f32_dtype(): + with pytest.raises(EmitError, match="f32"): + get_emitter("x86").emit(F.FusedRegion(epilogue=("relu",)), dtype="f16") + + +# ── 2. F4 gating (host-free-safe) ───────────────────────────────────────────── + +def test_oracle_trusts_correct_or_fallback_x86(): + # A correct kernel (or a reference fallback on a compiler-less host) is trusted. + F.clear_verification_cache() + for r in (F.FusedRegion(epilogue=("relu",)), + F.FusedRegion(epilogue=("bias", "gelu")), + F.FusedRegion(epilogue=(), reduction="softmax")): + assert F.verify_synthesized_region(r, runner=get_runner("x86"), force=True) is True + + +def test_oracle_rejects_wrong_x86_kernel(): + # Prove the gate BITES for x86: a runner that returns a wrong result under the + # real-execution tag must be rejected (not silently trusted). + F.clear_verification_cache() + + class _WrongX86(x86.X86CRunner): + def run_fused_region(self, region, A, B, bias=None, *a, **k): + return np.full((A.shape[0], B.shape[1]), 999.0, np.float32), "x86_native" + + assert F.verify_synthesized_region( + F.FusedRegion(epilogue=("relu",)), runner=_WrongX86(), force=True) is False + + +# ── 3. Real execution (needs a C compiler; runs on the Zen 5 box) ───────────── + +_CHAINS = [ + F.FusedRegion(epilogue=("relu",)), + F.FusedRegion(epilogue=("bias", "gelu")), + F.FusedRegion(epilogue=("silu",)), + F.FusedRegion(epilogue=("sigmoid",)), + F.FusedRegion(epilogue=("tanh",)), + F.FusedRegion(epilogue=("bias",), reduction="softmax"), + F.FusedRegion(epilogue=(), reduction="rmsnorm"), + F.FusedRegion(epilogue=("relu",), reduction="layer_norm"), + F.FusedRegion(epilogue=("gelu",), prologue=("relu",)), +] + + +@pytest.mark.skipif(not _HAVE_CC, reason="no C compiler (clang/cc/gcc) on host") +@pytest.mark.parametrize("region", _CHAINS, ids=lambda r: f"{r.epilogue}/{r.reduction}/{r.prologue}") +def test_x86_kernel_runs_and_matches_numpy(region): + runner = get_runner("x86") + rng = np.random.default_rng(0) + A = rng.standard_normal((8, 12)).astype(np.float32) + B = rng.standard_normal((12, 16)).astype(np.float32) + bias = rng.standard_normal((16,)).astype(np.float32) if region.has_bias else None + out, execution = runner.run_fused_region(region, A, B, bias) + assert execution == "x86_native" # a real compiled kernel ran on this box + assert np.allclose(out, region.reference(A, B, bias), atol=1e-3) + + +@pytest.mark.skipif(not _HAVE_CC, reason="no C compiler (clang/cc/gcc) on host") +def test_x86_residual_path_matches_numpy(): + region = F.FusedRegion(epilogue=("gelu",), residual=True) + rng = np.random.default_rng(1) + A = rng.standard_normal((8, 12)).astype(np.float32) + B = rng.standard_normal((12, 16)).astype(np.float32) + R = rng.standard_normal((8, 16)).astype(np.float32) + out, execution = get_runner("x86").run_fused_region(region, A, B, None, residual=R) + assert execution == "x86_native" + assert np.allclose(out, region.reference(A, B, None, R), atol=1e-3) From ada3a817dfb79457294bebe9337482119f687a5e Mon Sep 17 00:00:00 2001 From: gstoner Date: Mon, 6 Jul 2026 19:30:58 -0600 Subject: [PATCH 2/2] C1: guard x86 runner against NULL bias/residual deref (PR #286 review) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit P1 from review: a residual (or bias) FusedRegion invoked without the required buffer — e.g. verify_synthesized_region calls run_fused_region(region, A, B, bias) with no residual — left res_arr None, and the emitted C dereferences residual[(long)m*N + n] / bias[n]. The null pointer segfaulted the process (SIGSEGV, reproduced: rc=-11) BEFORE the except could fall back — uncatchable. Fix: validate required buffers in Python before the ctypes launch. If the region declares a bias/residual op but the corresponding buffer is absent, route through region.reference (which raises a clean, catchable ValueError naming the missing operand) instead of launching the kernel with a null it will deref. Test: test_x86_missing_required_buffer_declines_not_segfault runs the risky call in a CHILD process and asserts rc==0 (no SIGSEGV) + a clean ValueError, so a regression fails an assert rather than crashing the session. x86 19/19, rocm 9/9, emit contract green; ruff + mypy clean. Co-Authored-By: Claude Opus 4.8 (1M context) --- python/tessera/compiler/emit/x86_llvm.py | 9 +++++++ tests/unit/test_x86_plugin.py | 34 ++++++++++++++++++++++++ 2 files changed, 43 insertions(+) diff --git a/python/tessera/compiler/emit/x86_llvm.py b/python/tessera/compiler/emit/x86_llvm.py index b0a93a2cb..66d17277f 100644 --- a/python/tessera/compiler/emit/x86_llvm.py +++ b/python/tessera/compiler/emit/x86_llvm.py @@ -228,6 +228,15 @@ def run_fused_region(self, region: Any, A: Any, B: Any, bias: Any = None, *args: Any, residual: Any = None, **kwargs: Any) -> tuple[Any, str]: import numpy as np + # Required-buffer guard BEFORE any launch: the emitted C dereferences + # bias[n] / residual[...] whenever the region declares them, so a missing + # buffer would pass a NULL pointer and segfault PAST Python's ``except`` + # (an uncatchable SIGSEGV). Route such an ill-formed call through the numpy + # reference instead, which raises a clean, catchable ValueError naming the + # missing operand — never launch the kernel with a null it will deref. + if (region.has_bias and bias is None) or \ + (region.has_residual and residual is None): + return region.reference(A, B, bias, residual), "reference" try: Af = np.ascontiguousarray(A, np.float32) Bf = np.ascontiguousarray(B, np.float32) diff --git a/tests/unit/test_x86_plugin.py b/tests/unit/test_x86_plugin.py index 315ebd078..6d0a41f87 100644 --- a/tests/unit/test_x86_plugin.py +++ b/tests/unit/test_x86_plugin.py @@ -120,6 +120,40 @@ def test_x86_kernel_runs_and_matches_numpy(region): assert np.allclose(out, region.reference(A, B, bias), atol=1e-3) +def test_x86_missing_required_buffer_declines_not_segfault(): + # A residual/bias region invoked WITHOUT the required buffer must NOT launch + # the kernel — the emitted C dereferences residual[...] / bias[n], so a null + # would SIGSEGV past Python's except. The runner routes through the reference, + # which raises a clean ValueError. Run in a CHILD process so a regression + # (segfault) surfaces as a failed assert, not a crashed test session. + import subprocess + import sys + import textwrap + code = textwrap.dedent( + """ + import numpy as np + import tessera.compiler.fusion as F + import tessera.compiler.emit.x86_llvm as x86 + r = x86.X86CRunner() + A = np.zeros((8, 12), np.float32) + B = np.zeros((12, 16), np.float32) + for region in (F.FusedRegion(epilogue=("relu",), residual=True), + F.FusedRegion(epilogue=("bias", "relu"))): + try: + r.run_fused_region(region, A, B, None) + raise SystemExit("expected ValueError, got a result") + except ValueError: + pass + print("ok") + """ + ) + p = subprocess.run([sys.executable, "-c", code], capture_output=True, text=True) + assert p.returncode == 0, ( + f"missing-buffer guard failed (rc={p.returncode}, -11=SIGSEGV): " + f"{p.stderr[-300:]}") + assert "ok" in p.stdout + + @pytest.mark.skipif(not _HAVE_CC, reason="no C compiler (clang/cc/gcc) on host") def test_x86_residual_path_matches_numpy(): region = F.FusedRegion(epilogue=("gelu",), residual=True)