diff --git a/ggml/src/ggml-cuda/fattn-new-mma.cu b/ggml/src/ggml-cuda/fattn-new-mma.cu
index b2f67e4e..8c60b563 100644
--- a/ggml/src/ggml-cuda/fattn-new-mma.cu
+++ b/ggml/src/ggml-cuda/fattn-new-mma.cu
@@ -2255,10 +2255,15 @@ void ggml_cuda_flash_attn_ext_mma_new(ggml_backend_cuda_context & ctx, ggml_tens
         return;
     }
     if (Q->ne[0] == 512 && K->ne[0] == 512 && V->ne[0] == 512) {
-        if (gqa_ratio == 8) {
+        // head_dim 512: ncols2=16 exceeds the max dynamic shared memory on some GPUs (e.g. Ada,
+        // where cudaFuncSetAttribute returns invalid argument), so route gqa_ratio % 8 == 0
+        // (covers 8 and 16) through the ncols2=8 kernel. It iterates over Q-head groups
+        // (iter_z = ceil(gqa_ratio/ncols2)), so 16 heads run as two passes of 8. This unblocks
+        // head_dim-512 models with a 16:1 GQA ratio such as Gemma 4 12B's global layers.
+        if (gqa_ratio % 8 == 0) {
             ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1<512, 512, 8>(ctx, dst);
         }
-        else if (gqa_ratio == 4) {
+        else if (gqa_ratio % 4 == 0) {
             ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1<512, 512, 4>(ctx, dst);
         }
         else {