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.