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) {