Skip to content

ggml-cuda: use mma-f16 flash-attn for D=256 on RDNA3.5 #40

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
roberteg16 wants to merge 1 commit into
base: gfx11
Choose a base branch
from
Open

ggml-cuda: use mma-f16 flash-attn for D=256 on RDNA3.5 #40

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions ggml/src/ggml-cuda/fattn-mma-f16.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1759,7 +1759,12 @@ static __global__ void flash_attn_ext_f16(
#endif // __CUDA_ARCH__ == GGML_CUDA_CC_TURING

#if defined(AMD_WMMA_AVAILABLE)
// DKQ=256 is only tuned/validated on RDNA3.5; other AMD WMMA archs keep the DKQ<=128 limit.
#if defined(RDNA3_5)
if (ncols1*ncols2 < 16 || ncols2 == 1 || DKQ > 256) {
#else
if (ncols1*ncols2 < 16 || ncols2 == 1 || DKQ > 128) {
#endif // defined(RDNA3_5)
NO_DEVICE_CODE;
return;
}
Expand Down
7 changes: 7 additions & 0 deletions ggml/src/ggml-cuda/fattn.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -499,6 +499,13 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
if (can_use_vector_kernel && Q->ne[1] <= 2) {
return BEST_FATTN_KERNEL_VEC;
}
// The mma-f16 kernel also lowers to RDNA WMMA instructions here, and at D=256 it packs
// ncols2 GQA query heads into a full 16-wide tile where the wmma kernel leaves it mostly
// empty. Needs GQA (ncols2 >= 2) and a filled tile (ncols1*ncols2 >= 16). Restricted to
// RDNA3.5, the only AMD arch this was tuned/validated on.
if (GGML_CUDA_CC_IS_RDNA3_5(cc) && Q->ne[0] == 256 && gqa_opt_applies && Q->ne[1]*gqa_ratio_eff >= 16) {
return BEST_FATTN_KERNEL_MMA_F16;
}
return BEST_FATTN_KERNEL_WMMA_F16;
}

Expand Down
Loading