Misc. bug: HIP backend performs poorly on AMD Ryzen AI MAX 395 (Strix Halo gfx1151)

[lhl]

Name and Version

❯ build/bin/llama-cli --version
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 Graphics, gfx1151 (0x1151), VMM: no, Wave Size: 32
version: 5392 (c753d7be)
built with cc (GCC) 15.0.1 20250418 (Red Hat 15.0.1-0) for x86_64-redhat-linux

Operating systems

Linux

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

llama-bench

Command line

llama.cpp-cpu/build/bin/llama-bench -m ~/models/llama-2-7b.Q4_0.gguf
llama.cpp-vulkan/build/bin/llama-bench -m ~/models/llama-2-7b.Q4_0.gguf
llama.cpp-hip/build/bin/llama-bench -m ~/models/llama-2-7b.Q4_0.gguf

Problem description & steps to reproduce

Recently I've been testing a Strix Halo (gfx1151) system and was a bit surprised by how poorly the HIP backend ran. All tests were run with llama-bench built on HEAD (b5392) with the standard TheBloke/Llama-2-7B-GGUF (Q4_0):

Backend	pp512 (t/s)	tg128 (t/s)
CPU	304.42 ± 2.05	28.65 ± 0.03
HIP	348.62 ± 0.35	48.70 ± 0.02
Vulkan	881.38 ± 2.11	52.82 ± 0.04

The HIP version performs far below what you'd expect in terms of tok/TFLOPS efficiency for prompt processing vs other RDNA3 architectures:

• gfx1103 Radeon 780M iGPU gets 14.51 tok/TFLOP. At that efficiency you'd expect the about 850 tok/s that the Vulkan backend delivers.
• gfx1100 Radeon 7900 XTX gets 25.12 tok/TFLOP. At that efficiency you'd expect almost 1500 tok/s, almost double what the Vulkan backend delivers, and >4X what the current HIP backend delivers.
• HIP pp512 barely beats out CPU backend numbers. I don't have an explanation for this.
• Just for a reference of how bad the HIP performance is, an 18CU M3 Pro has ~12.8 FP16 TFLOPS (4.6X less compute than Strix Halo) and delivers about the same pp512. Lunar Lake Arc 140V has 32 FP16 TFLOPS (almost 1/2 Strix Halo) and has a pp512 of 657 tok/s (1.9X faster)
• With the Vulkan backend pp512 is about the same as an M4 Max and tg128 is about equivalent to an M4 Pro

With monitoring, I've confirmed that both HIP and Vulkan reach the max graphics clock. This is a system running Linux 6.15.0-0.rc3, so should be up to date with the latest AMDGPU drivers.

These results are from a standard llama-bench run. I've tried -fa 1 and a rocWMMA build but they don't make much difference so excluded for clarity.

One interesting observation that may help track a potential regression, when I compile w/ gfx1100 support and run with HSA_OVERRIDE_GFX_VERSION=11.0.0 , the pp512 basically doubles to 598.84 ± 1.41 (this eventually leads to MES/kernel errors so obviously is not recommended for use, just an interesting observation that might help in tracking down the issue).

First Bad Commit

No response

Relevant log output

[IMbackK]

The performance of pp is mostly down to the performance of mulmat. Since we dont have any kernels that can make use of wmma or mfma, we convert to fp16 and use rocblas to perform the matrix multiplication.

From what i have heard the kernels rocblas includes for gfx115x are simply bad, which is why when you pretend to be gfx1100 performance improves.

you can test this theory with test-backend-ops perf -o MUL_MAT and something like rocblas-bench -f gemm_ex --transposeA T --transposeB N -m 4096 -n 512 -k 4096 --alpha 1 --a_type f16_r --lda 4096 --b_type f16_r --ldb 4096 --beta 0 --c_type f16_r --ldc 4096 --d_type f16_r --ldd 4096 --compute_type f16_r both of those will likely show very poor performance.

