Post-hackathon contribution offer: RDNA4 (gfx1201) LLM inference findings — kernels, ISA probes, profiler gaps, host tuning

[The-Monk]

Context: our AMD hackathon entry (Hyperloom Track 2, July 2026 — currently in judging) is built on a heavily tuned llama.cpp HIP backend for RDNA4 (2× Radeon AI PRO R9700, gfx1201), plus supporting profiling and benchmarking tooling. The submission is frozen while judging runs, but afterwards we would rather upstream what is generally useful than leave it in a fork. Before opening PRs, we would like maintainer direction on what is welcome and where it should land.

What we have

HIP kernel work (llama.cpp backend, RDNA4) — likely belongs in llama.cpp/ggml rather than ROCm proper; listed for visibility since it is all-AMD:

• Borrow-free SWAR unpacks for 1-bit/ternary weight formats replacing v_cmp/v_cndmask chains and the HIP __byte_perm emulation tax: on mainline llama.cpp HIP, Q1_0 tg128 29.5 → 64.6 t/s (+119%), pp2048 +30%, bit-exact (branch: The-Monk/llama.cpp hip-q1q2-swar).
• Native MXFP6 MMQ tile (e3m2 embedded into e4m3 at tile load, fp8 dot path reused): pp512 +40%, quality-exact.
• fp8 (e4m3) decode via hardware dot2 with lossless q8_1 activations: +2% throughput and better accuracy than the fp8×fp8 dot4 route.
• MUL_MAT_ID row-guard out-of-bounds fix (silent MoE decode corruption under rows-per-block > 1 geometry).
• DPP-only warp reductions (quad_perm / row_mirror / permlanex16 involution chain) replacing ds_bpermute butterflies in SSM kernels — zero LDS-pipe traffic.

ISA-level findings (documentation candidates):

• __byte_perm semantics divergence HIP vs CUDA for selectors ≥ 8 (mode-less wrap vs sign-replicate) — probe and details already posted on [Documentation]: ROCm Documentation Gap: WMMA Output Lane Mapping for gfx12 (RDNA4) #6025.
• WMMA C-accessor lane contract on gfx12 (column-major accessor indexing + the intra-warp-lane requirement) — also on [Documentation]: ROCm Documentation Gap: WMMA Output Lane Mapping for gfx12 (RDNA4) #6025.
• gfx1201 integer-dot inventory from a full LLVM mnemonic sweep: v_wmma_i32_16x16x32_iu4 is genuinely full-rate (measured 305 vs 148 Tmac/s iu8 = 2.07×), swmmac iu4 variants present, sdot8 (i4×i4) absent, no 2-bit dot/wmma variants exist at all. Rate microbenches included. As far as we can tell nothing upstream (composable_kernel included) wires the gfx12 16x16x32_iu4 shape yet — int4 matrix looks like an RDNA-exclusive lineage worth documenting.

rocprofiler-sdk / rocprofv3 gaps on gfx1201 (bug reports with repros ready):

• GL2C base counters return zeros/sentinels unless power_dpm_force_performance_level=profile_standard; the size-split EA variants (GL2C_EA_RDREQ_32B/64B/128B) are all-zero regardless of powerplay mask; base values under-report by a roughly constant ~4× (partial instance sampling suspected) — ratios are usable, absolutes are not.
• PC sampling: only host_trap works on gfx1201 (stochastic unsupported), minimum interval 512, occasional segfault at teardown.
• FETCH_SIZE-class derived counters all-zero.
• Profiler attach measurably suppresses graph-level kernel fusions in the profiled app — a perturbation worth a documentation note.

Host tuning + benchmarking methodology (documentation candidates): recommended amdgpu kernel args for inference hosts (gpu_recovery, lockup_timeout rationale), clock-pinning discipline (rocm-smi --setperflevel high is 10–12% slower than auto for inference on R9700 but ~3× tighter variance — pin only for A/B discrimination, publish auto numbers), NUMA/topology discipline, GPU heat-soak effects on back-to-back runs, and a benchmark-provenance tool (per-run commit + config + counter capture) we would be happy to contribute.

Questions for maintainers

1. Which of these would you take as PRs or issues, and against which repos (ROCm docs, rocprofiler-sdk, composable_kernel, elsewhere)?
2. For the profiler gaps: are issues-with-repros the right channel, or is there a preferred triage path?
3. Would there be interest in RDNA support for AMD-AGI/Hyperloom (a llama.cpp runner + a gfx1201 instrument map)? Happy to open a design issue there first if so.

