hip: AMD Instinct MI50 / MI60 / Radeon VII (gfx906, wave64) as an opt-in build

[JeanP00l]

What

The HIP backend (STRATA_ENABLE_HIP) is wave32 RDNA only, and the engine refuses a wave64 card. AMD Instinct MI50 / MI60 and Radeon VII (gfx906) are cheap 16/32 GB cards with ~1 TB/s HBM2. This PR adds an opt-in build for them: -DSTRATA_HIP_GFX906=ON -DCMAKE_HIP_ARCHITECTURES=gfx906.

It compiles the CUDA sources as HIP through a small compat layer (include/strata/platform/hip_compat/). A CUDA warp is a logical half of the 64-lane wavefront: 32-wide shuffles, and a ballot of its own half. With the option off, nothing changes for CUDA or the RDNA backend. Every new path is under STRATA_HIP_GFX906, a macro neither existing build defines.

Why a separate option and not STRATA_ENABLE_HIP + gfx906

The RDNA backend's kernels and compat layer assume wave32 throughout. Treating the CUDA warp as half a wavefront let the CUDA path run unchanged on gfx906 first, and then the hot kernels got wave64 layouts of their own. If you'd rather fold gfx906 into STRATA_ENABLE_HIP / cmake/hip_backend.cmake, I'm happy to rework it that way. Tell me which you prefer.

Correctness fixes found on the way (each changed results or hung)

• __fmul_rn / __fadd_rn / __fsub_rn contracted into FMA under HIP. They no longer do, so the elementwise and quantize_act parity tests pass bitwise.
• ggml MMQ got NVIDIA compute capability 9.0 for gfx906 and hung the prompt path. It now gets the AMD arch code.
• HIP's __byte_perm is a scratch-memory array. The IQ4 table lookup now uses llama.cpp's v_perm_b32 sequence: agent-request decode on 2x MI50 went from 29.9 to 39.2 tok/s.
• gfx906's wall clock is 25 MHz, not 100: fixed in the stage profiler.
• The CUDART_VERSION toolkit check is skipped under HIP. gdn_rec_parity (cp.async, CUDA occupancy API) is not built for gfx906.

wave64 kernels (gfx906 only; each has an env switch back to the CUDA layout for A/B)

Each was checked bitwise against the layout it replaces: parity tests, native_expert_bench on real GGUF rows, and the greedy text of a fixed request.

kernel	change	effect	switch
grouped native experts	mode 7: signs as dp4a(g ^ m, u) - dp4a(m, u) (gfx906 has no byte SIMD; __vcmpne4/__vsub4 were ~12 word ops each), grids and activations in LDS, one weight load per group	verify window 60.0 -> 51.3 ms	STRATA_EXP_MODE=2
MMVQ	one wavefront per row, 64-lane butterfly	+2.3% decode	STRATA_MMVQ_WAVE=0
router top-10	one wavefront, register argmax	bit-identical	-
verify-window routing (256-expert router)	one BF16 projection + one top-k per window	-3 ms per window	STRATA_ROUTE_PER_TOKEN=1
gr_down_multi (opt-in)	split along K; sums in another order than the single-token read	+2.9%	STRATA_GR_SPLIT=1
hyper-connection read	loads in flight, 8 lanes per row	51.3 -> 50.2 ms	STRATA_GR_FAST=0

docs/AMD_HIP.md gets a gfx906 section: the build recipe on the community ROCm 7.14 image (AMD dropped gfx906 from current ROCm), the table above, and measurements.

Measured (on the port this is cut from)

2x MI50 16 GB (85 W each, PCIe 3.0 x16), Xeon E5-2666 v3, 32 GB DDR4. Coder IQ1_M, 128K context, --kv int8 --kv-resident 32768, --layer-split 27, --prefill 4096, MTP --spec 4.

• All 12,288 experts are in VRAM.
• A 17K-token agent prompt is read in 43.5 s; decode on it runs at 57.7-59.3 tok/s.
• 6/6 long agent requests and a 900 s two-client load test x6 (407 requests) ran with no stalls or faults.
• For reference, llama.cpp on the same model and machine: 26.9 tok/s tg128.

These numbers also include #639 (layer-split weight trim) and #640 (mapped arena for small RAM), which are separate PRs.

Verified for this branch

• Builds all with STRATA_HIP_GFX906=ON (gfx906, ROCm 7.14).
• Builds all with STRATA_ENABLE_HIP=ON for gfx1100, so the RDNA backend is intact.
• Builds all with CUDA 13.0 for sm_86, so the CUDA path is intact.

Status

Re-run on 2x MI50 with this branch: ctest 35/35 on the cards and the greedy-text comparison with the port branch are in the comment below; builds checked for gfx906, gfx1100 and CUDA sm_86 at the branch head.

🤖 Generated with Claude Code

[JeanP00l]

Re-run on 2x MI50 (gfx906). Running ctest on the cards found four things, now fixed on the branch:

