Skip to content

Commit 0869128

Browse files
Add atomic_fence to triage
1 parent c1f8a74 commit 0869128

File tree

2 files changed

+5
-0
lines changed

2 files changed

+5
-0
lines changed

dpctl/tensor/libtensor/include/kernels/sorting/sort_utils.hpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -111,6 +111,8 @@ sycl::event map_back_impl(sycl::queue &exec_q,
111111
cgh.parallel_for<KernelName>(
112112
sycl::range<1>(nelems), [=](sycl::id<1> id) {
113113
const IndexTy linear_index = flat_index_data[id];
114+
sycl::atomic_fence(sycl::memory_order::acq_rel,
115+
sycl::memory_scope::device);
114116
reduced_index_data[id] = (linear_index % row_size);
115117
});
116118
});

dpctl/tensor/libtensor/include/kernels/sorting/topk.hpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -122,6 +122,9 @@ sycl::event write_out_impl(sycl::queue &exec_q,
122122
const IndexTy res_ind = index_data[src_idx];
123123
const argTy v = arg_tp[res_ind];
124124

125+
sycl::atomic_fence(sycl::memory_order::acq_rel,
126+
sycl::memory_scope::device);
127+
125128
const std::size_t dst_idx = data_id;
126129
vals_tp[dst_idx] = v;
127130
inds_tp[dst_idx] = (res_ind % axis_nelems);

0 commit comments

Comments
 (0)