Please ship official gfx90a builds of the fmha_v3_fwd ASM kernels (48 of 56 are assembler-verified portable)

[davetha]

Related: #4525 — a small PR adding gfx90a to GFX_CU_NUM_MAP, which today makes any explicit GPU_ARCHS=gfx90a build fail before a kernel is generated. Separate and independent of this request, but it is the other half of getting AITER usable on CDNA2, and it is worth landing regardless of what happens here.

The ask

Please ship official gfx90a builds of the fmha_v3_fwd ASM kernels, alongside the existing gfx942/gfx950 ones.

MI210 and MI250 are widely available — ex-datacenter cards are now common in homelabs, universities and small research groups, and MI250 still fills large HPC installations. Today those cards have no vendor-supported fast attention path: AITER ships the kernels, gfx90a is excluded by an architecture-string comparison, and the fallback is PyTorch SDPA.

An official, AMD-validated gfx90a build would close that for a large installed base. Everything below is evidence that it is worth doing and that the kernels will build.

Why it looks like an oversight rather than a hardware limit

aiter/ops/mha.py describes these as "hand-written gfx9 ASM". gfx90a is gfx9. The gate is a literal architecture-string comparison against gfx942/gfx950, and the failure is silent — the fast path simply never engages, with no warning that a faster kernel exists and was skipped.

The work behind this request

This comes out of a documented optimization effort on 2× MI210 — ~56 technical write-ups, every performance claim backed by an A/B with an assertion proving the code path actually ran. The parts relevant here:

1. Instruction-level portability analysis of the whole hsa/gfx942 tree. Every code object disassembled, every instruction classified as portable, renameable, or architecturally absent. Not a spot check.

2. An assembler-proof repatch tool. Disassembles each instruction, applies the gfx942→gfx90a mnemonic differences, and re-assembles for gfx90a. An instruction is accepted only if the gfx90a encoding is the same length; if anything fails to assemble the kernel is reported NOT PORTABLE and no file is written.

This exists because the naive approach is actively dangerous. An earlier hand-guessed byte patch substituted D3E1 → D3CD — bf16 MFMA for f16 MFMA. The kernel ran at full speed and computed the wrong thing for weeks before anyone checked. Re-assembling through the real assembler makes that class of error impossible: a wrong mnemonic either fails to assemble or changes encoding length.

3. Correctness before timing. Every benchmark number below came from a backend that passed a correctness check immediately before being timed, for exactly the reason above.

4. End-to-end integration and measurement under vLLM, not just microbenchmarks — including the vLLM-side work needed to reach AITER on gfx90a at all, since its dispatch gate calls on_mi3xx() while documenting itself as gfx9.

5. A negative finding published alongside the positive one. Of 242 kernels in the ASM tree, only ~48 turned out to matter for our workloads. That is recorded as prominently as the win.

Portability result

The substitutions are all renames of the same operation across ISA versions:

"v_mfma_f32_16x16x16_bf16" -> "v_mfma_f32_16x16x16bf16_1k"
"v_mfma_f32_32x32x8_bf16"  -> "v_mfma_f32_32x32x8bf16_1k"
"v_mfma_f32_4x4x4_bf16"    -> "v_mfma_f32_4x4x4bf16_1k"
"v_mfma_f32_16x16x16_f16"  -> "v_mfma_f32_16x16x16f16"
"v_mfma_f32_32x32x8_f16"   -> "v_mfma_f32_32x32x8f16"
"v_mfma_f32_4x4x4_f16"     -> "v_mfma_f32_4x4x4f16"

Anything matching _(fp8|bf8)_ | _xf32 | smfmac | _i8 has no gfx90a equivalent and disqualifies the kernel.

48 of 56 fmha_v3_fwd kernels — every bf16 one — assemble cleanly for gfx90a. The 8 that do not are the FP8 kernels, correctly so: CDNA2 has no FP8 ALU.

How the verification works

The substitution table alone would not be trustworthy — a table can be incomplete. Two further checks carry the weight. Core of convert():

    dis = run([OBJDUMP, "-d", "--mcpu=gfx942", src]).stdout
    # ... parse each line into (text, addr, encoding words)

    todo = []
    for txt, addr, words in insns:
        mn = txt.split()[0]
        if NOT_PORTABLE.search(mn):
            return "NOTPORT", f"gfx942-only op: {mn}"
        if mn in SUBST:
            if SUBST[mn] is None:
                return "NOTPORT", f"no gfx90a equivalent: {mn}"
            todo.append((txt, addr, words, SUBST[mn]))

    # (1) verify EVERY instruction is valid gfx90a -- catches gfx942-only ops
    #     that are NOT in the substitution table
    probe = "\n".join(substituted(t) for t, _, _ in insns)
    if assemble(probe) is None:
        return "NOTPORT", f"does not assemble for gfx90a: {bad}"

    encs = assemble("\n".join(newtxts))

    for (txt, addr, words, newmn), enc in zip(todo, encs):
        old = b"".join(bytes.fromhex(w.zfill(8))[::-1] for w in words)
        # (2) refuse any substitution that changes encoding length
        if len(enc) != len(old):
            return "NOTPORT", f"size change on {newmn} @ {addr:#x}"
        off = foff + (addr - vaddr)
        # (3) confirm the bytes on disk are what the disassembler claimed
        if bytes(data[off:off + len(old)]) != old:
            return "ERROR", f"byte mismatch at {addr:#x}"
        data[off:off + len(enc)] = enc

    struct.pack_into("<I", data, EFLAGS_OFF, (flags & ~0xFF) | GFX90A)

