From 1bb2b9fcbe85c35c4e3b787e1c868ef05ca702f3 Mon Sep 17 00:00:00 2001
From: uvos <carl@uvos.xyz>
Date: Sat, 10 Oct 2026 08:14:51 +0200
Subject: [PATCH] CUDA/HIP: fix race in flash_attn_ext_f16_process_tile when
 nbatch_combine != DKQ/2 (#30103)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Suggested-by: Johannes Gäßler <johannesg@5d6.de>
---
 ggml/src/ggml-cuda/fattn-mma-f16.cuh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh
index 7f0f6abfd..296aa2b34 100644
--- a/ggml/src/ggml-cuda/fattn-mma-f16.cuh
+++ b/ggml/src/ggml-cuda/fattn-mma-f16.cuh
@@ -1772,7 +1772,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
                 }
             }
         }
-        if (np > 1) {
+        if (np > 1 || nbatch_combine != DKQ/2) {
             __syncthreads();
         }
     }