Repository: ggml-org/llama.cpp
Pull request: HIP: disable rocwmma on gfx12 by default until rocm 7.0

Currently rocwmma dose not support gfx12 in any released version of rocm, causeing build failures when llamacpp is built with GGML_HIP_ROCWMMA_FATTN enabled against gfx12.

This causes issues for one of our downstreams (koboldcpp) who want to build a fat binary with gfx8-gfx12 support and simultaneously use rocwmma where applicable. This detail has also caused various issues https://github.com/ggml-org/llama.cpp/issues/14193 https://github.com/ggml-org/llama.cpp/issues/13110 to be filed against llamacpp by confused users.

This pr thus disables rocwmma on gfx12 even when GGML_HIP_ROCWMMA_FATTN is set if rocm < 7.0 is used.
I have it on good authority that gfx12 rocwmma support will land in rocm 7.0 which is scheduled for release later this year.

As the rocwmma git repo has contained gfx12 support for some time, it has been possible to compile llamacpp against git rocwmma to get a working build for some time, to facilitate this, this pr also adds a GGML_HIP_FORCE_ROCWMMA_FATTN_GFX12 option to override the check for rocm 7.0 and try to use rocwmma regardless.

Changed files:
filename: ggml/CMakeLists.txt
patch: @@ -172,6 +172,7 @@ option(GGML_HIP                             "ggml: use HIP"
 option(GGML_HIP_GRAPHS                      "ggml: use HIP graph, experimental, slow"         OFF)
 option(GGML_HIP_NO_VMM                      "ggml: do not try to use HIP VMM"                 ON)
 option(GGML_HIP_ROCWMMA_FATTN               "ggml: enable rocWMMA for FlashAttention"         OFF)
+option(GGML_HIP_FORCE_ROCWMMA_FATTN_GFX12   "ggml: enable rocWMMA FlashAttention on GFX12"    OFF)
 option(GGML_VULKAN                          "ggml: use Vulkan"                                OFF)
 option(GGML_VULKAN_CHECK_RESULTS            "ggml: run Vulkan op checks"                      OFF)
 option(GGML_VULKAN_DEBUG                    "ggml: enable Vulkan debug output"                OFF)

filename: ggml/src/ggml-cuda/common.cuh
patch: @@ -207,9 +207,9 @@ typedef float2 dfloat2;
 #define FP16_MMA_AVAILABLE
 #endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA
 
-#if defined(GGML_HIP_ROCWMMA_FATTN) && (defined(CDNA) || defined(RDNA3) || defined(RDNA4))
+#if defined(GGML_HIP_ROCWMMA_FATTN) && (defined(CDNA) || defined(RDNA3) || (defined(GGML_HIP_ROCWMMA_FATTN_GFX12) && defined(RDNA4)))
 #define FP16_MMA_AVAILABLE
-#endif // defined(GGML_HIP_ROCWMMA_FATTN) && (defined(CDNA) || defined(RDNA3) || defined(RDNA4))
+#endif // defined(GGML_HIP_ROCWMMA_FATTN) && (defined(CDNA) || defined(RDNA3) || (defined(GGML_HIP_ROCWMMA_FATTN_GFX12) && defined(RDNA4)))
 
 #if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING
 #define NEW_MMA_AVAILABLE

filename: ggml/src/ggml-hip/CMakeLists.txt
patch: @@ -113,6 +113,10 @@ if (GGML_HIP_ROCWMMA_FATTN)
     add_compile_definitions(GGML_HIP_ROCWMMA_FATTN)
 endif()
 
+if (GGML_HIP_FORCE_ROCWMMA_FATTN_GFX12 OR ${hip_VERSION} VERSION_GREATER_EQUAL 7.0)
+    add_compile_definitions(GGML_HIP_ROCWMMA_FATTN_GFX12)
+endif()
+
 if (NOT GGML_CUDA_FA)
     add_compile_definitions(GGML_CUDA_NO_FA)
 endif()