From 755422313aeb20a79755597ee1d972e83981b7ef Mon Sep 17 00:00:00 2001 From: sb32445 Date: Fri, 2 Oct 2026 09:17:12 +0200 Subject: [PATCH] cuda: smaller KV tile for the 8-column MMA flash attention config on Ampere/Ada (head size 256) With nbatch_fa 64 the kernel flash_attn_ext_f16<256,256,1,8> (one query, GQA 6) fits only 2 blocks per SM because of shared memory and reaches 58% of the DRAM bandwidth at 120k keys. nbatch_fa 32 fits 5 blocks per SM (84% DRAM). RTX 4070, Bonsai 2 27B, q4_0 K/V, one query, per layer: 32k keys 122 -> 87 us, 120k keys about 540 -> 370 us. Decode without speculative decoding: 32k 54.4 -> 56.4 tok/s, 120k 41.2 -> 46.0 tok/s. No change with 2-3 queries (different config), no change at depth 0. test-backend-ops FLASH_ATTN_EXT: 3006/3006 passed. Co-Authored-By: Claude Sonnet 5.5 --- ggml/src/ggml-cuda/fattn-mma-f16.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh index 0783e0d86329..463badb281d9 100644 --- a/ggml/src/ggml-cuda/fattn-mma-f16.cuh +++ b/ggml/src/ggml-cuda/fattn-mma-f16.cuh @@ -66,7 +66,7 @@ static constexpr __host__ __device__ fattn_mma_config ggml_cuda_fattn_mma_get_co GGML_CUDA_FATTN_MMA_CONFIG_CASE(192, 128, 32, 128, 2, 32, 96, 64, 64, 2, true); GGML_CUDA_FATTN_MMA_CONFIG_CASE(192, 128, 64, 128, 2, 32, 96, 64, 64, 2, true); - GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 8, 64, 4, 64, 128, 128, 128, 2, true); + GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 8, 64, 4, 32, 128, 128, 128, 2, true); GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 16, 64, 4, 32, 128, 128, 128, 2, true); GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 32, 128, 2, 32, 128, 128, 128, 2, true); GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 64, 128, 2, 32, 128, 128, 128, 2, true);