[SYCL] Fix explicit copy operation for host device (#2627)

Fix the linear index calculation in handler::copy() for
host-to-host-device and host-device-to-host cases.

Signed-off-by: Sergey Semenov <[email protected]>

Author: sergey-semenov

sycl/include/CL/sycl/handler.hpp

@@ -641,9 +641,7 @@ class __SYCL_EXPORT handler {
     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];
+      const size_t LinearIndex = detail::getLinearIndex(Index, Range);
       using TSrcNonConst = typename std::remove_const<TSrc>::type;
       (reinterpret_cast<TSrcNonConst *>(Dst))[LinearIndex] = Src[Index];
     });
@@ -678,9 +676,7 @@ class __SYCL_EXPORT handler {
     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];
+      const size_t LinearIndex = detail::getLinearIndex(Index, Range);
       Dst[Index] = (reinterpret_cast<const TDst *>(Src))[LinearIndex];
     });
   }

sycl/test/basic_tests/handler/handler_mem_op.cpp

@@ -34,7 +34,9 @@ template <typename T> struct point {
 
 template <typename T> void test_fill(T Val);
 template <typename T> void test_copy_ptr_acc();
+template <typename T> void test_3D_copy_ptr_acc();
 template <typename T> void test_copy_acc_ptr();
+template <typename T> void test_3D_copy_acc_ptr();
 template <typename T> void test_copy_shared_ptr_acc();
 template <typename T> void test_copy_shared_ptr_const_acc();
 template <typename T> void test_copy_acc_shared_ptr();
@@ -72,6 +74,14 @@ int main() {
     test_copy_ptr_acc<point<int>>();
     test_copy_ptr_acc<point<float>>();
   }
+  // handler.copy(ptr, acc) 3D
+  {
+    test_3D_copy_ptr_acc<int>();
+    test_3D_copy_ptr_acc<int>();
+    test_3D_copy_ptr_acc<point<int>>();
+    test_3D_copy_ptr_acc<point<int>>();
+    test_3D_copy_ptr_acc<point<float>>();
+  }
   // handler.copy(acc, ptr)
   {
     test_copy_acc_ptr<int>();
@@ -80,6 +90,14 @@ int main() {
     test_copy_acc_ptr<point<int>>();
     test_copy_acc_ptr<point<float>>();
   }
+  // handler.copy(acc, ptr) 3D
+  {
+    test_3D_copy_acc_ptr<int>();
+    test_3D_copy_acc_ptr<int>();
+    test_3D_copy_acc_ptr<point<int>>();
+    test_3D_copy_acc_ptr<point<int>>();
+    test_3D_copy_acc_ptr<point<float>>();
+  }
   // handler.copy(shared_ptr, acc)
   {
     test_copy_shared_ptr_acc<int>();
@@ -277,6 +295,28 @@ template <typename T> void test_copy_ptr_acc() {
   assert(DstValue == 99);
 }
 
+template <typename T> void test_3D_copy_ptr_acc() {
+  const range<3> Range{2, 3, 4};
+  const size_t Size = 2 * 3 * 4;
+  T Data[Size] = {0};
+  T Values[Size] = {0};
+  for (size_t I = 0; I < Size; ++I)
+    Values[I] = I;
+
+  {
+    buffer<T, 3> Buffer(Data, Range);
+    queue Queue;
+    Queue.submit([&](handler &Cgh) {
+      accessor<T, 3, access::mode::write, access::target::global_buffer>
+          Accessor(Buffer, Cgh, Range);
+      Cgh.copy(Values, Accessor);
+    });
+  }
+
+  for (int I = 0; I < Size; ++I)
+    assert(Data[I] == Values[I]);
+}
+
 template <typename T> void test_copy_acc_ptr() {
   const size_t Size = 10;
   T Data[Size] = {0};
@@ -345,6 +385,28 @@ template <typename T> void test_copy_acc_ptr() {
   assert(DstValue == 77);
 }
 
+template <typename T> void test_3D_copy_acc_ptr() {
+  const range<3> Range{2, 3, 4};
+  const size_t Size = 2 * 3 * 4;
+  T Data[Size] = {0};
+  T Values[Size] = {0};
+  for (size_t I = 0; I < Size; ++I)
+    Data[I] = I;
+
+  {
+    buffer<T, 3> Buffer(Data, Range);
+    queue Queue;
+    Queue.submit([&](handler &Cgh) {
+      accessor<T, 3, access::mode::read, access::target::global_buffer>
+          Accessor(Buffer, Cgh, Range);
+      Cgh.copy(Accessor, Values);
+    });
+  }
+
+  for (size_t I = 0; I < Size; ++I)
+    assert(Data[I] == Values[I]);
+}
+
 template <typename T> void test_copy_shared_ptr_acc() {
   const size_t Size = 10;
   T Data[Size] = {0};