Summary
This issue documents a systematic investigation into Volta-class GPU (SM70/SM72) flash attention performance in llama.cpp, motivated by 1Cat-vLLM's FLASH_ATTN_V100 backend which achieves significant improvements on V100 via TurboMind SM70 WMMA kernels and split-D attention design.
Benchmarking was performed on Jetson AGX Xavier (SM72, aarch64, CUDA 11.4) with Qwen3.6-35B-A3B MoE D=256 GGUF models.
Key Findings
Current Config is Already Reasonably Tuned
Attempting custom Volta flash attention configs (doubling nbatch_K2/nbatch_V2, adjusting occupancy) showed no improvement — in fact, decode throughput regressed ~5% due to increased shared memory pressure reducing occupancy.
| Scenario |
Current (Ampere fallback) |
Custom Volta config |
| Pure decode |
20.85 t/s |
19.68 t/s |
| 256 gen |
20.85 t/s |
21.08 t/s |
MMA Architecture Gap is Fundamental
The ~2× throughput gap between Volta (m8n8k4, no ldmatrix, no cp.async) and Turing/Ampere (m16n8k8, ldmatrix, cp.async) is architectural and cannot be closed by config tuning alone.
No Code Change Required Now
The existing dispatch logic in fattn.cu (Volta branching at line 491) and Ampere config fallback in ggml_cuda_fattn_mma_get_config_volta() at fattn-mma-f16.cuh:123 are already producing near-optimal results for this kernel design.
Potential Future Optimization Paths
-
Volta-specific split-D WMMA attention kernel (following 1Cat-vLLM's FLASH_ATTN_V100 design): Split D=256 into 4×64-dim slices processed by separate warps, reducing per-warp shared memory from 32KB to 8KB. Estimated gain: ~35% on SM72.
-
MTP speculative decoding support for Qwen3.x: llama.cpp already supports --draft-model, but Qwen3.x's native MTP heads are different — they're model-integrated, not a separate draft model. Adding MTP head support would benefit all architectures.
-
Persistent thread block design: Keep warps alive across KV slices instead of re-launching kernels, with pipelined loads into shared memory.
Next Steps
If maintainers are interested, I can develop a Volta-specific flash attention CUDA kernel following approach #1. The work would include:
- New kernel file (e.g.,
fattn-volta.cu)
- WMMA integration via
nvcuda::wmma (already used in mma.cuh)
- Test coverage via
test-backend-ops
- Benchmark data on AGX Xavier (SM72) and, if provided access, V100 (SM70)
Benchmark Environment
- GPU: Jetson AGX Xavier (SM72, Volta)
- CUDA: 11.4, aarch64
- Model: Qwen3.6-35B-A3B MoE (Q4_K_M, 20GB GGUF)
- llama.cpp v0.3.0-dev
Summary
This issue documents a systematic investigation into Volta-class GPU (SM70/SM72) flash attention performance in llama.cpp, motivated by 1Cat-vLLM's FLASH_ATTN_V100 backend which achieves significant improvements on V100 via TurboMind SM70 WMMA kernels and split-D attention design.
Benchmarking was performed on Jetson AGX Xavier (SM72, aarch64, CUDA 11.4) with Qwen3.6-35B-A3B MoE D=256 GGUF models.
Key Findings
Current Config is Already Reasonably Tuned
Attempting custom Volta flash attention configs (doubling nbatch_K2/nbatch_V2, adjusting occupancy) showed no improvement — in fact, decode throughput regressed ~5% due to increased shared memory pressure reducing occupancy.
MMA Architecture Gap is Fundamental
The ~2× throughput gap between Volta (m8n8k4, no ldmatrix, no cp.async) and Turing/Ampere (m16n8k8, ldmatrix, cp.async) is architectural and cannot be closed by config tuning alone.
No Code Change Required Now
The existing dispatch logic in
fattn.cu(Volta branching at line 491) and Ampere config fallback inggml_cuda_fattn_mma_get_config_volta()atfattn-mma-f16.cuh:123are already producing near-optimal results for this kernel design.Potential Future Optimization Paths
Volta-specific split-D WMMA attention kernel (following 1Cat-vLLM's FLASH_ATTN_V100 design): Split D=256 into 4×64-dim slices processed by separate warps, reducing per-warp shared memory from 32KB to 8KB. Estimated gain: ~35% on SM72.
MTP speculative decoding support for Qwen3.x: llama.cpp already supports
--draft-model, but Qwen3.x's native MTP heads are different — they're model-integrated, not a separate draft model. Adding MTP head support would benefit all architectures.Persistent thread block design: Keep warps alive across KV slices instead of re-launching kernels, with pipelined loads into shared memory.
Next Steps
If maintainers are interested, I can develop a Volta-specific flash attention CUDA kernel following approach #1. The work would include:
fattn-volta.cu)nvcuda::wmma(already used inmma.cuh)test-backend-opsBenchmark Environment