From 44e49deceb16a077590e9f81865841c4b6421d1a Mon Sep 17 00:00:00 2001 From: Joshua Cranmer Date: Tue, 3 Sep 2019 07:52:44 -0400 Subject: [PATCH] [SYCL] Implement USM memcpy/memset on handlers. Signed-off-by: Joshua Cranmer --- sycl/doc/extensions/USM/USM.adoc | 4 +- sycl/include/CL/sycl/detail/cg.hpp | 49 ++++++++++++++++- .../include/CL/sycl/detail/memory_manager.hpp | 9 +++ sycl/include/CL/sycl/handler.hpp | 30 ++++++++++ sycl/source/detail/memory_manager.cpp | 26 +++++++++ sycl/source/detail/scheduler/commands.cpp | 18 ++++++ sycl/test/usm/memcpy.cpp | 55 +++++++++++++++++++ sycl/test/usm/memset.cpp | 51 +++++++++++++++++ 8 files changed, 239 insertions(+), 3 deletions(-) create mode 100644 sycl/test/usm/memcpy.cpp create mode 100644 sycl/test/usm/memset.cpp diff --git a/sycl/doc/extensions/USM/USM.adoc b/sycl/doc/extensions/USM/USM.adoc index 044d074720d84..4cf76a3b4d0d2 100644 --- a/sycl/doc/extensions/USM/USM.adoc +++ b/sycl/doc/extensions/USM/USM.adoc @@ -256,7 +256,7 @@ class handler { ... public: ... - event memcpy(void* dest, const void* src, size_t count); + void memcpy(void* dest, const void* src, size_t count); }; class queue { @@ -279,7 +279,7 @@ class handler { ... public: ... - event memset(void* ptr, int value, size_t count); + void memset(void* ptr, int value, size_t count); }; class queue { diff --git a/sycl/include/CL/sycl/detail/cg.hpp b/sycl/include/CL/sycl/detail/cg.hpp index 854c2cd89d46e..51a3e4e89620f 100644 --- a/sycl/include/CL/sycl/detail/cg.hpp +++ b/sycl/include/CL/sycl/detail/cg.hpp @@ -326,7 +326,9 @@ class CG { COPY_ACC_TO_ACC, FILL, UPDATE_HOST, - RUN_ON_HOST_INTEL + RUN_ON_HOST_INTEL, + COPY_USM, + FILL_USM }; CG(CGTYPE Type, std::vector> ArgsStorage, @@ -461,6 +463,51 @@ class CGUpdateHost : public CG { Requirement *getReqToUpdate() { return MPtr; } }; +// The class which represents "copy" command group for USM pointers. +class CGCopyUSM : public CG { + void *MSrc; + void *MDst; + size_t MLength; + +public: + CGCopyUSM(void *Src, void *Dst, size_t Length, + std::vector> ArgsStorage, + std::vector AccStorage, + std::vector> SharedPtrStorage, + std::vector Requirements, + std::vector Events) + : CG(COPY_USM, std::move(ArgsStorage), std::move(AccStorage), + std::move(SharedPtrStorage), std::move(Requirements), + std::move(Events)), + MSrc(Src), MDst(Dst), MLength(Length) {} + + void *getSrc() { return MSrc; } + void *getDst() { return MDst; } + size_t getLength() { return MLength; } +}; + +// The class which represents "fill" command group for USM pointers. +class CGFillUSM : public CG { + std::vector MPattern; + void *MDst; + size_t MLength; + +public: + CGFillUSM(std::vector Pattern, void *DstPtr, size_t Length, + std::vector> ArgsStorage, + std::vector AccStorage, + std::vector> SharedPtrStorage, + std::vector Requirements, + std::vector Events) + : CG(FILL_USM, std::move(ArgsStorage), std::move(AccStorage), + std::move(SharedPtrStorage), std::move(Requirements), + std::move(Events)), + MPattern(std::move(Pattern)), MDst(DstPtr), MLength(Length) {} + void *getDst() { return MDst; } + size_t getLength() { return MLength; } + int getFill() { return MPattern[0]; } +}; + } // namespace detail } // namespace sycl } // namespace cl diff --git a/sycl/include/CL/sycl/detail/memory_manager.hpp b/sycl/include/CL/sycl/detail/memory_manager.hpp index 784308cb09c83..d6e9c83007f96 100644 --- a/sycl/include/CL/sycl/detail/memory_manager.hpp +++ b/sycl/include/CL/sycl/detail/memory_manager.hpp @@ -120,6 +120,15 @@ class MemoryManager { static void unmap(SYCLMemObjI *SYCLMemObj, void *Mem, QueueImplPtr Queue, void *MappedPtr, std::vector DepEvents, bool UseExclusiveQueue, RT::PiEvent &OutEvent); + + static void copy_usm(void *SrcMem, QueueImplPtr Queue, size_t Len, + void *DstMem, std::vector DepEvents, + bool UseExclusiveQueue, RT::PiEvent &OutEvent); + + static void fill_usm(void *DstMem, QueueImplPtr Queue, size_t Len, + int Pattern, std::vector DepEvents, + RT::PiEvent &OutEvent); + }; } // namespace detail } // namespace sycl diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 887d99ee66af9..2b7917961ffdd 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -164,6 +164,8 @@ class handler { void *MSrcPtr = nullptr; // Pointer to the dest host memory or accessor(depends on command type). void *MDstPtr = nullptr; + // Length to copy or fill (for USM operations). + size_t MLength = 0; // Pattern that is used to fill memory object in case command type is fill. std::vector MPattern; // Storage for a lambda or function object. @@ -383,6 +385,18 @@ class handler { std::move(MSharedPtrStorage), std::move(MRequirements), std::move(MEvents))); break; + case detail::CG::COPY_USM: + CommandGroup.reset(new detail::CGCopyUSM( + MSrcPtr, MDstPtr, MLength, std::move(MArgsStorage), + std::move(MAccStorage), std::move(MSharedPtrStorage), + std::move(MRequirements), std::move(MEvents))); + break; + case detail::CG::FILL_USM: + CommandGroup.reset(new detail::CGFillUSM( + std::move(MPattern), MDstPtr, MLength, std::move(MArgsStorage), + std::move(MAccStorage), std::move(MSharedPtrStorage), + std::move(MRequirements), std::move(MEvents))); + break; case detail::CG::NONE: throw runtime_error("Command group submitted without a kernel or a " "explicit memory operation."); @@ -1133,6 +1147,22 @@ class handler { }); } } + + // Copy memory from the source to the destination. + void memcpy(void* Dest, const void* Src, size_t Count) { + MSrcPtr = const_cast(Src); + MDstPtr = Dest; + MLength = Count; + MCGType = detail::CG::COPY_USM; + } + + // Fill the memory pointed to by the destination with the given bytes. + void memset(void *Dest, int Value, size_t Count) { + MDstPtr = Dest; + MPattern.push_back((char)Value); + MLength = Count; + MCGType = detail::CG::FILL_USM; + } }; } // namespace sycl } // namespace cl diff --git a/sycl/source/detail/memory_manager.cpp b/sycl/source/detail/memory_manager.cpp index c003c0ec840d5..ad48cc183e8a6 100644 --- a/sycl/source/detail/memory_manager.cpp +++ b/sycl/source/detail/memory_manager.cpp @@ -10,6 +10,7 @@ #include #include #include +#include #include #include @@ -462,6 +463,31 @@ void MemoryManager::unmap(SYCLMemObjI *SYCLMemObj, void *Mem, DepEvents.empty() ? nullptr : &DepEvents[0], &OutEvent)); } +void MemoryManager::copy_usm(void *SrcMem, QueueImplPtr SrcQueue, size_t Len, + void *DstMem, std::vector DepEvents, + bool UseExclusiveQueue, RT::PiEvent &OutEvent) { + RT::PiQueue Queue = UseExclusiveQueue + ? SrcQueue->getExclusiveQueueHandleRef() + : SrcQueue->getHandleRef(); + + sycl::context Context = SrcQueue->get_context(); + std::shared_ptr USMDispatch = + getSyclObjImpl(Context)->getUSMDispatch(); + PI_CHECK(USMDispatch->enqueueMemcpy(Queue, + /* blocking */ false, DstMem, SrcMem, Len, DepEvents.size(), + &DepEvents[0], &OutEvent)); +} + +void MemoryManager::fill_usm(void *Mem, QueueImplPtr Queue, size_t Length, + int Pattern, std::vector DepEvents, + RT::PiEvent &OutEvent) { + sycl::context Context = Queue->get_context(); + std::shared_ptr USMDispatch = + getSyclObjImpl(Context)->getUSMDispatch(); + PI_CHECK(USMDispatch->enqueueMemset(Queue->getHandleRef(), + Mem, Pattern, Length, DepEvents.size(), &DepEvents[0], &OutEvent)); +} + } // namespace detail } // namespace sycl } // namespace cl diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index db5284962dc3f..b51ea10797df5 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -480,6 +480,12 @@ void ExecCGCommand::printDot(std::ostream &Stream) const { case detail::CG::COPY_PTR_TO_ACC: Stream << "CG type: copy ptr to acc\\n"; break; + case detail::CG::COPY_USM: + Stream << "CG type: copy usm\\n"; + break; + case detail::CG::FILL_USM: + Stream << "CG type: fill usm\\n"; + break; default: Stream << "CG type: unknown\\n"; break; @@ -766,6 +772,18 @@ cl_int ExecCGCommand::enqueueImp() { return PI_SUCCESS; } + case CG::CGTYPE::COPY_USM: { + CGCopyUSM *Copy = (CGCopyUSM *)MCommandGroup.get(); + MemoryManager::copy_usm(Copy->getSrc(), MQueue, Copy->getLength(), + Copy->getDst(), std::move(RawEvents), MUseExclusiveQueue, Event); + return CL_SUCCESS; + } + case CG::CGTYPE::FILL_USM: { + CGFillUSM *Fill = (CGFillUSM *)MCommandGroup.get(); + MemoryManager::fill_usm(Fill->getDst(), MQueue, Fill->getLength(), + Fill->getFill(), std::move(RawEvents), Event); + return CL_SUCCESS; + } case CG::CGTYPE::NONE: default: throw runtime_error("CG type not implemented."); diff --git a/sycl/test/usm/memcpy.cpp b/sycl/test/usm/memcpy.cpp new file mode 100644 index 0000000000000..8fd622670dc8d --- /dev/null +++ b/sycl/test/usm/memcpy.cpp @@ -0,0 +1,55 @@ +//==---- memcpy.cpp - USM memcpy test --------------------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// RUN: %clangxx -fsycl %s -o %t1.out -lOpenCL +// RUN: %CPU_RUN_PLACEHOLDER %t1.out + +#include + +using namespace cl::sycl; + +static constexpr int count = 100; + +int main() { + queue q([](exception_list el) { + for (auto &e : el) + std::rethrow_exception(e); + }); + float *src = (float*)malloc_shared(sizeof(float) * count, q.get_device(), + q.get_context()); + float *dest = (float*)malloc_shared(sizeof(float) * count, q.get_device(), + q.get_context()); + for (int i = 0; i < count; i++) + src[i] = i; + + event init_copy = q.submit([&](handler &cgh) { + cgh.memcpy(dest, src, sizeof(float) * count); + }); + + q.submit([&](handler &cgh) { + cgh.depends_on(init_copy); + cgh.single_task([=]() { + for (int i = 0; i < count; i++) + dest[i] *= 2; + }); + }); + q.wait_and_throw(); + + for (int i = 0; i < count; i++) { + assert(dest[i] == i * 2); + } + + // Copying to nullptr should throw. + q.submit([&](handler &cgh) { + cgh.memcpy(nullptr, src, sizeof(float) * count); + }); + try { + q.wait_and_throw(); + assert(false && "Expected error from copying to nullptr"); + } catch (runtime_error e) { + } +} diff --git a/sycl/test/usm/memset.cpp b/sycl/test/usm/memset.cpp new file mode 100644 index 0000000000000..bb67fdfcccea9 --- /dev/null +++ b/sycl/test/usm/memset.cpp @@ -0,0 +1,51 @@ +//==---- memset.cpp - USM memset test --------------------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// RUN: %clangxx -fsycl %s -o %t1.out -lOpenCL +// RUN: %CPU_RUN_PLACEHOLDER %t1.out + +#include + +using namespace cl::sycl; + +static constexpr int count = 100; + +int main() { + queue q([](exception_list el) { + for (auto &e : el) + std::rethrow_exception(e); + }); + uint32_t *src = (uint32_t*)malloc_shared(sizeof(uint32_t) * count, q.get_device(), + q.get_context()); + + event init_copy = q.submit([&](handler &cgh) { + cgh.memset(src, 0x15, sizeof(uint32_t) * count); + }); + + q.submit([&](handler &cgh) { + cgh.depends_on(init_copy); + cgh.single_task([=]() { + for (int i = 0; i < count; i++) + src[i] *= 2; + }); + }); + q.wait_and_throw(); + + for (int i = 0; i < count; i++) { + assert(src[i] == 0x2a2a2a2a); + } + + // Filling to nullptr should throw. + q.submit([&](handler &cgh) { + cgh.memset(nullptr, 0, sizeof(uint32_t) * count); + }); + try { + q.wait_and_throw(); + assert(false && "Expected error from writing to nullptr"); + } catch (runtime_error e) { + } +}