Skip to content

Commit 823c201

Browse files
oleksandr-pavlykndgrigorian
authored andcommitted
Add explicit read into registers
1 parent cd1243f commit 823c201

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
@@ -145,8 +145,9 @@ topk_full_merge_sort_impl(sycl::queue &exec_q,
145145
std::size_t src_idx = iter_gid * axis_nelems + axis_gid;
146146
std::size_t dst_idx = iter_gid * k + axis_gid;
147147

148-
auto res_ind = index_data[src_idx];
149-
vals_tp[dst_idx] = arg_tp[res_ind];
148+
const IndexTy res_ind = index_data[src_idx];
149+
const argTy v = arg_tp[res_ind];
150+
vals_tp[dst_idx] = v;
150151
inds_tp[dst_idx] = res_ind % axis_nelems;
151152
});
152153
});
@@ -418,8 +419,9 @@ sycl::event topk_merge_impl(
418419
const std::size_t src_idx = iter_gid * alloc_len + axis_gid;
419420
const std::size_t dst_idx = gid;
420421

421-
const auto res_ind = index_data[src_idx];
422-
vals_tp[dst_idx] = arg_tp[res_ind];
422+
const IndexTy res_ind = index_data[src_idx];
423+
const argTy v = arg_tp[res_ind];
424+
vals_tp[dst_idx] = v;
423425
inds_tp[dst_idx] = (res_ind % axis_nelems);
424426
});
425427
});
@@ -528,7 +530,8 @@ sycl::event topk_radix_impl(sycl::queue &exec_q,
528530
const std::size_t dst_idx = gid;
529531

530532
const IndexTy res_ind = tmp_tp[src_idx];
531-
vals_tp[dst_idx] = arg_tp[res_ind];
533+
const argTy v = arg_tp[res_ind];
534+
vals_tp[dst_idx] = v;
532535
inds_tp[dst_idx] = (res_ind % axis_nelems);
533536
});
534537
});

0 commit comments

Comments
 (0)