public_git/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-10-12 01:05:43 +02:00
Code Issues Packages Projects Releases Wiki Activity

HIP: Enable fattn-mma kernel on cdna for dkq > 256 for large batch sizes (llama/28907)

Browse Source 
* 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
This commit is contained in:
uvos
2026-10-06 10:38:05 +03:00
committed by Georgi Gerganov
parent 14d1aa70e6
commit 743f1ad286
2 changed files with 5 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/fattn-mma-f16.cuh
+1 -1
View File
@@ -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;
}
ggml/src/ggml-cuda/fattn.cu
+4 -1
View File
@@ -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.