From 743f1ad28645d8edda0ff4a5dc3b44b563896d2a Mon Sep 17 00:00:00 2001 From: uvos Date: Sun, 27 Sep 2026 12:06:56 +0200 Subject: [PATCH] HIP: Enable fattn-mma kernel on cdna for dkq > 256 for large batch sizes (llama/28907) * HIP: Enable fattn-mma kernel on cdna for dkq > 256 for large batch sizes * CI: hip-quality-check: ignore spills for very large mfma mma kernels --- ggml/src/ggml-cuda/fattn-mma-f16.cuh | 2 +- ggml/src/ggml-cuda/fattn.cu | 5 ++++- 2 files changed, 5 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh index abb99a354..fa0d347af 100644 --- a/ggml/src/ggml-cuda/fattn-mma-f16.cuh +++ b/ggml/src/ggml-cuda/fattn-mma-f16.cuh @@ -1875,7 +1875,7 @@ static __global__ void flash_attn_ext_f16( #endif // defined(AMD_WMMA_AVAILABLE) #if defined(AMD_MFMA_AVAILABLE) - if (ncols1*ncols2 < 16 || DKQ > 256) { + if (ncols1*ncols2 < 16 || (DKQ > 256 && ncols1*ncols2 < 64)) { NO_DEVICE_CODE; return; } diff --git a/ggml/src/ggml-cuda/fattn.cu b/ggml/src/ggml-cuda/fattn.cu index d1fcf58cb..ad4d1df2c 100644 --- a/ggml/src/ggml-cuda/fattn.cu +++ b/ggml/src/ggml-cuda/fattn.cu @@ -681,7 +681,7 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const } // AMD MFMA needs a certain minimum batch size to outscale the tile kernel for large head sizes. - if ((amd_mfma_available(cc) && Q->ne[0] <= 256) && Q->ne[0] != 40 && Q->ne[0] != 72) { + if (amd_mfma_available(cc) && Q->ne[0] != 40 && Q->ne[0] != 72) { if ((Q->ne[0] <= 64 && Q->ne[1] * gqa_ratio_eff > 8)) { return BEST_FATTN_KERNEL_MMA_F16; } @@ -691,6 +691,9 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const if ((Q->ne[0] <= 256 && Q->ne[1] * gqa_ratio_eff > 64)) { return BEST_FATTN_KERNEL_MMA_F16; } + if (Q->ne[0] > 256 && gqa_opt_applies && Q->ne[1] * gqa_ratio_eff > 128) { + return BEST_FATTN_KERNEL_MMA_F16; + } } // AMD WMMA is faster than the tile kernel if the wide tiles with high arithmetic intensity can be utilized.