From 124733e860728f4797e665b31cfb719ca27ec5b4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jirka=20Sv=C3=ADtil?= Date: Sat, 29 Aug 2026 10:29:13 +0000 Subject: [PATCH] cuda blocking correctness issue fix --- ggml/src/ggml-cuda/fattn-mma-f16.cuh | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh index 7f4cfd5511ff..e172393c4c35 100644 --- a/ggml/src/ggml-cuda/fattn-mma-f16.cuh +++ b/ggml/src/ggml-cuda/fattn-mma-f16.cuh @@ -1535,8 +1535,6 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile( } } - __syncthreads(); - // Write back combined meta data: #pragma unroll for (int imeta = 0; imeta < nmeta; ++imeta) { @@ -1556,10 +1554,9 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile( float2 * dstk_fixup_meta = dstk_fixup + (gridDim.x + blockIdx.x)*ncols; dstk_fixup_meta[(threadIdx.y/np)*cols_per_warp + threadIdx.x] = make_float2(KQ_cmn, KQ_crs); } - } else if (np > 1) { - // Warps with threadIdx.y % np == 0 execute a __syncthreads() in the if branch. - // Therefore, all other warps also need to execute a __syncthreads(). - // Otherwise the points at which warps synchronize with each other would become misaligned. + } + + if constexpr (np > 1) { __syncthreads(); }