For gfx1151, llama.cpp should be built with `-DGGML_HIP_ROCWMMA_FATTN=ON` for a big performance boost

[lhl]

In my testing https://github.com/lhl/strix-halo-testing/tree/main/llm-bench I've found that pp (especially with FA) can be enormously faster when built with rocWMMA -DGGML_HIP_ROCWMMA_FATTN=ON than without.

To build for ROCm 7 I apply some janky patches: https://github.com/lhl/strix-halo-testing/blob/main/llm-bench/apply-rocwmma-fix.sh

BTW, here's an example with just the plainest Llama2 model but this isn't unrepresentative of the gains from some other models as well:

❯ llama.cpp-hip/build/bin/llama-bench --mmap 0 -fa 1 -m /models/gguf/llama-2-7b.Q4_K_M.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 | fa | mmap |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | -: | ---: | --------------: | -------------------: |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | ROCm       |  99 |  1 |    0 |           pp512 |        592.28 ± 1.63 |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | ROCm       |  99 |  1 |    0 |           tg128 |         40.40 ± 0.01 |

build: 7ad67ba9 (6111)

❯ ROCBLAS_USE_HIPBLASLT=1 llama.cpp-hip/build/bin/llama-bench --mmap 0 -fa 1 -m /models/gguf/llama-2-7b.Q4_K_M.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 | fa | mmap |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | -: | ---: | --------------: | -------------------: |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | ROCm       |  99 |  1 |    0 |           pp512 |        548.72 ± 1.99 |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | ROCm       |  99 |  1 |    0 |           tg128 |         40.43 ± 0.01 |

build: 7ad67ba9 (6111)

❯ ROCBLAS_USE_HIPBLASLT=1 llama.cpp-rocwmma/build/bin/llama-bench --mmap 0 -fa 1 -m /models/gguf/llama-2-7b.Q4_K_M.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 | fa | mmap |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | -: | ---: | --------------: | -------------------: |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | ROCm       |  99 |  1 |    0 |           pp512 |       1006.80 ± 3.80 |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | ROCm       |  99 |  1 |    0 |           tg128 |         39.46 ± 0.01 |

build: 7ad67ba9 (6111)

❯ llama.cpp-rocwmma/build/bin/llama-bench --mmap 0 -fa 1 -m /models/gguf/llama-2-7b.Q4_K_M.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 | fa | mmap |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | -: | ---: | --------------: | -------------------: |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | ROCm       |  99 |  1 |    0 |           pp512 |        899.73 ± 2.58 |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | ROCm       |  99 |  1 |    0 |           tg128 |         39.45 ± 0.01 |

build: 7ad67ba9 (6111)

[lhl]

BTW, last night I submitted a PR to llama.cpp ggml-org/llama.cpp#15239 and and AMD eng submitted a similar one shortly after ggml-org/llama.cpp#15241 for fixing rocWMMA builds w/ ROCm 7.0 - some back and forth going on on the "proper" way to fix things, but until then, my patching script continues to work.

[slojosic-amd]

Hi @lhl

It is a great coincidence that we both wrote a very similar patch for this problem with the warp sync function mask. :) Just to avoid confusion, we have been using mine patch internally for the last 15-20 days during llama.cpp testing with rocm7.0 The day before yesterday I first submitted a PR to our fork https://github.com/ROCm/llama.cpp and after the PR was approved I submitted a new one to the official llama.cpp and it turned out that I submitted it shortly after your PR. It's really incredible that we tried to solve the same problem in an almost identical way, and we didn't have any contact before.
I hope my latest changes will be accepted by the llama.cpp code-owners.

Best regards

[lhl]

No worries, figured it was one of those crazy coincidences. I've been using my own patches for a month or so as well but figured someone else would be submitting something better (I tried a bunch of different ifdef techniques but couldn't figure out a way around abstracting the mask if using the new native rocmWMMA implementation).

I finally submitted my patch to upstream since I started writing some docs Strix Halo docs for sharing: https://strixhalo-homelab.d7.wtf/AI/AI-Capabilities-Overview for some enthusiasts and figured it'd be easier than telling people to submit patches (which, well, still the case atm: https://strixhalo-homelab.d7.wtf/AI/llamacpp-with-ROCm )

[slojosic-amd]

Thanks for understanding. Great work with this Strix Halo docs.
Btw, after deeper investigation, it doesn't have anything related to "new" rocWMMA. The real issue here is macro change in https://github.com/ROCm/clr/blob/rocm-6.4.x-with-7.0-preview/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h#L30
HIP_DISABLE_WARP_SYNC_BUILTINS was HIP_ENABLE_WARP_SYNC_BUILTINS in https://github.com/ROCm/clr/blob/rocm-6.3.x/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h#L30 so for ROCm 6.5+ we should #define HIP_DISABLE_WARP_SYNC_BUILTINS instead of #undef HIP_ENABLE_WARP_SYNC_BUILTINS in https://github.com/ggml-org/llama.cpp/blob/master/ggml/src/ggml-cuda/fattn-wmma-f16.cu#L18 to prevent conflict with native built in warp sync functions after including rocwmma.hpp Please check: ggml-org/llama.cpp@74c7a2f

[lhl]

Ah interesting, I noticed there were a bunch of weird things w/ HIP_VERSION as well - eg, the earlier TheRock nightlies were updated on HIP_VERSION so those hip.h so I had to manually alter the datatype re-definitions for quite a while as well. I also noticed I didn't have a version guard on my code, another one of those things where if there are AMD employees with access to easy CI for all the various ROCm versions, that would be great.

Straying off topic, but is there anyone at AMD interested/able to host something like this? https://kyuz0.github.io/amd-strix-halo-toolboxes/ - being able to run regular testing for different versions and builds would be awesome, but more than can be expected for a hobbyist I think. I ran some sweeps that tooks many days (before discovering mmap and ROCm model loading did not play nice for big models) would be even better https://github.com/lhl/strix-halo-testing/tree/main/llm-bench - each model architecture has it's own peculiarities, but sadly, with all the libs/frameworks changing rapidly, these tests will be outdated very quickly: https://github.com/lhl/strix-halo-testing/tree/main/llm-bench/llama-2-7b.Q4_0

(ideally context would extend to 8K or 16K - there are reports that for example, the RADV driver scales much better than AMDVLK at longer context, frequently used for coding models)

[slojosic-amd]

I am not sure... maybe you can ask question/start discussion in https://github.com/ROCm/TheRock/discussions

[danielholanda]

Definitely noticed a lot of the same confusion around HIP versions. One of the key things that actually allows this llamacpp-rocm automation to work is a patch to that as shown here.

Also, looks like this just merged. Woohoo!
ggml-org/llama.cpp#15273

[lhl]

I am not sure... maybe you can ask question/start discussion in https://github.com/ROCm/TheRock/discussions

Sounds like a good idea, just FYI in case you're interested in taking a look: https://github.com/ROCm/ROCm/discussions/5194

[danielholanda]

@lhl and @slojosic-amd
Thanks for the great comments here. I just enabled ROCWMMA as part of THIS PR and the results are 🔥
Can you please take a look and let me know if you would change anything before I merge this?