Misc. bug: Under-Performance of Linux ROCm 7.2 binaries

[frankyriventek]

Name and Version

b8180

Operating systems

Linux

Which llama.cpp modules do you know to be affected?

llama-bench, llama-server, llama-cli

Command line

$ ROCBLAS_USE_HIPBLASLT=0 ./llama-bench --mmap 0 -fa 1 -b 1024 -ub 2048 -d 32768,65536 -m ~/.cache/huggingface/models/unsloth/Qwen3-Coder-Next-GGUF/Q8_0/Qwen3-Coder-Next-Q8_0-00001-of-00003.gguf

Problem description & steps to reproduce

The test platform is an Strix Halo Laptop ( Zbook Ultra G1a, 128GB Unified Memory ) using Ubuntu 24.04 and kernel 6.19.4 with ROCm 7.2

When running a self-compiled version of llama.cpp b8180 with the stated command line I got the following results:

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
qwen3next 80B.A3B Q8_0	78.98 GiB	79.67 B	ROCm	99	1024	2048	1	pp512 @ d32768	343.99 ± 1.68
qwen3next 80B.A3B Q8_0	78.98 GiB	79.67 B	ROCm	99	1024	2048	1	tg128 @ d32768	24.21 ± 0.41
qwen3next 80B.A3B Q8_0	78.98 GiB	79.67 B	ROCm	99	1024	2048	1	pp512 @ d65536	251.15 ± 8.99
qwen3next 80B.A3B Q8_0	78.98 GiB	79.67 B	ROCm	99	1024	2048	1	tg128 @ d65536	22.09 ± 0.07

While when I run the same using the ROCm 7.2 binaries directly downloaded from the Releases This is what I got:

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
qwen3next 80B.A3B Q8_0	78.98 GiB	79.67 B	ROCm	99	1024	2048	1	pp512 @ d32768	93.26 ± 1.15
qwen3next 80B.A3B Q8_0	78.98 GiB	79.67 B	ROCm	99	1024	2048	1	tg128 @ d32768	22.42 ± 0.32
qwen3next 80B.A3B Q8_0	78.98 GiB	79.67 B	ROCm	99	1024	2048	1	pp512 @ d65536	49.72 ± 1.26
qwen3next 80B.A3B Q8_0	78.98 GiB	79.67 B	ROCm	99	1024	2048	1	tg128 @ d65536	18.38 ± 0.33

There is a known performance regression previously observed in ROCm 7+ builds (compared to ROCm 6.4.4) has been resolved via a workaround.

The issue was caused by a compiler regression (llvm/llvm-project#147700) affecting loop unrolling thresholds. I have applied the workaround (-mllvm --amdgpu-unroll-threshold-local=600) in the builds, restoring full performance.

This workaround will be removed once the upstream fix lands. For details, see the issue: kyuz0/amd-strix-halo-toolboxes#45

Still, I do not think it explains everything, these are my compilation flags:
cmake -S .
-B build
-DGGML_HIP=ON
-DAMDGPU_TARGETS=gfx1151
-DGPU_TARGETS=gfx1151
-DGGML_HIP_UMA=OFF
-DGGML_HIP_ROCWMMA_FATTN=OFF
-DGGML_HIP_GRAPHS=ON
-DROCM_PATH=/opt/rocm
-DHIP_PLATFORM=amd
-DCMAKE_HIP_FLAGS="--rocm-path=/opt/rocm -mllvm --amdgpu-unroll-threshold-local=600"
-DGGML_CUDA_ENABLE_UNIFIED_MEMORY=ON
-DGGML_CUDA_FA=ON
-DGGML_CUDA_FA_ALL_QUANTS=ON
-DGGML_CUDA_FORCE_MMQ=OFF
-DGGML_CUDA_FORCE_CUBLAS=OFF \
-DGGML_NATIVE=ON
-DGGML_OPENMP=ON
-DCMAKE_CXX_FLAGS="-I$HIP_INCLUDE_PATH"
-DCMAKE_C_FLAGS="-I$HIP_INCLUDE_PATH"
-DGGML_RPC=ON
-DLLAMA_BUILD_SERVER=ON
-DCMAKE_BUILD_TYPE=Release \

I hope this can help. Thanks a lot for this great software :)

Kind Regards
Franky

First Bad Commit

I've observed this bad performance since the first ROCm 7.2 build

Relevant log output

Logs

