Skip to content

Commit

Permalink
Merge pull request #6056 from masterleinad/partially_reverse_5504
Browse files Browse the repository at this point in the history
Partially reverse #5504: SYCL: Remove workaround for submit_barrier not being enqueued properly
  • Loading branch information
dalg24 committed Apr 14, 2023
2 parents 0d96f88 + 079268c commit 9004274
Showing 1 changed file with 7 additions and 1 deletion.
8 changes: 7 additions & 1 deletion core/src/SYCL/Kokkos_SYCL_Space.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -40,7 +40,13 @@ void DeepCopySYCL(void* dst, const void* src, size_t n) {
void DeepCopyAsyncSYCL(const Kokkos::Experimental::SYCL& instance, void* dst,
const void* src, size_t n) {
sycl::queue& q = *instance.impl_internal_space_instance()->m_queue;
auto event = q.memcpy(dst, src, n);
// FIXME_SYCL memcpy doesn't respect submit_barrier which means that we need
// to actually fence the execution space to make sure the memcpy is properly
// enqueued when using out-of-order queues.
#ifndef KOKKOS_ARCH_INTEL_GPU
q.wait_and_throw();
#endif
auto event = q.memcpy(dst, src, n);
q.ext_oneapi_submit_barrier(std::vector<sycl::event>{event});
}

Expand Down

0 comments on commit 9004274

Please sign in to comment.