RDNA4 (gfx1201 / Radeon AI PRO R9700) support

[The-Monk]

Description: what and why

Adds RDNA4 / gfx1201 (Radeon AI PRO R9700) as a first-class target. Opening this at AMD's request, following up on #1196.

Upstream validates against Instinct only: _GFX_TO_RUNNER maps consumer arches to None, so a Radeon box resolves to no runner and no dispatch identity. This makes gfx1201 resolve, run, and report.

The port is smaller than #1196 originally scoped — 4 modified files, +21 lines of core change, no refactors:

File	Change
common/gpu_identity.py	+7 — one row, r9700 → (gfx1201, 64). The gpu_identity refactor means _AMD_GPU_TYPES and the rocm-smi product tags both derive from this table, so the row also makes the board CLI-accepted and autodetectable. CU count read off the board (amd-smi NUM_COMPUTE_UNITS), not a spec sheet.
inference_optimizer/gpu_types.py	+5 — gfx1201 → r9700 in _GFX_TO_RUNNER.
orchestrator/phases/quantization_schemes.py	+9 — docstring only. Records why mxfp4 is refused here: the MX scaled-converts are gfx950/gfx1250-gated, so an assembler capability probe grades them REJECTED. No hardware datapath, not a missing kernel.
inference_optimizer/tests/test_rdna4_r9700_support.py	+73 — new, 8 tests.

Everything else is additive and touches no core path: examples/rdna/custom_r9700.sh (llama.cpp scriptable runner — the bypass path already resolves an operator script by the custom_{runner_type}.sh convention, and it writes the InferenceX-shaped inferencex_result.json your collectors already consume, so nothing downstream needed changing), a container recipe, a support probe, docs/rdna-port/, and an issue template.

Why llama.cpp and not vLLM/SGLang: below ~32 GB the serving stacks' VRAM and concurrency assumptions do not transfer, and llama.cpp is what this class of card is actually served with.

Two things left deliberately negative, and tested that way

• gfx1200 (Navi 44) is unmapped. Same ISA family, but nobody has measured it here, and a runner label is a claim of support.
• mxfp4 raises rather than silently degrading on gfx1201.

Three files in this diff are fork-only — flagging rather than pre-deciding

This branch doubles as the fork's landing page, so it carries framing that is correct there and wrong here. I've left it visible instead of quietly stripping it, because where the line falls is your call, not mine:

• README.md (+23) — a "community port, not an AMD product, not endorsed by AMD" banner. Obviously not something upstream's README should say about itself. Assume this one is dropped unless you want some form of it.
• docs/rdna-port/README.md — line 3 opens with the same disclaimer; line 149 describes the fork's own branch layout.
• AGENTS.md — line 249 carries the disclaimer in the SPDX/licensing note.

Say the word and I'll strip all three in a follow-up commit — it's a five-minute change and it does not touch the core diff. Separately: AGENTS.md and the README.md banner are the only root-level files here, so if you'd rather this PR carry nothing project-wide, dropping both leaves everything else intact.

Linked issue(s)

• Refs RDNA support (gfx11/gfx12): interest in a contributed runner + profiling backend? #1196 — RDNA4 (this work). Requested directly; opening as asked.
• Refs RDNA3 (gfx1100/gfx1102) support: methodology validated independently, runner map doesn't cover it #1041 — RDNA3 (gfx1100/gfx1102), @kyuubyN, who validated the methodology independently against their own engine. Closed by the stale bot, not resolved. RDNA3 isn't in this PR; the runner map doesn't cover it yet.
• Context: magpie compare produces no performance ranking on RDNA4/gfx1201 with backend: metrix Magpie#70 — magpie compare produces no performance ranking on gfx1201.

Tests: added/updated? commands run?

Added test_rdna4_r9700_support.py (8 tests), including the two negative cases above.

Measured on 2× Radeon AI PRO R9700 (gfx1201, ROCm 7.14, amdgpu 6.19.4) — verified, not asserted:

• Autodetect resolves r9700 → ('gfx1201', 64).
• Correctness vs the CPU reference: test-backend-ops MUL_MAT 1196/1196, MUL_MAT_ID 865/865.
• Runner benchmarks end-to-end: prefill ~4476 t/s, decode ~169 t/s on an 8B Q2_0 at pp512/tg64.
• The result JSON is accepted by Hyperloom's own is_valid_measurement(), not by my reading of it.
• Reproduced in a clean container against stock upstream llama.cpp — see examples/rdna/container/.
• Full suite: 756 existing tests pass, 0 failures.
• Gates on this branch rebased onto current main: ruff check clean, ruff format --check clean, reuse lint compliant.

