AMD GPU (ROCm/HIP) and especially Strix Halo discussion group

[darchidev]

Feature Request: AMD GPU (ROCm) Backend Support

Summary

Add support for AMD GPUs via ROCm/HIP, similar to existing Metal backend for Apple Silicon.

Hardware Tested

• AMD Ryzen AI Max+ 395 (16-core/32-thread APU)
• AMD Radeon 8060S Graphics (integrated GPU, 64GB VRAM)
• Ubuntu 26.04 LTS
• ROCm 6.x installed and working (verified with rocminfo)

Current Situation

• ds4 compiles on Linux with `-DDS4_NO_METAL` flag
• CPU backend works but is extremely slow for 284B parameter model
• llama.cpp with HIP works on same hardware (182 t/s on small model)

Expected Behavior

• Run ds4 with `--backend cpu` (or new `--backend gpu`) on AMD GPUs
• Full model support (DeepSeek V4 Flash IQ2/Q4)
• KV cache persistence support

Implementation Suggestion

HIP/ROCm API is similar to Metal. The main components needed:

1. `ds4_hip.h` / `ds4_hip.c` - HIP tensor operations (analogous to ds4_metal.h/m)
2. Convert Metal shaders to HIP/compute kernels
3. Add backend selection in ds4_engine_options

Additional Notes

• GPU detected as gfx1151, 64GB VRAM available
• ROCm stack is stable on this hardware
• Happy to test and provide feedback

[canapaio]

+1 whit 128GB vram
Debian 13
kernel 6.17

[darchidev]

Problem

The current DeepSeek V4 Flash model (~81GB compressed) doesn't fit in consumer AMD GPUs with limited VRAM.

Hardware tested

• AMD Ryzen AI Max+ 395 (Strix Point)
• GPU: Radeon 8060S (62GB VRAM)
• System RAM: 128GB

Implementation

I've started implementing AMD HIP/ROCm support in a fork:

• Fork: https://github.com/darchidev/ds4-hip
• Files added: ds4_hip.c, ds4_hip.h, plus Makefile / ds4.c changes

Current situation

Configuration	Model	VRAM needed	VRAM available	Result
Apple Silicon 128GB	Q2	~81GB (RAM only)	128GB RAM	✅ Works
AMD 62GB VRAM	Q2	~81GB	62GB	❌ Doesn't fit

GPU detection works, but the model weights (81GB) exceed available VRAM (62GB), so it falls back to CPU.

Requests

1. Are there smaller versions of the model? (e.g., 30–40GB that fits in 62GB VRAM)
2. Is it possible to have a Q4 version with more aggressive MoE quantization?
3. Advice on how to proceed with layer-by-layer decompression on GPU?

Any guidance would be very helpful. Thanks!

[ejpir]

I've started work on the HIP kernels for DS4. Currently it works and loads the full model, or you can use the zero-copy that is in Metal as well. Full-copy requires 93GB ram, zero-copy-fast way less as it loads when needed.

Machine	Backend	Quant	Prompt	Prefill	Generation
AMD Radeon 8060S / ROCm HIP	HIP zero-copy fast	Q2_K	1003 tokens	33.21 t/s	7.17 t/s
AMD Radeon 8060S / ROCm HIP	HIP zero-copy fast	Q2_K	1905 tokens	32.61 t/s	8.88 t/s
AMD Radeon 8060S / ROCm HIP	HIP full-copy	Q2_K	1003 tokens	33.72 t/s	9.37 t/s
AMD Radeon 8060S / ROCm HIP	HIP full-copy	Q2_K	1905 tokens	33.46 t/s	9.20 t/s

The next steps are this:

1. rocWMMA on Q2_K routed experts
2. rocWMMA on dense Q8 projections
3. hipBLASLt-quality custom kernels
4. Decode/inference optimization

Right now I'm focussing on pre-fill speed, next steps are:

1. Decoding
2. K/V cache.

Fork is here: https://github.com/ejpir/ds4-hip/tree/hip
Model here: https://huggingface.co/cyberneurova/CyberNeurova-DeepSeek-V4-Flash-abliterated-GGUF

When the next steps are finished, pre-fill and decoding should be faster, but it won't be near what is possible on Metal, as the memory bandwidth is half for AMD Strix Halo.

