[SYCL] Mark sycl::marray device copyable (intel#4427)

sycl::marray is device copyable if the element type is device copyable.
Also test was added to verify that the following types are device copyable:
 - sycl::marray<T> when T is device copyable;
 - sycl::vec<T> (all supported element types are device copyable);
 - sycl::id;
 - sycl::range.

Author: vladimirlaz

sycl/include/CL/sycl/types.hpp

@@ -46,12 +46,13 @@
 #error "SYCL device compiler is built without ext_vector_type support"
 #endif // __HAS_EXT_VECTOR_TYPE__
 
-#include <CL/sycl/aliases.hpp>
 #include <CL/sycl/access/access.hpp>
+#include <CL/sycl/aliases.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/helpers.hpp>
 #include <CL/sycl/detail/type_traits.hpp>
 #include <CL/sycl/half_type.hpp>
+#include <CL/sycl/marray.hpp>
 #include <CL/sycl/multi_ptr.hpp>
 
 #include <array>
@@ -2333,6 +2334,12 @@ struct is_device_copyable<std::tuple<T, Ts...>>
     : detail::bool_constant<is_device_copyable<T>::value &&
                             is_device_copyable<std::tuple<Ts...>>::value> {};
 
+// marray is device copyable if element type is device copyable
+template <typename T, std::size_t N>
+struct is_device_copyable<sycl::marray<T, N>,
+                          std::enable_if_t<is_device_copyable<T>::value>>
+    : std::true_type {};
+
 namespace detail {
 template <typename T, typename = void>
 struct IsDeprecatedDeviceCopyable : std::false_type {};

sycl/test/basic_tests/is_device_copyable.cpp

@@ -69,12 +69,20 @@ void test() {
   BCopyable IamBadButCopyable(1);
   C IamAlsoGood;
   DCopyable IamAlsoBadButCopyable{0};
+  marray<int, 5> MarrayForCopyableIsCopyable(0);
+  range<2> Range{1,2};
+  id<3> Id{1,2,3};
+  vec<int, 3> VecForCopyableIsCopyable{3};
   queue Q;
   Q.single_task<class TestA>([=] {
     int A = IamGood.i;
     int B = IamBadButCopyable.i;
     int C = IamAlsoBadButCopyable.i;
     int D = IamAlsoGood.i;
+    int E = MarrayForCopyableIsCopyable[0];
+    int F = Range[1];
+    int G = Id[2];
+    int H = VecForCopyableIsCopyable[0];
   });
 
   Q.single_task<class TestB>(FunctorA{});

sycl/test/basic_tests/is_device_copyable_neg.cpp

@@ -55,10 +55,12 @@ struct FunctorB : public D {
 void test() {
   A IamBad(1);
   B IamAlsoBad{0};
+  marray<B, 2> MarrayForNotCopyable;
   queue Q;
   Q.single_task<class TestA>([=] {
     int A = IamBad.i;
     int B = IamAlsoBad.i;
+    int MB = MarrayForNotCopyable[0].i;
   });
 
   FunctorA FA;
@@ -69,13 +71,16 @@ void test() {
 }
 
 // CHECK: static_assert failed due to requirement 'is_device_copyable<A, void>
-// CHECK: is_device_copyable_neg.cpp:59:5: note: in instantiation of function
+// CHECK: is_device_copyable_neg.cpp:60:5: note: in instantiation of function
 
 // CHECK: static_assert failed due to requirement 'is_device_copyable<B, void>
-// CHECK: is_device_copyable_neg.cpp:59:5: note: in instantiation of function
+// CHECK: is_device_copyable_neg.cpp:60:5: note: in instantiation of function
+
+// CHECK: static_assert failed due to requirement 'is_device_copyable<sycl::marray<B, 2>, void>
+// CHECK: is_device_copyable_neg.cpp:60:5: note: in instantiation of function
 
 // CHECK: static_assert failed due to requirement 'is_device_copyable<C, void>
-// CHECK: is_device_copyable_neg.cpp:65:5: note: in instantiation of function
+// CHECK: is_device_copyable_neg.cpp:67:5: note: in instantiation of function
 
 // CHECK: static_assert failed due to requirement 'is_device_copyable<D, void>
-// CHECK: is_device_copyable_neg.cpp:68:5: note: in instantiation of function
+// CHECK: is_device_copyable_neg.cpp:70:5: note: in instantiation of function