From: slaren <redacted>
Date: Sun, 3 Mar 2024 13:26:18 +0000 (+0100)
Subject: cuda : fix data race in soft max (llama/5853)
X-Git-Tag: upstream/0.0.1642~889
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=31f3807578072faff08c3b17c09df9906225b01d;p=pkg%2Fggml%2Fsources%2Fggml

cuda : fix data race in soft max (llama/5853)
---

diff --git a/src/ggml-cuda.cu b/src/ggml-cuda.cu
index 9382d96a..7d027a30 100644
--- a/src/ggml-cuda.cu
+++ b/src/ggml-cuda.cu
@@ -6947,6 +6947,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;
         }