[SYCL] Support scalar accessor in handler::copy(acc,ptr) and copy(ptr,acc) (#1634)

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

Author: v-klochkov

sycl/include/CL/sycl/handler.hpp

@@ -522,6 +522,79 @@ class __SYCL_EXPORT handler {
     return true;
   }
 
+#ifndef __SYCL_DEVICE_ONLY__
+  /// Copies the content of memory object accessed by Src into the memory
+  /// pointed by Dst.
+  ///
+  /// \param Src is a source SYCL accessor.
+  /// \param Dst is a pointer to destination memory.
+  template <typename TSrc, typename TDst, int Dim, access::mode AccMode,
+            access::target AccTarget, access::placeholder IsPH>
+  detail::enable_if_t<(Dim > 0)>
+  copyAccToPtrHost(accessor<TSrc, Dim, AccMode, AccTarget, IsPH> Src,
+                   TDst *Dst) {
+    range<Dim> Range = Src.get_range();
+    parallel_for<class __copyAcc2Ptr<TSrc, TDst, Dim, AccMode, AccTarget, IsPH>>
+        (Range, [=](id<Dim> Index) {
+      size_t LinearIndex = Index[0];
+      for (int I = 1; I < Dim; ++I)
+        LinearIndex += Range[I] * Index[I];
+      (reinterpret_cast<TSrc *>(Dst))[LinearIndex] = Src[Index];
+    });
+  }
+
+  /// Copies 1 element accessed by 0-dimensional accessor Src into the memory
+  /// pointed by Dst.
+  ///
+  /// \param Src is a source SYCL accessor.
+  /// \param Dst is a pointer to destination memory.
+  template <typename TSrc, typename TDst, int Dim, access::mode AccMode,
+            access::target AccTarget, access::placeholder IsPH>
+  detail::enable_if_t<Dim == 0>
+  copyAccToPtrHost(accessor<TSrc, Dim, AccMode, AccTarget, IsPH> Src,
+                   TDst *Dst) {
+    single_task<class __copyAcc2Ptr<TSrc, TDst, Dim, AccMode, AccTarget, IsPH>>
+        ([=]() {
+      *Dst = readFromFirstAccElement(Src);
+    });
+  }
+
+  /// Copies the memory pointed by Src into the memory accessed by Dst.
+  ///
+  /// \param Src is a pointer to source memory.
+  /// \param Dst is a destination SYCL accessor.
+  template <typename TSrc, typename TDst, int Dim, access::mode AccMode,
+            access::target AccTarget, access::placeholder IsPH>
+  detail::enable_if_t<(Dim > 0)>
+  copyPtrToAccHost(TDst *Src,
+                   accessor<TSrc, Dim, AccMode, AccTarget, IsPH> Dst) {
+    range<Dim> Range = Dst.get_range();
+    parallel_for<class __copyPtr2Acc<TSrc, TDst, Dim, AccMode, AccTarget, IsPH>>
+        (Range, [=](id<Dim> Index) {
+      size_t LinearIndex = Index[0];
+      for (int I = 1; I < Dim; ++I)
+        LinearIndex += Range[I] * Index[I];
+      Dst[Index] = (reinterpret_cast<TDst *>(Src))[LinearIndex];
+    });
+  }
+
+  /// Copies 1 element pointed by Src to memory accessed by 0-dimensional
+  /// accessor Dst.
+  ///
+  /// \param Src is a pointer to source memory.
+  /// \param Dst is a destination SYCL accessor.
+  template <typename TSrc, typename TDst, int Dim, access::mode AccMode,
+            access::target AccTarget, access::placeholder IsPH>
+  detail::enable_if_t<Dim == 0>
+  copyPtrToAccHost(TDst *Src,
+                   accessor<TSrc, Dim, AccMode, AccTarget, IsPH> Dst) {
+    single_task<class __copyPtr2Acc<TSrc, TDst, Dim, AccMode, AccTarget, IsPH>>
+        ([=]() {
+      writeToFirstAccElement(Dst, *Src);
+    });
+  }
+#endif // __SYCL_DEVICE_ONLY__
+
   constexpr static bool isConstOrGlobal(access::target AccessTarget) {
     return AccessTarget == access::target::global_buffer ||
            AccessTarget == access::target::constant_buffer;
@@ -1206,7 +1279,7 @@ class __SYCL_EXPORT handler {
 
   // Explicit copy operations API
 
-  /// Copies the contents of memory object accessed by Src into the memory
+  /// Copies the content of memory object accessed by Src into the memory
   /// pointed by Dst.
   ///
   /// Source must have at least as many bytes as the range accessed by Dst.
@@ -1228,7 +1301,7 @@ class __SYCL_EXPORT handler {
     copy(Src, RawDstPtr);
   }
 
-  /// Copies the contents of memory pointed by Src into the memory object
+  /// Copies the content of memory pointed by Src into the memory object
   /// accessed by Dst.
   ///
   /// Source must have at least as many bytes as the range accessed by Dst.
@@ -1251,14 +1324,13 @@ class __SYCL_EXPORT handler {
     copy(RawSrcPtr, Dst);
   }
 
-  /// Copies the contents of memory object accessed by Src into the memory
+  /// Copies the content of memory object accessed by Src into the memory
   /// pointed by Dst.
   ///
   /// Source must have at least as many bytes as the range accessed by Dst.
   ///
   /// \param Src is a source SYCL accessor.
   /// \param Dst is a pointer to destination memory.
-  // TODO: support 0-dimensional and atomic accessors.
   template <typename T_Src, typename T_Dst, int Dims, access::mode AccessMode,
             access::target AccessTarget,
             access::placeholder IsPlaceholder = access::placeholder::false_t>
@@ -1270,17 +1342,8 @@ class __SYCL_EXPORT handler {
 #ifndef __SYCL_DEVICE_ONLY__
     if (MIsHost) {
       // TODO: Temporary implementation for host. Should be handled by memory
-      // manger.
-      range<Dims> Range = Src.get_range();
-      parallel_for< class __copyAcc2Ptr< T_Src, T_Dst, Dims, AccessMode,
-                                         AccessTarget, IsPlaceholder>>
-                                         (Range, [=](id<Dims> Index) {
-        size_t LinearIndex = Index[0];
-        for (int I = 1; I < Dims; ++I)
-          LinearIndex += Range[I] * Index[I];
-        ((T_Src *)Dst)[LinearIndex] = Src[Index];
-      });
-
+      // manager.
+      copyAccToPtrHost(Src, Dst);
       return;
     }
 #endif
@@ -1297,14 +1360,13 @@ class __SYCL_EXPORT handler {
     MAccStorage.push_back(std::move(AccImpl));
   }
 
-  /// Copies the contents of memory pointed by Src into the memory object
+  /// Copies the content of memory pointed by Src into the memory object
   /// accessed by Dst.
   ///
   /// Source must have at least as many bytes as the range accessed by Dst.
   ///
   /// \param Src is a pointer to source memory.
   /// \param Dst is a destination SYCL accessor.
-  // TODO: support 0-dimensional and atomic accessors.
   template <typename T_Src, typename T_Dst, int Dims, access::mode AccessMode,
             access::target AccessTarget,
             access::placeholder IsPlaceholder = access::placeholder::false_t>
@@ -1317,17 +1379,8 @@ class __SYCL_EXPORT handler {
 #ifndef __SYCL_DEVICE_ONLY__
     if (MIsHost) {
       // TODO: Temporary implementation for host. Should be handled by memory
-      // manger.
-      range<Dims> Range = Dst.get_range();
-      parallel_for< class __copyPtr2Acc< T_Src, T_Dst, Dims, AccessMode,
-                                         AccessTarget, IsPlaceholder>>
-                                         (Range, [=](id<Dims> Index) {
-        size_t LinearIndex = Index[0];
-        for (int I = 1; I < Dims; ++I)
-          LinearIndex += Range[I] * Index[I];
-
-        Dst[Index] = ((T_Dst *)Src)[LinearIndex];
-      });
+      // manager.
+      copyPtrToAccHost(Src, Dst);
       return;
     }
 #endif
@@ -1344,7 +1397,7 @@ class __SYCL_EXPORT handler {
     MAccStorage.push_back(std::move(AccImpl));
   }
 
-  /// Copies the contents of memory object accessed by Src to the memory
+  /// Copies the content of memory object accessed by Src to the memory
   /// object accessed by Dst.
   ///
   /// Dst must have at least as many bytes as the range accessed by Src.

sycl/test/basic_tests/handler/handler_mem_op.cpp

@@ -251,6 +251,20 @@ template <typename T> void test_copy_ptr_acc() {
   for (size_t I = 0; I < Size; ++I) {
     assert(Data[I] == Values[I]);
   }
+
+  // Check copy from memory to 0-dimensional accessor.
+  T SrcValue = 99;
+  T DstValue = 0;
+  {
+    buffer<T, 1> DstBuf(&DstValue, range<1>(1));
+    queue Queue;
+    Queue.submit([&](handler &Cgh) {
+      accessor<T, 0, access::mode::discard_write, access::target::global_buffer>
+          DstAcc(DstBuf, Cgh);
+      Cgh.copy(&SrcValue, DstAcc);
+    });
+  }
+  assert(DstValue == 99);
 }
 
 template <typename T> void test_copy_acc_ptr() {
@@ -272,6 +286,38 @@ template <typename T> void test_copy_acc_ptr() {
   for (size_t I = 0; I < Size; ++I) {
     assert(Data[I] == Values[I]);
   }
+
+  // Check copy from 0-dimensional accessor to memory
+  T SrcValue = 99;
+  T DstValue = 0;
+  {
+    buffer<T, 1> SrcBuf(&SrcValue, range<1>(1));
+    queue Queue;
+    Queue.submit([&](handler &Cgh) {
+      accessor<T, 0, access::mode::read, access::target::global_buffer>
+          SrcAcc(SrcBuf, Cgh);
+      Cgh.copy(SrcAcc, &DstValue);
+    });
+  }
+  assert(DstValue == 99);
+
+  // Check copy from 0-dimensional placeholder accessor to memory
+  SrcValue = 77;
+  DstValue = 0;
+  {
+    buffer<T, 1> SrcBuf(&SrcValue, range<1>(1));
+    accessor<T, 0, access::mode::read, access::target::global_buffer,
+             access::placeholder::true_t>
+        SrcAcc(SrcBuf);
+    {
+      queue Queue;
+      Queue.submit([&](handler &Cgh) {
+        Cgh.require(SrcAcc);
+        Cgh.copy(SrcAcc, &DstValue);
+      });
+    }
+  }
+  assert(DstValue == 77);
 }
 
 template <typename T> void test_copy_shared_ptr_acc() {