C1: x86 codegen plugin + ROCm runner→F4 gate + oracle accuracy budget - #286
Merged
Merged
Conversation
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) <noreply@anthropic.com>
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 4764a426d8
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
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) <noreply@anthropic.com>
Owner
Author
|
Fixed the P1 NULL-deref segfault in Confirmed the crash first: a residual/bias Fix: validate required buffers in Python before the ctypes launch — if the region declares a bias/residual op but the buffer is absent, route through Regression guard: |
This was referenced Jul 7, 2026
gstoner
added a commit
that referenced
this pull request
Jul 7, 2026
) * docs: reconcile status surfaces with the merged C/D/E backend wave (PRs #286–#302) The plan doc carried "landed" annotations for Workstreams C/D/E, but the all-up status surfaces had drifted behind. Refresh them (prose only; counts stay in the generated dashboards per Decision #26): - MASTER_AUDIT.md: frontmatter + body date → 2026-07-07; added a 2026-07-07 reconcile banner (A–E spine wave); refreshed the NVIDIA (Phase G) row — no longer "one op × one arch": now a generic CUDA lane over all four fusable region kinds + hand-emitted tensor-core mma.sync GEMM/flash-attn Tier-2 lanes the arbiter selects, via a shipped PTX launch bridge, all F4-gated and hardware-proven; updated the Runtime/backend + ROCm rows to note the generic plugin lanes + arbiter + x86 Zen 5 execution. - backend/nvidia/NVIDIA_AUDIT.md: frontmatter date → 2026-07-07 (body was already reconciled in #298). - COMPILER_REFACTOR_PLAN.md: the stale Phase-0 E3 row (⬜⬜⬜) → landed (E3 escape hatch lives in the D1 `force` path + PR #298 test); added the E3-landed annotation to the bullet. - README.md: refreshed the NVIDIA support-snapshot row; reframed "Where the compiler is going" from future-direction to largely-built (the three-tier framework + arbiter now live in python/tessera/compiler/emit/, NVIDIA/ROCm/x86 executing). - docs/README.md: same north-star reframe + a new NVIDIA sm_120 executing-lane status row. Doc-only. Frontmatter contract 8/8, generated-doc drift gate in sync (19), docs lint passed. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * docs: scope NVIDIA prose to the execution matrix (Codex P2) The README's Current Support Snapshot says generated audits are the authority when prose disagrees, but runtime_execution_matrix.md records only sm_120's shipped `nvidia_mma` GEMM. My reconcile prose claimed the new generic CUDA + tensor-core lanes as executable, which disagrees with the cited dashboard. Root cause: those lanes run through the arbiter/emit subsystem (emit/nvidia_cuda.py + candidate.py + in-process nvcc), a different path from the execution-matrix executor registry — so they're hardware-proven by the plugin/perf/conformance test gates (test_nvidia_plugin.py, test_nvidia_perf_ratchet.py, test_conformance_execute_compare_nvidia.py), not recorded in the matrix, and a regen won't add them. Fix (README, MASTER_AUDIT, docs/README): scope the execution-matrix-authoritative rows to `nvidia_mma`, attribute the arbiter/emit + tensor-core lanes to their real proof surface (the test gates), and add "promote the arbiter/emit lanes into the execution matrix" as an explicit still-open item — which is the genuine gap. Doc-only. Frontmatter 8/8, drift gate in sync (19), docs lint passed. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Reconciles the plan's §C1 with shipped reality, then brings two backends online
in the target-agnostic synthesizer framework (Workstream B/C0) + adds the
accuracy-budget seed that lets f16 lead kernels be gated honestly.
Doc — §C1 reconciled + status table
The plan sketched a 7-field
TargetPluginstruct; what B2/B4a + C0 actuallyshipped is three registered seams (emitter / compile_fn / runner). §C1 now
says so and maps the sketched fields onto reality (
emit_kernel/compile_fn/spec_policyexist;shape_table+cost_model→ A4/D1 arbiter;intrinsic_set→ a compile_fn build-flag;
async_model→ GPU-only, deferred). AOCL-DLP split toC1b. Added a phase-status table, softened the "Phase 0 gates everything"
wording, fixed a typo/date.
C1 — x86 (Zen 5) plugin
emit/x86_llvm.pyMirrors
apple_msl.pywith three real seams:FusedRegion→ portable f32 C (matmul + prologue/epilogue/residual/reduction). Own op→C-snippet tables matching each
fusion_corenumpyreference (the shared
EpilogueOp.emit(target)raises for non-Metal by design).Rejects non-Fused / DYNAMIC / non-f32 via
EmitError(Decision Apple GPU Tier-3: conv2d via MPSGraph convolution2D #21).cc/clang -O3 -march=native -shared→.so(real AOT).Shape-anonymous: M/N/K are runtime args, so one artifact serves all shapes.
ctypesdlopen + launch →(out, "x86_native"), else numpy"reference". F4-verified on this Zen 5 box across relu/gelu/silu/sigmoid/tanh/bias + softmax/rmsnorm/layer_norm + prologue + residual.
ROCm —
emit/rocm_hip.py(runner-only)Wires the shipped gfx1151 kernels into the universal F4 oracle — the
cross-backend differential-equivalence superpower applied to the lead's real
kernels, without claiming a generic emit lane (that's C3).
run_fused_attentionruns the compiled FA-2 lane on-device (
"rocm_hip") and is gated against thenumpy reference; other region kinds decline honestly. No emitter/compile_fn —
ROCm's kernels are shipped, not synthesized.
Oracle accuracy budget (D2 seed)
ROCm's lanes are f16 storage, so a fixed 1e-3 f32 tolerance misreads f16
rounding (~2.5e-3) as a miscompile. Added
KernelRunner.accuracy_atol; the fourverify_synthesized_*widen tolerance tomax(atol, runner.accuracy_atol).Non-breaking — Apple/x86 (f32/exact) declare no budget → unchanged. ROCm
declares 5e-3: loose enough for f16, tight enough that an O(1) miscompile is
still caught (explicitly tested).
Verification
emitter/cache 36/36; fusion/synthesis sweep 144 pass.
mypy python/tessera0 issues (348 files); ruff clean; audit-docs 8/8;generated-doc drift gate in sync (19).
🤖 Generated with Claude Code