Skip to content

Commit 51ead2b

Browse files
Add explicit read into registers
1 parent ce5914f commit 51ead2b

File tree

1 file changed

+8
-5
lines changed
  • dpctl/tensor/libtensor/include/kernels/sorting

1 file changed

+8
-5
lines changed

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

Lines changed: 8 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -149,8 +149,9 @@ topk_full_merge_sort_impl(sycl::queue &exec_q,
149149
std::size_t src_idx = iter_gid * axis_nelems + axis_gid;
150150
std::size_t dst_idx = iter_gid * k + axis_gid;
151151

152-
auto res_ind = index_data[src_idx];
153-
vals_tp[dst_idx] = arg_tp[res_ind];
152+
const IndexTy res_ind = index_data[src_idx];
153+
const argTy v = arg_tp[res_ind];
154+
vals_tp[dst_idx] = v;
154155
inds_tp[dst_idx] = res_ind % axis_nelems;
155156
});
156157
});
@@ -425,8 +426,9 @@ sycl::event topk_merge_impl(
425426
const std::size_t src_idx = iter_gid * alloc_len + axis_gid;
426427
const std::size_t dst_idx = gid;
427428

428-
const auto res_ind = index_data[src_idx];
429-
vals_tp[dst_idx] = arg_tp[res_ind];
429+
const IndexTy res_ind = index_data[src_idx];
430+
const argTy v = arg_tp[res_ind];
431+
vals_tp[dst_idx] = v;
430432
inds_tp[dst_idx] = (res_ind % axis_nelems);
431433
});
432434
});
@@ -538,7 +540,8 @@ sycl::event topk_radix_impl(sycl::queue &exec_q,
538540
const std::size_t dst_idx = gid;
539541

540542
const IndexTy res_ind = tmp_tp[src_idx];
541-
vals_tp[dst_idx] = arg_tp[res_ind];
543+
const argTy v = arg_tp[res_ind];
544+
vals_tp[dst_idx] = v;
542545
inds_tp[dst_idx] = (res_ind % axis_nelems);
543546
});
544547
});

0 commit comments

Comments
 (0)