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

2024-03-03 14:26:18 +01:00 · 2024-03-03 14:26:18 +01:00 · 93a84a143b
parent bd26876267
commit 93a84a143b
1 changed files with 1 additions and 0 deletions
--- a/ggml-cuda.cu
+++ b/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;
        }