Misc. bug: Performance regression: CUDA sparse flash attention decode 1.6x slower (b11047 -> b11062)

[goodbadwolf]

Name and Version

version: 0.4.1-dev (build 11062, commit 3cf03257f)
built with GNU 13.3.0 for Linux x86_64

Compared against build 11047, commit b23701f77.

Operating systems

Linux

Which llama.cpp modules do you know to be affected?

ggml (CUDA backend)

Command line

test-backend-ops perf -b CUDA0 -o FLASH_ATTN_EXT

Built with:

-DGGML_CUDA=ON -DGGML_CUDA_FA=ON -DGGML_CUDA_FA_ALL_QUANTS=ON -DGGML_CUDA_GRAPHS=ON -DCMAKE_CUDA_ARCHITECTURES=120

Problem description & steps to reproduce

FLASH_ATTN_EXT on the sparse path costs 64 percent more at batch one after
#28770, median over the 11 sparse configurations
test-backend-ops generates. The worst is 113 percent. Dense shapes are untouched. RTX PRO 6000 Blackwell
Workstation Edition, sm_120, driver 615.71.09, CUDA 13.4.

Run the command above at 3cf03257f and at b23701f77 and compare the configurations whose n_kv_max is not 0.
Use console output, since --output csv drops the timings.

Every sparse configuration, us/run:

hsk	nh	kv	n_kv_max	b11047	b11062	delta
512	1	49152	2048	25.80	55.01	+113%
512	1	32768	512	18.29	37.58	+106%
576	1	49152	2048	29.17	58.83	+102%
576	1	32768	512	21.78	41.60	+91%
512	1	16384	512	13.58	23.27	+71%
256	2	32768	2048	22.55	36.97	+64%
576	1	16384	512	17.12	27.14	+59%
256	2	16384	2048	15.03	22.53	+50%
512	1	4096	512	10.21	12.64	+24%
576	1	4096	512	13.59	16.18	+19%
256	2	4096	2048	16.45	11.81	-28%

Dense shapes over the same range moved 0.1 percent median on 58 decode configurations and 0.3 percent on 29 batched
ones. Nothing outside the sparse path changed.

Where the time goes. Dividing the added time by the mask width gives a constant, within each head size and across
a twelve-fold range of kv:

hsk	nh	kv	n_kv_max	added us	ns per mask column
512	1	4096	512	+2.43	0.593
512	1	16384	512	+9.69	0.591
512	1	32768	512	+19.29	0.589
512	1	49152	2048	+29.21	0.594
576	1	4096	512	+2.59	0.632
576	1	16384	512	+10.02	0.612
576	1	32768	512	+19.82	0.605
576	1	49152	2048	+29.66	0.603

The cost tracks kv and ignores n_kv_max, including across the pair at n_kv_max 512 and 2048 that read 0.589 and
0.594. The attention kernel's work is bounded by n_kv_max, so it cannot produce this. The mask scan reads exactly
kv columns per list, so it can. I have not profiled the two kernels separately, so this is inference from the
scaling rather than a direct attribution.

It points at the new inner loop in flash_attn_mask_to_sparse_indices:

for (int q = 0; q < q1 - q0 && !selected; ++q) {
    selected = i < ne30 && isfinite(__half2float(mask[q*s31 + i]));
}

where it used to be a single test. The trip count is a runtime value and the exit is data dependent, so the
#pragma unroll over values_per_lane around it no longer yields straight line code. At ncols1 == 1 the body
always runs once and the union reduces to the old single-query test, so specialising that case, or templating the
kernel on ncols1, would recover it without touching the batched path the PR was written for.

The one configuration that improved, kv=4096 at n_kv_max=2048, fits the same reading. The new per-list count
bounding kb0_stop is a real win, and at that shape it is larger than the scan cost.

Scope. test-backend-ops only generates sparse configurations at nb=1, so this says nothing about the batched
sparse path. The models reaching it are the ones calling ggml_flash_attn_ext_set_n_kv_max, currently qwen4exp and
deepseek4. A server decoding at batch one, meaning one slot and no draft model, sits on the regressed shape for
every token.