Self-Compiled llama.cpp:
ROCBLAS_USE_HIPBLASLT=0 llama-bench --mmap 0 -fa 1 -b 1024 -ub 2048 -d 32768,65536 -m ~/.cache/huggingface/models/unsloth/Qwen3-Coder-Next-GGUF/Q8_0/Qwen3-Coder-Next-Q8_0-00001-of-00003.gguf  
ggml_cuda_init: found 1 ROCm devices:
  Device 0: Radeon 8060S Graphics, gfx1151 (0x1151), VMM: no, Wave Size: 32
| model                          |       size |     params | backend    | ngl | n_batch | n_ubatch | fa |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | -------: | -: | --------------: | -------------------: |
| qwen3next 80B.A3B Q8_0         |  78.98 GiB |    79.67 B | ROCm       |  99 |    1024 |     2048 |  1 |  pp512 @ d32768 |        343.99 ± 1.68 |
| qwen3next 80B.A3B Q8_0         |  78.98 GiB |    79.67 B | ROCm       |  99 |    1024 |     2048 |  1 |  tg128 @ d32768 |         24.21 ± 0.41 |
| qwen3next 80B.A3B Q8_0         |  78.98 GiB |    79.67 B | ROCm       |  99 |    1024 |     2048 |  1 |  pp512 @ d65536 |        251.15 ± 8.99 |
| qwen3next 80B.A3B Q8_0         |  78.98 GiB |    79.67 B | ROCm       |  99 |    1024 |     2048 |  1 |  tg128 @ d65536 |         22.09 ± 0.07 |

build: d979f2b17 (8180)

Release b180 ROCm 7.2 binary:
ROCBLAS_USE_HIPBLASLT=0 ./llama-bench --mmap 0 -fa 1 -b 1024 -ub 2048 -d 32768,65536 -m ~/.cache/huggingface/models/unsloth/Qwen3-Coder-Next-GGUF/Q8_0/Qwen3-Coder-Next-Q8_0-00001-of-00003.gguf  
ggml_cuda_init: found 1 ROCm devices:
  Device 0: Radeon 8060S Graphics, gfx1151 (0x1151), VMM: no, Wave Size: 32
load_backend: loaded ROCm backend from /home/franky/Downloads/llama-b8180-bin-ubuntu-rocm-7.2-x64/llama-b8180/libggml-hip.so
load_backend: loaded RPC backend from /home/franky/Downloads/llama-b8180-bin-ubuntu-rocm-7.2-x64/llama-b8180/libggml-rpc.so
load_backend: loaded CPU backend from /home/franky/Downloads/llama-b8180-bin-ubuntu-rocm-7.2-x64/llama-b8180/libggml-cpu-zen4.so
| model                          |       size |     params | backend    | ngl | n_batch | n_ubatch | fa |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | -------: | -: | --------------: | -------------------: |
| qwen3next 80B.A3B Q8_0         |  78.98 GiB |    79.67 B | ROCm       |  99 |    1024 |     2048 |  1 |  pp512 @ d32768 |         93.26 ± 1.15 |
| qwen3next 80B.A3B Q8_0         |  78.98 GiB |    79.67 B | ROCm       |  99 |    1024 |     2048 |  1 |  tg128 @ d32768 |         22.42 ± 0.32 |
| qwen3next 80B.A3B Q8_0         |  78.98 GiB |    79.67 B | ROCm       |  99 |    1024 |     2048 |  1 |  pp512 @ d65536 |         49.72 ± 1.26 |
| qwen3next 80B.A3B Q8_0         |  78.98 GiB |    79.67 B | ROCm       |  99 |    1024 |     2048 |  1 |  tg128 @ d65536 |         18.38 ± 0.33 |

build: d979f2b17 (8180)

[IMbackK]

Our releases are built with these flags already see:

llama.cpp/.github/workflows/release.yml

Line 585 in 05728db

-DCMAKE_HIP_FLAGS="-mllvm --amdgpu-unroll-threshold-local=600" \

so if you see a regression with the latest releases its something else.

What artifact from the b8180 build are you using, exactly, the docker image or the plain build?

[IMbackK]

probably duplicate of #19480, please state if it affects other models too, then i will reopen.

[frankyriventek]

I was using the plain build:
https://github.com/ggml-org/llama.cpp/releases/download/b8180/llama-b8180-bin-ubuntu-rocm-7.2-x64.tar.gz