• 6a12d43 native_expert_grouped: the AMD layouts only handle the IQ packs' gate/up and down types. Unsloth's Q4_K / Q5_K gate/up and Q5_1 / Q8_0 down exited ("down type 7"). They now take the CUDA layout (native_grouped_parity).
• 6a12d43 The K-split gr_down broke gr_parity's contract (multi-token read == T single-token calls, bitwise), because its fixed-order finish sums in another order. It is now opt-in (STRATA_GR_SPLIT=1). The latency-hidden norm/up (same sums, same order) stay on by default.
• 6a12d43 Two tests did not know about gfx906:
  • mmvq_multi_parity's negative control has nothing to find, because the wave64 MMVQ is the same exact kernel at every column count. A second entry, mmvq_multi_parity_cuda_layout, runs it on the CUDA layout, where the control has power.
  • kv_hybrid_parity: the tensor-core prompt attention refuses on gfx906, as on RDNA.
• 17beb33 The verify window's routing for the Coder's 256-expert router (one multi-column BF16 projection + one top-k), with a new route_window_parity. gfx906 build only, so CUDA/RDNA are unchanged. -3 ms per window, same text.
• a914b13 Fixes the CUDA/RDNA compile break from 6a12d43, which my re-check of those builds caught.

Builds at a914b13 (all targets, compile only):

• STRATA_HIP_GFX906=ON (gfx906, HIP 7.14)
• STRATA_ENABLE_HIP=ON (gfx1100)
• CUDA 13.0 (sm_86)

ctest on the cards: 35 of 35 pass, including route_window_parity and mmvq_multi_parity_cuda_layout. iq_parity and ple_parity were not run: they need fixture files (tools/iq_fixture.py, bench/micro/ple_*.bin) that the build does not make.

Greedy-text comparison with the port branch. Coder IQ1_M, 32K context, --layer-split 27, --spec 4, --adapt-every 100000 --suffix-draft 0; a greedy 400-token request ("BST") and a greedy 1,500-token code request; md5 prefixes of the generated text.

build	settings	BST	1,500 tokens	ms / verify window
port branch (in daily use)	STRATA_STAGE_TRIM=0 STRATA_GR_SPLIT=0	1938073e3b0b	79449b452413	53.1
this PR	STRATA_IQ_MT_MIN=1	1938073e3b0b	79449b452413	52.3
this PR	defaults	1938073e3b0b	43f4429cdc95 (3 runs, the last at a914b13)	52.0-52.2
port branch	its defaults	1938073e3b0b	5be6772870f2 (3 runs)	50.3

Same settings, same text, bitwise. The port's defaults differ from this PR in three things, each of which changes which expert a GPU or the CPU computes, or the summation order:

• the port branch also has layer split: with explicit split points each GPU loads only its own layers' dense weights #639 (more experts in VRAM);
• the port uses the K-split gr_down by default;
• the port uses STRATA_IQ_MT_MIN=1 by default (upstream's default is 2).

Every run is deterministic: repeats gave the same hashes. The QSA scorer this PR leaves as upstream's (the scalar HIP kernel, not my FP32 variant) gave the same text as mine. With #639 and #640 merged in, the 128K configuration decodes at 61 tok/s (numbers in #639). A community benchmark in the COMMUNITY_BENCHMARKS.md format follows as its own PR.

Marking ready.

[Niko1221]

Merged for the next release, as its own opt-in build (STRATA_HIP_GFX906), with your commits. Changes on top: 0.1.39's segmented expert cache (#533) kept its cuMem* code out of HIP builds through STRATA_USE_HIP, which your build does not define, so expert_cache.cpp included <cuda.h>. Its virtual-memory parts are now off in the gfx906 build too. mrope.hpp's hip_runtime include is limited to the gfx906 build, and tests.list is left out. The gfx906 engine compiles and links for gfx906 with TheRock's ROCm on Windows. The CUDA release build's device code is unchanged, and the identity gate on an RTX 5070 is 10/10 identical on four models. The RDNA Windows build compiles. We have no gfx906 card, so a re-run of your ctest and greedy comparison on the merged branch would be welcome.

[JeanP00l]

Post-merge gfx906 re-run on the current upstream tree (6f32ec0), using the pinned llama.cpp revision from the checkout:

• Build: all gfx906 targets compile with HIP/ROCm 7.14 on MI50. The HIP compatibility wrapper uses the function form of cudaFuncSetAttribute from hip/gfx906: cudaFuncSetAttribute must be a function, not a macro (ROCm 7.2.1 build fix) #808, which handles kernel template arguments containing commas.
• CTest: 36 of 37 tests passed. ple_parity could not complete because its required Q2_0 GGUF fixture is not present. The IQ fixture tests were not registered because Python NumPy is unavailable in the build container.
• Greedy smoke: with the Qwen3.8-Flash-Next Coder IQ1_M pack, explicit layer split 27, MTP 4, 128K context, INT8 KV and 32K resident KV, a fixed temperature-0 request produced identical assistant code on the merge candidate and the existing gfx906 build. The code SHA-256 was 114cf38b24ff3651c135481a4791a9f30597e7c71229d432ce5e2dd3d3956f63 on both.

This is a short smoke request, not the earlier 1,500-token greedy comparison or the long 128K/multi-client acceptance run. The only CTest failure was the unavailable external PLE model fixture; it did not report a parity mismatch.