Ultimately its up to AMD to provide a half way competent implementation of BLAS.

For silly reasons, AMD has two sets indipendent of generators for matrix multiplication shaders one in tensile and one in tensilelt, the later being used by hipblaslt. You can try if these are any better by building rocblas with BUILD_WITH_HIPBLASLT and using export ROCBLAS_USE_HIPBLASLT=1 to make rocblas cross dispatch to hipblaslt which will then use tensilelt kernels. They might be better.

[lhl]

OK, this makes sense to me if llama.cpp isn't using hipBLASLt.

So I just gave build/bin/test-backend-ops perf -o MUL_MAT a run and got:

MUL_MAT(type_a=q4_0,type_b=f32,m=4096,n=8,k=14336,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0):               6206 runs -   163.23 us/run - 939.52 MFLOP/run -   5.76 TFLOPS

This roughly matches w/ my mamf-finder rocBLAS/hipblas results https://llm-tracker.info/_TOORG/Strix-Halo#mamf-finder give 5.0 TFLOPS (<9% theoretical)

Tried  3375 shapes => the best outcomes were:
mean:   5.0 TFLOPS @ 4096x9216x1024 (MxNxK)
median: 5.0 TFLOPS @ 12288x3072x1024 (MxNxK)
max:    5.1 TFLOPS @ 11264x3072x1024 (MxNxK)
Elapsed time: 1 day, 10:40:32

mamf-finder performs much better w/ hipBLASLt (>60% theoretical, is inline w/ MI300X efficiency):

Tried  3375 shapes => the best outcomes were:
mean:   35.1 TFLOPS @ 15360x3072x1024 (MxNxK)
median: 35.1 TFLOPS @ 15360x3072x1024 (MxNxK)
max:    36.9 TFLOPS @ 6144x3072x3072 (MxNxK)

Elapsed time: 6:04:34