Not tested, and labelled as such in the docs: gfx1200, anything above 2 cards, RDNA3. The 2-card host reports peers: false with host↔GPU capped at PCIe Gen3 (an on-board switch's upstream port downgrades the link; the leaf link reads Gen5 and hides it). A 4-card rig will hit walls this one can't see.

Breaking changes: no

Additive. No existing target's behaviour changes; the Instinct rows and paths are untouched.

---

One thing this port surfaces: the profile-configuration feedback loop

Raising this as a question, not a change — it needs your call on a cross-system contract, so I haven't touched it.

The optimizer's whole value here is finding a profile configuration that beats the shipped baseline. On RDNA it does. But storing that result so it survives is where the current KB doesn't quite reach:

1. Keying is fine. The canonical id is inference:{model}:{hardware}:{framework_name}:{model_type}:{architectures}:{framework_version}:{precision} — hardware is a dimension, so an r9700 recipe is already a distinct row and can't collide with an Instinct one. Good.
2. Staleness detection is not. StackFingerprint is vllm_version / aiter_commit / rocm_version — vLLM/Instinct-shaped by construction, and its docstring says it mirrors arbor's exactly "to keep binary readability." On the llama.cpp scriptable path there is no vLLM and no aiter, so two of three fields are empty.
3. And the framework version collapses. detect_framework_version resolves via _FRAMEWORK_VERSION_MODULES = {sglang, vllm, atom} by importing the package. llama.cpp is a binary, not an importable module, so it falls through to DEFAULT_FRAMEWORK_VERSION_SLUG = "unknown_version".

Net effect: a better-than-baseline configuration measured on gfx1201 can be written, but never invalidated or attributed to the build that produced it. Every llama.cpp recipe lands in the same unknown_version bucket with a fingerprint that can't answer "is this stale?" — so the loop that should make the optimizer compound across runs and across users doesn't close on this path.

The ask: is there an intended way to record an engine identity for a non-importable, non-serving runner? A llama_cpp_commit (or a generic engine_version) on StackFingerprint would do it, but that field mirrors arbor and I'm not going to unilaterally change a contract shared with a system I can't see. If you point me at the right shape I'll implement it and measure it on the R9700s — and if the answer is "out of scope for now," that's a fine answer too; the port stands without it.

The broader version of the same point: contributors with hardware AMD doesn't have will find configurations that beat the provided defaults. Right now there's no path for those to flow back. examples/rdna/rdna-support-probe.sh and the RDNA hardware-report issue template in this PR are a deliberately low-tech stand-in — one command, paste-ready output, CONFIRMED / PARTIAL / UNSUPPORTED, with a required "what you did NOT test" field. That's a bug tracker being used as a database, and it would be better as a real one.

---

Related work — RDNA4 findings this port draws on

All measured on gfx1201; linked so the numbers above have provenance rather than being claims.

The stack this runner drives

• The-Monk/The-Rock8 — RDNA4 inference appliance (container + fp8 GGUFs), the reference the runner's numbers come from.
• The-Monk/llama.cpp @ roc8 — the gfx1201 kernel work: int4 decode, 2:4-sparse SWMMAC prefill, e4m3 dot2, and a MUL_MAT_ID out-of-bounds fix.
• PrismML-Eng/llama.cpp#116 — HIP-path Q1_0/Q2_0 vec_dots (+37% / +16% decode on RDNA4), opt-in quant dedup, opt-in hipBLASLt prefill routes.

Upstream fixes filed from this hardware

• ROCm/composable_kernel#3759 — 3 correctness bugs in ck_tile 2:4-sparse SWMMAC on gfx1201, with a fail→pass repro. Open.
• ROCm/clr#283 — __byte_perm sign-replicate selector modes on the device path. Open.
• ROCm discussion #6618 — the fuller inventory: ISA findings, the gfx1201 rocprofv3 gaps that docs/rdna-port/README.md documents (host_trap-only PC sampling, FETCH_SIZE zeros, GL2C perf-level gating), and host tuning.

The profiler gaps are worth a look for anyone porting Instinct-shaped tooling: profile_standard un-gates the counters but throttles the card to ~1593 MHz vs ~2330, so it silently suppresses anything benchmarking alongside it.

Hardware I don't own: @kyuubyN reported gfx1100/gfx1102 results in #1041, and a field-tester independently validated the ISA census and cache behaviour on an RX 7600 XT (gfx1102, ROCm 7.2.4). Both are RDNA3 — out of scope for this PR, and the reason the support matrix in docs/rdna-port/README.md marks RDNA3 as methodology validated, runner map not covered.

/cc @arsalanfrq-a11y @l33tm4st3r @kyuubyN