Feat/rdna4 gfx1201 rocwmma v2

[bestbug456]

Hello (again!), while I was looking around I found out #599 and #521 . Both of them where reporting no support for WMMA v2 introduced with the new RDNA4. Looking around I found the documentation on it at https://gpuopen.com/learn/wmma-guide-amd-rdna-4-gpus-part-2/ and decided to give it a try (of course assisted by an LLM), with this change I was able to reach a whopping 0.3 tkn/s in prefill and generation. Down a more structured PR message

This PR adds a WMMA v2 fused double-K matmul kernel for RDNA4 (gfx1200/1201/1202) and improves the ROCm build system.

Changes

GPU Kernel

• Wide-K WMMA v2 kernel (matmul_q8_0_f32_batch_wmma_v2_4w_kernel) — fuses two consecutive 16-K WMMA mma_sync calls per loop iteration so each thread fetches 32 K-elements, saturating the 128-bit memory bus on RDNA4.
• Dispatch guards — runtime detection (g_wmma_v1, g_is_gfx12) selects the correct kernel: WMMA v1 builtins on RDNA3/3.5 (gfx11xx), WMMA v2 rocWMMA on RDNA4 (gfx12xx).
• Numerical correctness guaranteed by associativity: A[0:32]B[0:32] + A[32:64]B[32:64] is identical to two separate WMMA accumulates.

Build System

• Auto-detect ROCm GPU arch via rocminfo — make rocm now automatically detects RDNA3 or RDNA4 and compiles with the correct --offload-arch. Falls back to gfx1151 if rocminfo is unavailable.
• make rx9070 target for explicit gfx1201 builds.
• make strix-halo unchanged — still targets gfx1151 explicitly.
• Reverted ds4.c changes from this branch to match main.

Tests

• tests/test_rocm_make.sh — 20 smoke tests covering:
  • make help lists all ROCm targets
  • rocminfo detection works
  • make rocm and make rx9070 build successfully
  • All 5 binaries (ds4, ds4-server, ds4-bench, ds4-eval, ds4-agent) are produced and executable
  • Binaries are linked against libhipblas (not cublas)
  • Auto-detected arch is used in the build command

Files changed

• rocm/ds4_rocm_matmul.cuh — WMMA v2 dispatch guards
• rocm/ds4_rocm_q8.cuh — new wide-K WMMA v2 kernel
• rocm/ds4_rocm_runtime.cuh — g_wmma_v1 / g_is_gfx12 runtime detection
• Makefile — auto-detect via rocminfo, rx9070 target, help text
• tests/test_rocm_make.sh — new smoke test suite
• ds4.c — reverted to main

Testing

$ make clean && make rocm # builds successfully on gfx1201
$ ./tests/test_rocm_make.sh # PASS=20 FAIL=0