Performance of llama.cpp on AMD ROCm (HIP) #15021
Replies: 71 comments 131 replies
RX 7800 XT (Sapphire Pulse 280W)ggml_cuda_init: found 1 ROCm devices:
build: 00131d6 (6031) ggml_vulkan: Found 1 Vulkan devices:
build: baad948 (6056) Notes:
|
|
ggml_cuda_init: GGML_CUDA_FORCE_MMQ: no
build: e725a1a (6104) |
|
Happy to replicate: ggml_cuda_init: found 1 ROCm devices:
build: 9c35706 (6060) On Linux |
RX 7600 XTggml_cuda_init: found 1 ROCm devices:
build: 9c35706 (6060) Running on Linux 6.12.32, mainline amdgpu, ROCm 6.4.1. ggml_vulkan: Found 1 Vulkan devices:
build: 9c35706 (6060) |
|
@olegshulyakov , the 7600 XT actually has a 128 bit memory bus. |
|
AMD MI60. Happy to contribute.
I will post FA=1 and vulkan results once I have time during the weekend. |
MI100Using ./llama-bench -m llama-2-7b.Q4_0.gguf -ngl 99 -fa 0,1 -sm none -mg 0
build: 9c35706 (6060) I'm running Ubuntu 24.04.2 and ROCm 6.4.1 |
|
I expected it to be faster than RX 7800 XT because of HBM2... Have you tried to launch with a single device only? |
|
Bandwidth utilization is still fairly low on the gcn/cdna parts (gcn/cdna = same thing for tg). |
AMD Instinct MI300Xroot@0-4-9-gpu-mi300x1-192gb-devcloud-atl1:~/llama.cpp# ./build/bin/llama-bench -m llama-2-7b.Q4_0.gguf -ngl 99 -fa 0,1
build: 2bf3fbf (6069) Ref: #14640 |
|
I'm just referring to the rocWMMA flag from the build instructions: https://github.com/ggml-org/llama.cpp/blob/master/docs/build.md#hip
It should work for CDNA too but we have only tested with our RDNA3 cards (7900 XTX) and saw huge performance jumps in PP with FA on: #10879 (reply in thread) Please try it out because 1/3rd the performance in PP with FA on is just... strange at best |
|
With root@0-4-9-gpu-mi300x1-192gb-devcloud-atl1:~/llama.cpp# ./build/bin/llama-bench -m llama-2-7b.Q4_0.gguf -ngl 99 -fa 0,1
build: ee3a9fc (6090) |
|
Performance is in the middle between RTX 4090 and RTX 5090. |
|
So rocWMMA does work for CDNA in FA :) |
|
not very well |
Pro V620Why does FA slow down the V620 so much? Been a question I've been trying to answer for a while now.
build: 03d4698 (6074) Linux, ROCm 6.4.1 ( will try upgrading soon) |
|
@samteezy Can you please run per each device PRO V620/Pro WX 3200 and ROCm only backend? |
|
@olegshulyakov The numbers come out the same. Forcing root@llama:~# /root/llama-builds/llama.cpp/bin/llama-bench -m /mnt/models/llama-2-7b.Q4_0.gguf -ngl 99 -fa 0,1 -mg 0 -sm none
build: 5c0eb5e (6075) |
Powercolor Hellhound RX 7900 XTX (400W power limit)Opensuse tumbleweed system with rocm packages from
build: 5c0eb5e (6075) Sapphire Nitro 7900 XTX (400W power limit)In a different PC unfortunately because these GPUs are too chonky to fit in a regular case
build: 9c35706 (6060) |
|
So what are recommend settings/what to do to get best performance on 7900 XTX ? I have SAPPHIRE NITRO+ AMD Radeon RX 7900 XTX Vapor-X 24GB and without changing anything I get WAY worse results.
build: b9be58d (1005) And this is what LACT showed while running it 
|
|
At these high speeds with fast gpus the cpu gets important for results, your results and his would be within the expected variance for differing cpu performance. Not that it really matters mutch as the cpu gets almost irrelevant when a model of a size to fill the device is used. Benchmarking llama 7B Q4_0 is not really that great, as it dosent reflect actual usage much, this hurts the most on cdna devices which scale better than you would expect performance wise when increasing the number parameters. |
|
Go for ROCm 7.0, it has officially been released now. And compile llama.cpp with ROCWMMA enabled, see #15021 (comment) . You should get much better results. |
|
ROCm 7 is released? that's great news! I'll try it out as soon at is lands in my distro's package manager. That said I don't think it'll give any performance improvements for the 7900 XTX. The reason 9070 XT gets a boost with rocm7 is because pre rocm7 WMMA is not implemented for RDNA4. But we'll see. BTW FYI in my benchmarks, the powercolor 7900XTX was paired with a ryzen 2700 with 64GB RAM, and the sapphire 7900XTX was paired with a 5700X3d also 64GB RAM. I did not see an appreciable performance difference between the two setups in LLM inference. |
|
I see. To me that 3573 vs 3053 seems big difference. I'm running this on AMD Threadripper 1920X (12 core) with PCIe 3.0 x16. What benchmark I could use to compare more accurate real world inference performance? Now when I booted with |
Powercolor Red Devil 7900XTXAdrenalin 25.8.1 just came out, so time to test again
ggml_cuda_init: GGML_CUDA_FORCE_MMQ: no
ggml_cuda_init: GGML_CUDA_FORCE_CUBLAS: no
ggml_cuda_init: found 1 ROCm devices:
Device 0: AMD Radeon RX 7900 XTX, gfx1100 (0x1100), VMM: no, Wave Size: 32
build: 2572689 (6099) Still lower than the historical highs on May 26th (3599 and 3743), and a loss and a win against July 22nd (3529 and 3598). |
|
RX 7900 XTX (ASUS TUF)
build: 6c7e9a5 (6118) |
RX 6800 (16GB 203W)ROCm 6.3.4 on Ubuntu 24.04 in a Docker container
build: 79c1160 (6123) Bonus benchmarksI ran these to compare ROCm versions on various models. Obviously the results are specific to my RX 6800 and shouldn't be used to make any judgments about ROCm performance in general, especially on RDNA3 and later gpus. I use 6.3.4 because I don't care about LLama 3 8B. Note how fast the new MoE models are - gpt-oss-20B even at Q6_K_XL is faster than this 7B Q4_0 model. (Do make sure that you have a fixed version because the original gpt-oss releases had some issues - I used https://huggingface.co/unsloth/gpt-oss-20b-GGUF). ROCm 6.3.4
ROCm 6.4.3
|
RX 6800 (16GB 203W)ROCm 7.2.4 arch 7.0.10-arch1-1 |
|
RX 7900 XTX (ASUS TUF a bit overclocked for 100 mhz for core and VRAM) ./build/bin/llama-bench -m /home/vk/Downloads/llama-2-7b.Q4_0.gguf -ngl 99 -fa 0,1
build: 648ebcd (6146) |
|
RX 6900 XT AMD Reference Card (Stock clocks) Debian Testing llama.cpp version: gguf-v0.17.1-386-gfd1234cb
|
|
@tdjb Results are pretty low, can you re-test using llama.cpp standalone without Docker? |
|
Just a quick test, installed the llama.cpp build from Debian sid (was surprised to even find a build to be honest), which appears to be b5882 and the results came in quite similar. I tried the benchmark on both of my devices, as one is on a slower PCIe 4x slot, the results below are from the faster run. Why do you think the 6900 XT should perform better?
Happy to run further tests. |
|
It should be 10% better on my understanding, according to specs: RX 6900 XT and RX 6800 XT |
|
the rdna2 results are mostly surprisingly high, given the hardware capabilities, not low. |
|
GigaByte R9700
|
|
So 7900 XTX has better performance, that's sad. Also weird that Vulkan performs way worse on pp than ROCm but better on tg. |
|
We'll see when someone a bit more experienced gives it a shot. My benchmarks are about as vanilla as it gets. Threw it in an unraid server (12700k and 128GB DDR4-2133), made docker images and ran benchmarks. Many of the 7900 XTX results are baremetal, have factory overclock or are manually overclocked, installed additional drivers, and/or have raised power limits. I bet someone will beat my benchmarks shortly. |
there is zero reason to expect 9070(xt) to perform better than the xtx |
Radeon RX 9070 (non-XT)
build: 65349f2 (6183) I tried to enable the use of rocwmma with Still surprising that these numbers are better than the 9070 XT. |
|
ROCm 7.2.3 /workspace/llama.cpp/build/bin/llama-bench -m /data2/llm/models/llama-2-7b.Q4_0.gguf -ngl 99 -fa 0,1 -p 512 -n 128 --device ROCm0
build: 046e284 (9085) 2 card on PCI 4
build: 046e284 (9085) |
|
Is the "AMD Radeon Graphics, gfx1201 (0x1201) -> AI PRO R9700 (ASRock)-> PCI 4x4" by any chance on chipset PCIe or via M.2 ? I found Chipset PCIe VERY unstable and rocm-bandwidth-test -A (or -a) very inconsistent and wildly wayyyy below it's theoretical speed. I've ended up using bifurcation on my B650M and running PCIe 4.0 x8/x8 (wigh iGPU disabled as it was also messing the bandwidth somehow). |
|
rocm Version : 7.13.0a20260508-1
|
Radeon 890M
Running the more detailed test of: Getting results of:
build: 338ea1e (9127) |
|
A data point from the image-diffusion side, since the rocWMMA flash-attn numbers here are (understandably) all LLM token-generation. Diffusion attention is a different shape — non-GQA ( Setup: AMD Radeon AI PRO R9700 (Navi 48 / RDNA4 / gfx1201, 32 GB),
So ~6% on the attention kernel, ~2.5% on total wall, output verified correct at the same seed. Two things worth noting for other AMD/diffusion users:
The small gain is consistent with this workload being memory-bandwidth-bound at 1024² (O(n²) over 4096 tokens), not matmul-FLOP-bound — matrix cores help, just not a lot here. The same card sees a bigger relative benefit at smaller, more compute-bound configs. Happy to run more diffusion points (resolutions/steps/other RDNA cards) if it's useful to the rocWMMA tuning work in #16827 — it's a fixed-context, non-GQA stress case the LLM benchmarks don't cover. |
|
it should not displatch the tile kernel for the non rocwmma case for most shapes, it should dispatch the mma kernel: llama.cpp/ggml/src/ggml-cuda/fattn.cu Line 519 in 354ebac for CDNA this is the fastest kernel and for RDNA4 this kernel should also be fastest, at least for the shapes used in llms. |
|
Quick follow-up to my earlier diffusion numbers — the rocWMMA flash-attn win on RDNA4 turns out to be strongly config-dependent, which lines up nicely with the LLM results others have posted in this thread. Same setup as before (R9700 / gfx1201 / ROCm 7.0.2 / rocWMMA 2.0, FLUX.1 Krea Q4,
So at small / compute-bound configs the matrix-core path is a big win (~2.4×), but as resolution grows the attention becomes memory-bandwidth-bound (O(n²) K/V streaming) and the gain shrinks to a few percent. Same shape as the LLM side here — large on long-context prefill, modest on short prompts / decode. Net for image-diffusion users on RDNA4: definitely worth building with rocWMMA fattn on (ROCm 7 / rocWMMA 2.0) — just be aware the upside is very workload-dependent. For reference, with FA on the R9700 lands right next to an A6000 (CUDA) at the same diffusion recipe. |
|
you should probably not be useing rocwmma fattn at all on gfx12, its slower than the default fattn implementation on this arch at least for llms |
|
I get better results with rocwmma on my R9700 and Gemma 4 31B for prompt processing. |
|
OS: CachyOS
build: ba5b911 (9518) Just for gigs, Vulkan, it is quite a bit faster.
|
|
Why should AMD be ashamed? AMD only maintains HIP/ROCm to help customers move their code to international standard APIs, and away from legacy APIs like CUDA. If anything, the people who keep writing CUDA should be ashamed. There is no reason to keep using a non-standard API that runs on a minority of hardware. The only thing I wish AMD would do is to replace HIP's compiler backend with one that outputs standard SPIR-V IR instead of directly compiling it to AMD targets. No more needing to ship CUDA and HIP libraries, everything is already built into your OS's driver stack. Hell, if they played their cards right, they could probably pull the right strings to get it merged as a Mesa state tracker. |
|
I refer more to Vulkan backend performing better than AMD's own stack. You would think AMD knows their hardware stack best and yet a "generic" Vulkan backend is outperforming it. |
|
This is sort of a long term issue that several CEOs never personally waded into and fixed, even though its a company-wide culture issue: teams that came from outside of AMD don't work well with teams that are from AMD. When AMD bought ATI, they had incompatible corporate cultures.... and this was never fixed. The various CPU and platform teams seem to work well with each other, like, drivers arrive in time and generally work when a new CPU comes out, or Linux kernel patches make it into mainline before product launch, etc, and stuff just works. However, now that GPUs are now fully integrated into AMD's product line, the GPU driver team doesn't really communicate with either hardware team, and there is more than one GPU driver team. The guys who make Adrenaline happen aren't the guys who make ROCm/HIP happen, and also aren't the guys that work on Mesa/DRI/DRM/etc. On Windows, they're two different toolchains, and ROCm/HIP uses a different internal API to talk to the driver (basically a continuation of the Catalyst/FGLRX era, and not how the driver works today); and even though ROCm/HIP uses LLVM as a shader compiler, its also a different toolchain than Mesa/DRI, which also uses LLVM. Only in the past 2-3 years Lisa Su has actually managed to make any sort of headway into fixing it; the rest of those years shes been in for was trying to pull the company up from the nosedive it was in. Various teams have been merged and juggled around, plus more people have been hired on. The cross-pollination between teams is unfucking a lot of the long standing weirdness AMD GPU software had. |
|
You seem to know a lot about AMD! Were you from the team before? |
|
No, but I've come across various employees that work on drivers at AMD due to my time in open source software and needing to get bugs on their side fixed. |
PowerColor AMD Radeon AI PRO R9700Using ./build/bin/llama-bench -m models/llama-2-7b.Q4_0.gguf -ngl 99 -fa 0,1 -sm none -mg 1 (which seems to be the fastest of the 4 cards I have)
build: 6e9007a (9641) I'm running Ubuntu 24.04.4 and ROCm 7.2.4 |
|
Envious of that prompt processing speed! |
|
Updating Mi210 it seems overall better than before. Thermals are just ok for this run at 53C idel. ggml_cuda_init: found 2 ROCm devices (Total VRAM: 131040 MiB):
build: 6e9007a (9641) |
|
Going to wave this flag, just so not missed. :) I have an AMD Instinct MI25. At 16GB, it could in theory run decent-sized models. In the present hyper-overpriced world, the MI25 can be had for under $100. This is older hardware. Not (remotely) what you would use for commercial research. Not what you would use for production. But if you care less about time, do not want to burn tokens, and do not want to send data into the cloud - looks like a useful alternative. Seems this old card wants ROCm 5.4.3 (which I have). Seems the present llama.cpp does not work with this old card. The infallible AI (Microsoft Copilot) eventually came to the conclusion that this is impossible. As a recently retired software guy, I have other things to do, so care less if the LLM runs a bit longer. Before wandering off to try other approach, going to put the question in front of llama.cpp folk, to be sure this is a choice. :) |
just a heads up, amd has re-added gfx900, gfx90c(vega igpu) and gfx906 to the builds so the next rocm major version will re-add support, you can grab the preview builds for gfx900 already. ROCm/TheRock#3564 |
|
https://github.com/phil2sat/llamacpp-rocmcpp-turboquant/releases/tag/b1001 [58629] 1.16.199.754 I slot print_timing: id 0 | task 0 | prompt eval time = 20048.29 ms / 4713 tokens ( 4.25 ms per token, 235.08 tokens per second) nice for gemma-4-26b and MI25 @110W capped |
|
Have to admit, I spent the last several years working in a cave (at an air-gapped defense contractor), so have not spent much time in this domain, before. My aim is to spin up and use local LLMs for more extended experiments with various models. Did just get to that point, with current llama.cpp built for Vulcan. Made an attempt at a benchmark reported: here With llama.cpp in "router" mode, I can exercise any of the models in the huggingface cache through a web browser. This is very cool. Also I have a 16-year-old grandson who is learning to program, and he has access. At present, I have llama.cpp using Vulcan working with the MI25. But perhaps could get better performance with ROCm, if I can get a version of ROCm that works with gfx900. To summarize: From the prior responses, seems I could get there by moving to ArchLinux, as ROCm there kept gfx900 support. Seems AMD recently re-added gfx900 support to ROCm, which would land in 7.13, for which there is a preview release. Could install the preview release, or wait for the general release. If I read this right, the "turboquant" work would net more "speed and context" out of the MI25. Could build from: Did I miss anything? |
|
youre right you loose but gain, its always a matter of what you win and loose, most of the time you loose stupid features that your card doesnt support, like models that need 200+gb vram, lol. aktually the best you can get in my opinion is gemma-4 26b. its stock ROCm nothing patched and it works. so gfx900 is definitively inside, its some time ago that i yelled at them, why remove stuff that works and finally they get it. archlinux since a couple of years had gfx900 support inside hipblas. if you read the https://forum.level1techs.com/t/mi25-stable-diffusions-100-hidden-beast/ but its some damn old infos. i also could maybe give you my torch build scripts for the gfx900. it opens the door to the whole world of Image-Generation and other crazy stuff https://huggingface.co/Phil2Sat/Qwen-Image-Edit-Rapid-AIO-GGUF. but python is another story. a PRO TIP, MI25 hates llama.cpp Flash Attention and huge batch size. its faster without but you could try tris: thats the fastest i can get out of it an im testin, compiling, benching always. all my benches are real world benchmarks!!! |
|
Personally, myself... I like using cram, instead of cache-ram... Especially if I'm testing 24/7... |
|
im used https://github.com/lemonade-sdk/llamacpp-rocm
officialy bin https://github.com/ggml-org/llama.cpp/releases/download/b9994/llama-b9994-bin-ubuntu-rocm-7.2-x64.tar.gz works bad
|
RX 9070 XT (gfx1201) — Vulkan vs ROCm/HIP on Fedora 44Hardware
Software
Note: Fedora is not an officially validated distro for ROCm. ROCm 7.1.1 is the version available in Fedora 44 official repositories (ROCm 7.2 added official gfx1201 support, but 7.1.1 works on Fedora without Build flagsVulkan: ROCm/HIP: Benchmark resultsModel: Vulkan (3 runs):
ROCm/HIP (3 runs):
ObservationsToken generation (tg128): Vulkan is consistently ~23% faster than ROCm/HIP (220 vs 179 t/s). Both are very stable across runs. Prompt processing (pp512): ROCm is remarkably stable across runs (<1% variance). Vulkan shows high variance between runs — the first run after a cold start is fast (~7500 t/s), subsequent runs drop to ~5600 t/s. This is likely related to radv shader pipeline cache warm-up behavior. gfx1201 on ROCm 7.1 (Fedora): The GPU is correctly identified without Fedora note: ROCm is not officially supported on Fedora, but the packages available in Fedora 44 official repos (rocm-runtime 7.1.1) work correctly for llama.cpp inference on gfx1201. This may be a useful data point for anyone running RDNA4 on a non-Ubuntu distro. Post structure and formatting assisted by Claude (Anthropic). All benchmarks, system configuration, and observations are my own. |
Update — Llama 2 7B Q4_0 (standard model)Running the standard benchmark model for comparability with other entries. Note: using Vulkan (3 runs) — llama-2-7b.Q4_0,
ROCm/HIP (3 runs) — llama-2-7b.Q4_0,
Observations (7B Q4_0):
For the full system configuration and software versions, see the original post above. |
|
For those who are wondering about gfx90c. ROCm Vulkan |
Strix Halo (AMD AI Max 395+, gfx1151) with ROCm 7.14The standard benchmark numbers that show somewhat ~10% improvement, compared to the current numbers. Built |
|
Strix Halo ROCm (HIP) numbers, measured head-to-head against Vulkan on the same machine. Hardware: Ryzen AI MAX+ 395 / Radeon 8060S (gfx1151) / 128 GB LPDDR5X UMA, Ubuntu 24.04.4, kernel 6.17. Method: llama.cpp master (2026-08), separate ROCm (HIP) + Vulkan (RADV) builds, llama-bench + llama-server, single request serial, ±5%/cell. Decode throughput (tok/s):
On this box ROCm trails Vulkan: bare −2~−25% (dense 27B only Full matrix + scripts: https://github.com/JJJJJovi/strix-halo-benchmark |
|
llamacpp version: b10590 llamacpp compiled with; Card Setting: GPU1 (Core up to 3000mz, vram up to 2650mz, power up to 350W) Card Setting: GPU1 + GPU0 (Core up to 3000mz, vram up to 2650mz, power up to 350W) -sm layer,tensor Tensor + fa 0 - doesn't work run with fa 1 only Interestingly, split tensor for such a small model makes it worse instead of helping. For comparison gemma 4 31B Q8: Tensor + fa 0 - doesn't work run with fa 1 only Very nice performance for dual gpu with tensor split (about +50% for generation is impressive, the boost in processing is less impressive - but it's still nice). |
and others
Uh oh!
There was an error while loading. Please reload this page.
Uh oh!
There was an error while loading. Please reload this page.
This is similar to the Performance of llama.cpp on Apple Silicon M-series, Performance of llama.cpp on Nvidia CUDA and Performance of llama.cpp with Vulkan, but for ROCm! I think it's good to consolidate and discuss our results here.
We'll be testing the Llama 2 7B model like the other thread to keep things consistent, and use Q4_0 as it's simple to compute and small enough to fit on a 4GB GPU. You can download it here.
Instructions
Either run the commands below or download one of our ROCm(HIP) releases. If you have multiple GPUs please run the test on a single GPU using
-sm none -mg YOUR_GPU_NUMBERunless the model is too big to fit in VRAM.Share your llama-bench results along with the git hash and ROCm info string in the comments. Feel free to try other models and compare backends, but only valid runs will be placed on the scoreboard.
If multiple entries are posted for the same device I'll prioritize newer commits with substantial ROCm updates, otherwise I'll pick the one with the highest overall score at my discretion. Performance may vary depending on driver, operating system, board manufacturer, etc. even if the chip is the same. For integrated graphics note that your memory speed and number of channels will greatly affect your inference speed!
ROCm Scoreboard for Llama 2 7B, Q4_0 (no FA)
ROCm Scoreboard for Llama 2 7B, Q4_0 (with FA)
More detailed test
The main idea of this test is to show a decrease in performance with increasing size.
All reactions