Cold-Air-Networks/mirror-llvm-project
2
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions 7

[libsycl] Implement memcpy API (#205369)

Browse source
This commit is contained in:
Sergey Semenov 2026-07-07 11:29:57 +02:00 committed by GitHub
parent 5318bd4123
commit 6e0de1cea1
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
15 changed files with 445 additions and 48 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

3
libsycl/docs/index.rst
View file

@ -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

2
libsycl/include/sycl/__impl/async_handler.hpp
View file

@ -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

21
libsycl/include/sycl/__impl/detail/obj_utils.hpp
View file

@ -22,6 +22,7 @@
#include <optional>
#include <type_traits>
#include <utility>
#include <vector>
_LIBSYCL_BEGIN_NAMESPACE_SYCL
@ -59,6 +60,26 @@ auto getSyclObjImpl(const SyclObject &Obj)
return ImplUtils::getSyclObjImpl(Obj);
}
template <typename SyclObject>
auto getSyclObjImpls(const std::vector<SyclObject> &Objs) {
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) {
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>
SyclObject createSyclObjFromImpl(Impl &&ImplObj) {
return ImplUtils::createSyclObjFromImpl<SyclObject>(

37
libsycl/include/sycl/__impl/queue.hpp
View file

@ -332,6 +332,43 @@ public:
std::forward<Rest>(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<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.
/// \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<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.
/// \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,

5
libsycl/src/detail/offload/offload_topology.cpp
View file

@ -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

2
libsycl/src/detail/platform_impl.cpp
View file

@ -45,6 +45,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>(

118
libsycl/src/detail/queue_impl.cpp
View file

@ -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,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<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 +101,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 +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<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) {
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<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());
}
}
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

13
libsycl/src/detail/queue_impl.hpp
View file

@ -96,7 +96,20 @@ public:
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);
EventImplPtr createEvent();
// Queue features.
ol_queue_handle_t MOffloadQueue = {};
const bool MIsInorder;

15
libsycl/src/queue.cpp
View file

@ -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<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,

115
libsycl/test/usm/memcpy.cpp Normal file
View file

@ -0,0 +1,115 @@
// REQUIRES: any-device
// RUN: %clangxx -fsycl %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 (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<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;
}

53
libsycl/unittests/mock/helpers.cpp
View file

@ -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;
});
@ -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<ol_event_handle_t>(
reinterpret_cast<unsigned char *>(&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<ol_event_handle_t>(
reinterpret_cast<unsigned char *>(&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;
});

10
libsycl/unittests/mock/helpers.hpp
View file

@ -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<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

13
libsycl/unittests/mock/mock.cpp
View file

@ -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);

1
libsycl/unittests/queue/CMakeLists.txt
View file

@ -1,4 +1,5 @@
add_sycl_unittest(QueueTests
memcpy.cpp
queue.cpp
sycl_kernel_launch.cpp
)

85
libsycl/unittests/queue/memcpy.cpp Normal file
View file

@ -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);
*(static_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);
}