CUDA: enable sparse FlashAttention on Volta

[CoREse]

Summary

Sparse flash attention (the n_kv_max path added in #27970) is currently gated behind turing_mma_available(), so on Volta it never runs: DeepSeek-V4's CSA layers attend over the whole compressed KV cache even though the mask only keeps min(n_swa, 128) + top_k(512) = 640 entries per query row. This PR enables the sparse path on Volta, where it is worth up to 14x for prompt processing and 6x for token generation at long context.

Two constraints have to be satisfied at once:

• the sparse gather uses one index row per launch, so ncols1 == 1;
• the Volta build emits no device code for launches with fewer than 32 columns (fattn-mma-f16.cuh: #ifdef VOLTA_MMA_AVAILABLE if (ncols1*ncols2 < 32) NO_DEVICE_CODE), so the existing sparse instances (<512, 512, 1, 8>, <576, 512, 1, 16>) abort there.

The only viable combination is <512, 512, 1, 32>: all GQA heads of one query row in a single 32-column launch. DeepSeek-V4 has 64 query heads on 1 KV head, so its CSA layers split into two such launches per query row. The instance is added to the existing ncols1_1-ncols2_32 file (which already carries 320/256 and 576/512).

Changes:

• ggml_cuda_fattn_volta_sparse_ok(): allow sparse on Volta for DKQ == DV == 512 with gqa_ratio % 32 == 0; the rest of the sparse preconditions are unchanged.
• ggml_cuda_flash_attn_ext_mma_f16_switch_ncols2(): route such shapes to <DKQ, DV, 1, 32> on Volta.
• ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1(): keep Volta out of the narrower sparse instances.
• ggml_cuda_get_best_fattn_kernel(): for a single token, prefer the sparse MMA launch over the dense tile kernel (which reads all of K for every token).
• Split the body of shall_use_sparse() into a cc-based helper so the kernel selection can use it too.
• may_use_sparse(), extern declaration, instance file and generate_cu_files.py updated for the new instance; only ncols1 == 1 is generated for 512/32.

Test

test-backend-ops, V100 (sm_70), CUDA 12.9:

• 3 new cases covering the DeepSeek-V4 CSA layout (64 query heads on 1 KV head, sinks, n_kv_max = 640, with V both separate and a view of K, batch 32 and 1) — pass.
• -o FLASH_ATTN_EXT otherwise unchanged: 2874/2874 pass. (Two pre-existing Volta failures are excluded and are unrelated to this PR — hsk=320 aborts in cudaFuncSetAttribute with "invalid argument", and max_bias=8 hits "no device code for arch 700" because ALiBi disables the GQA optimization and lands on a 16-column instance. Both reproduce on master.)

test-backend-ops perf, V100, DeepSeek-V4 CSA shape (hsk=hsv=512, 64 query heads on 1 KV head, n_kv_max=640):

kv	batch	dense (master)	sparse (this PR)
8192	256	12.87 ms	2.69 ms	4.8x
32768	256	50.86 ms	3.47 ms	14.6x
8192	1	148 us	72 us	2.1x
32768	1	517 us	85 us	6.1x

End to end, DeepSeek-V4-Flash (UD-Q8_K_XL, 4x V100-16GB, experts on CPU), 54,564-token prompt:

	master	this PR
prompt processing	147.1 t/s	178.4 t/s
token generation @ 55K ctx	12.88 t/s	13.65 t/s
GPU time per 1024-token ubatch	1069 ms + 54.1 ms per 1K of context	1193 ms + 2.65 ms per 1K of context

i.e. the attention cost stops growing with context length; what is left is the (CPU-side) expert matmul.

Quality, llama-perplexity --kl-divergence, wikitext-2, -c 32768 --chunks 2 (the scored half of each chunk is past the point where the sparse path engages):

• mean KLD 0.0046, top-1 agreement 97.4%, PPL 3.9660 vs 3.9689 on master (±0.048).
• This is not lost information: running the same <512, 512, 1, 32> instance without the sparse gather gives KLD 0.000000 / 100% agreement, and an instrumented compaction kernel confirms no row ever has more unmasked entries than n_kv_max, so nothing is dropped. The deviation is the model's own sensitivity to summation order in its top-k routing: an exact-f32, mathematically equivalent reordering elsewhere in the graph produces the same 0.0043. For reference, Q8_0 -> Q4_K on this model is KLD 0.0102.

Notes

• One extra template instance is compiled for every target architecture, not just Volta. Happy to gate it differently if that is a concern.
• Runs on Turing and newer are unaffected: all new code paths are behind cc == GGML_CUDA_CC_VOLTA.

---

🤖 Generated with Claude Code

https://claude.ai/code/session_01FBqdxhtLYCXWCqbGx6LJuh

[ggml-gh-bot]

Hi @CoREse, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• PR Template not respected: Please respect the template when creating a new pull request. Make sure to fill out all required sections.

• AI-generated content: While code is allowed to be generated by AI, please write the PR description and commit messages on your own without the help of AI.

---

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

[taronaeo]

https://github.com/ggml-org/llama.cpp/blob/master/AGENTS.md#prohibited-ai-usage-results-in-immediate-pr-closure

[yudagua]

+1 to reopen — 7×Volta rig here (Quadro GV100 + 6×Tesla PG503-216, 32GB HBM2, sm_70). Running DeepSeek-V4-Flash GGUF at 100K–262K context on 2.41.0: the CSA layers attend over the full compressed KV exactly as described, and PP at long context is painful (~100 tok/s). Sparse FA on Volta would be a massive win for this hardware class. Happy to beta-test — I have reproducible long-context workloads and can run instrumented builds if that helps move this forward. Also relevant: our gridDim.y investigation (#27901/#28175) shows DSV4's indexer softmax hits the same 65535 grid cap at ≥262K ctx, so sparse-FA + the norm clamp would jointly unlock long-context Volta for DSV4.