I'm using ChatGPT5.5 to do the work for me ;)

[puppinoo]

Amazing work.
This sounds desperate but I try: Any chance this will work using older gfx906 gpus and just 64gb or ram (dual 32gb to be precise)?
I must have 2x MI50 32gb + a Machinist x99 MR9A PRO + Xeon E5-2680 V4 somewhere in the closet.

Thanks.

[janimo]

#79 (comment)

[MaxFerAlten]

here my implementation with original Antirez gguf on my beelink Mini PC GTR9 Pro 128gb https://github.com/MaxFerAlten/ds4-rocm

[MaxFerAlten]

rocm antirez branch
beelink mini pc gtr9 pro 128gb
ubuntu 24.04
rocm 7.2.3

ds4: prefill: 26.12 t/s, generation: 10.20 t/s

---

./ds4

env
-u DS4_CUDA_DIRECT_MODEL
-u DS4_CUDA_COPY_MODEL_CHUNKED
-u DS4_CUDA_MODEL_PREFETCH_SYNC
-u DS4_CUDA_WEIGHT_CACHE_LIMIT_GB
-u DS4_CUDA_WEIGHT_ARENA_CHUNK_MB
DS4_CUDA_NO_INDEXER_WMMA=1
DS4_CUDA_NO_DIRECT_IO=1
DS4_CUDA_WEIGHT_CACHE_VERBOSE=1
./ds4
-m /home/tendermachine/.lmstudio/models/antirez/deepseek-v4-gguf/DeepSeek-V4-Flash-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-chat-v2.gguf
--backend cuda
--ctx 2048
--nothink
-p "Rispondi solo: ROCm OK"

[antirez]

If you have well tested and checked improvements for the rocm support, tell me explicitly where, what changes, and so forth, so that I can understand how to act. Thanks!

[MaxFerAlten]

[MaxFerAlten]

To achieve this result, I only made a very small change
in the Makefile and in ./script by running a Python script. The reason is that I don't want to touch
your work on ds4_cuda.cu to avoid creating performance regressions in CUDA.

make clean
make rocm ROCM_ARCH=gfx1151

Makefile:

-GPU_CFLAGS ?= -O3 -fno-finite-math-only -pthread -D__HIP_PLATFORM_AMD__ -Wno-unused-command-line-argument --offload-arch=$(ROCM_ARCH)
+GPU_CFLAGS ?= -O3 -fno-finite-math-only -pthread -D__HIP_PLATFORM_AMD__ -I. -Wno-unused-command-line-argument --offload-arch=$(ROCM_ARCH)
 GPU_LDLIBS = -lm -pthread -L$(ROCM_PATH)/lib -lhipblas
 
 @echo "ROCM_ARCH: $(ROCM_ARCH)"
@@ -177,8 +177,17 @@ ds4_bench_cpu.o: ds4_bench.c ds4.h
 ds4_metal.o: ds4_metal.m ds4_gpu.h $(METAL_SRCS)
        $(CC) $(OBJCFLAGS) -c -o $@ ds4_metal.m
 
-ds4_cuda.o: ds4_cuda.cu ds4_gpu.h ds4_iq2_tables_cuda.inc $(EXTRA_DEPS)
-       $(GPU_CC) $(GPU_CFLAGS) -c -o $@ ds4_cuda.cu
+GPU_CUDA_SRC ?= ds4_cuda.cu
+
+ifeq ($(GPU_BACKEND),rocm)
+GPU_CUDA_SRC := .build/rocm/ds4_cuda.rocm.cu
+endif
+
+.build/rocm/ds4_cuda.rocm.cu: ds4_cuda.cu scripts/gen_ds4_cuda_rocm.py
+       ./scripts/gen_ds4_cuda_rocm.py ds4_cuda.cu .build/rocm/ds4_cuda.rocm.cu
+
+ds4_cuda.o: $(GPU_CUDA_SRC) ds4_gpu.h ds4_iq2_tables_cuda.inc $(EXTRA_DEPS)
+       $(GPU_CC) $(GPU_CFLAGS) -c -o $@ $(GPU_CUDA_SRC)
 
 tests/cuda_long_context_smoke: tests/cuda_long_context_smoke.o ds4_cuda.o
        $(GPU_CC) $(GPU_CFLAGS) -o $@ $^ $(GPU_LDLIBS)