(1) is the important one. The whole kernel — every instruction, substituted where applicable — is assembled for gfx90a before anything is written. So an unknown gfx942-only opcode that nobody thought to put in NOT_PORTABLE still fails the kernel. The table is an optimization; the assembler is the authority.

(3) guards against a disassembly/offset mismatch: the bytes about to be overwritten must equal what llvm-objdump reported at that address, or the file is rejected rather than corrupted.

Anything not reported OK is never written. hsa/gfx90a/ is treated as generated — rebuilt with the tool, never hand-edited.

One integration detail worth mentioning, since it is easy to miss: the kernel manifest has to be pruned in step with the blobs. The loader hard-fails on a missing code object (AITER_CHECK(file.is_open(), ...) in aiter_hip_common.h) and selection picks by shape from the CSV-derived table without checking the file exists — so a manifest that outlives its kernels turns an unsupported shape into a crash instead of a fallback.

Full tool: configs/repatch_gfx942_to_gfx90a.py (249 lines, MIT).

A full run, top to bottom

Real output, produced today against the aiter_meta/hsa/gfx942 tree shipped in a ROCm 7.14 container on an MI210 box.

$ python3 repatch_gfx942_to_gfx90a.py \
      /opt/python/.../aiter_meta/hsa/gfx942  /tmp/out  fmha_v3_fwd/

  NOTPORT  fmha_v3_fwd/MI308/fwd_hd128_fp8.co: gfx942-only op: v_mfma_f32_32x32x16_fp8_fp8
  NOTPORT  fmha_v3_fwd/MI308/fwd_hd128_fp8_causal.co: gfx942-only op: v_mfma_f32_32x32x16_fp8_fp8
  NOTPORT  fmha_v3_fwd/MI308/fwd_hd128_fp8_causal_group.co: gfx942-only op: v_mfma_f32_32x32x16_fp8_fp8
  NOTPORT  fmha_v3_fwd/MI308/fwd_hd128_fp8_group.co: gfx942-only op: v_mfma_f32_32x32x16_fp8_fp8
  NOTPORT  fmha_v3_fwd/MI300/fwd_hd128_fp8.co: gfx942-only op: v_mfma_f32_32x32x16_fp8_fp8
  NOTPORT  fmha_v3_fwd/MI300/fwd_hd128_fp8_causal.co: gfx942-only op: v_mfma_f32_32x32x16_fp8_fp8
  NOTPORT  fmha_v3_fwd/MI300/fwd_hd128_fp8_causal_group.co: gfx942-only op: v_mfma_f32_32x32x16_fp8_fp8
  NOTPORT  fmha_v3_fwd/MI300/fwd_hd128_fp8_group.co: gfx942-only op: v_mfma_f32_32x32x16_fp8_fp8
  PRUNED   fmha_v3_fwd/fmha_fwd.csv: dropped 4 of 28 rows (kernel not portable)

TALLY: {'OK': 48, 'NOTPORT': 8}
manifests: 24 rows kept, 4 dangling rows dropped

48 written, 8 refused. Every refusal is the same FP8 MFMA, split across the MI300/ and MI308/ product subdirectories. Nothing else in the tree tripped a check.

Verifying one of the 48 that came out:

$ readelf -h  <input>/fmha_v3_fwd/MI308/fwd_hd192x128_bf16_causal_rtz.co | grep Flags
  Flags:  0x54c, <unknown AMDGPU GPU type: 0x4c>, xnack any, sramecc any

$ readelf -h /tmp/out/fmha_v3_fwd/MI308/fwd_hd192x128_bf16_causal_rtz.co | grep Flags
  Flags:  0x53f, gfx90a, xnack any, sramecc any

and it disassembles as gfx90a, with the MFMA renamed and nothing else changed:

$ llvm-objdump -d --mcpu=gfx90a  /tmp/out/.../fwd_hd192x128_bf16_causal_rtz.co
  216 mfma instructions, all:  v_mfma_f32_32x32x8bf16_1k

$ llvm-objdump -d --mcpu=gfx942 <input>/.../fwd_hd192x128_bf16_causal_rtz.co
  v_mfma_f32_32x32x8_bf16

Same operation, two spellings across ISA versions — which is the whole basis of the claim, and is now visible rather than asserted.

Note the manifest line. fmha_fwd.csv lost the 4 rows whose kernels were refused, because the loader hard-fails on a missing code object and selection picks by shape without checking the file exists. A manifest that outlives its kernels turns an unsupported shape into a crash instead of a fallback.

What it is worth

Kernels in isolation:

