From 99db719bdf07f7d537431ca556a0512bf3d995b6 Mon Sep 17 00:00:00 2001 From: Paul Mairhofer Date: Mon, 19 Oct 2026 07:25:04 +0200 Subject: [PATCH] CUDA: fix shared-memory WAR hazard in flash_attn_ext_vec (#30573) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * CUDA: fix KQ write-after-read race in flash_attn_ext_vec Assisted-by: ChatGPT * CUDA: apply review suggestions for KQ sync comments Co-authored-by: Johannes Gäßler --------- Co-authored-by: Johannes Gäßler --- ggml/src/ggml-cuda/fattn-vec.cuh | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/fattn-vec.cuh b/ggml/src/ggml-cuda/fattn-vec.cuh index 57a285565..74f282047 100644 --- a/ggml/src/ggml-cuda/fattn-vec.cuh +++ b/ggml/src/ggml-cuda/fattn-vec.cuh @@ -317,7 +317,7 @@ static __global__ void flash_attn_ext_vec( #endif // V_DOT2_F32_F16_AVAILABLE } - ggml_cuda_syncwarp(); + ggml_cuda_syncwarp(); // Needed for KQ synchronization. #pragma unroll for (int k0 = 0; k0 < WARP_SIZE; k0 += V_cols_per_iter) { @@ -374,6 +374,8 @@ static __global__ void flash_attn_ext_vec( } #endif // V_DOT2_F32_F16_AVAILABLE } + + ggml_cuda_syncwarp(); // Needed for KQ synchronization. } if (sinks && blockIdx.y == 0) {