in scripts/gen_ds4_cuda_rocm.py:

`
#!/usr/bin/env python3
from pathlib import Path
import sys

src = Path(sys.argv[1])
dst = Path(sys.argv[2])

text = src.read_text()

old = "const float alpha = rsqrtf((float)head_dim);"
new = "const float alpha = 1.0f / sqrtf((float)head_dim);"

count = text.count(old)

if count != 2:
raise SystemExit(
f"Expected exactly 2 host alpha rsqrtf occurrences, found {count}. "
"Refusing to generate ROCm source."
)

text = text.replace(old, new)

dst.parent.mkdir(parents=True, exist_ok=True)
dst.write_text(text)

print(f"Generated {dst} with {count} ROCm host rsqrtf fixes")
`

[mgiustiniani]

#133 (comment)

[mgiustiniani]

After further investigation, above 16 GB of VRAM, the free and used VRAM values are reported incorrectly, and the application degrades.

[ejpir]

A little update from my side:

Current HIP zero-copy fast path is now roughly ~60 t/s prefill on 2048-token prefill chunks with the CyberNeurova Q2_K GGUF:

Machine	Backend	Quant	ctx	prefill tokens	Prefill	Generation
AMD Radeon 8060S / ROCm HIP	HIP zero-copy fast	Q2_K	2048	2048	~60.4 t/s	~7.0 t/s
AMD Radeon 8060S / ROCm HIP	HIP zero-copy fast	Q2_K	4096	2048	~54.6 t/s	~6.8 t/s
AMD Radeon 8060S / ROCm HIP	HIP zero-copy fast	Q2_K	6144	2048	~53.6 t/s	~6.7 t/s
AMD Radeon 8060S / ROCm HIP	HIP zero-copy fast	Q2_K	8192	2048	~53.1 t/s	~6.6 t/s

ROCm / HIP stack:

  ROCm package:   7.2.3
  HIP compiler:   HIP 7.2.53211-364a905
  AMD clang:      22.0.0git
  ROCm path:      /opt/rocm
  GPU arch:       gfx1151
  GPU:            AMD Radeon 8060S Graphics
  Wavefront:      32

ROCm packages:

  rocm-core       7.2.3-1.1
  hip-runtime-amd 7.2.3-1.1
  hipblas         7.2.3-1.1
  hipblaslt       7.2.3-1
  rocblas         7.2.3-1
  rocwmma header  2.1.0
  composable-kernel 7.2.3-1

There is also a new experimental opt-in MoE-down Q8_K path:

DS4_HIP_MOE_Q8K_DOWN=1

Machine	Backend	Quant	ctx	prefill tokens	Prefill	Generation
AMD Radeon 8060S / ROCm HIP	HIP zero-copy + Q8_K MoE-down opt-in	Q2_K	2048	2048	~61.5 t/s	~7.2 t/s
AMD Radeon 8060S / ROCm HIP	HIP zero-copy + Q8_K MoE-down opt-in	Q2_K	4096	2048	~61.5 t/s	~7.0 t/s
AMD Radeon 8060S / ROCm HIP	HIP zero-copy + Q8_K MoE-down opt-in	Q2_K	6144	2048	~56.2 t/s	~6.7 t/s
AMD Radeon 8060S / ROCm HIP	HIP zero-copy + Q8_K MoE-down opt-in	Q2_K	8192	2048	~53.3 t/s	~6.6 t/s

@puppinoo Memory wise it should work to load, but I haven't asked ChatGPT about what it would entail to support gfx906 yet. It'll be hard for me to verify as I don't have access to the hardware myself, but I'll have a look anyway when I'm done with pre-fill/decoding.

As a side note; I'm migrating the work from my branch into the same shape as the rocm branch here, so at some point both GGUFs can be supported.

[puppinoo]

Thanks. Much appeciated.

[NathZuckerman011]

If you have well tested and checked improvements for the rocm support, tell me explicitly where, what changes, and so forth, so that I can understand how to act. Thanks!

Hi Salvatore, I can't gift you one but if you need it I can give you access to my Beelink GTR9 Pro 24/7:

• 128GB Ram (up to 124 GB available for the GPU)
• Fedora Linux 43
• Kernel: 6.19.14-200.fc43.x86_64
• ROCm 6.4.4