[Experimental] Add ROCm/HIP support (100% vibe‑coded, for testing purposes)#48
[Experimental] Add ROCm/HIP support (100% vibe‑coded, for testing purposes)#48IIIIIllllIIIIIlllll wants to merge 4 commits into
Conversation
Enable ENGINE_ENABLE_HIP=ON to build with HIP/ROCm for AMD GPUs. - New BackendType::Hip in module.h - init_backend, backend_type, release_backend_graph_resources, query_backend_memory all support HIP - Reuses GGML's existing ggml-hip with vendor header mapping CUDA→HIP APIs - GGML_USE_CUDA is still defined for HIP (zero new kernel code needed) - Runtime distinguishes HIP devices via 'ROCm' name prefix - Skip audio.cpp's own .cu files (ISTFT, TorchRandom) on HIP builds - --backend hip in CLI help and server config
|
@IIIIIllllIIIIIlllll Sorry I missed your new updates. Great work! I've seen many Reddit users asking about running these models on ROCm, so your PR is really important for closing a major gap. I'm curious, though, can't the Python implementations run on ROCm? The assumption that “if it works for llama.cpp, it should work for audio.cpp” is generally correct. However, the audio-model ecosystem is much messier, while LLM architectures are relatively standardized. Some GGML operations used by audio models may be supported on CUDA but unavailable on Vulkan, Metal, or other backends. In some cases, I have to modify GGML to add the missing backend support. These issues often only appear at runtime. |
|
@IIIIIllllIIIIIlllll I called for help on reddit and hopefully more people can help with testing the model behaviors on ROCm. |
AMD GPUs have many compatibility issues and are not as simple and easy to use as CUDA. Therefore, your project is of great value to AMD users. |
|
@IIIIIllllIIIIIlllll Just to let you know, I’m currently upgrading the vendored GGML from v0.12 to upstream v0.16. It brings a decent performance boost for some models and better support for other backends. Another user asked about ROCm support today. Looking forward to your update and please keep me posted! |
- ggml-hip: route all cuBLAS-equivalent GEMM through hipBLASLt
(new GGML_HIP_HIPBLASLT option, default ON, falls back to hipBLAS
when the package is missing). hipBLASLt ships kernels for arches
rocBLAS lacks on Windows (notably gfx1103 / Radeon 780M). Handles
legacy hipblasDatatype_t vs hipDataType numbering on ROCm < 6.5,
6.x and 7.x. Pointer-array batched GEMM is emulated per batch
element. Verified against CPU reference on gfx1103/ROCm 6.4.
- Forward ENGINE_ENABLE_CUDA_GRAPHS to GGML_HIP_GRAPHS so graphs can
be disabled on memory-constrained iGPUs (per-graph VRAM reserves
caused OOM on the second request).
- Add scripts/build_windows_hip.ps1: one-shot Windows HIP build
(auto-detects ROCm, GPU targets, cmake, ninja; graphs OFF and
NO_VMM ON by default).
- Add Hip case to ace_step planner backend switch.
- Fix moss_tts_local parity tests for the TensorSource constructor
API (codec_{encode,decode,dequant}_parity build again).
- Update docs/HIP.md: hipBLASLt path, support matrix, Windows build
instructions, verification checklist.
|
Today I gave the popular Kimi K3 model a try, asking it to help me continue optimizing the ROCm backend. The process was painfully slow and tormented me for quite a while. Anyway, the ROCm backend can now be compiled on Windows — my 780M iGPU can run Index-TT2 without issues, and everything works fine on Linux as well. More testing… to be continued tomorrow. |
Test Summary – ROCm/HIP Backend (Qwen3-TTS)
GPU:AMD Radeon Graphics (gfx1151, Strix Halo), 120 GiB VRAM
ROCm:7.2.4
Build:cmake -S . -B build_hip -DENGINE_ENABLE_HIP=ON -DGPU_TARGETS=gfx1151 -DCMAKE_C_COMPILER="$(hipconfig -l)/clang" -DCMAKE_CXX_COMPILER="$(hipconfig -l)/clang++" -DCMAKE_BUILD_TYPE=Release
Model:Qwen3-TTS-12Hz-1.7B-CustomVoice
Input:
Generation time:9.8 s → RTF = 0.89× (faster than real‑time)
Output:
11 s audio file
test-rocm.wav
Detection:ggml_cuda_init correctly identifies ROCm device, backend_type() returns Hip, memory queries accurate
About this PR:
I noticed that llama.cpp already has ROCm support via its ggml-hip backend, and the key insight is that HIP maps almost all CUDA APIs to HIP equivalents at compile time through a vendor header. Since audio.cpp reuses GGML's CUDA backend code, I thought: “If it works for llama.cpp, there’s no reason it shouldn't work for audio.cpp – there might be tricky bugs, but at least I should prove the concept.”
To get started, I took the ROCm‑related CMakeLists.txt from llama.cpp and fed it to DeepSeek V4 PRO (an LLM), asking it to adapt the build logic for this codebase. The result was a clean CMake configuration that successfully compiled and ran on Linux with the HIP backend enabled. This is therefore a 100% vibe‑coded proof‑of‑concept – it works, but please treat it as a reference/test implementation rather than a fully polished production feature.
I'm not English speaker, so please forgive any unclear phrasing. I'm happy to clarify or adjust anything.