CUDA: XOR swizzle flash attn K,V smem fp16 tiles by ynankani · Pull Request #25635 · ggml-org/llama.cpp

ynankani · 2026-07-13T19:23:11Z

Overview

This PR adds XOR swizzling for K/V shared-memory tiles in the CUDA flash-attention fp16 MMA kernel (fattn-mma-f16.cuh), replacing row padding (+4) as the primary fix for shared-memory bank conflicts on cp.async stores and ldmatrix loads. On Turing+ (Ampere/Ada/Blackwell), when the K/V tile stride is a power-of-two ≥ 32, the kernel uses a pow2 stride with per-row XOR address remapping instead of padded row stride.

Additional information

Followed resource : https://lubits.ch/flash/Part-4

This change helped in improving perf for very high context, collected data for 65K depth

Performance DGX -SPARK

Model	Metric	ToT	Swizzle	Δ
Gemma-4-26B-A4B-NVFP4	pp2048 @ d65536	1616.61 ± 2.94	1644.18 ± 1.80	1.71%
Gemma-4-26B-A4B-NVFP4	tg128 @ d65536	29.30 ± 0.04	30.32 ± 0.03	3.48%
Qwen3.6-27B-Q4_K_M	pp2048 @ d65536	574.91 ± 1.19	583.02 ± 1.44	1.41%
Qwen3.6-27B-Q4_K_M	tg128 @ d65536	9.43 ± 0.01	9.76 ± 0.01	3.50%
Qwen3.6-27B-NVFP4	pp2048 @ d65536	648.32 ± 0.77	665.23 ± 0.60	2.61%
Qwen3.6-27B-NVFP4	tg128 @ d65536	8.29 ± 0.00	8.53 ± 0.01	2.89%
Qwen3.6-35B-A3B-NVFP4	pp2048 @ d65536	1729.35 ± 8.33	1758.57 ± 3.19	1.69%
Qwen3.6-35B-A3B-NVFP4	tg128 @ d65536	49.22 ± 0.20	52.39 ± 0.15	6.44%
Qwen3.6-35B-A3B-UD-Q4_K_M	pp2048 @ d65536	1662.27 ± 4.71	1680.99 ± 3.85	1.13%
Qwen3.6-35B-A3B-UD-Q4_K_M	tg128 @ d65536	46.24 ± 0.20	48.48 ± 0.63	4.84%

Performance RTX6000-PRO-BLACKWELL

Model	Metric	ToT (t/s)	Swizzle (t/s)	Δ
Qwen3.6-35B-A3B NVFP4 (W4A16)	pp2048 @ d65536	5105.62 ± 8.83	5213.21 ± 7.96	2.11%
Qwen3.6-35B-A3B NVFP4 (W4A16)	tg128 @ d65536	191.42 ± 3.47	193.00 ± 2.71	0.83%
Qwen3.6-35B-A3B-UD-Q4_K_M	pp2048 @ d65536	4757.63 ± 13.50	4947.80 ± 9.65	4.00%
Qwen3.6-35B-A3B-UD-Q4_K_M	tg128 @ d65536	197.44 ± 2.57	197.45 ± 2.75	0.01%
Qwen3.6-27B-NVFP4	pp2048 @ d65536	1805.03 ± 4.16	1868.25 ± 2.47	3.50%
Qwen3.6-27B-NVFP4	tg128 @ d65536	53.91 ± 0.20	53.71 ± 0.23	−0.37%
Qwen3.6-27B-Q4_K_M	pp2048 @ d65536	1518.74 ± 1.59	1630.54 ± 5.56	7.36%
Qwen3.6-27B-Q4_K_M	tg128 @ d65536	54.70 ± 0.21	55.91 ± 0.15	2.21%
Gemma-4-26B-A4B-NVFP4	pp2048 @ d65536	4997.70 ± 36.84	5213.21 ± 7.96	4.31%
Gemma-4-26B-A4B-NVFP4	tg128 @ d65536	144.94 ± 0.51	143.93 ± 0.87	−0.70%

Requirements

I have read and agree with the contributing guidelines
AI usage disclosure: Yes, paired with Claude on this

naamfung · 2026-07-15T10:17:16Z