First Bad Commit

3cf03257f (#28770, "CUDA: enable sparse fa for qwen4").

15 commits separate b23701f77 from 3cf03257f, and only 3cf03257f touches ggml/src/ggml-cuda/ at all. The rest
are Metal, Hexagon, server, chat and UI changes.

I measured #28536 separately, as the other CUDA flash attention
change in the neighbourhood. It costs 0.7 percent on these same configurations, so it is not involved.

[ServeurpersoCom]

I reproduced the regression on my RTX PRO 6000 and measured the same drop, this small patch restores the performance on master, could you give it a try?

diff --git a/ggml/src/ggml-cuda/fattn.cu b/ggml/src/ggml-cuda/fattn.cu
index 7098c8b4ca..1d63b7cd10 100644
--- a/ggml/src/ggml-cuda/fattn.cu
+++ b/ggml/src/ggml-cuda/fattn.cu
@@ -7,10 +7,11 @@

 #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
 // one list per group of ncols1 queries: a column is selected if any query of the group can see it
+template <int ncols1>
 __launch_bounds__(256, 1)
 static __global__ void flash_attn_mask_to_sparse_indices(
         const half * mask_ptr, int32_t * indices_ptr, int32_t * counts_ptr, const int ne30, const int n_queries,
-        const int ncols1, const int n_kv_max, const int64_t s31, const int64_t s33) {
+        const int n_kv_max, const int64_t s31, const int64_t s33) {
     ggml_cuda_pdl_sync();

     constexpr int values_per_lane = 8;
@@ -42,8 +43,11 @@ static __global__ void flash_attn_mask_to_sparse_indices(
         for (int item = 0; item < values_per_lane; ++item) {
             const int i = i0 + (warp*values_per_lane + item)*WARP_SIZE + lane;
             bool selected = false;
-            for (int q = 0; q < q1 - q0 && !selected; ++q) {
-                selected = i < ne30 && isfinite(__half2float(mask[q*s31 + i]));
+            if (i < ne30) {
+#pragma unroll
+                for (int q = 0; q < ncols1; ++q) {
+                    selected |= (ncols1 == 1 || q < q1 - q0) && isfinite(__half2float(mask[q*s31 + i]));
+                }
             }
             selected_warp[item] = __ballot_sync(0xFFFFFFFF, selected);
             warp_count += __popc(selected_warp[item]);
@@ -110,8 +114,10 @@ void ggml_cuda_flash_attn_ext_compact_mask(
     const dim3 blocks_num((n_queries + ncols1 - 1)/ncols1, mask->ne[3], 1);
     const dim3 block_dim(256, 1, 1);
     const ggml_cuda_kernel_launch_params launch_params(blocks_num, block_dim, 0, stream);
-    ggml_cuda_kernel_launch(flash_attn_mask_to_sparse_indices, launch_params,
-        (const half *) mask->data, indices, counts, int(mask->ne[0]), n_queries, ncols1, n_kv_max, s31, s33);
+    GGML_ASSERT(ncols1 == 1 || ncols1 == 8);
+    const auto kernel = ncols1 == 1 ? flash_attn_mask_to_sparse_indices<1> : flash_attn_mask_to_sparse_indices<8>;
+    ggml_cuda_kernel_launch(kernel, launch_params,
+        (const half *) mask->data, indices, counts, int(mask->ne[0]), n_queries, n_kv_max, s31, s33);
     CUDA_CHECK(cudaGetLastError());
 #endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
 }

a0d1bb6

[shenron0101]

I also managed to recreate the issue and created a small fix here.

Results

Mask kernel on the worst config: 65.2 → 35.5 us (old was 33.2).

Results

Mask kernel on the worst config: 65.2 -> 35.5 us (old was 33.2).

All 12 sparse configs, us/run:

hsk	nh	kv	n_kv_max	old	master	fixed	vs master
512	1	4096	512	14.70	19.66	17.26	-12%
576	1	4096	512	17.34	19.64	17.03	-13%
256	2	4096	2048	19.31	21.88	17.14	-22%
512	1	16384	512	19.93	30.89	22.87	-26%
576	1	16384	512	24.02	31.09	23.38	-25%
256	2	16384	2048	48.15	35.45	25.49	-28%
512	1	32768	512	27.43	46.02	31.24	-32%
576	1	32768	512	30.03	46.63	32.11	-31%
256	2	32768	2048	89.60	57.68	40.00	-31%
512	1	49152	2048	49.44	78.38	49.93	-36%
512	2	49152	2048	n/a	108.61	73.43	-32%
576	1	49152	2048	47.97	78.49	51.87	-34%

[ggerganov]

Can you verify that #29298 fixes the issue?

[ServeurpersoCom]

Can you verify that #29298 fixes the issue?

I tested #29298 on my RTX PRO 6000 and it does not fix it: sparse decode timings are identical to master and nsys still shows the mask scan at 46 us out of 58 at 49k, since the PR only changes the batched gating while the regression comes from the runtime query loop in flash_attn_mask_to_sparse_indices, which a0d1bb6 fixes (26 us, back to the b11047 level).

[ServeurpersoCom]

I can also measure where the gain actually comes from (compile time unroll, the early exit or the bounds check) to trim the patch down and open a clean PR

[am17an]

I will fix and push a change to #29298

[shenron0101]

Yup the speed changes from running the sparse attention kernel itself.

When llama.cpp runs flash attention, two things happen in sequence:

1. Decide whether to use sparse attention for this shape at all (a cost-estimate gate)
2. If yes, build the compact index list, then run the sparse attention kernel

#29298 touches step 1.

[ServeurpersoCom]

I measured each change separately and the compile time unroll alone accounts for the whole gain, the early exit and the bounds check barely matter on their own, since a known trip count turns the batch one scan into straight line code with its 8 loads issued together.

[am17an]

@ServeurpersoCom you can push your change to 29298, I verified it fixes it for me too.

[sssr5hsssdssqdse]

Thank you, I will review this.

…

On Thu, 24 Sept 2026 at 10:16, Aman Gupta ***@***.***> wrote: Closed #29281 <#29281> as completed. — Reply to this email directly, view it on GitHub <#29281?email_source=notifications&email_token=CJHCJ7732B3C77URTBTIX3T5QTQ7NA5CNFSNUABQM5UWIORPF5TWS5BNNB2WEL2JONZXKZKFOZSW45CON52GSZTJMNQXI2LPNYXTGMJXGQZTENRUGY3DJJTSMVQXG33OVJZXKYTTMNZGSYTFMSSWK5TFNZ2KYZTPN52GK4S7MNWGSY3L#event-31743264664>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/CJHCJ77SDP6G4N6XRFG6J335QTQ7NAVCNFSNUABFKJSXA33TNF2G64TZHM3DCMRTGU2DOOBUHNEXG43VMU5TKNJUGMYTQMRTG432C5QC> . Triage notifications, keep track of coding agent tasks and review pull requests on the go with GitHub Mobile for iOS <https://github.com/notifications/mobile/ios/CJHCJ73NY5OKSJDCTW3S5O35QTQ7NA5CNFSNUABQM5UWIORPF5TWS5BNNB2WEL2JONZXKZKFOZSW45CON52GSZTJMNQXI2LPNYXTGMJXGQZTENRUGY3DJJTSMVQXG33OVJZXKYTTMNZGSYTFMSSWK5TFNZ2KUZTPN52GK4S7NFXXG> and Android <https://github.com/notifications/mobile/android/CJHCJ72VMASVGRB7WYM722D5QTQ7NA5CNFSNUABQM5UWIORPF5TWS5BNNB2WEL2JONZXKZKFOZSW45CON52GSZTJMNQXI2LPNYXTGMJXGQZTENRUGY3DJJTSMVQXG33OVJZXKYTTMNZGSYTFMSSWK5TFNZ2K4ZTPN52GK4S7MFXGI4TPNFSA>. Download it today! You are receiving this because you are subscribed to this thread.Message ID: ***@***.***>