Speeds up CUDA flash attention with XOR swizzle
CUDA speeds up flash attention with XOR swizzle memory optimization.
CUDA: XOR swizzle flash attn K,V smem fp16 tiles (#25635) CUDA: XOR swizzle flash attn K,V smem fp16 tiles Signed-off-by: ynankani Fix use 64bit generic pointer instead of 32bit shared pointer Signed-off-by: ynankani fix shared memory race in FA on DGX Spark Handle corener case Signed-off-by: ynankani Add swizzle test cases and gate sync for swizzled path only Signed-off-by: ynankani gate CUDA PTX Signed-off-by: ynankani offset calculation specific for swizzle branch Signed-off-by: ynankani Reafctor code Signed-off-by: ynankani Refactor FA swizzle ldmatrix if/else into helpers (K row/col, V offset) Signed-off-by: ynankani rebase and update test case args Signed-off-by: ynankani Allow swizzle for non-pow2 shapes, for which nbatch_2%32==0 Signed-off-by: ynankani --------- Signed-off-by: ynankani Website: Attestations: macOS/iOS: macOS Apple Silicon (arm64) macOS Apple Silicon (arm64, KleidiAI enabled) DISABLED macOS Intel (x64) iOS XCFramework Linux: Ubuntu x64 (CPU) Ubuntu arm64 (CPU) Ubuntu s390x (CPU) Ubuntu x64 (Vulkan) Ubuntu arm64 (Vulkan) Ubuntu x64 (ROCm 7.14) Ubuntu x64 (OpenVINO) Ubuntu x64 (SYCL FP32) Ubuntu x64 (SYCL FP16) Android: Android arm64 (CPU) Windows: Windows…
- github.comllama.cpp b10728primary