From 67be2ce1015d070b3b2cd488bcb041eefb61de72 Mon Sep 17 00:00:00 2001
From: slaren <slarengh@gmail.com>
Date: Sun, 3 Mar 2024 14:26:18 +0100
Subject: [PATCH] cuda : fix data race in soft max (#5853)

---
 ggml-cuda.cu | 1 +
 1 file changed, 1 insertion(+)

diff --git a/ggml-cuda.cu b/ggml-cuda.cu
index 7ed97430f4fa4..04c6cb1b8fada 100644
--- a/ggml-cuda.cu
+++ b/ggml-cuda.cu
@@ -6904,6 +6904,7 @@ static __global__ void soft_max_f32(const float * x, const float * mask, const f
     // find the sum of exps in the block
     tmp = warp_reduce_sum(tmp);
     if (block_size > WARP_SIZE) {
+        __syncthreads();
         if (warp_id == 0) {
             buf_iw[lane_id] = 0.0f;
         }