Repository navigation
AMD Instinct MI50 / MI60 / Radeon VII (gfx906, wave64): a working port, numbers, and PRs to upstream it #641
Description
I've been running Strata on 2x AMD Instinct MI50 16 GB (gfx906, wave64) daily since 29.09, as the backend for a coding agent: Coder IQ1_M, 128K context, layer split across the two cards. It runs at 57-59 tok/s decode on a 17K-token agent prompt. For comparison, llama.cpp on the same model and machine does 26.9 tok/s tg128.
Each upstream release currently has to be merged into my port by hand, so I'd like to upstream it. I split it into small PRs:
- hip: AMD Instinct MI50 / MI60 / Radeon VII (gfx906, wave64) as an opt-in build #638 (draft): gfx906 as an opt-in build (
STRATA_HIP_GFX906=ON). It adds a compat layer, correctness fixes that changed results or hung (FMA contraction of__fmul_rn, ggml MMQ getting cc 9.0,__byte_permin scratch memory, the 25 MHz wall clock) and wave64 kernels. Nothing changes with the option off. It builds for CUDA sm_86, HIP gfx1100 and HIP gfx906. - layer split: with explicit split points each GPU loads only its own layers' dense weights #639 (draft): with explicit layer-split points, each GPU loads only its own layers' dense weights. +1,800 experts in VRAM on 2x16 GB. Not AMD-specific.
- STRATA_ARENA_MMAP=1: the expert arena as a read-only mapped file (Linux), for small-RAM machines #640 (draft):
STRATA_ARENA_MMAP=1, the native-pack expert arena as a read-only mapped file, for machines where the GPUs hold most experts but RAM is small. RAM available on 32 GB: ~1 GB -> 25 GB. It also covers a ROCclr pin-in-place behaviour worth knowing about. - serve: retry an engine start that fails, and keep the last known context after a failed restart #637: server fix. A dead engine was restarted next to its still-alive predecessor, ran out of VRAM, and left the server answering
400 ... context (0)forever.
Questions for the maintainers:
- Build option. Would you keep gfx906 as its own build option, or should it be folded into
STRATA_ENABLE_HIP/cmake/hip_backend.cmake(wave64 next to wave32)? I can do the second if you prefer. The current shape came from getting the CUDA path running unchanged first, with a CUDA warp as half a wavefront. - Setup. setup.py detection and an automatic build for gfx906 are not included. AMD dropped gfx906 from current ROCm, so the build uses a community ROCm 7.14 image. Is a "build from source" section in
docs/AMD_HIP.mdenough for now?
Also available, not submitted: a tensor split (each layer's halves on two cards, an own P2P all-reduce at 7.4-7.9 µs per exchange on gfx906). It works, but is not faster than the layer split on these cards for long prompts.
Next: the drafts become ready once the cleaned branches have re-run on the cards (ctest, a greedy-text comparison with the port), and I'll add a community benchmark in the docs/COMMUNITY_BENCHMARKS.md format.
🤖 Generated with Claude Code
Activity
Niko1221 commentedon Oct 3, 2026
Thanks, this is a well-prepared port and the numbers are impressive. Answers to your questions:
STRATA_HIP_GFX906) for now. It leaves the CUDA and RDNA builds untouched and is easier to review; we can fold it intohip_backend.cmakelater once it has settled.docs/AMD_HIP.mdis enough. No detection or automatic build in setup for now.We'll review in this order: #637 first (server only, with tests; please rebase it on 0.1.38, where
restart()already ends and waits for the old process, so the retry and the kept context are the parts that add to it), then #639 and #640 once you've re-run them on the cards, then #638 when it's out of draft with the ctest results and the greedy-text comparison. We have no gfx906 card, so for #638 we'll rely on your runs and check that the CUDA and RDNA builds and their results stay the same.JeanP00l commentedon Oct 3, 2026
Thanks! I am keeping
STRATA_HIP_GFX906as its own option and the docs as a "build from source" section, without setup changes.zqdeepblueskyzq commentedon Oct 3, 2026
谢谢你做出的贡献，我现在也在尝试用Radeon VII*2运行 strata，刚开始着手做这件事情，加油
JeanP00l commentedon Oct 3, 2026
Update: all four PRs are re-run on the cards and out of draft.
gr_down, now opt-in); both are fixed. CUDA sm_86 and RDNA gfx1100 builds pass at the branch head.COMMUNITY_BENCHMARKS.mdformat. Decode 50.1 / 47.8 / 45.7 tok/s at 4K / 32K / 128K, prompt about 520 tok/s, 6/6 needles.Niko1221 commentedon Oct 4, 2026
#638 and #677 are in for the next release, after #639 and #640 in 0.1.39. gfx906 stays its own opt-in build, documented in docs/AMD_HIP.md and the new docs/OLDER_GPUS.md (the support matrix for older cards). Closing this; new findings on the cards are welcome as issues or PRs.
JeanP00l commentedon Oct 4, 2026
Post-merge validation update on current upstream
main(6f32ec0) with the gfx906 port:cudaFuncSetAttributecompatibility fix proposed in hip/gfx906: cudaFuncSetAttribute must be a function, not a macro (ROCm 7.2.1 build fix) #808.ple_parityneeds a Q2_0 GGUF fixture that was not available; IQ fixture tests were not registered because NumPy was unavailable.STRATA_ARENA_MMAP=1, returned identical generated code on the candidate and existing gfx906 build (SHA-256114cf38b24ff3651c135481a4791a9f30597e7c71229d432ce5e2dd3d3956f63). This is a smoke check, not the original long greedy comparison or the multi-client 128K acceptance run.The old tensor-split experiment was not included in the active build targets: its implementation no longer compiles against current upstream APIs, and it was not part of the validated gfx906 serving configuration. The existing checkout and deployed build were left unchanged.