[SYCL] Implement USM memcpy/memset on handlers.

Signed-off-by: Joshua Cranmer <[email protected]>

Author: jcranmer-intel

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 {

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<std::vector<char>> 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<std::vector<char>> ArgsStorage,
+            std::vector<detail::AccessorImplPtr> AccStorage,
+            std::vector<std::shared_ptr<const void>> SharedPtrStorage,
+            std::vector<Requirement *> Requirements,
+            std::vector<detail::EventImplPtr> 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<char> MPattern;
+  void *MDst;
+  size_t MLength;
+
+public:
+  CGFillUSM(std::vector<char> Pattern, void *DstPtr, size_t Length,
+            std::vector<std::vector<char>> ArgsStorage,
+            std::vector<detail::AccessorImplPtr> AccStorage,
+            std::vector<std::shared_ptr<const void>> SharedPtrStorage,
+            std::vector<Requirement *> Requirements,
+            std::vector<detail::EventImplPtr> 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

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<RT::PiEvent> DepEvents,
                     bool UseExclusiveQueue, RT::PiEvent &OutEvent);
+
+  static void copy_usm(void *SrcMem, QueueImplPtr Queue, size_t Len,
+                       void *DstMem, std::vector<RT::PiEvent> DepEvents,
+                       bool UseExclusiveQueue, RT::PiEvent &OutEvent);
+
+  static void fill_usm(void *DstMem, QueueImplPtr Queue, size_t Len,
+                       int Pattern, std::vector<RT::PiEvent> DepEvents,
+                       RT::PiEvent &OutEvent);
+
 };
 } // namespace detail
 } // namespace sycl

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<char> 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<void *>(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

sycl/source/detail/memory_manager.cpp

@@ -10,6 +10,7 @@
 #include <CL/sycl/detail/event_impl.hpp>
 #include <CL/sycl/detail/memory_manager.hpp>
 #include <CL/sycl/detail/queue_impl.hpp>
+#include <CL/sycl/detail/usm_dispatch.hpp>
 
 #include <algorithm>
 #include <cassert>
@@ -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<RT::PiEvent> DepEvents,
+                             bool UseExclusiveQueue, RT::PiEvent &OutEvent) {
+  RT::PiQueue Queue = UseExclusiveQueue
+    ? SrcQueue->getExclusiveQueueHandleRef()
+    : SrcQueue->getHandleRef();
+
+  sycl::context Context = SrcQueue->get_context();
+  std::shared_ptr<usm::USMDispatcher> 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<RT::PiEvent> DepEvents,
+                             RT::PiEvent &OutEvent) {
+  sycl::context Context = Queue->get_context();
+  std::shared_ptr<usm::USMDispatcher> USMDispatch =
+    getSyclObjImpl(Context)->getUSMDispatch();
+  PI_CHECK(USMDispatch->enqueueMemset(Queue->getHandleRef(),
+    Mem, Pattern, Length, DepEvents.size(), &DepEvents[0], &OutEvent));
+}
+
 } // namespace detail
 } // namespace sycl
 } // namespace cl

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.");

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 <CL/sycl.hpp>
+
+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<class double_dest>([=]() {
+        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) {
+  }
+}

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 <CL/sycl.hpp>
+
+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<class double_dest>([=]() {
+        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) {
+  }
+}