[SYCL] fix softmax r2r result wrong issue (ggml-org#7811)

pengxin99 · web-flow · commit d5c938cd7716 · 2024-06-07T14:28:26.000+08:00
diff --git a/ggml-sycl.cpp b/ggml-sycl.cpp
@@ -9108,6 +9108,7 @@ static void soft_max_f32(const float * x, const float * mask, float * dst, const
     // find the sum of exps in the block
     tmp = warp_reduce_sum(tmp, item_ct1);
     if (block_size > WARP_SIZE) {
+        item_ct1.barrier(sycl::access::fence_space::local_space);
         if (warp_id == 0) {
             buf[lane_id] = 0.f;
         }

Original file line number	Diff line number	Diff line change
`@@ -9108,6 +9108,7 @@ static void soft_max_f32(const float * x, const float * mask, float * dst, const`
`9108`	`9108`	`// find the sum of exps in the block`
`9109`	`9109`	`tmp = warp_reduce_sum(tmp, item_ct1);`
`9110`	`9110`	`if (block_size > WARP_SIZE) {`
	`9111`	`+ item_ct1.barrier(sycl::access::fence_space::local_space);`
`9111`	`9112`	`if (warp_id == 0) {`
`9112`	`9113`	`buf[lane_id] = 0.f;`
`9113`	`9114`	`}`