Everything stays frozen until judging concludes — this post is only to get direction so contributions land where you want them.

[The-Monk]

@doplxyz — following up here since this is the "what we'd like to contribute" thread, and after the last week (your CK #3759 review + mutation matrix, the EA request-size work on #6613) it's clear you're the right person to ask.

Would you be willing to look through our two trees and carry upstream whatever you judge upstream-worthy — as author or co-author, entirely your call? I care that the fixes land, not whose name is on them. Concretely:

• The-Monk/llama.cpp rock8-release — the living RDNA4 tree: ternary Q1_0/Q2_0 decode kernels (mmvq + SWAR MMQ tile unpack), fp8 E4M3 paths + hipBLASLt prefill routes, a native MXFP6 MMQ tile (unpack-to-e4m3 at tile load, +40% prefill, PPL-exact), KV/dispatch fixes, and the measurement tooling: an ISA datapath grid for gfx1201, a calibrated kernel-roofline calculator, and the GL2C-EA byte instrument built on exactly the 256B/request model you established (our audit yesterday used bytes = RDREQ×256 as ground truth to correct several of our own roofline formulas — happy to share the full audit).
• The-Monk/The-Rock8 — the hackathon appliance this all shipped in. Read-only until judging ends, but everything in it is visible, and the post-judging changelog will unfreeze the rest.

Low-effort starting points if you want them: the __byte_perm sign-replicate fix just went up as ROCm/clr#283 (your kind of thing — the device test is exhaustive-16-bit); the MXFP6 tile probably belongs upstream in some form; and the boundary double-fetch behavior we measured (+1×256B line per row, matches your 512 MiB excess-count observation) might interest you for #6613.

We can support with anything you need: repro runs on the two R9700s, perplexity/correctness gates, benchmarks under the DPM/profiling constraints we've mapped. No obligation and no timeline — take whatever interests you and ignore the rest.

[doplxyz]

@The-Monk — thank you, and I want to give you a straight answer rather than a polite one.

I am not going to carry the trees. Not because of the work, and not because of authorship —
your offer on that is more generous than it needed to be. It is that carrying something upstream
means vouching for it, and the only way I know how to vouch for a change is the way I did on
CK #3759 and clr #283: build it, measure it on the hardware, try to break it, and only then say
anything. That method does not scale to a fork the size of rock8-release, and if I submitted
work I had not put through it, the first reviewer question I could not answer would cost you more
than my not submitting it. I also could not commit to the maintenance and the follow-up that
being the submitter implies, over a set of changes that large.

What I can do is the thing that has actually worked twice now: when an individual PR or issue
goes up in an area where a gfx1201 measurement is the missing piece, point me at it and I will
look at whether I can take it. I decide case by case and depending on what else is on my plate
— please read this as an open door, not as a queue I have committed to draining. I have one
RX 9070 XT (gfx1201, 16 GB, shared with my desktop), so even when I do take something, it is
bounded by what one card can answer: correctness, ISA behaviour, counter and profiler questions,
single-GPU repro. Multi-GPU work (your two R9700s) is outside what I can verify at all.

On the llama.cpp side specifically: rock8-release and The-Rock8 are llama.cpp-side work, which
I keep separate from the ROCm-stack work this thread is about, and they are frozen for judging
anyway. The one coupling between your tree and the ROCm stack that I happened to look at closely was
__byte_perm — I reviewed clr#283 yesterday. I have not
audited the trees for others, so that is not a claim that it is the only one.

Three technical notes. The first two are about how my #6613 numbers are being used in the message
above; the third is about a line in your opening post, not in that message.

1. bytes = RDREQ x 256 — fine for reads, do not carry it to writes. For the read side your
usage is consistent with what I measured: my fit is count = bytes/256 + 2, so against
RDREQ x 256 you are off by 512 B, which is nothing. The write side has a different intercept:
GL2C_EA_WRREQ_sum fits count = bytes/256 - 24,576 exactly at every point from 16 MiB to
512 MiB, and 24,576 x 256 B is exactly 6 MiB. The slope is the same 256 B; the offset is not.
So if any of your roofline formulas apply the read relation to the write side, the gap is
24,578 requests — the 24,576 plus the read fit's +2 — i.e. 6 MiB + 512 B, not the 6 MiB you
would get from the WRREQ intercept alone. I am only claiming this for RDREQ and WRREQ taken
separately; I have not measured a combined read+write total and would not want to guess how the
two intercepts behave when a kernel does both. All of this is from two kernels (a flat float4
streaming read, and a matching write) and those two counters, on gfx1201 with ROCm 7.2.4 and
7.14.0~pre3. The ranges are not the same on the two sides: the read fit is established on clean
repeats at 16, 32, 64, 128 and 256 MiB, while the write fit holds exactly at every point from
16 MiB to 512 MiB. I would not assume either intercept holds outside its own range.

2. The boundary double-fetch and my excess counts: please do not treat those as agreeing.
I would like that pairing to hold, but my own data cannot carry it, and the reason matters
more than the numbers. My excess is not reproducible run to run. Over 8 repeats of the same
256 MiB read, seven runs sat exactly on the fit at 1,048,578 with all split counters at zero, and
one run came in at 1,050,439. A separate inflated run reached 1,110,127. So the excess I have is
1,861 in one run, ~61,549 in another, and zero in the large majority of runs at the same size.
The 512 MiB figure you are matching against (2,099,030, i.e. 1,876 over fit) is a single run —
I never had a repeat set at that size. That spread is the problem: the match is being read off the
1,861-vs-1,876 closeness, but at the same 256 MiB size I have runs at 0, at 1,861, and at ~61,549,
and at 512 MiB I have exactly one run. So I have no reproducible correspondence between the two
sizes to match against in the first place. On top of that my kernel is a flat float4 read with
no natural row structure to map your "+1 line per row" onto, and I said explicitly in #6613 that I
cannot identify what produces the excess — I could not exclude display, blits, firmware, or the
collection machinery, and I could not prove it was not my own kernel either. Because I cannot
attribute the excess at all, I am not claiming this refutes your mechanism; I am saying my numbers
neither confirm nor refute it, and should not be cited as agreeing with it. They are the wrong
instrument for the question. If you want to test it against my setup, the two things that would
make it checkable are what you map to a "row" in a flat read, and how many rows your model
predicts at 256 MiB and at 512 MiB. A repeat set at 512 MiB on my side would also be needed before
any number there means anything, since what I have is one run.

3. One correction to the opening post: CK does have the gfx12 iu4 intrinsic. In the ISA
inventory bullet you wrote "As far as we can tell nothing upstream (composable_kernel included)
wires the gfx12 16x16x32_iu4 shape yet". You hedged it, and you were right to — the part that is
off is narrow. I have not measured the full-rate result myself, so I am neither confirming nor
questioning that half of the bullet; this is only about the "nothing upstream wires it" clause. But composable_kernel does carry that exact shape. I checked
ROCm/composable_kernel at 8fc1ac24e (on develop). include/ck_tile/core/arch/mma/wmma/wmma_gfx12.hpp
has amdgcn_mma specializations for both gfx12 iu4 shapes — __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12
at line 293 and __builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12 at line 325 — and
test/ck_tile/core/arch/mma/test_amdgcn_mma_layout.inc covers both (lines 559-560), plus the
sparse swmmac iu4 variants at 570-571. The 16x16x32 shape you name is the one at line 325.
So at the intrinsic layer it is present, not absent.

Where your bullet does match what I found is above that layer. On the two paths I looked at, I did
not find anything reaching those intrinsics:

• include/ck_tile/ops/gemm/warp/ — a recursive case-insensitive content grep for int4 and
  iu4 across all 18 headers there (~7.7k lines, including
  warp_gemm_attribute_wmma_impl.hpp and warp_gemm_dispatcher.hpp) returns zero matches.
  The impl traits files in that directory are ..._8bit_traits.hpp, ..._16bit_traits.hpp and
  ..._highprec_traits.hpp, with no 4-bit counterpart. So on this path I found no int4
  warp-GEMM attribute for a dispatcher to reach.
• the older include/ck/tensor_operation/gpu/warp/wmma_gemm.hpp — there is a
  GetWmma<int4_t, int4_t, int, 16, 16>(), but it is behind
  CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4 and returns the non-gfx12
  WmmaInstr::wmma_i32_16x16x16_iu4. Compare the iu8 case right above it, which does branch on
  __gfx125__; the int4 one has no gfx12 branch.

Those two paths and that grep are the whole of what I checked. The grep is a content search, not
a filename one, so within that directory the absence is solid; what I have not done is establish
that this directory is the only way in. The claim I am making is that I found no exposure on the
paths I looked at. I did not audit CK's other entry points and I am not saying
gfx12 iu4 is unreachable in CK, nor guessing how much work any gap would be. I am not planning to
write that wiring myself — but if the shape matters to you, this looks like a reasonable thing to
raise as its own issue against ROCm/composable_kernel, and that is the kind of thing I would look
at (same caveat as above about deciding case by case).

[The-Monk]

@doplxyz — that is the right answer, and the reasoning is one I'd rather have than a yes. Vouching for a change means having built it, measured it, and tried to break it; a submitter who can't answer the first reviewer question does the work more harm than no submitter at all. Understood on all three counts, and the open-door framing is noted as exactly that — I'll point at individual items sparingly, and only where a gfx1201 measurement is genuinely the missing piece, not as a queue. Multi-GPU is ours to carry; I won't send you anything that needs two cards.

Your three technical notes, in order.

1. WRREQ intercept — audited, and you were right to flag it.

I went through the current decode roofline calculator looking for exactly the error you describe. The result is that it never computes write bytes at all: the numerator is executed GEMV weight bytes (per-format traffic factor applied to tensor sizes) plus a mandatory non-GEMV term that is itself entirely reads — GDN state, ssm_conv, flash-attn KV, embedding row, quantize reads. So the read relation is not carried to the write side anywhere in it, and the 24,578-request gap doesn't reach us. That is luck as much as design: the calculator is read-only because decode is read-bound, not because we had your write fit when it was written.

Two things worth stating precisely rather than leaving implied. We use bytes = requests x 256 without your +2, so we are 512 B optimistic on the read side by construction — you already called that nothing, and I agree, but it is a known constant offset rather than an accident. And we do have a calib_write harness built against your bytes/256 - 24,576 model; it has not been run to a fit, so I have no write-side data of my own to offer, only your numbers. If we do run a repeat set I'll post it here rather than in a tree.

I'll also take the range caveat seriously: read fit established 16–256 MiB, write fit 16–512 MiB, neither assumed outside its own range, and no guess about how the two intercepts combine when a kernel does both.

2. The double-fetch pairing — withdrawn, and the problem is on my side too.

You're right that your excess can't carry it. But it's worse than unsupported from where I'm sitting, and you've asked the two questions that make it obvious.

What we map to a "row": the mechanism isn't geometric rows, it's block-size against the 256 B line. The only row-boundary effect our tooling measures is a per-format traffic factor, and it is non-trivial for exactly one format — f8e4m3, at 1.03–1.05, because its 34-byte blocks leave rows unaligned to line boundaries. Q2_0 measures 1.001 and q6_k 1.000, because those align.

How many rows the model predicts at 256 MiB and 512 MiB for your kernel:
zero. A flat float4 stream is 256 B-aligned with no block structure and no
row boundaries, so there is nothing for the mechanism to act on. Our own model
says your kernel should show no excess at all.

So I was matching a number that our model predicts shouldn't exist, against a number yours can't reproduce. It was wrong to cite as agreement and I've stopped citing it — not "unproven pending better data," just wrong. I'm not asking you for a 512 MiB repeat set on that basis; the pairing doesn't deserve one.

3. CK gfx12 iu4 — the correction is accepted, and thank you for the precision.

You're right, and the hedge in my bullet doesn't save it: wmma_gfx12.hpp does carry both shapes, including the 16x16x32 I named, with test coverage and the sparse swmmac variants alongside. I should have grepped the intrinsic layer before writing "nothing upstream wires it." The claim was wrong as stated.

Your narrower finding is the interesting one and matches what I'd expect from the CK side: the intrinsics exist, and nothing above them reaches. No int4 warp-GEMM attribute in ck_tile/ops/gemm/warp/, no 4-bit traits file next to the 8/16/highprec ones, and the older wmma_gemm.hpp int4 path gated behind CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4 returning the non-gfx12 instruction while the iu8 case right above it branches on __gfx125__. That asymmetry looks like the shape of an oversight rather than a decision.

On filing it: I'm happy to open the issue against composable_kernel with the inventory and that asymmetry laid out, referencing your line numbers and commit. Say the word if you'd rather file it yourself — you found it, and I don't want to take the write-up out from under you. Either way I'll leave the "full rate" half of my original bullet alone, since neither of us has measured it.

Thanks for taking the time to be exact about all three. It cost you more effort than a decline needed to.

[doplxyz]

@The-Monk — go ahead and file it yourself. You found the gap and did the write-up; no reason to hand that off. If it turns into something that needs a gfx1201 measurement, point me at it and I'll take a look, same as everything else in this thread.