From 6e0de1cea1c8c567304e5e539ba480cada3e28f5 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Tue, 7 Jul 2026 11:29:57 +0200 Subject: [PATCH] [libsycl] Implement memcpy API (#205369) --- libsycl/docs/index.rst | 3 + libsycl/include/sycl/__impl/async_handler.hpp | 2 + .../include/sycl/__impl/detail/obj_utils.hpp | 21 ++++ libsycl/include/sycl/__impl/queue.hpp | 37 ++++++ .../src/detail/offload/offload_topology.cpp | 5 +- libsycl/src/detail/platform_impl.cpp | 2 + libsycl/src/detail/queue_impl.cpp | 118 +++++++++++++----- libsycl/src/detail/queue_impl.hpp | 13 ++ libsycl/src/queue.cpp | 15 ++- libsycl/test/usm/memcpy.cpp | 115 +++++++++++++++++ libsycl/unittests/mock/helpers.cpp | 53 ++++++-- libsycl/unittests/mock/helpers.hpp | 10 ++ libsycl/unittests/mock/mock.cpp | 13 ++ libsycl/unittests/queue/CMakeLists.txt | 1 + libsycl/unittests/queue/memcpy.cpp | 85 +++++++++++++ 15 files changed, 445 insertions(+), 48 deletions(-) create mode 100644 libsycl/test/usm/memcpy.cpp create mode 100644 libsycl/unittests/queue/memcpy.cpp diff --git a/libsycl/docs/index.rst b/libsycl/docs/index.rst index a5d657f3e640..593cc4456153 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/async_handler.hpp b/libsycl/include/sycl/__impl/async_handler.hpp index 53bbcecbf48e..258214401997 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 + #include _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 dcb18f1c03de..8a3518be4169 100644 --- a/libsycl/include/sycl/__impl/detail/obj_utils.hpp +++ b/libsycl/include/sycl/__impl/detail/obj_utils.hpp @@ -22,6 +22,7 @@ #include #include #include +#include _LIBSYCL_BEGIN_NAMESPACE_SYCL @@ -59,6 +60,26 @@ auto getSyclObjImpl(const SyclObject &Obj) return ImplUtils::getSyclObjImpl(Obj); } +template +auto getSyclObjImpls(const std::vector &Objs) { + std::vector>> + Result; + Result.reserve(Objs.size()); + for (const SyclObject &Obj : Objs) + Result.push_back(getSyclObjImpl(Obj)); + return Result; +} + +template +auto getSyclObjHandles(const std::vector &Impls) { + std::vectorgetHandle())> Result; + Result.reserve(Impls.size()); + for (const SyclObjectImpl &Impl : Impls) + Result.push_back(Impl->getHandle()); + return Result; +} + template SyclObject createSyclObjFromImpl(Impl &&ImplObj) { return ImplUtils::createSyclObjFromImpl( diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp index d05592e5fa3f..96d764e6f184 100644 --- a/libsycl/include/sycl/__impl/queue.hpp +++ b/libsycl/include/sycl/__impl/queue.hpp @@ -332,6 +332,43 @@ public: std::forward(rest)...); } + /// 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) { + return memcpy(dest, src, numBytes, std::vector{}); + } + + /// 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. + /// \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) { + return memcpy(dest, src, numBytes, std::vector{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. + /// \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 &depEvents); + private: template event parallelForImpl(range numWorkItems, diff --git a/libsycl/src/detail/offload/offload_topology.cpp b/libsycl/src/detail/offload/offload_topology.cpp index 12ea9c78f108..c04024eef48c 100644 --- a/libsycl/src/detail/offload/offload_topology.cpp +++ b/libsycl/src/detail/offload/offload_topology.cpp @@ -86,9 +86,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 9e36b287fc59..96d52d1363a9 100644 --- a/libsycl/src/detail/platform_impl.cpp +++ b/libsycl/src/detail/platform_impl.cpp @@ -45,6 +45,8 @@ const std::vector &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( diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp index beba3082f06e..6eadecb42f8b 100644 --- a/libsycl/src/detail/queue_impl.cpp +++ b/libsycl/src/detail/queue_impl.cpp @@ -10,6 +10,7 @@ #include #include +#include #include #include @@ -64,27 +65,34 @@ QueueImpl::~QueueImpl() { backend QueueImpl::getBackend() const noexcept { return MDevice.getBackend(); } +static ol_device_handle_t getHostOLDevice() { + static ol_device_handle_t HostDevice = + *(getOffloadTopologies()[OL_PLATFORM_BACKEND_HOST].getDevices(0).begin()); + return HostDevice; +} + void QueueImpl::wait() { callAndThrow(olSyncQueue, MOffloadQueue); } -static bool checkEventsPlatformMatch(std::vector &Events, +static void checkEventsPlatformMatch(const std::vector &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 &&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 +101,7 @@ void QueueImpl::setKernelParameters(std::vector &&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 +112,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"); @@ -138,13 +132,77 @@ void QueueImpl::submitKernelImpl(DeviceKernelInfo &KernelInfo, void *ArgData, KernelInfo.getName().data() + ") failed with " + formatCodeString(Result)); + MCurrentSubmitInfo.LastEvent = createEvent(); +} + +// Returns the {DeviceHandle, IsHostDevice} pair associated with the ptr. +static std::pair 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 +QueueImpl::memcpy(void *Dest, const void *Src, std::size_t NumBytes, + const std::vector &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"); + } + + auto [DestOLDevice, IsDestOLDeviceHost] = getAllocDevice(Dest); + auto [SrcOLDevice, IsSrcOLDeviceHost] = getAllocDevice(Src); + + // 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) + throw sycl::exception( + sycl::make_error_code(sycl::errc::feature_not_supported), + "Host-to-host copy is not implemented yet"); + + handleEventDependencies(EventHandles); + callAndThrow(olMemcpy, MOffloadQueue, Dest, DestOLDevice, Src, SrcOLDevice, + NumBytes); + return createEvent(); +} + +void QueueImpl::handleEventDependencies(std::vector &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()); + } +} + +EventImplPtr QueueImpl::createEvent() { ol_event_handle_t NewEvent{}; ol_event_flags_t Flags{}; callAndThrow(olCreateEvent, MOffloadQueue, Flags, &NewEvent); - - MCurrentSubmitInfo.LastEvent = - EventImpl::createEventWithHandle(NewEvent, MDevice.getPlatformImpl()); + 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 ac5b303e95dc..de0371370821 100644 --- a/libsycl/src/detail/queue_impl.hpp +++ b/libsycl/src/detail/queue_impl.hpp @@ -96,7 +96,20 @@ public: void setKernelParameters(std::vector &&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 &DepEvents); + private: + void handleEventDependencies(std::vector &Deps); + EventImplPtr createEvent(); + // Queue features. ol_queue_handle_t MOffloadQueue = {}; const bool MIsInorder; diff --git a/libsycl/src/queue.cpp b/libsycl/src/queue.cpp index b57324219e46..56366507e8e7 100644 --- a/libsycl/src/queue.cpp +++ b/libsycl/src/queue.cpp @@ -35,18 +35,21 @@ 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, + const std::vector &depEvents) { + std::shared_ptr EventImplPtr = + impl->memcpy(dest, src, numBytes, detail::getSyclObjImpls(depEvents)); + assert(EventImplPtr); + return detail::createSyclObjFromImpl(EventImplPtr); +} + event queue::getLastEvent() { return detail::createSyclObjFromImpl(impl->getLastEvent()); } void queue::setKernelParameters(const std::vector &Events, const detail::UnifiedRangeView &Range) { - std::vector 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 000000000000..d45c3388dafe --- /dev/null +++ b/libsycl/test/usm/memcpy.cpp @@ -0,0 +1,115 @@ +// REQUIRES: any-device +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: %t.out + +#include + +#include +#include +#include + +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 +void test(queue &Q, MemcpyFuncT MemCpyFunc, + std::tuple AllocFs) { + constexpr std::size_t NAllocations = std::tuple_size_v; + static_assert(NAllocations > 1); + + std::vector> 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 +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(new int[DataSize], + [&](int *Ptr) { delete[] Ptr; }); + }; + auto USMDeleter = [&](int *Ptr) { free(Ptr, Q); }; + auto HostUSMAllocF = [&]() { + return std::shared_ptr(malloc_host(DataSize, Q), USMDeleter); + }; + auto DeviceUSMAllocF = [&]() { + return std::shared_ptr(malloc_device(DataSize, Q), USMDeleter); + }; + auto SharedUSMAllocF = [&]() { + return std::shared_ptr(malloc_shared(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 (const 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([&](queue &Q, event Dep, auto... OtherArgs) { + (void)Dep; + return Q.memcpy(OtherArgs...); + }); + runTestsForMemCpyFunc([&](queue &Q, event Dep, auto... OtherArgs) { + return Q.memcpy(OtherArgs..., Dep); + }); + runTestsForMemCpyFunc([&](queue &Q, event Dep, auto... OtherArgs) { + return Q.memcpy(OtherArgs..., std::vector({Dep})); + }); + + return 0; +} diff --git a/libsycl/unittests/mock/helpers.cpp b/libsycl/unittests/mock/helpers.cpp index 6bc89e1b5713..60704273e4c5 100644 --- a/libsycl/unittests/mock/helpers.cpp +++ b/libsycl/unittests/mock/helpers.cpp @@ -46,6 +46,10 @@ void mock::MockLiboffload::initDefault() { DefaultDevice = mock::createDummyHandleWithData( reinterpret_cast(&DefaultPlatform), sizeof(DefaultPlatform)); + HostPlatform = mock::createDummyHandle(); + HostDevice = mock::createDummyHandleWithData( + reinterpret_cast(&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(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; }); @@ -254,6 +264,19 @@ void mock::MockLiboffload::initDefault() { 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 { + if (!Queue) + return makeEmptyStrError(OL_ERRC_INVALID_NULL_HANDLE); + if (!Event) + return makeEmptyStrError(OL_ERRC_INVALID_NULL_POINTER); + std::ignore = Flags; + *Event = mock::createDummyHandleWithData( + reinterpret_cast(&Queue), sizeof(Queue)); + return OL_SUCCESS; + }); + ON_CALL(*this, olLaunchKernel) .WillByDefault([this](ol_queue_handle_t Queue, ol_device_handle_t Device, ol_symbol_handle_t Kernel, @@ -277,16 +300,28 @@ void mock::MockLiboffload::initDefault() { 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 { - if (!Queue) + 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 (!Event) + if (!DstPtr || !SrcPtr) return makeEmptyStrError(OL_ERRC_INVALID_NULL_POINTER); - std::ignore = Flags; - *Event = mock::createDummyHandleWithData( - reinterpret_cast(&Queue), sizeof(Queue)); + return OL_SUCCESS; + }); + + ON_CALL(*this, olGetMemInfo) + .WillByDefault([this](const void *Ptr, ol_mem_info_t PropName, + size_t PropSize, void *PropValue) -> ol_result_t { + if (!Ptr) + return makeEmptyStrError(OL_ERRC_INVALID_NULL_POINTER); + // Other properties are not used by the runtime yet + EXPECT_EQ(PropName, OL_MEM_INFO_DEVICE); + if (PropSize < sizeof(ol_device_handle_t)) + return makeEmptyStrError(OL_ERRC_INVALID_SIZE); + assignAs(PropValue, DefaultDevice); return OL_SUCCESS; }); diff --git a/libsycl/unittests/mock/helpers.hpp b/libsycl/unittests/mock/helpers.hpp index 793c687b7276..8d2fa09d8d06 100644 --- a/libsycl/unittests/mock/helpers.hpp +++ b/libsycl/unittests/mock/helpers.hpp @@ -116,12 +116,20 @@ public: 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 @@ private: std::unordered_map 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 a4e66548c7f4..5dfdb278a678 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 2c66d3f11045..ea98c6d0d48f 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 000000000000..d2fbbb0ea2df --- /dev/null +++ b/libsycl/unittests/queue/memcpy.cpp @@ -0,0 +1,85 @@ +#include + +#include +#include + +#include +#include + +#include +#include + +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(1); + int *DstPtr = reinterpret_cast(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); + *(static_cast(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); +}