Skip to content

Commit

Permalink
Add explicit read into registers
Browse files Browse the repository at this point in the history
  • Loading branch information
oleksandr-pavlyk committed Dec 23, 2024
1 parent ce5914f commit f1b2045
Showing 1 changed file with 8 additions and 5 deletions.
13 changes: 8 additions & 5 deletions dpctl/tensor/libtensor/include/kernels/sorting/topk.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -149,8 +149,9 @@ topk_full_merge_sort_impl(sycl::queue &exec_q,
std::size_t src_idx = iter_gid * axis_nelems + axis_gid;
std::size_t dst_idx = iter_gid * k + axis_gid;

auto res_ind = index_data[src_idx];
vals_tp[dst_idx] = arg_tp[res_ind];
const IndexTy res_ind = index_data[src_idx];
const argTy v = arg_tp[res_ind];
vals_tp[dst_idx] = v;
inds_tp[dst_idx] = res_ind % axis_nelems;
});
});
Expand Down Expand Up @@ -425,8 +426,9 @@ sycl::event topk_merge_impl(
const std::size_t src_idx = iter_gid * alloc_len + axis_gid;
const std::size_t dst_idx = gid;

const auto res_ind = index_data[src_idx];
vals_tp[dst_idx] = arg_tp[res_ind];
const IndexTy res_ind = index_data[src_idx];
const argTy v = arg_tp[res_ind];
vals_tp[dst_idx] = v;
inds_tp[dst_idx] = (res_ind % axis_nelems);
});
});
Expand Down Expand Up @@ -538,7 +540,8 @@ sycl::event topk_radix_impl(sycl::queue &exec_q,
const std::size_t dst_idx = gid;

const IndexTy res_ind = tmp_tp[src_idx];
vals_tp[dst_idx] = arg_tp[res_ind];
const v = arg_tp[res_ind];
vals_tp[dst_idx] = v;
inds_tp[dst_idx] = (res_ind % axis_nelems);
});
});
Expand Down

0 comments on commit f1b2045

Please sign in to comment.