[StevenChenSE]

Second attempt of #260, rebased on current main (engine 0.1.30) — reworked per the review. Single commit, 10 files.

Addressing the review

review point	what changed
opt-in (STRATA_WMMA_GEMM=1)	done — nothing runs unless asked. STRATA_WMMA_GEMM=1 enables the dense WMMA GEMMs (fp16 + bf16), STRATA_PA_WMMA=1 the WMMA prompt attention; STRATA_WMMA_BF16=0 still excludes just bf16. With the switches unset, WMMA does not run at all — the default hipBLASLt-table-to-fallback path is unchanged and output is bit-identical to today's build. When opted in, WMMA dispatches ahead of the table; measurement below shows that arrangement is also the faster one on this card.
runtime gate on gfx11 only	done — both paths query hipGetDeviceProperties().gcnArchName and return false without a gfx11 prefix, so a gfx1201, a mixed-arch build, or any other device never executes WMMA code regardless of what was compiled. Additionally the STRATA_WMMA_GFX11 compile definition is now set only for gfx11 CMAKE_HIP_ARCHITECTURES configurations.
fp16 rounds toward zero	acknowledged and bounded: the matrix core accumulates in fp32 with RZ while hipBLAS rounds to nearest. With the switches off this is unobservable by construction; with them on, the parity test bounds the deviation at 4 ULP (every measured delta ≤ 1 ULP and toward zero; bf16 compares exactly), documented in docs/AMD_HIP.md and tests/hip/prefill_wmma_gemm_parity.cpp.
CMakeLists.txt BOM / comment encoding	preserved — patched at byte level; BOM re-verified present after the edit (efbbbf), and cmake/hip_backend.cmake (no BOM upstream) likewise untouched in encoding.
with/without numbers	below — including against a freshly calibrated hipBLASLt table, not just the version-mismatched fallback.

What the PR contains

