Skip to content

Commit af56d6c

Browse files
committed
Wait on radix sort event and write out event
1 parent 5191d68 commit af56d6c

File tree

1 file changed

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

1 file changed

+4
-4
lines changed

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -498,10 +498,10 @@ sycl::event topk_radix_impl(sycl::queue &exec_q,
498498
exec_q, iter_nelems, axis_nelems, workspace, tmp_tp, proj_op,
499499
ascending, {iota_ev});
500500

501+
radix_sort_ev.wait();
502+
501503
// Write out top k of the temporary
502504
sycl::event write_topk_ev = exec_q.submit([&](sycl::handler &cgh) {
503-
cgh.depends_on(radix_sort_ev);
504-
505505
using KernelName = topk_radix_map_back_krn<argTy, IndexTy>;
506506

507507
cgh.parallel_for<KernelName>(iter_nelems * k, [=](sycl::id<1> id) {
@@ -519,9 +519,9 @@ sycl::event topk_radix_impl(sycl::queue &exec_q,
519519
});
520520
});
521521

522-
sycl::event cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
523-
cgh.depends_on(write_topk_ev);
522+
write_topk_ev.wait();
524523

524+
sycl::event cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
525525
const sycl::context &ctx = exec_q.get_context();
526526

527527
using dpctl::tensor::alloc_utils::sycl_free_noexcept;

0 commit comments

Comments
 (0)