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

Merged

JohannesGaessler merged 11 commits into

ggml-org:masterfrom

ynankani:ynankani/swizzle_flash_attn

Aug 31, 2026

Commits on Aug 30, 2026

CUDA: XOR swizzle flash attn K,V smem fp16 tiles
ynankani
committed
e13512a
Fix use 64bit generic pointer instead of 32bit shared pointer
ynankani
committed
8cb5912
fix shared memory race in FA on DGX Spark
ynankani
committed
73e5d1e
Handle corener case
ynankani
committed
50b0d08
Add swizzle test cases and gate sync for swizzled path only
ynankani
committed
24da808
gate CUDA PTX
ynankani
committed
bc56f68
offset calculation specific for swizzle branch
ynankani
committed
c9a90f7
Reafctor code
ynankani
committed
2a57183
Refactor FA swizzle ldmatrix if/else into helpers (K row/col, V offset)
ynankani
committed
de01a0c
rebase and update test case args
ynankani
committed
8ccc1da

Commits on Aug 31, 2026

Allow swizzle for non-pow2 shapes, for which nbatch_2%32==0
ynankani
committed
088d425