ROCm: run on discrete RDNA4 (gfx1201) — opt-in WMMA fallback + configurable VRAM reserve

[cm999club]

First report (as far as I know) of ds4 running on a discrete RDNA4 GPU: AMD Radeon AI PRO R9700 32GB, gfx1201. Two small opt-in changes were needed; default gfx1151 builds are unchanged. Related discussion: #16.

Problems

1. Build fails on gfx1201: the hand-written Q8 batch GEMM (rocm/ds4_rocm_q8.cuh) calls __builtin_amdgcn_wmma_f32_16x16x16_f16_w32, a gfx11-style WMMA builtin. RDNA4 changed the WMMA intrinsic signatures, so hipcc dies with Cannot select: intrinsic %llvm.amdgcn.wmma.f32.16x16x16.f16. A proper gfx12 WMMA port needs different fragment layouts, so this PR takes the conservative route: an opt-in DS4_ROCM_NO_WMMA guard that routes batch prefill to the existing sharedx kernel.

2. SSD streaming unusable above ~12GB expert cache on a 32GB card: cuda_stream_resident_free_reserve_bytes() hardcodes a 16 GiB free-VRAM reserve. That is reasonable on unified-memory APUs (Strix Halo shares the pool with the OS), but on a discrete 32GB card it wastes half the VRAM: decode fails with streaming expert cache cannot keep 6.75 MiB for layer=N expert=M while preserving 16.00 GiB free. New env var DS4_ROCM_FREE_RESERVE_MB overrides it; default is unchanged.

Plus a make rdna4 convenience target (= strix-halo with ROCM_ARCH=gfx1201 and the guard).

Test environment

• Radeon AI PRO R9700 32GB (gfx1201, wave32), monitor on iGPU so full VRAM free
• Ryzen 9 9950X3D, 128GB DDR5, NVMe SSD
• Ubuntu 26.04, ROCm 7.2.2 (hipcc/AMD clang 22), rocWMMA 2.2.0, hipblas/hipblaslt
• Model: DeepSeek-V4-Flash-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-chat-v2-imatrix.gguf (q2-imatrix, ~81GB)

Commands run

make rdna4                      # new target, builds clean
make rocm                       # default gfx1151 target still builds clean (compile regression check)
./ds4-eval --self-test-extractors   # passed

DS4_ROCM_FREE_RESERVE_MB=2048 ./ds4 --rocm --ssd-streaming \
  --ssd-streaming-cache-experts 22GB --ctx 8192 --nothink -p "..."

ds4_test / ds4_agent_test link against $(NVCC) unconditionally, so I could not run them on a CUDA-less AMD box — that seems to affect every ROCm setup, not just this one.

Results (one-shot CLI, cold expert cache included)

Config	prefill t/s	generation t/s
default reserve (16GiB), 12GB expert cache	3.67	2.30
DS4_ROCM_FREE_RESERVE_MB=2048, 22GB cache	6.24	4.43
DS4_ROCM_FREE_RESERVE_MB=2048, 26GB cache	6.33	4.70

Long-form generations (300+ tokens, Italian and English) are coherent; the server (/v1/chat/completions and /v1/messages) works, including as a local backend for coding agents. Also hit the mgiustiniani VRAM-reporting caveat from #16 indirectly: with other GPU tenants evicted the allocator behaves as expected at 22–26GB cache.

Happy to run more tests or numeric comparisons against baseline if useful. This PR was developed with an AI coding agent driving the terminal, human-approved throughout — in the spirit of the project.

🤖 Generated with Claude Code

https://claude.ai/code/session_01VSVkwdZNYXKAfgFgdzcPDc

[fred-chen]

This I personally feel is a necessary fix, with it I could run a worker on discrete R9700 GPUs and still be able to run other workers on Cuda then the coordinator on Strix Halo, to get a bigger model running on my rig. 👍

A small suggestion is to also implement the UNIX domain socket (UDS) to replace TCP socket for worker-worker, worker-coordinator communication if all workers and the coordinator is on the same host, it could reduce the latency of IPC.