[SYCL] Fix handler::copy(acc, ptr) and handler::copy(ptr, acc) (#1684)

This patch fixes the cases when the memory referenced by a pointer
is either 'const void *' or 'void *'.
It also fixes handler::copy(ptr, acc) where 'acc' accesses elements
having type with const qualifier.

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

Author: v-klochkov

sycl/include/CL/sycl/handler.hpp

@@ -539,7 +539,8 @@ class __SYCL_EXPORT handler {
       size_t LinearIndex = Index[0];
       for (int I = 1; I < Dim; ++I)
         LinearIndex += Range[I] * Index[I];
-      (reinterpret_cast<TSrc *>(Dst))[LinearIndex] = Src[Index];
+      using TSrcNonConst = typename std::remove_const<TSrc>::type;
+      (reinterpret_cast<TSrcNonConst *>(Dst))[LinearIndex] = Src[Index];
     });
   }
 
@@ -555,7 +556,8 @@ class __SYCL_EXPORT handler {
                    TDst *Dst) {
     single_task<class __copyAcc2Ptr<TSrc, TDst, Dim, AccMode, AccTarget, IsPH>>
         ([=]() {
-      *Dst = readFromFirstAccElement(Src);
+      using TSrcNonConst = typename std::remove_const<TSrc>::type;
+      *(reinterpret_cast<TSrcNonConst *>(Dst)) = readFromFirstAccElement(Src);
     });
   }
 
@@ -566,15 +568,15 @@ class __SYCL_EXPORT handler {
   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) {
+  copyPtrToAccHost(TSrc *Src,
+                   accessor<TDst, 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];
+      Dst[Index] = (reinterpret_cast<const TDst *>(Src))[LinearIndex];
     });
   }
 
@@ -586,11 +588,11 @@ class __SYCL_EXPORT handler {
   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) {
+  copyPtrToAccHost(TSrc *Src,
+                   accessor<TDst, Dim, AccMode, AccTarget, IsPH> Dst) {
     single_task<class __copyPtr2Acc<TSrc, TDst, Dim, AccMode, AccTarget, IsPH>>
         ([=]() {
-      writeToFirstAccElement(Dst, *Src);
+      writeToFirstAccElement(Dst, *(reinterpret_cast<const TDst *>(Src)));
     });
   }
 #endif // __SYCL_DEVICE_ONLY__

sycl/test/basic_tests/handler/handler_mem_op.cpp

@@ -252,6 +252,20 @@ template <typename T> void test_copy_ptr_acc() {
     assert(Data[I] == Values[I]);
   }
 
+  // Check copy from 'const void *' memory to accessor.
+  {
+    buffer<T, 1> Buffer(Size);
+    queue Queue;
+    Queue.submit([&](handler &Cgh) {
+      auto Acc = Buffer.template get_access<access::mode::discard_write>(Cgh);
+      Cgh.copy(reinterpret_cast<const void *>(Values), Acc);
+    });
+
+    auto Acc = Buffer.template get_access<access::mode::read>();
+    for (int I = 0; I < Size; ++I)
+      assert(Acc[I] == Values[I]);
+  }
+
   // Check copy from memory to 0-dimensional accessor.
   T SrcValue = 99;
   T DstValue = 0;
@@ -287,6 +301,21 @@ template <typename T> void test_copy_acc_ptr() {
     assert(Data[I] == Values[I]);
   }
 
+  // Check copy from 'const T' accessor to 'void *' memory.
+  {
+    buffer<const T, 1> Buffer(reinterpret_cast<const T *>(Data), range<1>(Size));
+    T Dst[Size] = {0};
+    queue Queue;
+    Queue.submit([&](handler &Cgh) {
+      auto Acc = Buffer.template get_access<access::mode::read>(Cgh);
+      Cgh.copy(Acc, reinterpret_cast<void *>(Dst));
+    });
+    Queue.wait();
+
+    for (int I = 0; I < Size; ++I)
+      assert(Dst[I] == Data[I]);
+  }
+
   // Check copy from 0-dimensional accessor to memory
   T SrcValue = 99;
   T DstValue = 0;