[Experimental] Add ROCm/HIP support (100% vibe‑coded, for testing purposes)#48
Draft
IIIIIllllIIIIIlllll wants to merge 10 commits into
Draft
[Experimental] Add ROCm/HIP support (100% vibe‑coded, for testing purposes)#48IIIIIllllIIIIIlllll wants to merge 10 commits into
IIIIIllllIIIIIlllll wants to merge 10 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
Owner
|
@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. |
0xShug0
added
enhancement
Owner
|
@IIIIIllllIIIIIlllll I called for help on reddit and hopefully more people can help with testing the model behaviors on ROCm. |
Author
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. |
Owner
|
@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.
Author
|
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. |
Author
|
ROCm-7.14 Test for MOSS-TTS-Local-Transformer-v1.5 Time consuming: Result: Test for Index-TTS2 Time consuming: Result: Test for Qwen3-TTS-12Hz-1.7B-CustomVoice Time consuming: Result: Test for Qwen3-ASR-0.6B & Qwen3-ASR-1.7B Input: qwen3-tts-custom-voice.wav Command: ./audiocpp_cli --task asr --family qwen3_asr --model /home/mark/App/models/Qwen3-ASR-1.7B --backend hip --audio /home/mark/qwen3-tts-custom-voice.wav --text "" --text-out transcript.txt --log Qwen3-ASR-0.6B: ~12s audio transcribed in 1.17s |
Owner
|
@IIIIIllllIIIIIlllll Thank you for testing! IndexTTS2 looks very slow on ROCm. I’m starting to wonder whether the CPU might actually be faster...Let me think about how to debug this without ROCm machine. I’m ready to merge your PR whenever you’re ready. You can just test whichever models you’re interested in. Please don’t feel pressured to cover all models --- the tests you’ve done are already extremely valuable. |
Author
|
@0xShug0 Thank you for your kind words! The current issue is that the ROCm backend hasn't received any user feedback yet, and my own hardware doesn't meet the requirements for testing. If someone could compile and run it on Windows with a regular consumer-grade GPU, that would pretty much solve the problem. On a side note, I haven't found any optimization solutions for IndexTTS2, but ROCm should still be much faster than CPU in the end :( |
Owner
|
@IIIIIllllIIIIIlllll I added a unit test It should print something like (in your case backend should be hip)
If you share the output from the unit test I may be able to find the root cause of the IndexTTS performance issue. |
Author
|
@0xShug0 I just execute 'conv_lowering_matrix_test', here is the result.
|
Owner
|
@IIIIIllllIIIIIlllll Could you try this patch? pinned to commit 615971e |
Author
|
@0xShug0 I think I'm not mistaken.
|
Contributor
|
👋 Hey, I'm trying a rocm port on my side as well. I started from scratch and I'm stealing some ideas from this PR So far these are my results
I'm comparing generated wav files as well and diffs are in the margin of randomness. It's still the beginning, I have several ideas to add here and there in the codebase, I'm still navigating it to understand it. As soon as I have something decent I can push and we can try to combine our PRs (if it makes sense) |
Author
|
@fedeizzo Welcome aboard – it's great to have someone more professional on this! If you're willing to take over, that would be fantastic. Feel free to close my PR anytime. I'm not comfortable with these AI-generated codes since I lack the necessary expertise. |
|
I ran the conv_lowering_matrix_test on GFX1100 (7900XTX):
|
|
I had to install this package on arch linux I have a Radeon RX 9060 XT I got some errors with the build and got junie to help me out. https://github.com/hartmark/audio.cpp/tree/fix/hip-build-rdna3 Anyone have any scripts to download models and kickstart some transcribing? 🐌 |
Owner
|
@hartmark Thank you for testing! You can download all models from https://huggingface.co/audio-cpp/audio.cpp-gguf/tree/main. @IIIIIllllIIIIIlllll @fedeizzo @wizardeur Should we first prepare a merge-ready PR for ROCm support, along with documentation tracking the current testing status? It looks like some models already work out of the box. We can then use follow-up PRs to update the test results and add performance optimizations. |
It's a bit of a jungle for me as I have mostly just messing around with stable diffusion and LLMs. What model is a good first step to test on 16GB vram? |
Owner
Try OmniVoice for TTS and Nemotron-3.5-ASR-Streaming-0.6B for ASR. |
|
got a crash using the 16bit omnivoice models Generating with supertonic works however: |
Owner
|
Unfortunately, I don’t have access to a ROCm machine to debug this. This is exactly why runtime testing is essential when adding support for a new backend. |
Author
Let me try. |
Author
|
@hartmark Sorry to disappoint you. My test showed that the original safetensor model works. I'm now checking if it's a problem with the GGUF model. Could you provide a download link for your GGUF model? Thanks! :) |
|
I downloaded the GGUF model from the repo that @0xShug0 linked: For context on my environment, I'm currently running on a custom branch (fix/hip-build-rdna3) to resolve some HIP build issues specific to my AMD GFX1200 (RDNA 4) architecture. To get it compiling, a couple of fixes were applied:
Let me know if you think these compiler workarounds might be causing the issue with the GGUF model, or if there's anything else you'd like me to test! |
Bias tensors were loaded with native storage, so BF16 GGUF checkpoints fed BF16 tensors directly into elementwise add ops. The ggml CUDA/HIP binary broadcast path only supports F32/F16 src1 and asserted in binbcast.cu during decoder graph warmup. Norm weights and snake alpha were already forced to F32; do the same for all bias loads. Widening BF16 to F32 is lossless, so numerics and parity are unchanged.
Co-authored-by: Junie <junie@jetbrains.com>
Author
|
@hartmark Thank you! I have fixed both issues: omnivoice-bf16.gguf failure and compilation problems. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
5 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
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.