CUDA: XOR swizzle flash attn K,V smem fp16 tiles #25635
JohannesGaessler
merged
11 commits into
ggml-org:master
from
ynankani:ynankani/swizzle_flash_attn
ORippler
commented
on 2026-07-21
gaugarg-nv
commented
on 2026-07-30
ORippler
commented
on 2026-08-14
gaugarg-nv
approved these changes
on 2026-08-24
am17an
approved these changes
on 2026-08-25
JohannesGaessler
requested changes
on 2026-08-27
CUDA: XOR swizzle flash attn K,V smem fp16 tiles
e13512ab
Fix use 64bit generic pointer instead of 32bit shared pointer
8cb5912b
fix shared memory race in FA on DGX Spark
73e5d1e4
Handle corener case
50b0d08f
Add swizzle test cases and gate sync for swizzled path only
24da8081
gate CUDA PTX
bc56f686
offset calculation specific for swizzle branch
c9a90f7c
Reafctor code
2a571830
Refactor FA swizzle ldmatrix if/else into helpers (K row/col, V offset)
de01a0cc
rebase and update test case args
8ccc1daa
Allow swizzle for non-pow2 shapes, for which nbatch_2%32==0
088d425c
JohannesGaessler
approved these changes
on 2026-08-31
gaugarg-nv
approved these changes
on 2026-08-31
Reviewers
Assignees
Labels
testing
ggml
CUDA
Milestone
No milestone
Login to write a write a comment.
Login via GitHub