Developer Tools

llama.cpp b9589 fixes CUDA data races for stable GPU inference

Missing syncthreads caused random crashes on NVIDIA GPUs during SSM scan.

Deep Dive

llama.cpp, the widely-used open-source C++ library for running large language models locally, has released version b9589 with a targeted fix for CUDA GPU users. The update resolves data races in the ssm_scan_f32 kernel — a component used for state-space model operations. The root cause was a lack of __syncthreads() barriers before reusing the cub::BlockLoad temp storage in shared memory, which could lead to random crashes or incorrect outputs under concurrent thread execution. The patch ensures all threads finish reading shared memory before it is overwritten in the next loop iteration, aligning with NVIDIA's CUB documentation. This small but critical change eliminates intermittent failures for users running models like Mamba or other SSM architectures on NVIDIA GPUs.

The b9589 release also includes a wide range of platform-specific builds: macOS (Apple Silicon and Intel), Linux (x64, arm64, s390x with Vulkan, ROCm, OpenVINO, SYCL), Windows (x64, arm64, CUDA 12/13, Vulkan, HIP), and Android (arm64). Developers using llama.cpp for production or research will benefit from improved reliability, especially when running long inference sessions. The fix is available immediately via the GitHub release, and users are encouraged to upgrade to avoid potential silent data corruption.

Key Points
  • Fixes data races in ssm_scan_f32 CUDA kernel by adding missing __syncthreads barriers before reusing shared memory cub_temp_storage.
  • Ensures stable state-space model inference on NVIDIA GPUs (CUDA 12/13) without random crashes.
  • Release b9589 includes builds for 15+ platforms including macOS, Linux, Windows, Android, and openEuler.

Why It Matters

Critical stability fix for llama.cpp users running SSM models on NVIDIA GPUs, preventing silent errors in local AI inference.

📬 Get the top 10 AI stories daily