Skip to content

ggml-cuda: tune D=256 tile flash-attn config for RDNA3.5 #41

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

Merged
mgehre-amd merged 1 commit into from
Jul 10, 2026
Merged

ggml-cuda: tune D=256 tile flash-attn config for RDNA3.5 #41

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
17 changes: 15 additions & 2 deletions ggml/src/ggml-cuda/fattn-tile.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -309,8 +309,19 @@ static constexpr __host__ __device__ uint32_t ggml_cuda_fattn_tile_get_config_am
return 0;
}

static constexpr __host__ __device__ uint32_t ggml_cuda_fattn_tile_get_config_amd_rdna3_5(const int DKQ, const int DV, const int ncols) {
// With rocWMMA FlashAttention off, D=256 prefill runs on the tile kernel; on RDNA3.5 a smaller
// K tile with higher occupancy is faster than the shared RDNA values. Other cases fall back.
GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 32, 256, 4, 64, 64)

return ggml_cuda_fattn_tile_get_config_amd_rdna(DKQ, DV, ncols);
}

static __host__ uint32_t ggml_cuda_fattn_tile_get_config(const int DKQ, const int DV, const int ncols, const int cc) {
if (GGML_CUDA_CC_IS_AMD(cc)) {
if (GGML_CUDA_CC_IS_RDNA3_5(cc)) {
return ggml_cuda_fattn_tile_get_config_amd_rdna3_5(DKQ, DV, ncols);
}
if (GGML_CUDA_CC_IS_RDNA(cc)) {
return ggml_cuda_fattn_tile_get_config_amd_rdna(DKQ, DV, ncols);
}
Expand All @@ -324,11 +335,13 @@ static __host__ uint32_t ggml_cuda_fattn_tile_get_config(const int DKQ, const in

static constexpr __device__ uint32_t ggml_cuda_fattn_tile_get_config(const int DKQ, const int DV, const int ncols) {
#ifdef GGML_USE_HIP
#ifdef RDNA
#ifdef RDNA3_5
return ggml_cuda_fattn_tile_get_config_amd_rdna3_5(DKQ, DV, ncols);
#elif defined(RDNA)
return ggml_cuda_fattn_tile_get_config_amd_rdna(DKQ, DV, ncols);
#else
return ggml_cuda_fattn_tile_get_config_amd(DKQ, DV, ncols);
#endif // RDNA
#endif // RDNA3_5
#else
#ifdef FAST_FP16_AVAILABLE
return ggml_cuda_fattn_tile_get_config_nvidia_fp16(DKQ, DV, ncols);
Expand Down
Loading