I have not observed it with another MoE model, mostly with Qwen3-Coder-Next at 8bits. I'm doing more tests if I can see more.

The major point is that I compiled the same build b8180 myself with the flags stated above, and then I got ~4x-5x more performance on PP and 10%-15% more performance on TG , with the same source of b8180.

[IMbackK]

could you a/b test DGGML_HIP_ROCWMMA_FATTN=OFF, DGGML_HIP_ROCWMMA_FATTN=ON afaik thats the only difference in build flags

[frankyriventek]

I got a confirmation with a dense model, Qwen3.5-27B:

• With my compilation flags:

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
qwen35 ?B Q8_0	26.62 GiB	26.90 B	ROCm,ROCm	99	1024	2048	1	pp512 @ d32768	159.53 ± 21.94
qwen35 ?B Q8_0	26.62 GiB	26.90 B	ROCm,ROCm	99	1024	2048	1	tg128 @ d32768	6.42 ± 0.05

• With the Release:

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
qwen35 ?B Q8_0	26.62 GiB	26.90 B	ROCm	99	1024	2048	1	pp512 @ d32768	44.03 ± 0.16
qwen35 ?B Q8_0	26.62 GiB	26.90 B	ROCm	99	1024	2048	1	tg128 @ d32768	5.80 ± 0.09

I'll do the test with the DGGML_HIP_ROCWMMA_FATTN=ON and report back.

[IMbackK]

Please try something other than qwen3.5. qwen3.5 and qwen3next are both very unusual models, in that they use delta net attention, in trying to figure you if its a problem in that path.

[frankyriventek]

Hi there, I've been able to do a proper A/B testing and I can confirm the problem is GGML_HIP_ROCWMMA_FATTN=ON . I recall from a previous discussion that the implementation for ROCWMMA was not performing ( #16827 ) but the developers rather preferred to implement a proper fix on the generic WMMA kernels rather than jusr ROCWMMA, which can explain the behaviour.

The data. I recompiled build b8180 with your build release flags and tested several models.

With GGML_HIP_ROCWMMA_FATTN=ON the results are:

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
qwen3next 80B.A3B Q8_0	78.98 GiB	79.67 B	ROCm	99	1024	2048	1	pp512 @ d32768	100.52 ± 0.32
qwen3next 80B.A3B Q8_0	78.98 GiB	79.67 B	ROCm	99	1024	2048	1	tg128 @ d32768	22.62 ± 0.13

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
gemma3 27B Q8_0	29.62 GiB	27.01 B	ROCm	99	1024	2048	1	pp512 @ d32768	83.86 ± 0.17
gemma3 27B Q8_0	29.62 GiB	27.01 B	ROCm	99	1024	2048	1	tg128 @ d32768	4.50 ± 0.01

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
nemotron_h_moe 31B.A3.5B Q8_0	31.27 GiB	31.58 B	ROCm	99	1024	2048	1	pp512 @ d32768	198.08 ± 0.28
nemotron_h_moe 31B.A3.5B Q8_0	31.27 GiB	31.58 B	ROCm	99	1024	2048	1	tg128 @ d32768	38.35 ± 1.59

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
gpt-oss 120B MXFP4 MoE	59.02 GiB	116.83 B	ROCm	99	1024	2048	1	pp512 @ d32768	204.26 ± 1.21
gpt-oss 120B MXFP4 MoE	59.02 GiB	116.83 B	ROCm	99	1024	2048	1	tg128 @ d32768	18.75 ± 0.17

With GGML_HIP_ROCWMMA_FATTN=OFF there are improvements across all ,the results are:

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
qwen3next 80B.A3B Q8_0	78.98 GiB	79.67 B	ROCm	99	1024	2048	1	pp512 @ d32768	342.63 ± 1.55
qwen3next 80B.A3B Q8_0	78.98 GiB	79.67 B	ROCm	99	1024	2048	1	tg128 @ d32768	25.97 ± 0.14

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
gemma3 27B Q8_0	29.62 GiB	27.01 B	ROCm	99	1024	2048	1	pp512 @ d32768	238.99 ± 1.31
gemma3 27B Q8_0	29.62 GiB	27.01 B	ROCm	99	1024	2048	1	tg128 @ d32768	5.95 ± 0.01

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
nemotron_h_moe 31B.A3.5B Q8_0	31.27 GiB	31.58 B	ROCm	99	1024	2048	1	pp512 @ d32768	562.33 ± 2.50
nemotron_h_moe 31B.A3.5B Q8_0	31.27 GiB	31.58 B	ROCm	99	1024	2048	1	tg128 @ d32768	44.60 ± 0.14

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
gpt-oss 120B MXFP4 MoE	59.02 GiB	116.83 B	ROCm	99	1024	2048	1	pp512 @ d32768	233.61 ± 1.35
gpt-oss 120B MXFP4 MoE	59.02 GiB	116.83 B	ROCm	99	1024	2048	1	tg128 @ d32768	33.41 ± 0.65

