[llvm] [libsycl] Implement memcpy API (PR #205369)
Sergey Semenov via llvm-commits
llvm-commits at lists.llvm.org
Thu Jul 2 08:45:24 PDT 2026
https://github.com/sergey-semenov updated https://github.com/llvm/llvm-project/pull/205369
>From f991dd8d23e776866451aac7de61e2f147b1efc6 Mon Sep 17 00:00:00 2001
From: Sergey Semenov <sergey.semenov at intel.com>
Date: Fri, 8 May 2026 14:08:53 -0700
Subject: [PATCH 1/2] [libsycl] Implement memcpy API
---
libsycl/docs/index.rst | 3 +
.../include/sycl/__impl/detail/obj_utils.hpp | 23 ++++
libsycl/include/sycl/__impl/queue.hpp | 30 +++++
.../src/detail/offload/offload_topology.cpp | 5 +-
libsycl/src/detail/platform_impl.cpp | 2 +
libsycl/src/detail/queue_impl.cpp | 108 ++++++++++++----
libsycl/src/detail/queue_impl.hpp | 12 ++
libsycl/src/queue.cpp | 23 +++-
libsycl/test/usm/memcpy.cpp | 116 ++++++++++++++++++
9 files changed, 287 insertions(+), 35 deletions(-)
create mode 100644 libsycl/test/usm/memcpy.cpp
diff --git a/libsycl/docs/index.rst b/libsycl/docs/index.rst
index 4e92a219163ca..cbf104731d33a 100644
--- a/libsycl/docs/index.rst
+++ b/libsycl/docs/index.rst
@@ -109,6 +109,9 @@ TODO for added SYCL classes
* ``queue``:
* to implement USM methods
+
+ * ``memcpy``: enable the host-to-host case (blocked by liboffload limitations)
+
* to implement synchronization methods
* to implement submit & copy with accessors (low priority)
* get_info & properties
diff --git a/libsycl/include/sycl/__impl/detail/obj_utils.hpp b/libsycl/include/sycl/__impl/detail/obj_utils.hpp
index dcb18f1c03de4..ece4472550394 100644
--- a/libsycl/include/sycl/__impl/detail/obj_utils.hpp
+++ b/libsycl/include/sycl/__impl/detail/obj_utils.hpp
@@ -22,6 +22,7 @@
#include <optional>
#include <type_traits>
#include <utility>
+#include <vector>
_LIBSYCL_BEGIN_NAMESPACE_SYCL
@@ -59,6 +60,28 @@ auto getSyclObjImpl(const SyclObject &Obj)
return ImplUtils::getSyclObjImpl(Obj);
}
+template <typename ObjT, typename FuncT>
+auto transformVec(const std::vector<ObjT> &Vec, FuncT Func) {
+ std::vector<
+ std::remove_const_t<std::remove_reference_t<decltype(Func(Vec[0]))>>>
+ Result;
+ Result.reserve(Vec.size());
+ for (const ObjT &Obj : Vec)
+ Result.push_back(Func(Obj));
+ return Result;
+}
+
+template <typename SyclObject>
+auto getSyclObjImpls(const std::vector<SyclObject> &Objs) {
+ return transformVec(Objs, getSyclObjImpl<SyclObject>);
+}
+
+template <typename SyclObjectImpl>
+auto getSyclObjHandles(const std::vector<SyclObjectImpl> &Impls) {
+ return transformVec(
+ Impls, [&](const SyclObjectImpl &Impl) { return Impl->getHandle(); });
+}
+
template <typename SyclObject, typename Impl>
SyclObject createSyclObjFromImpl(Impl &&ImplObj) {
return ImplUtils::createSyclObjFromImpl<SyclObject>(
diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index 87f1a4d7e6c14..22dec053ac421 100644
--- a/libsycl/include/sycl/__impl/queue.hpp
+++ b/libsycl/include/sycl/__impl/queue.hpp
@@ -331,6 +331,36 @@ class _LIBSYCL_EXPORT queue {
std::forward<Rest>(rest)...);
}
+ /// Submits a memory copy operation from one USM or host pointer to another.
+ ///
+ /// \param dest is the pointer to copy to.
+ /// \param src is the pointer to copy from.
+ /// \param numBytes is the number of bytes to copy.
+ /// \return an event that represents the status of the operation.
+ event memcpy(void *dest, const void *src, std::size_t numBytes);
+
+ /// Submits a memory copy operation from one USM or host pointer to another.
+ ///
+ /// \param dest is the pointer to copy to.
+ /// \param src is the pointer to copy from.
+ /// \param numBytes is the number of bytes to copy.
+ /// \param depEvent is an event that represents a dependency for the
+ /// operation.
+ /// \return an event that represents the status of the operation.
+ event memcpy(void *dest, const void *src, std::size_t numBytes,
+ event depEvent);
+
+ /// Submits a memory copy operation from one USM or host pointer to another.
+ ///
+ /// \param dest is the pointer to copy to.
+ /// \param src is the pointer to copy from.
+ /// \param numBytes is the number of bytes to copy.
+ /// \param depEvents is a vector of events that represent dependencies for the
+ /// operation.
+ /// \return an event that represents the status of the operation.
+ event memcpy(void *dest, const void *src, std::size_t numBytes,
+ const std::vector<event> &depEvents);
+
private:
template <typename KernelName, int Dims, typename... Rest>
event parallelForImpl(range<Dims> numWorkItems,
diff --git a/libsycl/src/detail/offload/offload_topology.cpp b/libsycl/src/detail/offload/offload_topology.cpp
index ab4c57ecf37eb..9458876f14ac5 100644
--- a/libsycl/src/detail/offload/offload_topology.cpp
+++ b/libsycl/src/detail/offload/offload_topology.cpp
@@ -88,9 +88,8 @@ void discoverOffloadDevices() {
if (Res != OL_SUCCESS)
return true;
- // Ignore host and unknown backends
- if (OL_PLATFORM_BACKEND_HOST == OlBackend ||
- OL_PLATFORM_BACKEND_UNKNOWN == OlBackend)
+ // Ignore unknown backends
+ if (OL_PLATFORM_BACKEND_UNKNOWN == OlBackend)
return true;
// Ignore the device if the backend index exceeds the number of backends
diff --git a/libsycl/src/detail/platform_impl.cpp b/libsycl/src/detail/platform_impl.cpp
index 932f282619b0e..7f8d0e19350c0 100644
--- a/libsycl/src/detail/platform_impl.cpp
+++ b/libsycl/src/detail/platform_impl.cpp
@@ -43,6 +43,8 @@ const std::vector<PlatformImplUPtr> &PlatformImpl::getPlatforms() {
auto &PlatformCache = getPlatformCache();
for (const auto &Topo : getOffloadTopologies()) {
+ if (Topo.getBackend() == OL_PLATFORM_BACKEND_HOST)
+ continue;
size_t PlatformIndex = 0;
for (const auto &OffloadPlatform : Topo.getPlatforms()) {
PlatformCache.emplace_back(std::make_unique<PlatformImpl>(
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index beba3082f06e2..a10737c13d624 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -10,6 +10,7 @@
#include <detail/device_impl.hpp>
#include <detail/event_impl.hpp>
+#include <detail/global_objects.hpp>
#include <detail/program_manager.hpp>
#include <algorithm>
@@ -64,27 +65,35 @@ QueueImpl::~QueueImpl() {
backend QueueImpl::getBackend() const noexcept { return MDevice.getBackend(); }
+static ol_device_handle_t getHostOLDevice() {
+ auto HostDeviceRange =
+ getOffloadTopologies()[OL_PLATFORM_BACKEND_HOST].getDevices(0);
+ assert(HostDeviceRange.size() == 1);
+ return *HostDeviceRange.begin();
+}
+
void QueueImpl::wait() { callAndThrow(olSyncQueue, MOffloadQueue); }
-static bool checkEventsPlatformMatch(std::vector<EventImplPtr> &Events,
+static void checkEventsPlatformMatch(const std::vector<EventImplPtr> &Events,
const PlatformImpl &QueuePlatform) {
// liboffload limitation to olWaitEvents. We can't do any extra handling for
// cross context/platform events without host task support now.
// "The input events can be from any queue on any device provided by the
// same platform as `Queue`."
- return std::all_of(Events.cbegin(), Events.cend(),
- [&QueuePlatform](const EventImplPtr &Event) {
- return &Event->getPlatformImpl() == &QueuePlatform;
- });
+ if (!std::all_of(Events.cbegin(), Events.cend(),
+ [&QueuePlatform](const EventImplPtr &Event) {
+ return &Event->getPlatformImpl() == &QueuePlatform;
+ })) {
+ throw sycl::exception(
+ sycl::make_error_code(sycl::errc::feature_not_supported),
+ "libsycl doesn't support cross-context/platform event dependencies "
+ "yet.");
+ }
}
void QueueImpl::setKernelParameters(std::vector<EventImplPtr> &&Events,
const detail::UnifiedRangeView &Range) {
- if (!checkEventsPlatformMatch(Events, MDevice.getPlatformImpl()))
- throw sycl::exception(
- sycl::make_error_code(sycl::errc::feature_not_supported),
- "libsycl doesn't support cross-context/platform event dependencies "
- "now.");
+ checkEventsPlatformMatch(Events, MDevice.getPlatformImpl());
// TODO: this conversion and storing of only offload events is possible only
// while we don't have host tasks (or features based on host tasks, like
@@ -93,11 +102,7 @@ void QueueImpl::setKernelParameters(std::vector<EventImplPtr> &&Events,
// implemented on offload level (no data now).
assert(MCurrentSubmitInfo.DepEvents.empty() &&
"Kernel submission must clean up dependencies.");
- MCurrentSubmitInfo.DepEvents.reserve(Events.size());
- for (auto &Event : Events) {
- assert(Event && "Event impl object can't be nullptr");
- MCurrentSubmitInfo.DepEvents.push_back(Event->getHandle());
- }
+ MCurrentSubmitInfo.DepEvents = getSyclObjHandles(Events);
setKernelLaunchArgs(Range, MCurrentSubmitInfo.Range);
}
@@ -108,17 +113,7 @@ void QueueImpl::submitKernelImpl(DeviceKernelInfo &KernelInfo, void *ArgData,
KernelInfo, MDevice);
assert(Kernel);
- // TODO: liboffload supports only in-order queues and no cross context waiting
- // is available now that means that this code is excessive but correct. I
- // don't want to skip it and rely on default liboffload behaviour that is
- // applicable for in-order queue only. Once OOO queues are added this waiting
- // must be disabled for in-order queues. Once host tasks are added - cross
- // context dependencies should be enabled and checked as well.
- if (!MCurrentSubmitInfo.DepEvents.empty()) {
- callAndThrow(olWaitEvents, MOffloadQueue,
- MCurrentSubmitInfo.DepEvents.data(),
- MCurrentSubmitInfo.DepEvents.size());
- }
+ handleEventDependencies(MCurrentSubmitInfo.DepEvents);
assert(ArgData && "At least one argument must exist");
assert(ArgSize && "Arguments size must be greater than 0");
@@ -146,5 +141,66 @@ void QueueImpl::submitKernelImpl(DeviceKernelInfo &KernelInfo, void *ArgData,
EventImpl::createEventWithHandle(NewEvent, MDevice.getPlatformImpl());
}
+// Returns the {DeviceHandle, IsHostDevice} pair associated with the ptr.
+static std::pair<ol_device_handle_t, bool> getAllocDevice(const void *ptr) {
+ // TODO consider caching this information to avoid querying it every time.
+ ol_device_handle_t Device{};
+ [[maybe_unused]] ol_result_t Result =
+ callNoCheck(olGetMemInfo, ptr, OL_MEM_INFO_DEVICE,
+ sizeof(ol_device_handle_t), &Device);
+ if (detail::isFailed(Result)) {
+ // If liboffload could not find the allocation, assume it is a host one.
+ if (Result->Code == OL_ERRC_NOT_FOUND) {
+ return {getHostOLDevice(), true};
+ }
+ checkAndThrow(Result);
+ }
+
+ assert(Device);
+ return {Device, false};
+}
+
+std::shared_ptr<EventImpl>
+QueueImpl::memcpy(void *Dest, const void *Src, std::size_t NumBytes,
+ const std::vector<EventImplPtr> &DepEvents) {
+ if (!Dest || !Src) {
+ throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
+ "Nullptr argument in memcpy operation");
+ }
+ checkEventsPlatformMatch(DepEvents, MDevice.getPlatformImpl());
+
+ auto [DestOLDevice, IsDestOLDeviceHost] = getAllocDevice(Dest);
+ auto [SrcOLDevice, IsSrcOLDeviceHost] = getAllocDevice(Src);
+
+ // FIXME Currently, liboffload does not let us specify a queue for
+ // host-to-host cases, which means that the operation would be synchronous and
+ // any implicit dependencies wouldn't be respected for an in-order queue.
+ if (IsDestOLDeviceHost && IsSrcOLDeviceHost)
+ throw sycl::exception(
+ sycl::make_error_code(sycl::errc::feature_not_supported),
+ "Host-to-host copy is not implemented yet");
+
+ auto EventHandles = getSyclObjHandles(DepEvents);
+ handleEventDependencies(EventHandles);
+
+ callAndThrow(olMemcpy, MOffloadQueue, Dest, DestOLDevice, Src, SrcOLDevice,
+ NumBytes);
+ ol_event_handle_t NewEvent{};
+ ol_event_flags_t Flags{};
+ callAndThrow(olCreateEvent, MOffloadQueue, Flags, &NewEvent);
+ return EventImpl::createEventWithHandle(NewEvent, MDevice.getPlatformImpl());
+}
+
+void QueueImpl::handleEventDependencies(std::vector<ol_event_handle_t> &Deps) {
+ // TODO: liboffload supports only in-order queues and no cross context waiting
+ // is available now that means that this code is excessive but correct. I
+ // don't want to skip it and rely on default liboffload behaviour that is
+ // applicable for in-order queue only. Once OOO queues are added this waiting
+ // must be disabled for in-order queues. Once host tasks are added - cross
+ // context dependencies should be enabled and checked as well.
+ if (!Deps.empty()) {
+ callAndThrow(olWaitEvents, MOffloadQueue, Deps.data(), Deps.size());
+ }
+}
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/queue_impl.hpp b/libsycl/src/detail/queue_impl.hpp
index ac5b303e95dc6..d81503921f020 100644
--- a/libsycl/src/detail/queue_impl.hpp
+++ b/libsycl/src/detail/queue_impl.hpp
@@ -96,7 +96,19 @@ class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
void setKernelParameters(std::vector<EventImplPtr> &&Events,
const detail::UnifiedRangeView &Range);
+ /// Submits a memory copy operation from one USM or host pointer to another.
+ ///
+ /// \param Dest is the pointer to copy to.
+ /// \param Src is the pointer to copy from.
+ /// \param NumBytes is the number of bytes to copy.
+ /// \param DepEvents is a vector of dependencies for the operation.
+ /// \return an event impl object that represents the status of the operation.
+ EventImplPtr memcpy(void *Dest, const void *Src, std::size_t NumBytes,
+ const std::vector<EventImplPtr> &DepEvents);
+
private:
+ void handleEventDependencies(std::vector<ol_event_handle_t> &Deps);
+
// Queue features.
ol_queue_handle_t MOffloadQueue = {};
const bool MIsInorder;
diff --git a/libsycl/src/queue.cpp b/libsycl/src/queue.cpp
index b57324219e46b..3807de0f296a6 100644
--- a/libsycl/src/queue.cpp
+++ b/libsycl/src/queue.cpp
@@ -35,18 +35,29 @@ bool queue::is_in_order() const { return impl->isInOrder(); }
void queue::wait() { impl->wait(); }
+event queue::memcpy(void *dest, const void *src, std::size_t numBytes) {
+ return memcpy(dest, src, numBytes, std::vector<event>{});
+}
+
+event queue::memcpy(void *dest, const void *src, std::size_t numBytes,
+ event depEvent) {
+ return memcpy(dest, src, numBytes, std::vector<event>{depEvent});
+}
+event queue::memcpy(void *dest, const void *src, std::size_t numBytes,
+ const std::vector<event> &depEvents) {
+ std::shared_ptr<detail::EventImpl> EventImplPtr =
+ impl->memcpy(dest, src, numBytes, detail::getSyclObjImpls(depEvents));
+ assert(EventImplPtr);
+ return detail::createSyclObjFromImpl<event>(EventImplPtr);
+}
+
event queue::getLastEvent() {
return detail::createSyclObjFromImpl<event>(impl->getLastEvent());
}
void queue::setKernelParameters(const std::vector<event> &Events,
const detail::UnifiedRangeView &Range) {
- std::vector<detail::EventImplPtr> DepEventImplRefs;
- DepEventImplRefs.reserve(Events.size());
- for (const auto &Event : Events) {
- DepEventImplRefs.push_back(detail::getSyclObjImpl(Event));
- }
- return impl->setKernelParameters(std::move(DepEventImplRefs), Range);
+ return impl->setKernelParameters(detail::getSyclObjImpls(Events), Range);
}
void queue::submitKernelImpl(detail::DeviceKernelInfo &KernelInfo,
diff --git a/libsycl/test/usm/memcpy.cpp b/libsycl/test/usm/memcpy.cpp
new file mode 100644
index 0000000000000..2d738c4a11d5f
--- /dev/null
+++ b/libsycl/test/usm/memcpy.cpp
@@ -0,0 +1,116 @@
+// REQUIRES: any-device
+// RUN: %clangxx %sycl_options %s -o %t.out
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <cstddef>
+#include <numeric>
+#include <tuple>
+
+using namespace sycl;
+
+constexpr std::size_t DataSize = 1024;
+constexpr std::size_t NumBytes = DataSize * sizeof(int);
+
+// Tests memcpy API by making multiple memory allocations using AllocFs and
+// performing a sequence of copies from one allocation to the next,
+// using MemCpyFunc to specify dependencies.
+// Assumes that the first and the last allocations are accessible on host.
+template <typename MemcpyFuncT, typename... AllocFuncssT>
+void test(queue &Q, MemcpyFuncT MemCpyFunc,
+ std::tuple<AllocFuncssT...> AllocFs) {
+ constexpr std::size_t NAllocations = std::tuple_size_v<decltype(AllocFs)>;
+ static_assert(NAllocations > 1);
+
+ std::vector<std::shared_ptr<int>> PtrPipeline;
+ PtrPipeline.reserve(NAllocations);
+ std::apply([&](auto &&...Fs) { ((PtrPipeline.push_back(Fs())), ...); },
+ AllocFs);
+
+ std::iota(PtrPipeline[0].get(), PtrPipeline[0].get() + DataSize, 0);
+
+ // No dependencies for the first operation.
+ event LastEvent =
+ Q.memcpy(PtrPipeline[1].get(), PtrPipeline[0].get(), NumBytes);
+ for (int I = 1; I < NAllocations - 1; ++I) {
+ LastEvent = MemCpyFunc(Q, LastEvent, PtrPipeline[I + 1].get(),
+ PtrPipeline[I].get(), NumBytes);
+ }
+
+ Q.wait();
+
+ int *ResultPtr = PtrPipeline.back().get();
+ for (int I = 0; I < DataSize; ++I)
+ assert(ResultPtr[I] == I);
+}
+
+template <bool ExplicitDeps, typename MemCpyFuncT>
+void runTestsForMemCpyFunc(MemCpyFuncT MemCpyFunc) {
+ // TODO: add in_order property here when ExplicitDeps == false once it is
+ // implemented. For now all sycl::queue objects are in-order due to liboffload
+ // limitation.
+ sycl::queue Q;
+ auto HostAllocF = [&]() {
+ return std::shared_ptr<int>(new int[DataSize],
+ [&](int *Ptr) { delete[] Ptr; });
+ };
+ auto USMDeleter = [&](int *Ptr) { free(Ptr, Q); };
+ auto HostUSMAllocF = [&]() {
+ return std::shared_ptr<int>(malloc_host<int>(DataSize, Q), USMDeleter);
+ };
+ auto DeviceUSMAllocF = [&]() {
+ return std::shared_ptr<int>(malloc_device<int>(DataSize, Q), USMDeleter);
+ };
+ auto SharedUSMAllocF = [&]() {
+ return std::shared_ptr<int>(malloc_shared<int>(DataSize, Q), USMDeleter);
+ };
+ auto RunTest = [&](auto... AllocFuncs) {
+ test(Q, MemCpyFunc, std::tuple(AllocFuncs...));
+ };
+
+ // These cases only do a single memcpy, so they would end up using just the
+ // overload without dependencies regardless of what MemCpyFunc is.
+ // TODO Pass a default-constructed event as a dependency in these cases
+ // instead when those are implemented.
+ if constexpr (!ExplicitDeps) {
+ // TODO Remove try-catch once host-to-host copies are supported.
+ try {
+ RunTest(HostAllocF, HostAllocF);
+ assert(false);
+ } catch (sycl::exception e) {
+
+ assert(e.code() == make_error_code(errc::feature_not_supported));
+ }
+ RunTest(HostAllocF, HostUSMAllocF);
+ RunTest(HostAllocF, SharedUSMAllocF);
+
+ RunTest(HostUSMAllocF, HostUSMAllocF);
+ RunTest(HostUSMAllocF, HostAllocF);
+ RunTest(HostUSMAllocF, SharedUSMAllocF);
+
+ RunTest(SharedUSMAllocF, SharedUSMAllocF);
+ RunTest(SharedUSMAllocF, HostAllocF);
+ RunTest(SharedUSMAllocF, HostUSMAllocF);
+ }
+
+ RunTest(HostAllocF, DeviceUSMAllocF, HostAllocF);
+ RunTest(SharedUSMAllocF, DeviceUSMAllocF, SharedUSMAllocF);
+ RunTest(HostUSMAllocF, DeviceUSMAllocF, HostUSMAllocF);
+ RunTest(HostAllocF, DeviceUSMAllocF, DeviceUSMAllocF, HostAllocF);
+}
+
+int main() {
+ runTestsForMemCpyFunc<false>([&](queue &Q, event Dep, auto... OtherArgs) {
+ (void)Dep;
+ return Q.memcpy(OtherArgs...);
+ });
+ runTestsForMemCpyFunc<true>([&](queue &Q, event Dep, auto... OtherArgs) {
+ return Q.memcpy(OtherArgs..., Dep);
+ });
+ runTestsForMemCpyFunc<true>([&](queue &Q, event Dep, auto... OtherArgs) {
+ return Q.memcpy(OtherArgs..., std::vector<event>({Dep}));
+ });
+
+ return 0;
+}
>From 3a1d1a3385d7ba60b377f0e9a8dcee47d9622198 Mon Sep 17 00:00:00 2001
From: Sergey Semenov <sergey.semenov at intel.com>
Date: Wed, 1 Jul 2026 07:34:37 -0700
Subject: [PATCH 2/2] Add unit tests and apply comments
---
libsycl/include/sycl/__impl/async_handler.hpp | 2 +
.../include/sycl/__impl/detail/obj_utils.hpp | 26 +++---
libsycl/include/sycl/__impl/queue.hpp | 11 ++-
libsycl/src/detail/queue_impl.cpp | 34 +++++---
libsycl/src/detail/queue_impl.hpp | 1 +
libsycl/src/queue.cpp | 8 --
libsycl/test/usm/memcpy.cpp | 3 +-
libsycl/unittests/mock/helpers.cpp | 24 +++++-
libsycl/unittests/mock/helpers.hpp | 10 +++
libsycl/unittests/mock/mock.cpp | 13 +++
libsycl/unittests/queue/CMakeLists.txt | 1 +
libsycl/unittests/queue/memcpy.cpp | 85 +++++++++++++++++++
12 files changed, 176 insertions(+), 42 deletions(-)
create mode 100644 libsycl/unittests/queue/memcpy.cpp
diff --git a/libsycl/include/sycl/__impl/async_handler.hpp b/libsycl/include/sycl/__impl/async_handler.hpp
index 53bbcecbf48e4..258214401997e 100644
--- a/libsycl/include/sycl/__impl/async_handler.hpp
+++ b/libsycl/include/sycl/__impl/async_handler.hpp
@@ -19,6 +19,8 @@
#ifndef _LIBSYCL___IMPL_ASYNC_HANDLER_HPP
#define _LIBSYCL___IMPL_ASYNC_HANDLER_HPP
+#include <sycl/__impl/detail/config.hpp>
+
#include <functional>
_LIBSYCL_BEGIN_NAMESPACE_SYCL
diff --git a/libsycl/include/sycl/__impl/detail/obj_utils.hpp b/libsycl/include/sycl/__impl/detail/obj_utils.hpp
index ece4472550394..8a3518be41698 100644
--- a/libsycl/include/sycl/__impl/detail/obj_utils.hpp
+++ b/libsycl/include/sycl/__impl/detail/obj_utils.hpp
@@ -60,26 +60,24 @@ auto getSyclObjImpl(const SyclObject &Obj)
return ImplUtils::getSyclObjImpl(Obj);
}
-template <typename ObjT, typename FuncT>
-auto transformVec(const std::vector<ObjT> &Vec, FuncT Func) {
- std::vector<
- std::remove_const_t<std::remove_reference_t<decltype(Func(Vec[0]))>>>
- Result;
- Result.reserve(Vec.size());
- for (const ObjT &Obj : Vec)
- Result.push_back(Func(Obj));
- return Result;
-}
-
template <typename SyclObject>
auto getSyclObjImpls(const std::vector<SyclObject> &Objs) {
- return transformVec(Objs, getSyclObjImpl<SyclObject>);
+ std::vector<std::remove_const_t<
+ std::remove_reference_t<decltype(getSyclObjImpl(Objs[0]))>>>
+ Result;
+ Result.reserve(Objs.size());
+ for (const SyclObject &Obj : Objs)
+ Result.push_back(getSyclObjImpl(Obj));
+ return Result;
}
template <typename SyclObjectImpl>
auto getSyclObjHandles(const std::vector<SyclObjectImpl> &Impls) {
- return transformVec(
- Impls, [&](const SyclObjectImpl &Impl) { return Impl->getHandle(); });
+ std::vector<decltype(Impls[0]->getHandle())> Result;
+ Result.reserve(Impls.size());
+ for (const SyclObjectImpl &Impl : Impls)
+ Result.push_back(Impl->getHandle());
+ return Result;
}
template <typename SyclObject, typename Impl>
diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index 785ea3acf8f45..96d764e6f184a 100644
--- a/libsycl/include/sycl/__impl/queue.hpp
+++ b/libsycl/include/sycl/__impl/queue.hpp
@@ -333,14 +333,18 @@ class _LIBSYCL_EXPORT queue {
}
/// Submits a memory copy operation from one USM or host pointer to another.
+ /// USM pointers must be accessible on the device associated with the queue.
///
/// \param dest is the pointer to copy to.
/// \param src is the pointer to copy from.
/// \param numBytes is the number of bytes to copy.
/// \return an event that represents the status of the operation.
- event memcpy(void *dest, const void *src, std::size_t numBytes);
+ event memcpy(void *dest, const void *src, std::size_t numBytes) {
+ return memcpy(dest, src, numBytes, std::vector<event>{});
+ }
/// Submits a memory copy operation from one USM or host pointer to another.
+ /// USM pointers must be accessible on the device associated with the queue.
///
/// \param dest is the pointer to copy to.
/// \param src is the pointer to copy from.
@@ -349,9 +353,12 @@ class _LIBSYCL_EXPORT queue {
/// operation.
/// \return an event that represents the status of the operation.
event memcpy(void *dest, const void *src, std::size_t numBytes,
- event depEvent);
+ event depEvent) {
+ return memcpy(dest, src, numBytes, std::vector<event>{depEvent});
+ }
/// Submits a memory copy operation from one USM or host pointer to another.
+ /// USM pointers must be accessible on the device associated with the queue.
///
/// \param dest is the pointer to copy to.
/// \param src is the pointer to copy from.
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index a10737c13d624..8b0da48a5a93a 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -66,10 +66,9 @@ QueueImpl::~QueueImpl() {
backend QueueImpl::getBackend() const noexcept { return MDevice.getBackend(); }
static ol_device_handle_t getHostOLDevice() {
- auto HostDeviceRange =
- getOffloadTopologies()[OL_PLATFORM_BACKEND_HOST].getDevices(0);
- assert(HostDeviceRange.size() == 1);
- return *HostDeviceRange.begin();
+ static ol_device_handle_t HostDevice =
+ *(getOffloadTopologies()[OL_PLATFORM_BACKEND_HOST].getDevices(0).begin());
+ return HostDevice;
}
void QueueImpl::wait() { callAndThrow(olSyncQueue, MOffloadQueue); }
@@ -137,8 +136,7 @@ void QueueImpl::submitKernelImpl(DeviceKernelInfo &KernelInfo, void *ArgData,
ol_event_flags_t Flags{};
callAndThrow(olCreateEvent, MOffloadQueue, Flags, &NewEvent);
- MCurrentSubmitInfo.LastEvent =
- EventImpl::createEventWithHandle(NewEvent, MDevice.getPlatformImpl());
+ MCurrentSubmitInfo.LastEvent = createEvent();
}
// Returns the {DeviceHandle, IsHostDevice} pair associated with the ptr.
@@ -163,16 +161,22 @@ static std::pair<ol_device_handle_t, bool> getAllocDevice(const void *ptr) {
std::shared_ptr<EventImpl>
QueueImpl::memcpy(void *Dest, const void *Src, std::size_t NumBytes,
const std::vector<EventImplPtr> &DepEvents) {
+ checkEventsPlatformMatch(DepEvents, MDevice.getPlatformImpl());
+ auto EventHandles = getSyclObjHandles(DepEvents);
+ if (NumBytes == 0) {
+ handleEventDependencies(EventHandles);
+ return createEvent();
+ }
+
if (!Dest || !Src) {
throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
"Nullptr argument in memcpy operation");
}
- checkEventsPlatformMatch(DepEvents, MDevice.getPlatformImpl());
auto [DestOLDevice, IsDestOLDeviceHost] = getAllocDevice(Dest);
auto [SrcOLDevice, IsSrcOLDeviceHost] = getAllocDevice(Src);
- // FIXME Currently, liboffload does not let us specify a queue for
+ // TODO: Currently, liboffload does not let us specify a queue for
// host-to-host cases, which means that the operation would be synchronous and
// any implicit dependencies wouldn't be respected for an in-order queue.
if (IsDestOLDeviceHost && IsSrcOLDeviceHost)
@@ -180,15 +184,10 @@ QueueImpl::memcpy(void *Dest, const void *Src, std::size_t NumBytes,
sycl::make_error_code(sycl::errc::feature_not_supported),
"Host-to-host copy is not implemented yet");
- auto EventHandles = getSyclObjHandles(DepEvents);
handleEventDependencies(EventHandles);
-
callAndThrow(olMemcpy, MOffloadQueue, Dest, DestOLDevice, Src, SrcOLDevice,
NumBytes);
- ol_event_handle_t NewEvent{};
- ol_event_flags_t Flags{};
- callAndThrow(olCreateEvent, MOffloadQueue, Flags, &NewEvent);
- return EventImpl::createEventWithHandle(NewEvent, MDevice.getPlatformImpl());
+ return createEvent();
}
void QueueImpl::handleEventDependencies(std::vector<ol_event_handle_t> &Deps) {
@@ -202,5 +201,12 @@ void QueueImpl::handleEventDependencies(std::vector<ol_event_handle_t> &Deps) {
callAndThrow(olWaitEvents, MOffloadQueue, Deps.data(), Deps.size());
}
}
+
+EventImplPtr QueueImpl::createEvent() {
+ ol_event_handle_t NewEvent{};
+ ol_event_flags_t Flags{};
+ callAndThrow(olCreateEvent, MOffloadQueue, Flags, &NewEvent);
+ return EventImpl::createEventWithHandle(NewEvent, MDevice.getPlatformImpl());
+}
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/queue_impl.hpp b/libsycl/src/detail/queue_impl.hpp
index d81503921f020..de03713708217 100644
--- a/libsycl/src/detail/queue_impl.hpp
+++ b/libsycl/src/detail/queue_impl.hpp
@@ -108,6 +108,7 @@ class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
private:
void handleEventDependencies(std::vector<ol_event_handle_t> &Deps);
+ EventImplPtr createEvent();
// Queue features.
ol_queue_handle_t MOffloadQueue = {};
diff --git a/libsycl/src/queue.cpp b/libsycl/src/queue.cpp
index 3807de0f296a6..56366507e8e74 100644
--- a/libsycl/src/queue.cpp
+++ b/libsycl/src/queue.cpp
@@ -35,14 +35,6 @@ bool queue::is_in_order() const { return impl->isInOrder(); }
void queue::wait() { impl->wait(); }
-event queue::memcpy(void *dest, const void *src, std::size_t numBytes) {
- return memcpy(dest, src, numBytes, std::vector<event>{});
-}
-
-event queue::memcpy(void *dest, const void *src, std::size_t numBytes,
- event depEvent) {
- return memcpy(dest, src, numBytes, std::vector<event>{depEvent});
-}
event queue::memcpy(void *dest, const void *src, std::size_t numBytes,
const std::vector<event> &depEvents) {
std::shared_ptr<detail::EventImpl> EventImplPtr =
diff --git a/libsycl/test/usm/memcpy.cpp b/libsycl/test/usm/memcpy.cpp
index 2d738c4a11d5f..d8c870516810d 100644
--- a/libsycl/test/usm/memcpy.cpp
+++ b/libsycl/test/usm/memcpy.cpp
@@ -78,8 +78,7 @@ void runTestsForMemCpyFunc(MemCpyFuncT MemCpyFunc) {
try {
RunTest(HostAllocF, HostAllocF);
assert(false);
- } catch (sycl::exception e) {
-
+ } catch (const sycl::exception &e) {
assert(e.code() == make_error_code(errc::feature_not_supported));
}
RunTest(HostAllocF, HostUSMAllocF);
diff --git a/libsycl/unittests/mock/helpers.cpp b/libsycl/unittests/mock/helpers.cpp
index 6bc89e1b5713f..973aedbdcd975 100644
--- a/libsycl/unittests/mock/helpers.cpp
+++ b/libsycl/unittests/mock/helpers.cpp
@@ -46,6 +46,10 @@ void mock::MockLiboffload::initDefault() {
DefaultDevice = mock::createDummyHandleWithData<ol_device_handle_t>(
reinterpret_cast<unsigned char *>(&DefaultPlatform),
sizeof(DefaultPlatform));
+ HostPlatform = mock::createDummyHandle<ol_platform_handle_t>();
+ HostDevice = mock::createDummyHandleWithData<ol_device_handle_t>(
+ reinterpret_cast<unsigned char *>(&HostPlatform),
+ sizeof(HostPlatform));
return OL_SUCCESS;
});
@@ -55,6 +59,8 @@ void mock::MockLiboffload::initDefault() {
mock::releaseDummyHandle(DefaultPlatform);
mock::releaseDummyHandle(DefaultDevice);
+ mock::releaseDummyHandle(HostPlatform);
+ mock::releaseDummyHandle(HostDevice);
return OL_SUCCESS;
});
@@ -91,7 +97,9 @@ void mock::MockLiboffload::initDefault() {
if (PropSize != sizeof(ol_platform_backend_t))
return makeEmptyStrError(OL_ERRC_INVALID_SIZE);
assignAs<ol_platform_backend_t>(PropValue,
- OL_PLATFORM_BACKEND_LEVEL_ZERO);
+ Platform == HostPlatform
+ ? OL_PLATFORM_BACKEND_HOST
+ : OL_PLATFORM_BACKEND_LEVEL_ZERO);
return OL_SUCCESS;
}
@@ -159,6 +167,8 @@ void mock::MockLiboffload::initDefault() {
assert(DefaultDevice);
std::ignore = Callback(DefaultDevice, UserData);
+ assert(HostDevice);
+ std::ignore = Callback(HostDevice, UserData);
return OL_SUCCESS;
});
@@ -276,7 +286,17 @@ void mock::MockLiboffload::initDefault() {
}
return OL_SUCCESS;
});
-
+ ON_CALL(*this, olMemcpy)
+ .WillByDefault([this](ol_queue_handle_t Queue, void *DstPtr,
+ ol_device_handle_t DstDevice, const void *SrcPtr,
+ ol_device_handle_t SrcDevice,
+ size_t Size) -> ol_result_t {
+ if (!Queue || !DstDevice || !DstDevice)
+ return makeEmptyStrError(OL_ERRC_INVALID_NULL_HANDLE);
+ if (!DstPtr || !SrcPtr)
+ return makeEmptyStrError(OL_ERRC_INVALID_NULL_POINTER);
+ return OL_SUCCESS;
+ });
ON_CALL(*this, olCreateEvent)
.WillByDefault([this](ol_queue_handle_t Queue, ol_event_flags_t Flags,
ol_event_handle_t *Event) -> ol_result_t {
diff --git a/libsycl/unittests/mock/helpers.hpp b/libsycl/unittests/mock/helpers.hpp
index 793c687b7276b..8d2fa09d8d06f 100644
--- a/libsycl/unittests/mock/helpers.hpp
+++ b/libsycl/unittests/mock/helpers.hpp
@@ -116,12 +116,20 @@ class MockLiboffload {
const ol_kernel_launch_size_args_t *LaunchSizeArgs,
const ol_kernel_launch_prop_t *Properties, size_t NumArgs,
void **ArgPtrs, const size_t *ArgSizes));
+ MOCK_METHOD(ol_result_t, olMemcpy,
+ (ol_queue_handle_t Queue, void *DstPtr,
+ ol_device_handle_t DstDevice, const void *SrcPtr,
+ ol_device_handle_t SrcDevice, size_t Size));
+ MOCK_METHOD(ol_result_t, olGetMemInfo,
+ (const void *Ptr, ol_mem_info_t PropName, size_t PropSize,
+ void *PropValue));
ol_result_t makeEmptyStrError(ol_errc_t Code) {
auto [Iterator, Flag] =
Errors.emplace(std::make_pair(Code, ol_error_struct_t{Code, ""}));
return &Iterator->second;
}
+ ol_device_handle_t getHostOLDevice() { return HostDevice; }
private:
void initDefault();
@@ -129,6 +137,8 @@ class MockLiboffload {
std::unordered_map<ol_errc_t, ol_error_struct_t> Errors;
ol_platform_handle_t DefaultPlatform{};
ol_device_handle_t DefaultDevice{};
+ ol_platform_handle_t HostPlatform{};
+ ol_device_handle_t HostDevice{};
};
#ifndef _LIB_EXPORT
diff --git a/libsycl/unittests/mock/mock.cpp b/libsycl/unittests/mock/mock.cpp
index a4e66548c7f40..5dfdb278a6783 100644
--- a/libsycl/unittests/mock/mock.cpp
+++ b/libsycl/unittests/mock/mock.cpp
@@ -95,6 +95,19 @@ ol_result_t olLaunchKernel(ol_queue_handle_t Queue, ol_device_handle_t Device,
NumArgs, ArgPtrs, ArgSizes);
}
+ol_result_t olMemcpy(ol_queue_handle_t Queue, void *DstPtr,
+ ol_device_handle_t DstDevice, const void *SrcPtr,
+ ol_device_handle_t SrcDevice, size_t Size) {
+ return mock::getMockLiboffload().olMemcpy(Queue, DstPtr, DstDevice, SrcPtr,
+ SrcDevice, Size);
+}
+
+ol_result_t olGetMemInfo(const void *Ptr, ol_mem_info_t PropName,
+ size_t PropSize, void *PropValue) {
+ return mock::getMockLiboffload().olGetMemInfo(Ptr, PropName, PropSize,
+ PropValue);
+}
+
ol_result_t olCreateEvent(ol_queue_handle_t Queue, ol_event_flags_t Flags,
ol_event_handle_t *Event) {
return mock::getMockLiboffload().olCreateEvent(Queue, Flags, Event);
diff --git a/libsycl/unittests/queue/CMakeLists.txt b/libsycl/unittests/queue/CMakeLists.txt
index 2c66d3f110451..ea98c6d0d48fe 100644
--- a/libsycl/unittests/queue/CMakeLists.txt
+++ b/libsycl/unittests/queue/CMakeLists.txt
@@ -1,4 +1,5 @@
add_sycl_unittest(QueueTests
+ memcpy.cpp
queue.cpp
sycl_kernel_launch.cpp
)
diff --git a/libsycl/unittests/queue/memcpy.cpp b/libsycl/unittests/queue/memcpy.cpp
new file mode 100644
index 0000000000000..54d8a08295c92
--- /dev/null
+++ b/libsycl/unittests/queue/memcpy.cpp
@@ -0,0 +1,85 @@
+#include <mock/helpers.hpp>
+
+#include <sycl/__impl/device.hpp>
+#include <sycl/__impl/queue.hpp>
+
+#include <detail/device_impl.hpp>
+#include <detail/queue_impl.hpp>
+
+#include <gmock/gmock.h>
+#include <gtest/gtest.h>
+
+using namespace sycl;
+using namespace ::testing;
+
+TEST(Queue, Memcpy) {
+ constexpr int NumBytes = 32;
+ constexpr int NMemcpies = 4;
+ constexpr int NMemcpyAttempts = 5;
+
+ mock::MockWrapper Mock;
+ queue Q;
+
+ bool IsSrcHostPtr = false;
+ bool IsDstHostPtr = false;
+ int *SrcPtr = reinterpret_cast<int *>(1);
+ int *DstPtr = reinterpret_cast<int *>(2);
+ ol_device_handle_t OLDev =
+ detail::getSyclObjImpl(Q.get_device())->getOLHandle();
+
+ EXPECT_CALL(Mock.get(), olGetMemInfo(_, OL_MEM_INFO_DEVICE,
+ sizeof(ol_device_handle_t), _))
+ .Times(NMemcpyAttempts * 2)
+ .WillRepeatedly([&](const void *Ptr, ol_mem_info_t PropName,
+ size_t PropSize, void *PropValue) -> ol_result_t {
+ EXPECT_TRUE(Ptr == SrcPtr || Ptr == DstPtr);
+ bool IsHostPtr = Ptr == SrcPtr ? IsSrcHostPtr : IsDstHostPtr;
+ if (IsHostPtr)
+ return mock::getMockLiboffload().makeEmptyStrError(OL_ERRC_NOT_FOUND);
+ *(reinterpret_cast<ol_device_handle_t *>(PropValue)) = OLDev;
+ return OL_SUCCESS;
+ });
+ EXPECT_CALL(Mock.get(), olMemcpy(_, DstPtr, _, SrcPtr, _, NumBytes))
+ .Times(NMemcpies)
+ .WillRepeatedly([&](ol_queue_handle_t Queue, void *DstPtr,
+ ol_device_handle_t DstDevice, const void *SrcPtr,
+ ol_device_handle_t SrcDevice,
+ size_t Size) -> ol_result_t {
+ EXPECT_NE(Queue, nullptr);
+ ol_device_handle_t HostDevice =
+ mock::getMockLiboffload().getHostOLDevice();
+ EXPECT_EQ(DstDevice, IsDstHostPtr ? HostDevice : OLDev);
+ EXPECT_EQ(SrcDevice, IsSrcHostPtr ? HostDevice : OLDev);
+ return OL_SUCCESS;
+ });
+
+ EXPECT_CALL(Mock.get(), olCreateEvent(_, _, _)).Times(NMemcpies);
+
+ event Event = Q.memcpy(DstPtr, SrcPtr, NumBytes);
+
+ EXPECT_CALL(Mock.get(), olWaitEvents(_, _, 1));
+ Q.memcpy(DstPtr, SrcPtr, NumBytes, Event);
+
+ IsSrcHostPtr = true;
+ Q.memcpy(DstPtr, SrcPtr, NumBytes);
+ IsSrcHostPtr = false;
+ IsDstHostPtr = true;
+ Q.memcpy(DstPtr, SrcPtr, NumBytes);
+
+ IsSrcHostPtr = true;
+ try {
+ Q.memcpy(DstPtr, SrcPtr, NumBytes);
+ } catch (const exception &e) {
+ EXPECT_EQ(e.code(), make_error_code(errc::feature_not_supported));
+ }
+}
+
+TEST(Queue, MemcpyZeroBytes) {
+ mock::MockWrapper Mock;
+ queue Q;
+ EXPECT_CALL(Mock.get(), olWaitEvents(_, _, 1)).Times(1);
+ EXPECT_CALL(Mock.get(), olGetMemInfo(_, _, _, _)).Times(0);
+ EXPECT_CALL(Mock.get(), olMemcpy(_, _, _, _, _, _)).Times(0);
+ event Event = Q.memcpy(nullptr, nullptr, 0);
+ Q.memcpy(nullptr, nullptr, 0, Event);
+}
\ No newline at end of file
More information about the llvm-commits
mailing list