ROCm Stages A–D: RDNA WMMA path → emit/assemble → launch bridge → on-hardware execute-and-compare (gfx1151) - #83
Merged
Merged
Conversation
The Tile→ROCm matmul lowering always emitted `tessera_rocm.mfma` regardless
of arch, which is a silent miscompile on RDNA (gfx11xx/gfx12xx) — RDNA has no
MFMA matrix core, only WMMA. First step of the Strix Halo bring-up ladder
(STRIX_HALO_EXECUTION_PLAN.md Stage A), grounded on the box: the Radeon 8060S
enumerates as gfx1100 under WSL + ROCm 7.2.4.
- Add `tessera_rocm.wmma` op (mirrors `tessera_rocm.mfma`; 16x16x16, wave32).
- `lower-tile-to-rocm{arch=...}` now arch-selects the matrix op: RDNA (gfx11xx/
gfx12xx) → WMMA, CDNA (gfx9xx) → MFMA. The no-FP8-on-RDNA gate is preserved
(fp8 matmul on gfx1100/gfx1151 is still a hard, named error).
- `lower-tessera-target-to-rocdl` lowers `tessera_rocm.wmma` to the
`llvm.amdgcn.wmma.contract` artifact marker (mirrors the mfma marker).
- New lit fixture rocm/wmma_rdna3_matmul.mlir: gfx1100 tile.mma → wmma (not
mfma) and the full path → amdgcn.wmma marker. CDNA MFMA path unchanged
(tile_matmul_to_rocm.mlir).
Docs: STRIX_HALO_EXECUTION_PLAN.md + ROCM_AUDIT.md updated with box-landed
findings — gfx1100 (not gfx1151) under WSL, cleared toolchain gates (rocminfo
enumerates w/o HSA_OVERRIDE, hipcc compiles WMMA for gfx1100, lit 11→12), the
HIP loader↔tsrRegisterGpuLauncher gap (Stage C), and the stale `--rocm-target`
fixture flag.
Verified on the box: ROCm lit suite 12/12, 101 rocm unit tests pass, doc drift
gate green, tessera-opt + tessera-rocm-opt build clean (LLVM/MLIR 22.1.8).
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: 2bb7178c97
ℹ️ 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".
Comment on lines
+123
to
+124
| StringRef matrixOp = | ||
| isWmmaArch(arch) ? "tessera_rocm.wmma" : "tessera_rocm.mfma"; |
There was a problem hiding this comment.
When lower-tile-to-rocm{arch=gfx1100} or gfx1151 sees a tile.mma with f32 operands, this new arch switch now emits tessera_rocm.wmma because only FP8 is gated earlier. The checked target selector only maps RDNA WMMA inputs through the per-arch K tables in python/tessera/compiler/rocm_mma.py (RDNA3 has fp16/bf16/int8; gfx1200 adds fp8) and rejects dtypes absent from that table, so an FP32 matmul is lowered to a non-existent WMMA instruction instead of failing or taking a fallback.
Useful? React with 👍 / 👎.
Comment on lines
+42
to
+43
| static bool isWmmaArch(llvm::StringRef arch) { | ||
| return arch.starts_with("gfx11") || arch.starts_with("gfx12"); |
There was a problem hiding this comment.
For arch=gfx1250 or gfx1251, this gfx12 prefix check selects the WMMA op, but the lowering still unconditionally attaches shape = "m16n16k16" below. The repo target model documents gfx125x as the v2 WMMA ABI with doubled K (16x16x32 for f16/bf16 and 16x16x64/128 for fp8 in python/tessera/compiler/rocm_mma.py and _WMMA_VARIANTS), so gfx125x matmuls get an invalid artifact shape.
Useful? React with 👍 / 👎.
…/gfx1100) Stage B of the Strix Halo bring-up. rocdl_emit.py (the AMD analog of ptx_emit.py) already emitted llvm.amdgcn.wmma.* LLVM IR and llc-assembled it to real v_wmma_* AMDGCN as host-free pre-work; with the box here it now runs for real (LLVM 22.1.8 AMDGPU llc) and is grounded on the actual hardware. - The part is RDNA 3.5 = gfx1151 (true ISA). Verified on the box: both `hipcc --offload-arch=gfx1151` and `llc -mcpu=gfx1151` compile, so gfx1151 is the codegen target. WSL/ROCm 7.2.4 *enumerates* the device as gfx1100 (RDNA 3 profile) — also assembles — so the rung-3 GEMM tests are parametrized over both (gfx1151 primary + gfx1100). Grounded in the RDNA 3.5 ISA Ref Guide (AMD doc 70649, 23-Jul-2024). - Add llc_object(): the WMMA GEMM lowers to a real relocatable AMD GPU ELF (EM_AMDGPU) — the plan's "compiles A to a real object" gate (test_rung3_gemm_assembles_to_amdgpu_elf_object, gfx1100+gfx1151). - _find_llc() now also finds the apt.llvm.org llc (/usr/lib/llvm-22/bin/llc), so the rung-3 llc tests run by default on the box (no PATH tweak). - test_rocdl_emit.py: 96 passed, 0 skipped. Docs: STRIX_HALO_EXECUTION_PLAN.md + ROCM_AUDIT.md — Stage B marked done/verified, gfx1151-vs-gfx1100 clarified, ISA doc URL added to cross-refs, and the MLIR --tessera-emit-rocdl crash recorded (pass `tessera-to-linalg` unregistered in tessera-opt; Stage B rides the direct LLVM-IR emitter, the ptx_emit pattern). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Owner
Author
|
Stage B added to this branch (commit
Next: Stage C (register a HIP launcher into |
gstoner
changed the title
… bridge A real GPU kernel now executes on the gfx1100 device through Tessera's C-ABI launch bridge (tsrLaunchKernel -> registered tsrGpuLauncherFn -> HIP launch), the ROCm analog of the Apple G7 proof and the first non-Apple backend to do so. - tests/unit/test_runtime_abi_rocm_launch_bridge.py: a hipcc-compiled harness registers a launcher for (target="rocm", "tessera_rocm_gemm_f32") that runs a real __global__ GEMM (hipMalloc/H2D/launch/sync/D2H) over the params' buffers+dims, compiles a `rocm` artifact, launches via tsrLaunchKernel, and verifies the GPU output equals A@B. An unregistered kernel name still returns UNIMPLEMENTED (the bridge never silently succeeds). Mirrors the Apple G7 test. - Runtime CMake HIP-include fix: with -DTESSERA_ENABLE_HIP=ON, hip_backend.cpp was compiled without the HIP include path (fatal error: hip/hip_runtime.h). tessera_runtime now links hip::host (or falls back to $ROCM_PATH/include), so libtessera_runtime.a builds with HIP enabled. - WSL device quirk handled: hipGetDeviceCount reports 0 under WSL even though kernels launch+compute correctly, so the harness gates on a real HIP probe (malloc + sync round-trip) and skip-cleans (SKIP_NO_DEVICE) without a GPU. The launcher is also two-step compiled (hipcc compiles a `.a` as source otherwise). Honesty ceiling: Stage C proves the launch-bridge + execution mechanics with a naive GEMM; it does not yet route the Stage A/B WMMA kernel, and the launcher lives in the test harness (as the Apple G7 proof does), not an auto-registered backend lib. Stage D (WMMA execute-and-compare -> flip backend_kernel) is next. Docs: STRIX_HALO_EXECUTION_PLAN.md + ROCM_AUDIT.md updated (Stage C done). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Owner
Author
|
Stage C added (commit pushed): first non-Apple kernel executes through Tessera's C-ABI launch bridge. A hipcc harness registers a |
gstoner
changed the title
…rget Clarify (per hardware owner): the Radeon 8060S (Strix Halo) is RDNA 3.5 = gfx1151; WSL/ROCm 7.2.4 currently enumerates it as gfx1100 only as a temporary limitation — AMD's WSL enablement will report the native gfx1151. So gfx1151 is the codegen target and gfx1100 is a current-WSL transitional alias (tests cover both; the Stage C launcher uses hipcc's device-default arch and auto-adapts). Also note gfx1150 is the distinct Strix Point iGPU (Radeon 890M), not this part. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…ough the bridge The real RDNA WMMA matrix instruction now runs on the gfx1100/gfx1151 device and produces a numerically correct GEMM, routed through Tessera's C-ABI launch bridge — the first on-hardware execute-and-compare of a Tessera matmul on non-Apple silicon. - tests/unit/test_rocm_wmma_execute_compare.py: a hipcc harness whose launched kernel uses __builtin_amdgcn_wmma_f32_16x16x16_f16_w32 (the same v_wmma_f32_16x16x16_f16 rocdl_emit.py emits) to compute a 16x16x16 f32<-f16 GEMM, routed through tsrLaunchKernel, compared to a host reference. The operand/accumulator fragment layout matches rocdl_emit.py's grounded mapping (col = lane&15, row = 2*e + lane>>4). maxerr ~3e-8 standalone, <1e-2 through the bridge (f16 rounding). f32<-f16 first (bf16 has documented gfx115x bugs). Honest status — NOT promoted to hardware_verified / backend_kernel complete. This clears the *numerical-proof* half of the backend_manifest hardware_verified contract (execute_compare_fixture), but that status also requires a *shipped* runtime_symbol (an auto-registered ROCm runtime launcher); today the kernel + launcher live in the test harness (like the Apple G7 proof), so flipping the status would be Decision #25 inflation. backend_kernel stays 474/0. The flip is gated on shipping the launcher (ROCM_AUDIT.md Next Work #4) and becomes mechanical once the symbol ships. Docs: STRIX_HALO_EXECUTION_PLAN.md + ROCM_AUDIT.md (Stage D proof + honest gate). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
gstoner
changed the title
Owner
Author
|
Stage D added (commit Deliberately NOT flipped to |
…compare fixture Provides the *shipped* half of the backend_manifest `hardware_verified` contract for ROCm matmul (the numerical-proof half landed in Stage D). - src/.../runtime/hip/tessera_rocm_gemm.cpp: extern "C" tessera_rocm_wmma_gemm_f16(A,B,D,M,N,K) — HIPRTC-compiles the RDNA WMMA kernel for the device's arch at load time (gfx1100 under WSL / gfx1151 after AMD's WSL enablement) and launches a 16x16x16 f32<-f16 GEMM. Built by the ordinary host compiler; only needs HIP runtime + HIPRTC at link time (no hipcc-as-compiler). - CMake: new `tessera_rocm_gemm` SHARED lib target (HIP-gated), links hip::host + hiprtc. Exports the symbol (verified with nm). - tests/unit/test_rocm_wmma_runtime_symbol.py: dlopens the shipped .so via ctypes, runs the symbol, and compares to a numpy reference GEMM (maxerr < 1e-2); plus a shape-rejection check. Skip-clean without the lib / a GPU. 2 passed. This is the production counterpart of the Stage C/D harness launcher. Both halves of `hardware_verified` (shipped runtime_symbol + execute_compare_fixture) now exist; the remaining step is the drift-gated manifest/coverage flip (rocm matmul row artifact_only -> hardware_verified) + dashboard regen. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
gstoner
force-pushed
the
rocm/gfx1100-wmma-bringup
branch
from
June 22, 2026 23:26
8929906 to
3810314
Compare
…_f16 symbol The generated runtime_abi dashboard picks up the two new extern "C" tessera_rocm_* C-ABI symbols introduced by the shipped GEMM runtime (tessera_rocm_gemm.cpp). Regenerated to clear the drift gate. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Both halves of the backend_manifest hardware_verified contract now ship for
ROCm matmul, so flip the row artifact_only -> hardware_verified:
* runtime_symbol = tessera_rocm_wmma_gemm_f16 (libtessera_rocm_gemm.so,
HIPRTC-compiled RDNA WMMA kernel)
* execute_compare_fixture = tests/unit/test_rocm_wmma_runtime_symbol.py
(dlopens the shipped symbol, f32<-f16 16x16x16 WMMA vs numpy, maxerr <1e-2)
Implemented via a new _ROCM_HARDWARE_VERIFIED override table (the ROCm analog
of _APPLE_GPU_KERNELS); when an op is in it the generic MFMA artifact row is
replaced by the hardware_verified row. Honest dtype scope (Decision #25): the
row claims fp16 + WMMA only (no CDNA MFMA shape/descriptor), shape_envelope
documents the single-tile limit.
No audit inflation: per-primitive backend_kernel stays 474 open / 0 complete
(primitive_is_complete(matmul) is still False — x86/apple/nvidia/cpu rows are
not hardware_verified). Only the rocm target row is hardware-verified.
runtime_execution_matrix gets NO rocm row yet — that maps to runtime.launch()
dispatch, and the shipped symbol is dlopened directly (not routed through
launch()); ROCm stays in _UNIMPLEMENTED_TARGETS until an auto-registered
executor lands.
Regenerated rocm_target_map (matmul -> hardware_verified|fp16; artifact 32->31)
and op_target_conformance. Doc drift gate green (17 in sync); manifest /
numerical-check / s-series universal-gate / execution-matrix tests pass.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…ch()
Adds the runtime.launch() dispatch lane for target='rocm' matmul artifacts:
* runtime.py: _execute_rocm_wmma_artifact + a cached lib loader
(_load_rocm_gemm_runtime) + host probe (_rocm_wmma_runtime_available).
The executor dlopens the shipped libtessera_rocm_gemm.so, binds the
tessera_rocm_wmma_gemm_{f16,bf16} C-ABI symbols, runs the GEMM on the AMD
GPU and returns the f32 result. Registered as 'rocm_wmma' in _executor_table.
* execution_matrix.py: KNOWN_EXECUTORS['rocm_wmma'] + an executable
(rocm, rocm_wmma) native_gpu ExecutionRow; 'rocm' dropped from
_UNIMPLEMENTED_TARGETS (named sub-arches stay — the shipped symbol HIPRTC-
compiles for whatever arch the device enumerates, so the generic rocm lane
covers execution).
Proven end-to-end on the Strix Halo box (gfx1100/WSL): launch() of a hand-built
rocm matmul artifact runs a real RDNA WMMA GEMM, maxerr ~5e-7 vs numpy
(test_rocm_launch_execute.py). Host-gated like the Apple GPU lane: the loader/
probe return None/False without an AMD GPU, and launch() reports unimplemented.
The matrix row is host-independent (dashboard renders everywhere); only a real
ROCm box dispatches there.
The jit auto-stamp of @jit(target='rocm') artifacts is intentionally NOT
included: JitFn.is_executable reads compile_bundle.execution_kind (compile-time)
which a host runtime probe can't honestly drive, so making it execute via @jit
would conflate compile-time and host-runtime state. launch() is the wired lane
(matches how Apple G7 earned its matrix row before full jit integration).
Test fixups from the prior manifest flip (matmul rocm row -> WMMA
hardware_verified, no CDNA MFMA shape/descriptor): redirect the MFMA-shape
assertion to gemm and the mma_descriptor assertion to batched_gemm (both still
on the artifact_only MFMA path). Regenerated runtime_execution_matrix +
test_coverage; doc gate 17 in sync; mypy clean.
Pre-existing env failures on this box (stale tessera-opt: apple_value_target_ir
+ emit_nvvm) are unrelated and unchanged by this commit.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Extend the shipped libtessera_rocm_gemm.so kernel from a single 16x16x16 f16
tile to a general tiled/K-looped GEMM over both half-precision storage dtypes:
* tessera_rocm_gemm.cpp: one HIPRTC source template (element type + WMMA
builtin substituted at load) tiled over (M,N) in 16x16 output tiles with a
K-loop; ragged M/N/K zero-padded on load, bounds-checked on store. Grid =
(ceil(N/16), ceil(M/16)) waves. Ships a second symbol
tessera_rocm_wmma_gemm_bf16 alongside _f16.
* Verified on the Strix Halo box (gfx1100/WSL): f16 over 16^3 / 64x48x32 /
17^3 / 128x96x64 and bf16 over 16^3 / 64x48x32 / 100x33x80 all match numpy.
Runtime + manifest follow the kernel:
* runtime._execute_rocm_wmma_artifact already forwards arbitrary M/N/K and
picks the symbol by storage dtype — now reaches the bf16 lane; rc=1 message
updated (bad shape = non-positive M/N/K).
* backend_manifest matmul rocm row: dtypes ('fp16',) -> ('fp16','bf16'),
shape_envelope -> general tiled GEMM. rocm_target_map: matmul ->
hardware_verified | fp16,bf16.
* execute_compare_fixture extended: parametrized f16/bf16 across ragged +
K-looped shapes; the old 17^3-rejection test now asserts non-positive-dim
rejection (17^3 is a valid tiled shape).
* launch() execute test covers bf16 + tiled shapes.
runtime_abi: the templated device kernel name drops the _f16_kernel pseudo-entry
(never a real host ABI symbol — it's a __global__) and adds the real _bf16 host
symbol; rocm host-symbol count stays 2. Doc gate 17 in sync; mypy clean;
manifest / numerical-check / fixture / launch tests green.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…2.3x) Move the WMMA GEMM off correctness-first naive tiling onto a *measured* perf ladder grounded in the AMD Gluon v0->v9 tutorial (ROCM_PATTERNS §B1/§B2). Kernel (tessera_rocm_gemm.cpp): each 32-lane wave now computes an MTxNT grid of 16x16 WMMA output tiles (register blocking) — a loaded A fragment is reused across NT B-tiles, a B fragment across MT A-tiles, cutting global-load traffic per output element. MT/NT are compile-time so the tile loops unroll and accumulators stay in registers. Measurement: added a device-timed bench entry point tessera_rocm_wmma_gemm_f16_bench (hipEvent, buffers reused, kernel-only — no H2D/D2H in the timed loop) + benchmarks/rocm/benchmark_rocm_wmma_gemm.py emitting the stable JSON schema (Decision #12) with a --ladder sweep. Skip-clean without a GPU; never fabricates numbers. Result on gfx1100/WSL (Ryzen AI Max+ 395 / Radeon 8060S), f16, best-of-3: 1x1 naive: 3.36 / 4.00 TFLOP/s (1024/2048) 2x4 SHIP: 7.87 / 9.46 TFLOP/s -> ~2.3x The Gluon lesson reproduced empirically: 2x2 REGRESSED below naive (occupancy/ register pressure); the non-square 2x4 won — tile shape is the lever. Shipped tiling kProdMT=2,kProdNT=4. Correctness unchanged (execute-compare fixture green at 2x4). runtime_abi: +tessera_rocm_wmma_gemm_f16_bench (318 symbols). Open rungs (documented, not yet wired): LDS staging, K-loop software pipelining, arch-aware LDS layout — heed Gluon's v6 double-buffer regression. Doc gate 17 in sync; rocm/manifest/execution-matrix tests green. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
… blocking still wins) Implement the rung-2 LDS-staged, multi-wave-workgroup WMMA GEMM and measure it honestly against rung-1 register blocking. Kernel (tessera_rocm_gemm.cpp): a WM x WN-wave workgroup cooperatively stages the A/B 16-wide K-panels for its (WM*MT*16)x(WN*NT*16) macro-tile into LDS once per K-step (constexpr LDS sizes), syncs, then every wave reads its WMMA fragments from LDS and does MT x NT register-blocked WMMA. Refactored the device launch/timing into shared helpers (runDeviceGemm + timedKernelLaunches) used by both rungs. New shipped symbols: tessera_rocm_wmma_gemm_f16_lds (correctness path) + ..._bench_lds (device-timed). Numerically correct across shapes incl. ragged 257x129x200 (fixture test_shipped_rocm_wmma_lds_matches_numpy). Verdict on gfx1100 (Ryzen AI Max+ 395), best-of-3 f16 TFLOP/s: size rung1 reg 2x4 best rung2 LDS 512^3 3.47 3.20 1024^3 8.09 7.85 2048^3 8.88 9.38 (+6%) 4096^3 11.40 8.46 Single-buffer LDS staging is a wash-to-regression here — the Strix Halo unified-memory story: global bandwidth is shared with the CPU and is NOT the bottleneck LDS staging targets, so the __syncthreads + occupancy cost isn't repaid. So PRODUCTION STAYS rung-1 register blocking (2x4). The Gluon v6 lesson generalized: measure the obvious optimization, don't assume it. Rung 2 ships as a correctness-guarded symbol because it's the substrate for rung-3 software pipelining and should pay off on discrete RDNA/CDNA. benchmark_rocm_wmma_gemm.py gains --lds (reproduces the comparison). runtime_abi: +2 symbols (320 total). Doc gate 17 in sync; rocm/manifest/ execution-matrix tests green. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The WMMA bring-up promoted rocm/matmul to hardware_verified (backed by a
real C-ABI runtime_symbol + skip-clean execute_compare_fixture) and
updated the mfma_shape / mma_descriptor guards, but missed
test_only_apple_gpu_claims_hardware_verified_today — which hard-coded
apple_gpu as the only legal hardware_verified target.
Extend it (as its own docstring anticipated for "when real NVIDIA/ROCm
proof lands"): a hardware_verified row is allowed iff it carries both
evidence fields AND is either an Apple GPU encode-session op or a Strix
Halo ROCm WMMA op (_ROCM_HARDWARE_VERIFIED_OPS = {matmul}). NVIDIA remains
disallowed (no on-silicon proof yet). Other targets / unexpected ops still
fail the guard.
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.
Summary
First rung of the Strix Halo bring-up (STRIX_HALO_EXECUTION_PLAN.md Stage A), grounded on the now-present box (Ryzen AI Max+ 395, Ubuntu 24.04/WSL2, ROCm 7.2.4, LLVM/MLIR 22.1.8).
Bug fixed:
lower-tile-to-rocmemittedtessera_rocm.mfmafor every arch — a silent miscompile on RDNA, which has no MFMA matrix core (only WMMA).Changes
tessera_rocm.wmmaop (mirrorstessera_rocm.mfma; 16×16×16, wave32).lower-tile-to-rocm{arch=...}arch-selects the matrix op: RDNA (gfx11xx/gfx12xx) → WMMA, CDNA (gfx9xx) → MFMA. The no-FP8-on-RDNA gate is preserved (fp8 matmul on gfx1100/gfx1151 stays a hard, named error).lower-tessera-target-to-rocdllowerstessera_rocm.wmma→llvm.amdgcn.wmma.contractmarker.rocm/wmma_rdna3_matmul.mlir(gfx1100 → wmma, full path → amdgcn marker; CDNA path unchanged).tsrRegisterGpuLaunchergap (Stage C); the stale--rocm-targetfixture flag.Verification (on the box)
tessera_rocm.wmma; gfx942 →tessera_rocm.mfma; fp8-on-gfx1100 → hard errortessera-opt+tessera-rocm-optbuild cleanNext (not in this PR)
Stage B (assemble via hipcc
--offload-arch=gfx1100), Stage C (register HIP launcher intotsrRegisterGpuLauncher+ HIPRTC), Stage D (execute-and-compare → first non-Applebackend_kernelproof).🤖 Generated with Claude Code