Skip to content

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

Closed
Closed
Misc. bug: HIP backend performs poorly on AMD Ryzen AI MAX 395 (Strix Halo gfx1151)#13565
Labels
AMD GPUIssues specific to AMD GPUsperformanceSpeed related topics

Description

@lhl
lhl
opened on May 15, 2025

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

Activity
  1. lhl
    added 
    bug-unconfirmed
     on May 15, 2025
  2. IMbackK commented on May 15, 2025 
    @IMbackK
    IMbackK
    on May 15, 2025
    Contributor
    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.
  3. IMbackK
    added 
    performanceSpeed related topics
     
    AMD GPUIssues specific to AMD GPUs
     and removed 
    bug-unconfirmed
     on May 15, 2025
  4. lhl commented on May 16, 2025 
    @lhl
    lhl
    on May 16, 2025
    Author
    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...
  5. lhl commented on May 16, 2025 
    @lhl
    lhl
    on May 16, 2025
    Author
    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/
  6. lhl commented on May 16, 2025 
    @lhl
    lhl
    on May 16, 2025
    Author
    

    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.
  7. lhl commented on May 16, 2025 
    @lhl
    lhl
    on May 16, 2025
    Author
    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)

  8. lhl commented on May 17, 2025 
    @lhl
    lhl
    on May 17, 2025
    Author
    BTW, besides testing the improvements using ROCBLAS_USE_HIPBLASLT=1 I also filed a bug about the gfx1100/gfx1151 kernel regressions: https://github.com/ROCm/ROCm/issues/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)

  9. IMbackK commented on May 17, 2025 
    @IMbackK
    IMbackK
    on May 17, 2025
    Contributor
    Seams like kernel selection is just really buggy in gfx115x
  10. Djip007 commented on May 17, 2025 
    @Djip007
    Djip007
    on May 17, 2025
    Contributor
    Nice!
    
have you made some bench with bf16 on CPU? and BF16/FP16 on GPU?
  11. IMbackK commented on May 18, 2025 
    @IMbackK
    IMbackK
    on May 18, 2025
    Contributor
    since this isent really a issue on the llamacpp side, i will close this here
  12. IMbackK
    closed this as completedon May 18, 2025
  13. fabiantax
    added 4 commits that reference this issue on Mar 4, 2026
    Add strix-halo-optimizer skill for AMD Ryzen AI Max+ 395
    af2af6e
    ggml-cuda: auto-detect AMD APU unified memory and improve iGPU support
    5c7ab57
    Add strix-halo-optimizer skill for AMD Ryzen AI Max+ 395
    60a4387
    ggml-cuda: auto-detect AMD APU unified memory and improve iGPU support
    4e5622c
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Metadata
Metadata
Assignees
No one assigned
    Labels
    AMD GPUIssues specific to AMD GPUsperformanceSpeed related topics
    Type
    No type
    Projects
    No projects
      Milestone
      No milestone
      Relationships
      None yet
      Development
      No branches or pull requests
      Issue actions