Conversation
Turns emit/rocm_hip.py from a runner-only plugin into a FULL three-seam plugin (parallel to the x86 C lane), giving ROCm a generic compiled kernel lane for the fusable middle ground on gfx1151: - RocmHipEmitter: FusedRegion → HIP source — a one-thread-per-row __global__ kernel + a host-pointer C-ABI wrapper (H2D/launch/D2H, same shape as the shipped libtessera_rocm_gemm.so symbols). Rejects non-Fused / DYNAMIC / non-f32 via EmitError (Decision #21). - _rocm_hip_compile_fn: hipcc --offload-arch=<gfx> -O3 -shared → .so (real AOT; arch via $TESSERA_ROCM_ARCH / runtime._rocm_chip() / gfx1151 default). - RocmHipRunner.run_fused_region: dlopen + launch on gfx1151 → (out, "rocm_hip"), else numpy reference. Same NULL-buffer guard as x86 (missing bias/residual declines to the reference instead of a null-deref SIGSEGV). - run_fused_attention keeps the shipped compiled FA-2 lane (Tier-3 hand-tuned). Shared scalar body: extracted emit/_fused_scalar_body.py (pointwise/reduction snippets + the per-row compute body); x86 refactored onto it so the x86 C kernel and the ROCm HIP kernel synthesize the IDENTICAL body and stay locked to the one fusion_core numpy reference. x86 emit is byte-behavior-preserving (19/19). Lead-safety: the generic scalar HIP kernel is a correctness-first middle-ground candidate, NOT a replacement for the crown-jewel WMMA/MFMA GEMM (C3 tail = wire those Generate* passes through the loop as arbiter candidates; still open). Tests: rocm 12/12 (incl. live gfx1151 F4 gates for 7 FusedRegion chains via hipcc + the shipped flash lane) + the subprocess NULL-guard; x86 19/19; emit contract updated (unknown-emitter test now uses a genuinely-unregistered target; rocm emitter/compile now assert-registered). mypy python/tessera 0 (349); ruff clean; audit-docs 8/8; 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: a82cff7810
ℹ️ 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".
| M, K = Af.shape | ||
| _, N = Bf.shape |
There was a problem hiding this comment.
For calls where B.shape[0] != A.shape[1], this discards B's actual leading dimension and the generated HIP wrapper sizes szB from K (taken from A) and N; if B has fewer than K rows, hipMemcpy reads past the NumPy buffer, and if it has more rows the kernel silently computes a truncated matmul. The previous ROCm path delegated to region.reference, which raises on shape mismatches, so please mirror the Apple runner's K2 != K validation before building/launching the HIP kernel.
Useful? React with 👍 / 👎.
1 participant
Gives ROCm a generic compiled kernel lane on gfx1151 — the GPU analogue of the
x86 C plugin — by turning
emit/rocm_hip.pyfrom runner-only into a fullthree-seam plugin.
The three seams (C3 generic lane)
RocmHipEmitter—FusedRegion→ HIP source: a one-thread-per-row__global__kernel + a host-pointer C-ABI wrapper (H2D / launch / D2H, sameshape as the shipped
libtessera_rocm_gemm.sosymbols). Rejects non-Fused /DYNAMIC / non-f32 via
EmitError._rocm_hip_compile_fn—hipcc --offload-arch=<gfx> -O3 -shared→.so(arch via
$TESSERA_ROCM_ARCH/runtime._rocm_chip()/ gfx1151 default).RocmHipRunner.run_fused_region— dlopen + launch on gfx1151 →(out, "rocm_hip"), else numpy reference. Same NULL-buffer guard as x86 (amissing bias/residual declines to the reference instead of a null-deref SIGSEGV).
run_fused_attentionkeeps the shipped compiled FA-2 lane (Tier-3 hand-tuned).Shared scalar body
Extracted
emit/_fused_scalar_body.py(pointwise/reduction snippets + the per-rowcompute body). x86 refactored onto it, so the x86 C kernel and the ROCm HIP kernel
synthesize the identical body and stay locked to the one
fusion_corenumpyreference. x86 emit is behavior-preserving (19/19).
Proven on-device (gfx1151)
The generic HIP kernel compiles with hipcc, runs on the GPU, and passes the F4
oracle across 7 FusedRegion chains (relu / gelu / silu / bias+softmax / rmsnorm /
layer_norm / prologue), f32-exact.
Lead-safety
The generic scalar HIP kernel is a correctness-first middle-ground candidate,
not a replacement for the crown-jewel WMMA/MFMA GEMM. The C3 tail — wiring
those
Generate*MLIR passes through the loop as arbiter candidates — is stillopen (tracked in the status table); the D1 arbiter picks per measured latency +
accuracy budget.
Verification
hipcc+ theshipped flash lane + the subprocess NULL-guard); x86 19/19; emit contract
updated & green.
mypy python/tessera0 (349 files); ruff clean; audit-docs 8/8;generated-doc drift gate in sync.
🤖 Generated with Claude Code