So I think forcing llama.cpp to use hipBLASLt will be the way to do. Do the BLAS libraries get statically compiled for llama.cpp (I believe so since I'm able to compile for example -DGGML_HIP_ROCWMMA_FATTN=ON in a docker (rocWMMA) and have it work?) or do I need to pass additional flags? Due to the way my test system is setup, I have to do most of my work w/ the latest ROCm/HIP libraries within a docker container...

[lhl]

BTW, just an FYI, I've done a bunch more tests and ran theoretical vs actual perf analyses including comparisons to gfx1100 and gfx1103 hardware efficiency for the ROCm paths here, and testing some other models and backends (eg, the HIP backend outperforms Vulkan for Qwen3 MoE, but mostly because there are other bottlenecks there): https://www.reddit.com/r/LocalLLaMA/comments/1kmi3ra/amd_strix_halo_ryzen_ai_max_395_gpu_llm/

[lhl]

The performance of pp is mostly down to the performance of mulmat. Since we dont have any kernels that can make use of wmma or mfma, we convert to fp16 and use rocblas to perform the matrix multiplication.

BTW, also, I do have rocWMMA built and have a built with -DGGML_HIP_ROCWMMA_FATTN=ON (which works amazingly well on long context - no performance loss on pp8K/tg8K and even pp128K) but it doesn't in general make things faster. Is there a way to force more of the code paths to use rocWMMA (or would that require major surgery HIP backend)? Just mentioning it in case there are easy wins available.

[lhl]

OK, well I thought I would have to rebuild rocBLAS and hit a bunch of issues, but it turns out rocblas-bench was lying to me and ROCBLAS_USE_HIPBLASLT=1 is supported out of the box.

[root@94174dd39af9 llama.cpp]# llama.cpp-hip/build/bin/llama-bench -m /home/lhl/models/llama-2-7b.Q4_0.gguf
/share/libdrm/amdgpu.ids: No such file or directory
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 Graphics, gfx1151 (0x1151), VMM: no, Wave Size: 32
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | ROCm,RPC   |  99 |           pp512 |        347.79 ± 0.66 |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | ROCm,RPC   |  99 |           tg128 |         48.61 ± 0.04 |

build: c753d7be (5392)

[root@94174dd39af9 llama.cpp]# ROCBLAS_USE_HIPBLASLT=1 llama.cpp-hip/build/bin/llama-bench -m /home/lhl/models/llama-2-7b.Q4_0.gguf
/share/libdrm/amdgpu.ids: No such file or directory
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 Graphics, gfx1151 (0x1151), VMM: no, Wave Size: 32
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | ROCm,RPC   |  99 |           pp512 |        882.81 ± 3.21 |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | ROCm,RPC   |  99 |           tg128 |         48.63 ± 0.02 |

build: c753d7be (5392)

It looks like we are now matching the Vulkan speeds.

This seems to work w/ the rocWMMA build as well:

[root@94174dd39af9 llama.cpp]# ROCBLAS_USE_HIPBLASLT=1 llama.cpp-rocwmma/build/bin/llama-bench -m /home/lhl/models/llama-2-7b.Q4_0.gguf
/share/libdrm/amdgpu.ids: No such file or directory
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 Graphics, gfx1151 (0x1151), VMM: no, Wave Size: 32
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | ROCm,RPC   |  99 |           pp512 |        880.18 ± 1.86 |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | ROCm,RPC   |  99 |           tg128 |         48.66 ± 0.01 |

build: 3cc1f1f1 (5393)

[lhl]

BTW, besides testing the improvements using ROCBLAS_USE_HIPBLASLT=1 I also filed a bug about the gfx1100/gfx1151 kernel regressions: ROCm/ROCm#4748

This is still important since certain models (Qwen3-30B-A3B-UD-Q4_K_XL for example) will not be able to use hipBLASLt: rocBLAS warning: hipBlasLT failed, falling back to tensile.

Interestingly, even though it reports this error, pp512 still benches almost 50% faster...

# llama.cpp-hip/build/bin/llama-bench -m /home/lhl/models/Qwen3-30B-A3B-UD-Q4_K_XL.gguf
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 Graphics, gfx1151 (0x1151), VMM: no, Wave Size: 32
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| qwen3moe 30B.A3B Q4_K - Medium |  16.49 GiB |    30.53 B | ROCm,RPC   |  99 |           pp512 |        403.08 ± 1.27 |
| qwen3moe 30B.A3B Q4_K - Medium |  16.49 GiB |    30.53 B | ROCm,RPC   |  99 |           tg128 |         56.95 ± 0.01 |

build: c753d7be (5392)

# ROCBLAS_USE_HIPBLASLT=1 llama.cpp-hip/build/bin/llama-bench -m /home/lhl/models/Qwen3-30B-A3B-UD-Q4_K_XL.gguf
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 Graphics, gfx1151 (0x1151), VMM: no, Wave Size: 32
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
rocBLAS error: No hipBLASLt solution found
This message will be only be displayed once, unless the ROCBLAS_VERBOSE_HIPBLASLT_ERROR environment variable is set.

rocBLAS warning: hipBlasLT failed, falling back to tensile.
This message will be only be displayed once, unless the ROCBLAS_VERBOSE_TENSILE_ERROR environment variable is set.
| qwen3moe 30B.A3B Q4_K - Medium |  16.49 GiB |    30.53 B | ROCm,RPC   |  99 |           pp512 |        590.40 ± 3.71 |
| qwen3moe 30B.A3B Q4_K - Medium |  16.49 GiB |    30.53 B | ROCm,RPC   |  99 |           tg128 |         57.02 ± 0.03 |

build: c753d7be (5392)

[IMbackK]

Seams like kernel selection is just really buggy in gfx115x

[Djip007]

Nice!
have you made some bench with bf16 on CPU? and BF16/FP16 on GPU?

[IMbackK]

since this isent really a issue on the llamacpp side, i will close this here