Bug Report (for PR #25635)

Title: CUDA illegal memory access during warmup on Ampere (RTX 3060 Ti, CC 8.6)

Environment:

GPU: NVIDIA RTX 3060 Ti (Ampere, CC 8.6, 8 GB)
CUDA Toolkit: 12.8
OS: Windows / MSVC 19.44
Build flags: GGML_CUDA=ON, GGML_CUDA_FA=ON, GGML_CUDA_FA_ALL_QUANTS=ON, CMAKE_CUDA_ARCHITECTURES=860
Fork based on mainline b9899

Model: Qwen-AgentWorld/Qwen3.5-35B-A3B (hybrid architecture: recurrent Gated DeltaNet + SWA layers, GGUF Q8_0)

Reproduction:

llama-server -m <hybrid_model.gguf> --ctx-size 131072 --flash-attn on \
  --spec-type draft-mtp --spec-draft-n-max 4 -b 16384 -ub 256 \
  --cache-type-k q8_0 --cache-type-v q8_0

Symptom: Server crashes during warmup ("warming up the model with an empty run") with the following errors:

0.15.168.379 E CUDA error: an illegal memory access was encountered
0.15.168.386 E   current device: 0, in function ggml_cuda_kernel_can_use_pdl at D:\Programs\llama-cpp-repos\laamaafung\ggml\src\ggml-cuda\common.cuh:1602
0.15.168.387 E   cudaFuncGetAttributes(&attr, kernel)

(with GGML_CUDA_PDL=0):

0.17.265.491 E CUDA error: an illegal memory access was encountered
0.17.265.779 E   current device: 0, in function ggml_cuda_kernel_launch at D:\Programs\llama-cpp-repos\laamaafung\ggml\src\ggml-cuda\common.cuh:1639
0.17.266.020 E D:\Programs\llama-cpp-repos\laamaafung\ggml\src\ggml-cuda\ggml-cuda.cu:106: CUDA error
  cudaGetLastError()

Analysis:

The XOR swizzle path in fattn-mma-f16.cuh (added by this PR) triggers the crash when flash_attn_ext_f16 is launched.
GGML_CUDA_PDL=0 does not fix the crash, confirming the kernel execution itself fails (not the PDL attribute query).
Reverting only the fattn-mma-f16.cuh changes from this PR (restoring the +4 row padding) and removing fattn-swizzle.cuh resolves the crash on my fork.

Scope:

Observed on Ampere (CC 8.6).
Reproducible with and without --swa-full; independent of MTP.

ynankani · 2026-07-15T11:49:59Z

Bug Report (for PR #25635)

Title: CUDA illegal memory access during warmup on Ampere (RTX 3060 Ti, CC 8.6)

Environment:

GPU: NVIDIA RTX 3060 Ti (Ampere, CC 8.6, 8 GB)

CUDA Toolkit: 12.8

OS: Windows / MSVC 19.44

Build flags: GGML_CUDA=ON, GGML_CUDA_FA=ON, GGML_CUDA_FA_ALL_QUANTS=ON, CMAKE_CUDA_ARCHITECTURES=860

Fork based on mainline b9899

Model: Qwen-AgentWorld/Qwen3.5-35B-A3B (hybrid architecture: recurrent Gated DeltaNet + SWA layers, GGUF Q8_0)

Reproduction:
llama-server -m <hybrid_model.gguf> --ctx-size 131072 --flash-attn on \
  --spec-type draft-mtp --spec-draft-n-max 4 -b 16384 -ub 256 \
  --cache-type-k q8_0 --cache-type-v q8_0
Symptom: Server crashes during warmup ("warming up the model with an empty run") with the following errors:
0.15.168.379 E CUDA error: an illegal memory access was encountered
0.15.168.386 E   current device: 0, in function ggml_cuda_kernel_can_use_pdl at D:\Programs\llama-cpp-repos\laamaafung\ggml\src\ggml-cuda\common.cuh:1602
0.15.168.387 E   cudaFuncGetAttributes(&attr, kernel)
(with GGML_CUDA_PDL=0):
0.17.265.491 E CUDA error: an illegal memory access was encountered
0.17.265.779 E   current device: 0, in function ggml_cuda_kernel_launch at D:\Programs\llama-cpp-repos\laamaafung\ggml\src\ggml-cuda\common.cuh:1639
0.17.266.020 E D:\Programs\llama-cpp-repos\laamaafung\ggml\src\ggml-cuda\ggml-cuda.cu:106: CUDA error
  cudaGetLastError()
Analysis:

The XOR swizzle path in fattn-mma-f16.cuh (added by this PR) triggers the crash when flash_attn_ext_f16 is launched.

GGML_CUDA_PDL=0 does not fix the crash, confirming the kernel execution itself fails (not the PDL attribute query).

Reverting only the fattn-mma-f16.cuh changes from this PR (restoring the +4 row padding) and removing fattn-swizzle.cuh resolves the crash on my fork.

Scope:

Observed on Ampere (CC 8.6).

Reproducible with and without --swa-full; independent of MTP.

Thank you for reporting the issue. I will try to repro at my end and fix it

ynankani · 2026-07-16T12:02:37Z

@naamfung can you please check if the issue is resolved with the latest change on your system?

gaugarg-nv · 2026-07-20T17:10:15Z

@@ -66,7 +67,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, 128, 2,  64, 128, 128, 128, 2, true);


What is the perf impact of this change in isolation?

Collected ncu traces on different config swizzle+occupancy_update , swizzle_only, master, master+occupancy_update to isolate the impact of the change. Main idea for the change was to increase occupancy and was seeing some regression

Ncu values for Qwen3.6_35B collected on RTX6000-PRO
Decode:

cell config swizzle block occ% bc_ld bc_st dur

1 64,4 off 64 4.20 1141 370 162 us

2 128,2 off 128 8.30 478 376 161 us

3 64,4 on 64 4.20 396 188 168 us

4 128,2 on 128 8.30 376 385 163 us

Prefill:

cell config swizzle block occ% bc_ld bc_st dur

1 64,4 off 128 15.9 27,362 26,478 4.33 ms

2 128,2 off 128 15.9 27,433 26,566 4.34 ms

3 64,4 on 128 14.9 1,29,524 17,390 4.01 ms

4 128,2 on 128 14.9 1,29,963 17,654 4.02 ms

on DGX Spark:
Decode:

cell config swizzle block occ% bc_ld bc_st dur

1 64,4 off 64 4.17 764 91 0.76 ms

2 128,2 off 128 8.33 152 122 0.75 ms

3 64,4 on 64 4.17 99 48 0.63 ms

4 128,2 on 128 8.33 107 96 0.63 ms

Prefill:

cell config swizzle block occ% bc_ld bc_st dur

1 64,4 off 128 16.63 10722 8422 30.43 ms

2 128,2 off 128 16.63 10903 7359 29.35 ms

3 64,4 on 128 15.29 26097 11759 26.49 ms

4 128,2 on 128 15 27460 10403 27.36 ms

At least on DGX Spark, it seems (128,2) is worse than (64,4).
Also, bank conflicts are high with this change during prefill. Is this expected?

At least on DGX Spark, it seems (128,2) is worse than (64,4).

The change is made only for decode config, for prefill it is most probably noise.

Also, bank conflicts are high with this change during prefill. Is this expected?

it is for op_ld and op_st those values are negligible in end to end run. we should be looking at op_ldgsts(cp.async) which is the swizzle change touching for K/V fill(There is no excessive wavefront on K read and V read). Somehow in ncu i don't see that value.

Collected wavefront data on DGX Spark with CTK 13.3(JFYI i was not able to find the ldgsts bank conflict field so had to look into SASS for which took help from Claude to export the SASS and count the "L1 Wavefronts Shared, L1 Wavefronts Shared Ideal, and L1 Wavefronts Shared Excessive") and below are the observations :

Decode:

cell config swizzle cp.async wf cp.async ideal xideal

1 64,4 off 2.8M 1.1M 2.62x

2 128,2 off 2.8M 1.1M 2.61x

3 64,4 on 1.1M 1.1M 1.00x

4 128,2 on 1.1M 1.1M 1.03x

Prefill:

cell config swizzle cp.async wf cp.async ideal xideal

1 64,4 off 40.6M 14.7M 2.76x

2 128,2 off 40.6M 14.7M 2.76x

3 64,4 on 15.0M 14.4M 1.04x

4 128,2 on 15.0M 14.4M 1.04x

According to these numbers, please also change the config for ncols == 16 and apply the same changes for the configs with DKQ == 192 && DV == 128.

gaugarg-nv · 2026-07-20T17:12:24Z

+    for (int kv : { 4096, 8192, 16384,32768, 65536, }) {
+        for (int hs : { 64, 128, 256, }) {
+            for (int nr : { 1, 4, 8, }) {
+                for (int nb : { 1, 4096, }) {
+                    test_cases.emplace_back(new test_flash_attn_ext(hs, hs, 8, {nr, 1}, kv, nb, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16));
+                }


You should reduce the number of tests. This will take too long to run.

Seems these are perf tests, not eval. If ppl want, they can filter with -p . But we can also reduce them 🤷

gaugarg-nv · 2026-07-20T17:23:09Z

Can you do a perf sweep for different depth values?

naamfung · 2026-07-21T12:54:39Z

Sorry, I just saw the message. I'll merge and test it right away. Please wait.

naamfung · 2026-07-21T13:25:23Z

8119/8119 tests passed
Backend CUDA0: OK
2/2 backends passed
OK

naamfung · 2026-07-21T13:57:46Z

Startup is working without errors now, but no performance improvement is visible. Maybe the variation is just normal fluctuation? :) It's likely that further optimization potential is limited under the Ampere architecture with small tiles.

ynankani · 2026-07-21T14:20:57Z

Startup is working without errors now, but no performance improvement is visible. Maybe the variation is just normal fluctuation? :) It's likely that further optimization potential is limited under the Ampere architecture with small tiles.

Thank you for confirming the fix is working.
you are running Q8_0 kv and Q8_0 model, you need to rebuild with GGML_CUDA_FA_ALL_QUANTS=ON . Can you please try with that?

naamfung · 2026-07-21T14:29:09Z

Startup is working without errors now, but no performance improvement is visible. Maybe the variation is just normal fluctuation? :) It's likely that further optimization potential is limited under the Ampere architecture with small tiles.

Thank you for confirming the fix is working. you are running Q8_0 kv and Q8_0 model, you need to rebuild with GGML_CUDA_FA_ALL_QUANTS=ON . Can you please try with that?

The build has GGML_CUDA_FA_ALL_QUANTS=ON enabled, but on Ampere with small tiles, the performance gain may fall within statistical noise.

naamfung · 2026-07-21T14:55:48Z

Thanks for the fix. The 64-bit generic pointer commit (de97e22) successfully
resolves the illegal-memory-access crash on Ampere (CC 8.6). However, after
re-running benchmarks on an RTX 3060 Ti, I observed a significant regression
on tg128 that forced me to revert the PR again.

Environment:

GPU: NVIDIA GeForce RTX 3060 Ti, compute capability 8.6, 8 GiB VRAM
CUDA: 12.8, driver 610.62
Build: GGML_CUDA=ON, GGML_CUDA_FA=ON, GGML_CUDA_FA_ALL_QUANTS=ON,
CMAKE_CUDA_ARCHITECTURES=86-real, MSVC 19.44.35228.0
Model: Qwen-AgentWorld-35B-A3B (Q4_K - Medium, 35.51 B params, 15.89 GiB)
- GQA architecture, 4 KV heads
Tool: llama-bench, -p 2048 -n 128 -r 10 -fa 1

Results (10 repetitions, mean +/- std):

test	swizzle (ac30012d1)	baseline +4 pad (0a46778a8)	delta
pp2048	110.98 +/- 0.67 t/s	110.83 +/- 0.60 t/s	+0.14%
tg128	7.72 +/- 0.01 t/s	8.24 +/- 0.02 t/s	-6.31%

The pp2048 delta is within statistical noise (95% CIs fully overlap), but the
tg128 regression is clearly significant: the 95% CIs do not overlap
([7.70, 7.74] vs [8.20, 8.28]), and the gap is 25-50x the standard deviation.

This directionally matches the -0.37% tg128 regression you reported for
Qwen3.6-27B-NVFP4 on RTX6000-PRO-BLACKWELL, but the magnitude on Ampere +
Qwen GQA is much larger (-6.31% vs -0.37%).

Potential cause (speculative):

The 64-bit generic pointer ("l" constraint) consumes one extra register
compared to the 32-bit shared pointer ("r" constraint). On Ampere with
the smaller 16x8 fp16 tile (I=16, J=8), this may trigger register spilling
that degrades the fattn-mma kernel quality, particularly affecting the
decode-bound tg128 path while leaving the compute-bound pp2048 path
essentially unchanged.

The correctness fix is real and appreciated, but on this GPU + model
combination the net effect is negative for the generation workload that
matters most in practice. Happy to rerun any additional tests if helpful.

ORippler · 2026-07-21T16:55:31Z

+// XOR swizzle for K/V SMEM tiles to avoid bank conflicts without row padding (Turing+ only).
+// Stride must be a power-of-two >= 32 half2 columns,otherwise we keep +4 row padding.


Swizzling should be re-usable for GEMMS in mma.cuh, what are the steps needed to achieve/generalize this?

ORippler · 2026-07-21T17:10:37Z

+    for (int kv : { 4096, 8192, 16384,32768, 65536, }) {
+        for (int hs : { 64, 128, 256, }) {
+            for (int nr : { 1, 4, 8, }) {
+                for (int nb : { 1, 4096, }) {
+                    test_cases.emplace_back(new test_flash_attn_ext(hs, hs, 8, {nr, 1}, kv, nb, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16));
+                }


Seems these are perf tests, not eval. If ppl want, they can filter with -p . But we can also reduce them 🤷

ORippler · 2026-07-21T17:28:57Z

tg128 regression is clearly significant: the 95% CIs do not overlap
([7.70, 7.74] vs [8.20, 8.28]), and the gap is 25-50x the standard deviation.

Hmm you seem to be running both CPU offloading of some sort (16 GB weights on 8 GB GPU) + on a private fork. Can you:

Try reproducing locally on master vs this pr?
Execute the test-backend-ops perf -o FLASH_ATTN? This would help root-cause and isolate CPU-side-issues vs. GPU-side-issues

ynankani · 2026-07-21T17:45:09Z

tg128 regression is clearly significant: the 95% CIs do not overlap
([7.70, 7.74] vs [8.20, 8.28]), and the gap is 25-50x the standard deviation.

Hmm you seem to be running both CPU offloading of some sort (16 GB weights on 8 GB GPU) + on a private fork. Can you:

Try reproducing locally on master vs this pr?

Execute the test-backend-ops perf -o FLASH_ATTN? This would help root-cause and isolate CPU-side-issues vs. GPU-side-issues

I tried the second "Execute the test-backend-ops perf -o FLASH_ATTN" and i see perf benefit of 15% for hsk=256 and ~1-4% perf improvement for other shapes on 3060 Ti. @naamfung can you please confirm on your end as well?

naamfung · 2026-07-21T22:39:52Z

Please wait..

naamfung · 2026-07-22T00:06:44Z

Thanks for the pointer on GGML_CUDA_FA_ALL_QUANTS. I've rerun the benchmarks
on a clean upstream checkout (no private-fork modifications) to isolate the
PR's effect. Two builds were compared from the same base commit 3f08ef2,
differing only by this PR's two commits (4921f16a6 + de97e22f6).

Environment:

GPU: NVIDIA GeForce RTX 3060 Ti, compute capability 8.6, 8 GiB VRAM
CUDA: 12.8, driver 610.62
Build: GGML_CUDA=ON, GGML_CUDA_FA=ON, GGML_CUDA_FA_ALL_QUANTS=ON,
CMAKE_CUDA_ARCHITECTURES=86-real, MSVC 19.44.35228.0
Model: Qwen-AgentWorld-35B-A3B (Q4_K - Medium, 35.51 B params, 15.89 GiB)
Tool: llama-bench, -p 2048 -n 128 -r 10 -fa 1

Results (10 repetitions, mean +/- std):

test	baseline `3f08ef2`	PR de97e22f6	delta
pp2048	105.09 +/- 0.56 t/s	107.10 +/- 0.60	+1.91%
tg128	7.61 +/- 0.03 t/s	7.62 +/- 0.00	+0.13%

The pp2048 improvement is statistically significant (95% CIs of the mean do
not overlap). The tg128 delta is within noise.

Earlier I reported a -6.31% tg128 regression, but that was measured on a
private fork with additional patches.
On a clean upstream tree the regression disappears and a small pp2048 gain
is visible, so the prior regression was caused by interaction with local
modifications rather than this PR. Sorry for the noise on the earlier
report.

The 64-bit generic pointer fix (de97e22) also resolves the illegal-memory-
access crash on Ampere CC 8.6 that the earlier version of this PR triggered.

ynankani · 2026-07-22T10:54:13Z

Can you do a perf sweep for different depth values?

Performance sweep DGX Spark

Model	Kind	Depth	Master t/s	Swizzle t/s	Delta %
Qwen3.6-27B NVFP4	PP2048	0	982.53	998.12	1.6
Qwen3.6-27B NVFP4	TG128	0	10.03	10.12	0.9
Qwen3.6-27B NVFP4	PP2048	1024	967.83	984.68	1.7
Qwen3.6-27B NVFP4	TG128	1024	9.99	10.09	1
Qwen3.6-27B NVFP4	PP2048	2048	960.7	973.65	1.3
Qwen3.6-27B NVFP4	TG128	2048	9.96	10.06	1
Qwen3.6-27B NVFP4	PP2048	4096	946.79	958.67	1.3
Qwen3.6-27B NVFP4	TG128	4096	9.9	10	1
Qwen3.6-27B NVFP4	PP2048	8192	915.05	936.46	2.3
Qwen3.6-27B NVFP4	TG128	8192	9.74	9.88	1.4
Qwen3.6-27B NVFP4	PP2048	16384	863.01	881.53	2.1
Qwen3.6-27B NVFP4	TG128	16384	9.49	9.67	1.9
Qwen3.6-27B NVFP4	PP2048	32768	775.53	796.36	2.7
Qwen3.6-27B NVFP4	TG128	32768	9.06	9.27	2.3
Qwen3.6-27B NVFP4	PP2048	65536	644.1	663.42	3
Qwen3.6-27B NVFP4	TG128	65536	8.25	8.55	3.6
Gemma-4-26B NVFP4	PP2048	0	2856.53	2873.62	0.6
Gemma-4-26B NVFP4	TG128	0	38.31	38.64	0.9
Gemma-4-26B NVFP4	PP2048	1024	2772.28	2804.68	1.2
Gemma-4-26B NVFP4	TG128	1024	36.83	37.14	0.8
Gemma-4-26B NVFP4	PP2048	2048	2783.33	2796.37	0.5
Gemma-4-26B NVFP4	TG128	2048	36.7	37.02	0.9
Gemma-4-26B NVFP4	PP2048	4096	2725.8	2733.06	0.3
Gemma-4-26B NVFP4	TG128	4096	36.36	36.76	1.1
Gemma-4-26B NVFP4	PP2048	8192	2605.02	2618.86	0.5
Gemma-4-26B NVFP4	TG128	8192	35.82	36.29	1.3
Gemma-4-26B NVFP4	PP2048	16384	2391.07	2434.72	1.8
Gemma-4-26B NVFP4	TG128	16384	34.69	35.35	1.9
Gemma-4-26B NVFP4	PP2048	32768	2037.02	2076.93	2
Gemma-4-26B NVFP4	TG128	32768	32.67	33.57	2.8
Gemma-4-26B NVFP4	PP2048	65536	1595.6	1636.96	2.6
Gemma-4-26B NVFP4	TG128	65536	29.24	30.58	4.6
Qwen3.6-35B-A3B Q4_K_M	PP2048	0	2268.26	2293.37	1.1
Qwen3.6-35B-A3B Q4_K_M	TG128	0	57.53	57.84	0.5
Qwen3.6-35B-A3B Q4_K_M	PP2048	1024	2221	2254.72	1.5
Qwen3.6-35B-A3B Q4_K_M	TG128	1024	57.3	57.41	0.2
Qwen3.6-35B-A3B Q4_K_M	PP2048	2048	2223.46	2249.53	1.2
Qwen3.6-35B-A3B Q4_K_M	TG128	2048	56.97	57.08	0.2
Qwen3.6-35B-A3B Q4_K_M	PP2048	4096	2180.52	2212.14	1.5
Qwen3.6-35B-A3B Q4_K_M	TG128	4096	56.26	56.34	0.1
Qwen3.6-35B-A3B Q4_K_M	PP2048	8192	2130.81	2169.66	1.8
Qwen3.6-35B-A3B Q4_K_M	TG128	8192	54.79	55.05	0.5
Qwen3.6-35B-A3B Q4_K_M	PP2048	16384	2058.72	2071.08	0.6
Qwen3.6-35B-A3B Q4_K_M	TG128	16384	52.21	52.86	1.2
Qwen3.6-35B-A3B Q4_K_M	PP2048	32768	1895.27	1925.98	1.6
Qwen3.6-35B-A3B Q4_K_M	TG128	32768	48.1	49.12	2.1
Qwen3.6-35B-A3B Q4_K_M	PP2048	65536	1631.24	1669.68	2.4
Qwen3.6-35B-A3B Q4_K_M	TG128	65536	41.21	43.1	4.6
Qwen3.6-35B-A3B NVFP4	PP2048	0	2395.34	2432.05	1.5
Qwen3.6-35B-A3B NVFP4	TG128	0	75.74	76.22	0.6
Qwen3.6-35B-A3B NVFP4	PP2048	1024	2372.24	2396.53	1
Qwen3.6-35B-A3B NVFP4	TG128	1024	75.38	75.79	0.5
Qwen3.6-35B-A3B NVFP4	PP2048	2048	2358.05	2389.42	1.3
Qwen3.6-35B-A3B NVFP4	TG128	2048	74.99	75.14	0.2
Qwen3.6-35B-A3B NVFP4	PP2048	4096	2325.37	2345.71	0.9
Qwen3.6-35B-A3B NVFP4	TG128	4096	73.68	73.75	0.1
Qwen3.6-35B-A3B NVFP4	PP2048	8192	2261.95	2283.41	0.9
Qwen3.6-35B-A3B NVFP4	TG128	8192	71.09	71.75	0.9
Qwen3.6-35B-A3B NVFP4	PP2048	16384	2167.79	2189.21	1
Qwen3.6-35B-A3B NVFP4	TG128	16384	66.62	68.17	2.3
Qwen3.6-35B-A3B NVFP4	PP2048	32768	1997.6	2023.93	1.3
Qwen3.6-35B-A3B NVFP4	TG128	32768	60.02	62.25	3.7
Qwen3.6-35B-A3B NVFP4	PP2048	65536	1730	1756.92	1.6
Qwen3.6-35B-A3B NVFP4	TG128	65536	49.61	52.78	6.4
Qwen3.6-27B Q4_K_M	PP2048	0	835.22	828.23	-0.8
Qwen3.6-27B Q4_K_M	TG128	0	11.94	11.95	0.1
Qwen3.6-27B Q4_K_M	PP2048	1024	825.66	819.27	-0.8
Qwen3.6-27B Q4_K_M	TG128	1024	11.9	11.92	0.2
Qwen3.6-27B Q4_K_M	PP2048	2048	821.63	818.38	-0.4
Qwen3.6-27B Q4_K_M	TG128	2048	11.86	11.88	0.2
Qwen3.6-27B Q4_K_M	PP2048	4096	810.93	804.16	-0.8
Qwen3.6-27B Q4_K_M	TG128	4096	11.78	11.79	0.1
Qwen3.6-27B Q4_K_M	PP2048	8192	786.17	789.86	0.5
Qwen3.6-27B Q4_K_M	TG128	8192	11.57	11.62	0.4
Qwen3.6-27B Q4_K_M	PP2048	16384	753.6	753.15	-0.1
Qwen3.6-27B Q4_K_M	TG128	16384	11.21	11.33	1.1
Qwen3.6-27B Q4_K_M	PP2048	32768	667.1	681.07	2.1
Qwen3.6-27B Q4_K_M	TG128	32768	10.55	10.8	2.4
Qwen3.6-27B Q4_K_M	PP2048	65536	577.59	576.74	-0.1
Qwen3.6-27B Q4_K_M	TG128	65536	9.5	9.87	3.9

Performance sweep RTX6000-PRO

Model	Kind	Depth	Master t/s	Swizzle t/s	Delta %
Qwen3.6-27B NVFP4	PP2048	0	3735.767	3743.612	0.21
Qwen3.6-27B NVFP4	TG128	0	55.043	54.891	-0.278
Qwen3.6-27B NVFP4	PP2048	1024	3727.128	3767.222	1.076
Qwen3.6-27B NVFP4	TG128	1024	54.782	54.829	0.085
Qwen3.6-27B NVFP4	PP2048	2048	3662.326	3658.267	-0.111
Qwen3.6-27B NVFP4	TG128	2048	53.994	54.304	0.574
Qwen3.6-27B NVFP4	PP2048	4096	3554.905	3609.976	1.549
Qwen3.6-27B NVFP4	TG128	4096	52.891	53.054	0.308
Qwen3.6-27B NVFP4	PP2048	8192	3369.361	3402.983	0.998
Qwen3.6-27B NVFP4	TG128	8192	52.281	52.82	1.03
Qwen3.6-27B NVFP4	PP2048	16384	3087.817	3167.764	2.589
Qwen3.6-27B NVFP4	TG128	16384	50.989	51.643	1.282
Qwen3.6-27B NVFP4	PP2048	32768	2540.674	2613.858	2.881
Qwen3.6-27B NVFP4	TG128	32768	49.027	48.95	-0.156
Qwen3.6-27B NVFP4	PP2048	65536	1736.211	1813.591	4.457
Qwen3.6-27B NVFP4	TG128	65536	45.463	45.795	0.729
Gemma-4-26B NVFP4	PP2048	0	9857.254	9874.541	0.175
Gemma-4-26B NVFP4	TG128	0	164.337	163.588	-0.456
Gemma-4-26B NVFP4	PP2048	1024	9937.03	10014.22	0.777
Gemma-4-26B NVFP4	TG128	1024	161.638	161.186	-0.28
Gemma-4-26B NVFP4	PP2048	2048	9885.243	9912.221	0.273
Gemma-4-26B NVFP4	TG128	2048	160.594	160.256	-0.211
Gemma-4-26B NVFP4	PP2048	4096	9612.903	9713.084	1.042
Gemma-4-26B NVFP4	TG128	4096	159.312	158.821	-0.308
Gemma-4-26B NVFP4	PP2048	8192	9201.472	9369.367	1.825
Gemma-4-26B NVFP4	TG128	8192	157.447	156.803	-0.41
Gemma-4-26B NVFP4	PP2048	16384	8406.767	8612.412	2.446
Gemma-4-26B NVFP4	TG128	16384	153.912	153.121	-0.514
Gemma-4-26B NVFP4	PP2048	32768	6853.787	7091.345	3.466
Gemma-4-26B NVFP4	TG128	32768	148.665	148.688	0.015
Gemma-4-26B NVFP4	PP2048	65536	4893.351	5096.944	4.161
Gemma-4-26B NVFP4	TG128	65536	138.713	138.237	-0.343
Qwen3.6-35B-A3B Q4_K_M	PP2048	0	6717.521	6678.923	-0.575
Qwen3.6-35B-A3B Q4_K_M	TG128	0	232.623	232.392	-0.099
Qwen3.6-35B-A3B Q4_K_M	PP2048	1024	6623.178	6629.5	0.095
Qwen3.6-35B-A3B Q4_K_M	TG128	1024	238.249	236.712	-0.645
Qwen3.6-35B-A3B Q4_K_M	PP2048	2048	6587.27	6548.443	-0.589
Qwen3.6-35B-A3B Q4_K_M	TG128	2048	237.808	236.809	-0.42
Qwen3.6-35B-A3B Q4_K_M	PP2048	4096	6510.339	6468.767	-0.639
Qwen3.6-35B-A3B Q4_K_M	TG128	4096	231.487	229.493	-0.861
Qwen3.6-35B-A3B Q4_K_M	PP2048	8192	6334.852	6312.687	-0.35
Qwen3.6-35B-A3B Q4_K_M	TG128	8192	227.965	225.898	-0.907
Qwen3.6-35B-A3B Q4_K_M	PP2048	16384	6079.983	6143.106	1.038
Qwen3.6-35B-A3B Q4_K_M	TG128	16384	222.894	222.89	-0.001
Qwen3.6-35B-A3B Q4_K_M	PP2048	32768	5498.301	5565.342	1.219
Qwen3.6-35B-A3B Q4_K_M	TG128	32768	209.608	210.788	0.563
Qwen3.6-35B-A3B Q4_K_M	PP2048	65536	4503.546	4635.9	2.939
Qwen3.6-35B-A3B Q4_K_M	TG128	65536	191.038	192	0.503
Qwen3.6-35B-A3B NVFP4	PP2048	0	7889.913	8008.394	1.502
Qwen3.6-35B-A3B NVFP4	TG128	0	225.273	224.763	-0.226
Qwen3.6-35B-A3B NVFP4	PP2048	1024	7885.35	7877.043	-0.105
Qwen3.6-35B-A3B NVFP4	TG128	1024	231.535	227.721	-1.647
Qwen3.6-35B-A3B NVFP4	PP2048	2048	7818.813	7792.568	-0.336
Qwen3.6-35B-A3B NVFP4	TG128	2048	231.572	229.104	-1.066
Qwen3.6-35B-A3B NVFP4	PP2048	4096	7706.487	7663.216	-0.561
Qwen3.6-35B-A3B NVFP4	TG128	4096	226.107	228.745	1.166
Qwen3.6-35B-A3B NVFP4	PP2048	8192	7497.191	7528.131	0.413
Qwen3.6-35B-A3B NVFP4	TG128	8192	222.418	222.076	-0.154
Qwen3.6-35B-A3B NVFP4	PP2048	16384	7099.271	7115.332	0.226
Qwen3.6-35B-A3B NVFP4	TG128	16384	216.542	216.888	0.16
Qwen3.6-35B-A3B NVFP4	PP2048	32768	6340.828	6394.011	0.839
Qwen3.6-35B-A3B NVFP4	TG128	32768	204.84	204.294	-0.267
Qwen3.6-35B-A3B NVFP4	PP2048	65536	5040.035	5196.9	3.112
Qwen3.6-35B-A3B NVFP4	TG128	65536	187.883	187.137	-0.397
Qwen3.6-27B Q4_K_M	PP2048	0	2616.202	2634.481	0.699
Qwen3.6-27B Q4_K_M	TG128	0	63.458	63.637	0.282
Qwen3.6-27B Q4_K_M	PP2048	1024	2615.969	2625.161	0.351
Qwen3.6-27B Q4_K_M	TG128	1024	63.316	63.1	-0.342
Qwen3.6-27B Q4_K_M	PP2048	2048	2584.33	2574.847	-0.367
Qwen3.6-27B Q4_K_M	TG128	2048	62.668	62.587	-0.129
Qwen3.6-27B Q4_K_M	PP2048	4096	2536.773	2548.165	0.449
Qwen3.6-27B Q4_K_M	TG128	4096	61.709	61.756	0.075
Qwen3.6-27B Q4_K_M	PP2048	8192	2465.705	2478.911	0.536
Qwen3.6-27B Q4_K_M	TG128	8192	61.701	61.743	0.067
Qwen3.6-27B Q4_K_M	PP2048	16384	2331.064	2346.862	0.678
Qwen3.6-27B Q4_K_M	TG128	16384	60.533	60.623	0.148
Qwen3.6-27B Q4_K_M	PP2048	32768	2019.799	2054.841	1.735
Qwen3.6-27B Q4_K_M	TG128	32768	58.184	58.308	0.214
Qwen3.6-27B Q4_K_M	PP2048	65536	1487.298	1539.285	3.495
Qwen3.6-27B Q4_K_M	TG128	65536	53.718	54.154	0.813

I am collecting sweep on other hardware but facing high run to run variance. Will update here once i have perf sweep on other hardware .

ynankani · 2026-08-04T08:10:01Z

I see decent perf improvement with LLAMA3.1-8B_Q4_K_M and GPT-OSS-120B on DGX SPark

model	test	master t/s	swizzle t/s	delta
llama-3.1-8B	pp2048	3264.64 ± 23.77	3275.95 ± 8.38	0.35%
llama-3.1-8B	pp2048@d32768	1502.76 ± 9.99	1628.47 ± 12.91	8.37%
llama-3.1-8B	pp2048@d65536	983.54 ± 2.86	1074.51 ± 4.95	9.25%
llama-3.1-8B	tg128	42.20 ± 0.14	42.19 ± 0.08	−0.03%
llama-3.1-8B	tg128@d32768	23.89 ± 0.09	23.92 ± 0.09	0.11%
llama-3.1-8B	tg128@d65536	16.51 ± 0.00	16.53 ± 0.02	0.09%
gpt-oss-120b	pp2048	2004.26 ± 3.36	2016.92 ± 6.57	0.63%
gpt-oss-120b	pp2048@d32768	1395.58 ± 11.93	1438.45 ± 31.30	3.07%
gpt-oss-120b	pp2048@d65536	1014.04 ± 3.41	1137.23 ± 3.88	12.15%
gpt-oss-120b	tg128	59.33 ± 0.09	59.35 ± 0.18	0.04%
gpt-oss-120b	tg128@d32768	43.35 ± 0.10	45.43 ± 0.12	4.82%
gpt-oss-120b	tg128@d65536	34.09 ± 0.04	36.60 ± 0.06	7.37%

ynankani · 2026-08-04T08:16:10Z

Performance sweep Ampere

Model	Kind	Depth	Master t/s	Swizzle t/s	Delta %
Gemma-4-26B NVFP4	PP2048	0	2863.67	2831.2	-1.1
Gemma-4-26B NVFP4	TG128	0	86.94	87.31	0.4
Gemma-4-26B NVFP4	PP2048	1024	2965.68	2969.71	0.1
Gemma-4-26B NVFP4	TG128	1024	84.03	84.58	0.6
Gemma-4-26B NVFP4	PP2048	2048	2909.2	2904.02	-0.2
Gemma-4-26B NVFP4	TG128	2048	83.5	84.15	0.8
Gemma-4-26B NVFP4	PP2048	4096	2863.37	2877.13	0.5
Gemma-4-26B NVFP4	TG128	4096	83.03	83.51	0.6
Gemma-4-26B NVFP4	PP2048	8192	2683.69	2707.18	0.9
Gemma-4-26B NVFP4	TG128	8192	81.95	82.49	0.6
Gemma-4-26B NVFP4	PP2048	16384	2418.07	2500.35	3.4
Gemma-4-26B NVFP4	TG128	16384	79.13	80.93	2.3
Gemma-4-26B NVFP4	PP2048	32768	2026.74	2093.24	3.3
Gemma-4-26B NVFP4	TG128	32768	75.58	76.78	1.6
Gemma-4-26B NVFP4	PP2048	65536	1398.83	1429.87	2.2
Gemma-4-26B NVFP4	TG128	65536	63.81	66.86	4.8

Performance sweep Turing

Model	Depth	Metric	ToT median	PR median	Delta
Qwen3.5-4B-Q4_K_M	d0	pp2048	2,922.95	2,951.50	0.98%
Qwen3.5-4B-Q4_K_M	d0	tg128	101.67	103.16	1.46%
Qwen3.5-4B-Q4_K_M	d65536	pp2048	1,222.50	1,250.56	2.30%
Qwen3.5-4B-Q4_K_M	d65536	tg128	71.28	71.12	-0.23%
Meta-Llama-3.1-8B-Q4_K_M	d0	pp2048	2,043.31	2,075.57	1.58%
Meta-Llama-3.1-8B-Q4_K_M	d0	tg128	82.09	81.86	-0.27%
Meta-Llama-3.1-8B-Q4_K_M	d32768	pp2048	916.86	926.36	1.04%
Meta-Llama-3.1-8B-Q4_K_M	d32768	tg128	50.68	50.51	-0.34%

gaugarg-nv · 2026-08-04T12:13:28Z

@@ -1471,6 +1484,8 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
        const bool thread_should_write = T_C_KQ::J == 8 || T_C_KQ::get_j(threadIdx.x & 2) < 8;
 #endif // defined(TURING_MMA_AVAILABLE)

+        __syncthreads();


Are these __syncThreads needed only for swizzle path? If yes, can we make these conditional?

Yes, they are required for swizzled path due to "half2 * tile_K = Q_in_reg ? tile_Q" and padding is removed. Making this conditional.

Padding-removal alone should not change pipeline staging & data re-use though. Please clarify

Signed-off-by: ynankani <ynankani@nvidia.com>

ynankani · 2026-08-31T08:49:59Z

Performance sweep DGX Spark for JenniSD/DeepSeek-V2-Lite-Chat-Q4_K_M-GGUF

Test	Master	Allow nbatch_2%32=0	Change
pp2048 @ d0	4019.57 ± 16.39	4097.43 ± 33.85	1.90%
tg128 @ d0	113.19 ± 0.53	112.00 ± 0.46	-1.10%
pp2048 @ d32768	1165.39 ± 2.02	1236.21 ± 2.61	6.10%
tg128 @ d32768	71.45 ± 0.18	74.26 ± 0.22	3.90%
pp2048 @ d65536	674.13 ± 0.30	723.63 ± 0.57	7.30%
tg128 @ d65536	53.15 ± 0.13	57.00 ± 0.14	7.20%

ynankani · 2026-08-31T08:59:35Z

Performance sweep RTX-PRO-Blackwell 6000 for JenniSD/DeepSeek-V2-Lite-Chat-Q4_K_M-GGUF

Test	Master	Allow nbatch_2%32=0	Delta
pp2048 @ d0	11019.84	11100.41	0.73%
tg128 @ d0	327.18	324.07	-0.95%
pp2048 @ d32768	3518.76	3728.22	5.95%
tg128 @ d32768	239.89	238.48	-0.59%
pp2048 @ d65536	2040.49	2171.03	6.40%
tg128 @ d65536	197.9	195.21	-1.36%

ynankani · 2026-08-31T09:12:18Z

Thanks Johannes for the review. We can relax the condition to be more inclusive of other shapes. I will do a follow PR for refactoring and if find a way possible to get similar or more returns from memory pointer based swizzling as it will be cleaner approach.

JohannesGaessler · 2026-08-31T11:07:58Z

As discussed offsite, due to deadline constraints on the NVIDIA side I am fine with merging this PR as-is as long as there is a refactor in a follow-up. This is the performance I'm measuring when doing a sweep over batch size:

Performance

GPU	Model	Split mode	Microbatch size	Test	t/s b10707	t/s `088d425`	Speedup
RTX 3090	deepseek2 16B Q4_0	layer	1	pp1024@d32768	189.69	191.10	1.01
RTX 3090	deepseek2 16B Q4_0	layer	2	pp1024@d32768	294.89	296.54	1.01
RTX 3090	deepseek2 16B Q4_0	layer	4	pp1024@d32768	401.33	404.72	1.01
RTX 3090	deepseek2 16B Q4_0	layer	8	pp1024@d32768	474.49	489.46	1.03
RTX 3090	deepseek2 16B Q4_0	layer	16	pp1024@d32768	666.55	680.86	1.02
RTX 3090	deepseek2 16B Q4_0	layer	32	pp1024@d32768	778.61	797.69	1.02
RTX 3090	deepseek2 16B Q4_0	layer	64	pp1024@d32768	898.45	927.69	1.03
RTX 3090	deepseek2 16B Q4_0	layer	128	pp1024@d32768	911.17	875.52	0.96
RTX 3090	deepseek2 16B Q4_0	layer	256	pp1024@d32768	974.34	1024.77	1.05
RTX 3090	deepseek2 16B Q4_0	layer	512	pp1024@d32768	1081.46	1102.86	1.02
RTX 3090	gemma 2B Q4_0	layer	1	pp1024@d32768	346.67	351.72	1.01
RTX 3090	gemma 2B Q4_0	layer	2	pp1024@d32768	615.28	590.39	0.96
RTX 3090	gemma 2B Q4_0	layer	4	pp1024@d32768	991.04	1013.38	1.02
RTX 3090	gemma 2B Q4_0	layer	8	pp1024@d32768	1418.12	1419.10	1.00
RTX 3090	gemma 2B Q4_0	layer	16	pp1024@d32768	2673.65	2719.16	1.02
RTX 3090	gemma 2B Q4_0	layer	32	pp1024@d32768	3922.29	4101.52	1.05
RTX 3090	gemma 2B Q4_0	layer	64	pp1024@d32768	5118.77	5261.84	1.03
RTX 3090	gemma 2B Q4_0	layer	128	pp1024@d32768	5896.01	5542.83	0.94
RTX 3090	gemma 2B Q4_0	layer	256	pp1024@d32768	6261.00	7035.41	1.12
RTX 3090	gemma 2B Q4_0	layer	512	pp1024@d32768	7353.25	7736.17	1.05
RTX 3090	llama 8B Q4_0	layer	1	pp1024@d32768	92.84	92.31	0.99
RTX 3090	llama 8B Q4_0	layer	2	pp1024@d32768	178.41	176.74	0.99
RTX 3090	llama 8B Q4_0	layer	4	pp1024@d32768	297.95	269.61	0.90
RTX 3090	llama 8B Q4_0	layer	8	pp1024@d32768	398.07	370.72	0.93
RTX 3090	llama 8B Q4_0	layer	16	pp1024@d32768	838.28	785.35	0.94
RTX 3090	llama 8B Q4_0	layer	32	pp1024@d32768	1217.30	1190.95	0.98
RTX 3090	llama 8B Q4_0	layer	64	pp1024@d32768	1592.73	1365.97	0.86
RTX 3090	llama 8B Q4_0	layer	128	pp1024@d32768	1494.87	1458.58	0.98
RTX 3090	llama 8B Q4_0	layer	256	pp1024@d32768	1581.03	1570.17	0.99
RTX 3090	llama 8B Q4_0	layer	512	pp1024@d32768	2076.80	1882.33	0.91
RTX 4090	deepseek2 16B Q4_0	layer	1	pp1024@d32768	233.18	229.72	0.99
RTX 4090	deepseek2 16B Q4_0	layer	2	pp1024@d32768	382.40	379.40	0.99
RTX 4090	deepseek2 16B Q4_0	layer	4	pp1024@d32768	611.78	609.29	1.00
RTX 4090	deepseek2 16B Q4_0	layer	8	pp1024@d32768	960.92	978.25	1.02
RTX 4090	deepseek2 16B Q4_0	layer	16	pp1024@d32768	1260.35	1289.89	1.02
RTX 4090	deepseek2 16B Q4_0	layer	32	pp1024@d32768	1747.55	1807.47	1.03
RTX 4090	deepseek2 16B Q4_0	layer	64	pp1024@d32768	2165.00	2264.14	1.05
RTX 4090	deepseek2 16B Q4_0	layer	128	pp1024@d32768	2418.54	2533.24	1.05
RTX 4090	deepseek2 16B Q4_0	layer	256	pp1024@d32768	2942.81	3099.32	1.05
RTX 4090	deepseek2 16B Q4_0	layer	512	pp1024@d32768	3290.76	3474.35	1.06
RTX 4090	gemma 2B Q4_0	layer	1	pp1024@d32768	389.86	396.87	1.02
RTX 4090	gemma 2B Q4_0	layer	2	pp1024@d32768	707.04	671.01	0.95
RTX 4090	gemma 2B Q4_0	layer	4	pp1024@d32768	1340.48	1337.16	1.00
RTX 4090	gemma 2B Q4_0	layer	8	pp1024@d32768	2380.58	2365.00	0.99
RTX 4090	gemma 2B Q4_0	layer	16	pp1024@d32768	3975.73	3942.39	0.99
RTX 4090	gemma 2B Q4_0	layer	32	pp1024@d32768	6989.43	7031.07	1.01
RTX 4090	gemma 2B Q4_0	layer	64	pp1024@d32768	10145.70	10551.79	1.04
RTX 4090	gemma 2B Q4_0	layer	128	pp1024@d32768	12279.17	12701.43	1.03
RTX 4090	gemma 2B Q4_0	layer	256	pp1024@d32768	15364.57	15978.56	1.04
RTX 4090	gemma 2B Q4_0	layer	512	pp1024@d32768	16055.30	16534.94	1.03
RTX 4090	llama 8B Q4_0	layer	1	pp1024@d32768	104.40	104.39	1.00
RTX 4090	llama 8B Q4_0	layer	2	pp1024@d32768	200.40	200.57	1.00
RTX 4090	llama 8B Q4_0	layer	4	pp1024@d32768	391.65	391.59	1.00
RTX 4090	llama 8B Q4_0	layer	8	pp1024@d32768	749.72	751.79	1.00
RTX 4090	llama 8B Q4_0	layer	16	pp1024@d32768	1292.83	1282.36	0.99
RTX 4090	llama 8B Q4_0	layer	32	pp1024@d32768	2345.84	2345.94	1.00
RTX 4090	llama 8B Q4_0	layer	64	pp1024@d32768	3684.46	3809.77	1.03
RTX 4090	llama 8B Q4_0	layer	128	pp1024@d32768	4417.73	4596.19	1.04
RTX 4090	llama 8B Q4_0	layer	256	pp1024@d32768	5157.27	5387.12	1.04
RTX 4090	llama 8B Q4_0	layer	512	pp1024@d32768	5440.14	5638.00	1.04
RTX 5090	deepseek2 16B Q4_0	layer	1	pp1024@d32768	330.68	315.67	0.95
RTX 5090	deepseek2 16B Q4_0	layer	2	pp1024@d32768	552.32	539.30	0.98
RTX 5090	deepseek2 16B Q4_0	layer	4	pp1024@d32768	867.60	866.62	1.00
RTX 5090	deepseek2 16B Q4_0	layer	8	pp1024@d32768	1366.49	1376.25	1.01
RTX 5090	deepseek2 16B Q4_0	layer	16	pp1024@d32768	1482.22	1491.82	1.01
RTX 5090	deepseek2 16B Q4_0	layer	32	pp1024@d32768	2077.49	2103.47	1.01
RTX 5090	deepseek2 16B Q4_0	layer	64	pp1024@d32768	2448.73	2535.53	1.04
RTX 5090	deepseek2 16B Q4_0	layer	128	pp1024@d32768	2777.80	2909.15	1.05
RTX 5090	deepseek2 16B Q4_0	layer	256	pp1024@d32768	3901.43	4009.76	1.03
RTX 5090	deepseek2 16B Q4_0	layer	512	pp1024@d32768	4395.86	4555.75	1.04
RTX 5090	gemma 2B Q4_0	layer	1	pp1024@d32768	571.34	577.24	1.01
RTX 5090	gemma 2B Q4_0	layer	2	pp1024@d32768	992.75	922.86	0.93
RTX 5090	gemma 2B Q4_0	layer	4	pp1024@d32768	1922.69	1945.76	1.01
RTX 5090	gemma 2B Q4_0	layer	8	pp1024@d32768	3266.81	3285.40	1.01
RTX 5090	gemma 2B Q4_0	layer	16	pp1024@d32768	4885.39	5064.38	1.04
RTX 5090	gemma 2B Q4_0	layer	32	pp1024@d32768	8235.15	8474.15	1.03
RTX 5090	gemma 2B Q4_0	layer	64	pp1024@d32768	11538.44	11643.43	1.01
RTX 5090	gemma 2B Q4_0	layer	128	pp1024@d32768	14275.36	14359.23	1.01
RTX 5090	gemma 2B Q4_0	layer	256	pp1024@d32768	18827.35	19091.83	1.01
RTX 5090	gemma 2B Q4_0	layer	512	pp1024@d32768	20589.69	21024.50	1.02
RTX 5090	llama 8B Q4_0	layer	1	pp1024@d32768	166.43	166.36	1.00
RTX 5090	llama 8B Q4_0	layer	2	pp1024@d32768	333.24	333.43	1.00
RTX 5090	llama 8B Q4_0	layer	4	pp1024@d32768	642.64	642.76	1.00
RTX 5090	llama 8B Q4_0	layer	8	pp1024@d32768	1138.56	1134.76	1.00
RTX 5090	llama 8B Q4_0	layer	16	pp1024@d32768	1963.57	1964.46	1.00
RTX 5090	llama 8B Q4_0	layer	32	pp1024@d32768	3366.85	3475.25	1.03
RTX 5090	llama 8B Q4_0	layer	64	pp1024@d32768	4713.38	4775.74	1.01
RTX 5090	llama 8B Q4_0	layer	128	pp1024@d32768	5555.52	5602.75	1.01
RTX 5090	llama 8B Q4_0	layer	256	pp1024@d32768	6350.65	6536.92	1.03
RTX 5090	llama 8B Q4_0	layer	512	pp1024@d32768	7095.16	7254.80	1.02

Swizzling is faster on average but there are some cases where it is not. For now I think it is OK to merge like this but for the refactor please make whether or not to use swizzling part of the kernel configuration so that it can be controlled in a fine-grained way.

* CUDA: XOR swizzle flash attn K,V smem fp16 tiles Signed-off-by: ynankani <ynankani@nvidia.com> * Fix use 64bit generic pointer instead of 32bit shared pointer Signed-off-by: ynankani <ynankani@nvidia.com> * fix shared memory race in FA on DGX Spark * Handle corener case Signed-off-by: ynankani <ynankani@nvidia.com> * Add swizzle test cases and gate sync for swizzled path only Signed-off-by: ynankani <ynankani@nvidia.com> * gate CUDA PTX Signed-off-by: ynankani <ynankani@nvidia.com> * offset calculation specific for swizzle branch Signed-off-by: ynankani <ynankani@nvidia.com> * Reafctor code Signed-off-by: ynankani <ynankani@nvidia.com> * Refactor FA swizzle ldmatrix if/else into helpers (K row/col, V offset) Signed-off-by: ynankani <ynankani@nvidia.com> * rebase and update test case args Signed-off-by: ynankani <ynankani@nvidia.com> * Allow swizzle for non-pow2 shapes, for which nbatch_2%32==0 Signed-off-by: ynankani <ynankani@nvidia.com> --------- Signed-off-by: ynankani <ynankani@nvidia.com>

KV restore batching (ggml-org#27991), kv-cells seq-scan early stop (ggml-org#28011), MOE fusion to specdec + multi-token (ggml-org#27621), mm_ids_helper templated fast path (ggml-org#27978), qwen4exp recurrent state rollback (ggml-org#28123), n_layer_nextn load order (ggml-org#28159), FA K/V XOR-swizzle smem tiles (ggml-org#25635), --lazy-mode -lzm (ggml-org#27837/ggml-org#27969). TQ3/TurboQuant stack and vitriol-* integration auto-merged clean; no conflicts. Experiment E1 of mining-experiment-master-plan-2026-09-01.

* CUDA: XOR swizzle flash attn K,V smem fp16 tiles Signed-off-by: ynankani <ynankani@nvidia.com> * Fix use 64bit generic pointer instead of 32bit shared pointer Signed-off-by: ynankani <ynankani@nvidia.com> * fix shared memory race in FA on DGX Spark * Handle corener case Signed-off-by: ynankani <ynankani@nvidia.com> * Add swizzle test cases and gate sync for swizzled path only Signed-off-by: ynankani <ynankani@nvidia.com> * gate CUDA PTX Signed-off-by: ynankani <ynankani@nvidia.com> * offset calculation specific for swizzle branch Signed-off-by: ynankani <ynankani@nvidia.com> * Reafctor code Signed-off-by: ynankani <ynankani@nvidia.com> * Refactor FA swizzle ldmatrix if/else into helpers (K row/col, V offset) Signed-off-by: ynankani <ynankani@nvidia.com> * rebase and update test case args Signed-off-by: ynankani <ynankani@nvidia.com> * Allow swizzle for non-pow2 shapes, for which nbatch_2%32==0 Signed-off-by: ynankani <ynankani@nvidia.com> --------- Signed-off-by: ynankani <ynankani@nvidia.com>

Signed-off-by: ynankani <ynankani@nvidia.com>

github-actions Bot added testing Everything test related ggml changes relating to the ggml tensor library for machine learning CUDA Related to the CUDA backend labels Jul 13, 2026

ynankani force-pushed the ynankani/swizzle_flash_attn branch from 0fe9eeb to de97e22 Compare July 16, 2026 12:01

gaugarg-nv reviewed Jul 20, 2026

View reviewed changes

ORippler reviewed Jul 21, 2026

View reviewed changes

ynankani marked this pull request as ready for review July 23, 2026 15:05

ynankani requested review from a team and ggerganov as code owners July 23, 2026 15:05

ynankani force-pushed the ynankani/swizzle_flash_attn branch from de97e22 to 4093970 Compare July 29, 2026 15:07

gaugarg-nv reviewed Jul 30, 2026

View reviewed changes

Comment thread ggml/src/ggml-cuda/fattn-swizzle.cuh Outdated

gaugarg-nv reviewed Aug 4, 2026

View reviewed changes

ynankani force-pushed the ynankani/swizzle_flash_attn branch from c1fb1ce to 0166491 Compare August 4, 2026 16:29

ynankani added 11 commits August 30, 2026 16:43

CUDA: XOR swizzle flash attn K,V smem fp16 tiles

e13512a

Signed-off-by: ynankani <ynankani@nvidia.com>

Fix use 64bit generic pointer instead of 32bit shared pointer

8cb5912

Signed-off-by: ynankani <ynankani@nvidia.com>

fix shared memory race in FA on DGX Spark

73e5d1e

Handle corener case

50b0d08

Signed-off-by: ynankani <ynankani@nvidia.com>

Add swizzle test cases and gate sync for swizzled path only

24da808

Signed-off-by: ynankani <ynankani@nvidia.com>

gate CUDA PTX

bc56f68

Signed-off-by: ynankani <ynankani@nvidia.com>

offset calculation specific for swizzle branch

c9a90f7

Signed-off-by: ynankani <ynankani@nvidia.com>

Reafctor code

2a57183

Signed-off-by: ynankani <ynankani@nvidia.com>

Refactor FA swizzle ldmatrix if/else into helpers (K row/col, V offset)

de01a0c

Signed-off-by: ynankani <ynankani@nvidia.com>

rebase and update test case args

8ccc1da

Signed-off-by: ynankani <ynankani@nvidia.com>

Allow swizzle for non-pow2 shapes, for which nbatch_2%32==0

088d425

Signed-off-by: ynankani <ynankani@nvidia.com>

ynankani force-pushed the ynankani/swizzle_flash_attn branch from 987a04c to 088d425 Compare August 31, 2026 08:41

JohannesGaessler approved these changes Aug 31, 2026

View reviewed changes

gaugarg-nv approved these changes Aug 31, 2026

View reviewed changes

JohannesGaessler merged commit e4b9af0 into ggml-org:master Aug 31, 2026
26 of 30 checks passed

ORippler mentioned this pull request Sep 1, 2026

cuda: fuse MoE weighted expert reduction #25952

Merged

GenerelSchwerz mentioned this pull request Sep 3, 2026

cuda: optimize native quant FlashAttention integration (follow-up to #50) GenerelSchwerz/llama.cpp#55

Open

ggerganov mentioned this pull request Sep 4, 2026

ggml : bump version to 0.23.0 ggml-org/ggml#1618

Merged

JohannesGaessler self-assigned this Sep 4, 2026

BrewTestBot mentioned this pull request Sep 4, 2026

ggml 0.23.0 Homebrew/homebrew-core#302393

Merged

1 task

ynankani added a commit to ynankani/llama.cpp that referenced this pull request Sep 7, 2026

CUDA: Follow up of ggml-org#25635, refactor shared smem swizzle

5965e37

Signed-off-by: ynankani <ynankani@nvidia.com>

ynankani mentioned this pull request Sep 7, 2026

CUDA: Follow up of #25635, refactoring FA shared smem swizzle #28536

Draft

cell	config	swizzle	block	occ%	bc_ld	bc_st	dur
1	64,4	off	64	4.20	1141	370	162 us
2	128,2	off	128	8.30	478	376	161 us
3	64,4	on	64	4.20	396	188	168 us
4	128,2	on	128	8.30	376	385	163 us

cell	config	swizzle	block	occ%	bc_ld	bc_st	dur
1	64,4	off	128	15.9	27,362	26,478	4.33 ms
2	128,2	off	128	15.9	27,433	26,566	4.34 ms
3	64,4	on	128	14.9	1,29,524	17,390	4.01 ms
4	128,2	on	128	14.9	1,29,963	17,654	4.02 ms

cell	config	swizzle	block	occ%	bc_ld	bc_st	dur
1	64,4	off	64	4.17	764	91	0.76 ms
2	128,2	off	128	8.33	152	122	0.75 ms
3	64,4	on	64	4.17	99	48	0.63 ms
4	128,2	on	128	8.33	107	96	0.63 ms

cell	config	swizzle	block	occ%	bc_ld	bc_st	dur
1	64,4	off	128	16.63	10722	8422	30.43 ms
2	128,2	off	128	16.63	10903	7359	29.35 ms
3	64,4	on	128	15.29	26097	11759	26.49 ms
4	128,2	on	128	15	27460	10403	27.36 ms

cell	config	swizzle	cp.async wf	cp.async ideal	xideal
1	64,4	off	2.8M	1.1M	2.62x
2	128,2	off	2.8M	1.1M	2.61x
3	64,4	on	1.1M	1.1M	1.00x
4	128,2	on	1.1M	1.1M	1.03x

cell	config	swizzle	cp.async wf	cp.async ideal	xideal
1	64,4	off	40.6M	14.7M	2.76x
2	128,2	off	40.6M	14.7M	2.76x
3	64,4	on	15.0M	14.4M	1.04x
4	128,2	on	15.0M	14.4M	1.04x

		// XOR swizzle for K/V SMEM tiles to avoid bank conflicts without row padding (Turing+ only).
		// Stride must be a power-of-two >= 32 half2 columns,otherwise we keep +4 row padding.

Uh oh!

Conversation

ynankani commented Jul 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Overview

Additional information

Requirements

Uh oh!

naamfung commented Jul 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Bug Report (for PR #25635)

Uh oh!

ynankani commented Jul 15, 2026

Bug Report (for PR #25635)

Uh oh!

ynankani commented Jul 16, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

ynankani Jul 22, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

ynankani Jul 29, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

gaugarg-nv commented Jul 20, 2026

Uh oh!

naamfung commented Jul 21, 2026

Uh oh!

naamfung commented Jul 21, 2026

Uh oh!

naamfung commented Jul 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

ynankani commented Jul 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

naamfung commented Jul 21, 2026

Uh oh!

naamfung commented Jul 21, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

ORippler commented Jul 21, 2026

Uh oh!

ynankani commented Jul 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

naamfung commented Jul 21, 2026

Uh oh!

naamfung commented Jul 22, 2026

Uh oh!

ynankani commented Jul 22, 2026

Uh oh!

Uh oh!

ynankani commented Aug 4, 2026

Uh oh!

ynankani commented Aug 4, 2026

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

ORippler Aug 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

ynankani commented Jul 13, 2026 •

edited

Loading

naamfung commented Jul 15, 2026 •

edited

Loading

ynankani Jul 22, 2026 •

edited

Loading

ynankani Jul 29, 2026 •

edited

Loading

naamfung commented Jul 21, 2026 •

edited

Loading

ynankani commented Jul 21, 2026 •

edited

Loading

ynankani commented Jul 21, 2026 •

edited

Loading

ORippler Aug 13, 2026 •

edited

Loading