Skip to content

Commit 1b3ade0

Browse files
committed
Fix race condition in custom_inclusive_scan_over_group
By returning data from `local_mem_acc` after the group barrier, if memory is later overwritten, a race condition follows, which was especially obvious on CPU Now the value is stored in variable before the barrier and then returned
1 parent 4bd02b4 commit 1b3ade0

File tree

1 file changed

+2
-1
lines changed

1 file changed

+2
-1
lines changed

dpctl/tensor/libtensor/include/utils/sycl_utils.hpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -172,9 +172,10 @@ T custom_inclusive_scan_over_group(const GroupT &wg,
172172
}
173173
}
174174

175+
T accumulated_local_val = local_mem_acc[local_id];
175176
sycl::group_barrier(wg, sycl::memory_scope::work_group);
176177

177-
return local_mem_acc[local_id];
178+
return accumulated_local_val;
178179
}
179180

180181
// Reduction functors

0 commit comments

Comments
 (0)