Hope this helps :)

[IMbackK]

I tried reproducing this on gfx1100 and was not able.
As you mentioned the rocwmma path is on the way out, so the solution is maybe to just disabled it for gfx115x.

[frankyriventek]

Did you also tried with at least 32K depth ? With the default 4K you barely can see any difference .. it grows with the depth.

Also, I found that is needed for ROCM 7.2 to add ROCBLAS_USE_HIPBLASLT=0 , as hipblastlt library is less performant than regular rocblas.

If you do not see the difference, may be only noticeable to RDNA 3+

[IMbackK]

yes, hipblaslt is faster but depends on shape, gfx1100 is RDNA3

[frankyriventek]

I was saying the other way around.
WITH hipblast ( ROCBLAS_USE_HIPBLASLT=1 ):

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
qwen35moe 80B.A3B Q4_K - Medium	63.65 GiB	122.11 B	ROCm	99	1024	2048	1	pp512 @ d32768	132.76 ± 0.33
qwen35moe 80B.A3B Q4_K - Medium	63.65 GiB	122.11 B	ROCm	99	1024	2048	1	tg128 @ d32768	18.06 ± 0.01
qwen35moe 80B.A3B Q4_K - Medium	63.65 GiB	122.11 B	ROCm	99	1024	2048	1	pp512 @ d65536	104.11 ± 0.23
qwen35moe 80B.A3B Q4_K - Medium	63.65 GiB	122.11 B	ROCm	99	1024	2048	1	tg128 @ d65536	15.93 ± 0.03

WITHOUT hipblast ( ROCBLAS_USE_HIPBLASLT=0 ):

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
qwen35moe 80B.A3B Q4_K - Medium	63.65 GiB	122.11 B	ROCm	99	1024	2048	1	pp512 @ d32768	166.15 ± 0.58
qwen35moe 80B.A3B Q4_K - Medium	63.65 GiB	122.11 B	ROCm	99	1024	2048	1	tg128 @ d32768	17.17 ± 0.56
qwen35moe 80B.A3B Q4_K - Medium	63.65 GiB	122.11 B	ROCm	99	1024	2048	1	pp512 @ d65536	123.58 ± 0.28
qwen35moe 80B.A3B Q4_K - Medium	63.65 GiB	122.11 B	ROCm	99	1024	2048	1	tg128 @ d65536	15.77 ± 0.02

WITHOUT hipblast ( ROCBLAS_USE_HIPBLASLT=0 ) and adding DGGML_CUDA_FA=ON DGGML_CUDA_FA_ALL_QUANTS=ON to the release flags:

model	size	params	backend	ngl	n_batch	n_ubatch	fa	test	t/s
qwen35moe 80B.A3B Q4_K - Medium	63.65 GiB	122.11 B	ROCm	99	1024	2048	1	pp512 @ d32768	175.39 ± 1.07
qwen35moe 80B.A3B Q4_K - Medium	63.65 GiB	122.11 B	ROCm	99	1024	2048	1	tg128 @ d32768	18.01 ± 0.28
qwen35moe 80B.A3B Q4_K - Medium	63.65 GiB	122.11 B	ROCm	99	1024	2048	1	pp512 @ d65536	126.85 ± 2.20
qwen35moe 80B.A3B Q4_K - Medium	63.65 GiB	122.11 B	ROCm	99	1024	2048	1	tg128 @ d65536	15.81 ± 0.35

[IMbackK]

you misunderstood my post, here it is formatted differently:

1. Yes i tested at that depth
2. on gfx1100 hipblaslt is generally faster, but it depends on gemm shape
3. gfx1100 is RDNA3

hipblaslt being slower for you is not a surprise as gfx115x has fewer registers and lacks some instructions compared to gfx1100 so some solutions dont apply to it.

Its also possible that register pressure is affecting the rocwmma fattn kernel negatively