The llama.cpp b10728 release introduces a CUDA implementation of XOR swizzle for Flash Attention K,V shared memory fp16 tiles. This update also resolves a shared memory race condition in Flash Attention on DGX Spark hardware.
- Implements XOR swizzle flash attention with fp16 tiles for CUDA.
- Fixes use of 64-bit generic pointer instead of 32-bit shared pointer.
- Adds swizzle test cases and gates synchronization for the swizzled path only.
- Allows swizzling for non-power-of-2 shapes where nbatch_2%32==0.
- Refactors Flash Attention swizzle ldmatrix logic into helper functions.
The release provides binaries for macOS, Linux, Windows, Android, and openEuler across CPU, CUDA, Vulkan, ROCm, OpenVINO, and SYCL backends.