• src/prefill/wmma_gemm.{cu,h} — dense GEMM over the RDNA3 WMMA fragment layout (16x16x16, wave32 "doubled" inputs), fp16 + bf16, dispatched from Gemm::f16/Gemm::bf16.
• src/kernels/cuda/qsa_prompt_attn.cu — the prompt attention on WMMA instead of the ordered FP32 fallback.
• tests/hip/prefill_wmma_gemm_parity.cpp — ctest-registered parity vs a double-precision host reference, 440 shapes × 4 beta/ldy configurations × 2 dtypes (partial tiles, padding, guards, declined shapes); passes on this base (freshly rebuilt from this branch's sources).
• Minimal build wiring: one source in strata_prefill, the arch-gated compile definition, the ctest registration, and the CUDART_INF_F mapping hip_compat needs.
• A "RDNA3 WMMA kernels" section in docs/AMD_HIP.md.

With/without numbers

One RX 7900 XTX, 6-core host, PCIe 4.0 x16, engine 0.1.30, the documented measured configuration (--prefill 8192 --spec 4 --kv int8 --kv-resident 32768 --adapt-every 0 --pcie-frac 0, STRATA_PREFILL_MMQ=1), two interleaved repetitions per cell.

Against the strongest BLAS available here. The shipped tables (100100/100200) don't match this ROCm's hipBLASLt 1.4.1 (100401), so I calibrated a table for this runtime with the repository's own tools/hip/tune_hipblaslt and compared against it (4K, ranking consistent in both pairs):

arm	prefill (tok/s)	decode (tok/s)
WMMA (GEMM=1 PA=1)	1083.7 / 1308.4	32.1 / 41.3
calibrated hipBLASLt table (WMMA off)	852.2 / 1176.6	38.4 / 40.7
neither (hipblasGemmEx fallback)	658.6 / 552.8	37.3 / 37.3

WMMA leads the calibrated table by 10–27 % on prefill in both pairs, and the plain fallback by 55–90 %; decode is a wash at this sample size.

Switch sweep (the "off" arm here is the version-mismatched hipblasGemmEx fallback):

tier	switch state	prefill (tok/s)	decode (tok/s)
1K	off	376 / 455	35.2 / 45.4
1K	WMMA_GEMM=1	471 / 551	39.7 / 51.5
1K	GEMM=1 PA=1	619 / 471	50.1 / 50.3
32K	off	790 / 800	57.3 / 55.8
32K	GEMM=1 PA=1	1501 / 1484	55.1 / 57.1

Prefill: +31 % at the 1K median, +88 % at 32K vs the fallback; decode unchanged within this host's variance. Host and configuration named because they matter: expert offload over PCIe and a 6-core CPU are part of the setup.

Integration notes (without these the kernels compile but never run)

1. this tree's HIP macro is STRATA_USE_HIP — a private STRATA_BACKEND_HIP silently compiled the dispatch out;
2. the host pass of a HIP compile does not define __gfx1100__ (it defines __HIP_DEVICE_COMPILE__), so the build supplies STRATA_WMMA_GFX11;
3. hip_compat needs CUDART_INF_F (CUDA's math_constants.h spelling) mapped to HIP_INF_F.

Not claimed

• Decode gains: decode is unchanged within this host's variance in every comparison above.
• One host, one model pack, two repetitions per cell. With the switches off, behavior and output are unchanged from main.

[Niko1221]

Thanks! Two fixes needed: (1) STRATA_WMMA_GEMM=0 currently turns WMMA ON (getenv(...) != nullptr) - parse the value; (2) STRATA_WMMA_GFX11 is defined for every device pass when the arch list contains gfx11, so a gfx1100;gfx1201 build (setup writes those for mixed cards since 0.1.31) compiles gfx11 WMMA builtins for gfx1201 - guard the device code with __gfx11__/per-arch checks. We'll compile a mixed build on our box.

[StevenChenSE]

Rebased onto current main (0.1.31) and addressed both review points:

1. STRATA_WMMA_GEMM=0 now disables. Both switches are value-parsed (gemm.cu): a bare presence test previously turned WMMA on for any value including 0. Verified on the 7900 XTX: pp 446.6 tok/s with =0 vs 523.1 with =1 at 1K/4K.

2. Per-arch device guards. In a gfx1100;gfx1201 mixed build the TU compiles once per arch, so the RDNA3 v_wmma builtins are now guarded with gfx1100/gfx1101/gfx1102 on the device passes (wmma_gemm.cu mma(), and the four attention builtin call sites). A gfx1201 pass never compiles them; the kernels stay linkable for both arches (identity fallback in the non-gfx11 body, never launched thanks to the existing gcnArchName runtime gate).

Also: on re-testing after the rebase the switch check exposed that carrying upstream 0.1.31's trap-only mma16816 HIP path crashes our non-WMMA attention on HIP (prompt_attn_kernel hits __trap()). The PR keeps its fragment-preserving HIP scalar emulation and grafts upstream's Turing sm_75 (m16n8k8) and sm_80 (m16n8k16) CUDA paths alongside it. Parity: 22/22 on the WMMA test, fast-intrinsics exhaustive PASS.

[Niko1221]

Thanks! Not in 0.1.39. We no longer have an AMD card to test on, and the PR also changes the default HIP prompt attention (the scalar emulation in place of the upstream path), which must stay byte-identical with the switches off. Could you rebase on 0.1.39 and keep the default HIP path untouched? Please also post one run with the switches off showing identical output to 0.1.39 on gfx1100.

[xyzzing]

Measured gfx1100 data for this direction (RX 7900 XTX, engine 0.1.39, nightly ROCm 10.2 SDK + gfx1100 hipBLASLt-100500 table, production-mirrored config) — and a current-main-based implementation is now up as #786.

The FP32 fallback this PR targets is the single largest prefill phase on gfx1100: 25.6% of prefill GPU time at 128K, 26.9% at 32K. Porting the same v_wmma_f32_16x16x16_f16 to RDNA3 under the existing STRATA_HIP_WMMA=1 opt-in:

• parity (qsa_prompt_attn_parity on gfx1100, first time it runs off gfx12): 3/3 PASS, vs FP64 3.3–4.5e-06 (old kernel 1.8–2.1e-06, same class)
• kernel: 34.05 → 8.41 ms per chunk @ 32K×2048q (4.05×), 4.7–5.1× at smaller shapes
• paired e2e (5 pairs × 1k/32k/128k, arms alternating, same binary env-gated): prefill +22.8% median @128k (5/5 pairs), +17.9% @32k (5/5), decode flat (protection cap holds)

Two things we learned that may be useful for this PR's approach if it's ever rebased:

1. The RDNA3 fragment layout is not the gfx12 one. The unsuffixed builtin takes 16-half A/B operands with A[l&15][i], B[i][l&15], D[(l>>4)+2i][l&15] (AMD Matrix Instruction Calculator, rdna3 tables) — gfx12-style packing compiles fine and fails parity at ~30× output scale. Worth noting also that a plain __gfx1200__ || __gfx1201__ compile guard plus a runtime arch-prefix gate silently produced no-ops until the layout was fixed.
2. On the review points above: value-parsing the env and keeping the define out of other device passes are exactly right — hip: RDNA3 (gfx1100) WMMA kernel for the int8-KV QSA prompt attention (opt-in STRATA_HIP_WMMA, gfx12/gfx11) #786 does both, and its arch gate stays runtime-only (the host pass of a HIP TU sees no arch macros).

@StevenChenSE — the STRATA_PA_WMMA/STRATA_WMMA_GEMM split and the runtime gfx11 gating here were the right instinct; #786 covers only the prompt-attention half on current main, where (per Niko1221's note) the dense-GEMM side is partially overtaken by the per-card hipBLASLt tuning tables (our gfx1100 100500 table alone was +82% on 128K prefill).

[Niko1221]

Fixed in 0.1.40: the GEMM half of this PR is in, through our Strix Halo (gfx1151) branch. It is the RDNA3 / RDNA3.5 WMMA FP16 and BF16 dense GEMMs (opt-in STRATA_WMMA_GEMM=1), compiled into AMD builds only. We fixed two review points: an unlisted gfx11 part (gfx1103, gfx1152) now takes the hipBLASLt path, and the device check is cached per device. The attention half is not taken, because our own gfx11 prompt attention covers it (STRATA_HIP_WMMA=1). Thanks to StevenChenSE for the kernels. Closing.