path	vs alternative	peak
Prefill fmha_v3_fwd vs PyTorch SDPA	1.13–1.86×	89.9 TFLOP/s (50% of bf16 peak)
Decode pa_fwd_asm vs HIP kernel	0.99–1.72×	>1 TB/s (64% of HBM2e peak)

End to end under vLLM, Qwen3-14B bf16, attention backend the only variable:

prompt	conc 1	conc 8	conc 32
128 tokens	1.02×	1.00×	1.02×
4096 tokens	1.01×	1.23×	1.23×

Stated honestly: 1.23× serving throughput on long prompts under concurrency, and nothing on short prompts or single streams. Of that gain the ASM decode kernel accounts for ~1%, inside run-to-run noise — a 1.72× kernel proved indistinguishable from zero once the surrounding GEMMs and scheduling are included. The gain comes from the prefill path.

I would rather present it that way than quote 1.86× alone, because the kernel number is not what a user gets.

Scope — what is not portable

• The 8 FP8 fmha_v3_fwd kernels. CDNA2 has no FP8 ALU. Not portable, should stay gated.
• The fmoe tree. Separate finding, genuinely blocked: those kernels need global_atomic_pk_add_bf16, which gfx90a does not have at all. A rename cannot fix a missing instruction, and substituting global_atomic_pk_add_f16 would silently change the accumulation dtype. Mentioned so this is not read as "port the whole ASM tree" — it is specifically fmha_v3_fwd.

Why a binary repatcher is not the answer

I can generate working gfx90a code objects locally, and I do. But a repatched binary is a workaround: it is regenerated per AITER release, it is validated by nobody who owns the kernels, and it turns hsa/gfx90a/ into a build artifact users are tempted to hand-edit.

Building the 48 from source, in-tree, with AMD's validation behind them, is the durable fix — and the analysis above is offered as evidence the build will succeed, not as a substitute for it.

Everything is already public — take any of it

MIT licensed, no attribution needed, use it however is useful. Rather than offer to send things on request:

what	file
The repatch tool — disassemble, substitute, re-assemble, verify, write	configs/repatch_gfx942_to_gfx90a.py
Per-kernel portability classifier — why each kernel can or cannot run on gfx90a	configs/classify_gfx942_kernels.py
The five architecture-string gates that block the path once blobs exist	configs/enable_gfx90a_asm_paths.py
ELF/code-object inspector used to check e_flags and .text layout	configs/analyze_co_elf.py
MFMA scanner — which kernels in a fat binary actually use MFMA on gfx90a	configs/scan_fatbin_mfma.py
Kernel benchmarks + correctness harness (every number checked before timing)	benchmarks/asm-attention-gfx90a.md
End-to-end vLLM serving results, full method and raw numbers	benchmarks/vllm-aiter-asm-gfx90a.md
Instruction-level analysis of the fmoe tree — why that one is genuinely blocked	docs/49-the-fmoe-asm-tree-at-instruction-level.md

The last one is included deliberately. It documents a case where an earlier conclusion of mine was wrong about the mechanism, and was corrected. If you are weighing how much to trust the portability analysis above, that is the more useful document to read first.

I am also glad to run any build you produce on real MI210 hardware and report back with the same harness.

---

The request, plainly

Could you just release a gfx90a build of these kernels?

That is the whole ask. Not a code change, not a design discussion — a build target added to whatever produces the gfx942 and gfx950 blobs today, for the 48 bf16 fmha_v3_fwd kernels.

Everything above exists to answer the questions I would expect in response:

likely question	answer
Will they even build for gfx90a?	48 of 56 assemble cleanly; the 8 that do not are FP8 and should stay gated
Is it worth the build/validation time?	1.23× end-to-end serving throughput on long prompts under concurrency; 1.13–1.86× at kernel level
Is anyone actually on this hardware?	MI210/MI250 are widely available second-hand and MI250 is still in large HPC installs; none of them have a supported fast attention path today
Who will test it?	I will, on real MI210s, and report back with the same harness that produced the numbers above

I would much rather delete my repatcher and use an official kernel. A vendor-built, vendor-validated blob is better than anything I can generate downstream — it gets your correctness testing, it tracks your releases, and it does not silently rot when the ASM changes.

If a full gfx90a build target is more than you want to take on, even a statement of whether it is intended to be supported would help — right now the gate reads as an oversight rather than a decision, and downstream projects have no way to tell which it is.

[IMbackK]

same thing would work with gfx908 too, just also skipping bf16 kernels. gfx90a and gfx908 are almost 100% isa identical with very few differences, one of these differences is that it can process only half the size bf16 mfma tiles so if the larger v_mfma_fp32_*bf16_1k is used it wont work.

[davetha]

same thing would work with gfx908 too, just also skipping bf16 kernels. gfx90a and gfx908 are almost 100% isa identical with very few differences, one of these differences is that it can process only half the size bf16 mfma tiles so if the larger v_mfma_fp32_*bf16_1k is used it wont work.

I'd be happy make my code create binary patches for gfx908. I'm hoping the Aiter project will take this on.