cuda : fix data race in soft max (ggml-org#5853)

slaren · hodlen · commit 5481c462d863 · 2024-04-02T00:15:07.000+08:00
diff --git 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;
         }

Original file line number	Diff line number	Diff line change
`@@ -6904,6 +6904,7 @@ static __global__ void soft_max_f32(const float * x, const float * mask, const f`
`6904`	`6904`	`// find the sum of exps in the block`
`6905`	`6905`	`tmp = warp_reduce_sum(tmp);`
`6906`	`6906`	`if (block_size > WARP_SIZE) {`
	`6907`	`+ __syncthreads();`
`6907`	`6908`	`if (warp_id == 0) {`
`6908`	`6909`	`buf_iw[lane_id] = 0.0f;`
`6909`	`6910`	`}`