[SYCL] Make sycl::marray trivially copyable (#4885)

The sycl::marray class is not trivially copyable but marked as device copyable. It makes problems with any class that contain a member with the sycl::marray type: that class will not be trivially copyable and therefore device copyable regardless the sycl::marray one is. To solve this issue, the sycl::marray class is converted into a trivially copyable one. The constexpr attribute for the copy constructor is saved so the class can still be copied in constant expressions.

Author: Pavel Samolysov

sycl/include/CL/sycl/marray.hpp

@@ -66,17 +66,9 @@ template <typename Type, std::size_t NumElements> class marray {
       typename = typename std::enable_if<sizeof...(ArgTN) == NumElements>::type>
   constexpr marray(const ArgTN &... Args) : MData{Args...} {}
 
-  constexpr marray(const marray<Type, NumElements> &Rhs) {
-    for (std::size_t I = 0; I < NumElements; ++I) {
-      MData[I] = Rhs.MData[I];
-    }
-  }
+  constexpr marray(const marray<Type, NumElements> &Rhs) = default;
 
-  constexpr marray(marray<Type, NumElements> &&Rhs) {
-    for (std::size_t I = 0; I < NumElements; ++I) {
-      MData[I] = Rhs.MData[I];
-    }
-  }
+  constexpr marray(marray<Type, NumElements> &&Rhs) = default;
 
   // Available only when: NumElements == 1
   template <std::size_t Size = NumElements,
@@ -92,12 +84,7 @@ template <typename Type, std::size_t NumElements> class marray {
 
   const_reference operator[](std::size_t index) const { return MData[index]; }
 
-  marray &operator=(const marray<Type, NumElements> &Rhs) {
-    for (std::size_t I = 0; I < NumElements; ++I) {
-      MData[I] = Rhs.MData[I];
-    }
-    return *this;
-  }
+  marray &operator=(const marray<Type, NumElements> &Rhs) = default;
 
   // broadcasting operator
   marray &operator=(const Type &Rhs) {

sycl/include/CL/sycl/types.hpp

@@ -2414,10 +2414,13 @@ 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
+// marray is device copyable if element type is device copyable and it is also
+// not trivially copyable (if the element type is trivially copyable, the marray
+// is device copyable by default).
 template <typename T, std::size_t N>
-struct is_device_copyable<sycl::marray<T, N>,
-                          std::enable_if_t<is_device_copyable<T>::value>>
+struct is_device_copyable<
+    sycl::marray<T, N>, std::enable_if_t<is_device_copyable<T>::value &&
+                                         !std::is_trivially_copyable<T>::value>>
     : std::true_type {};
 
 // vec is device copyable on host, on device vec is trivially copyable

sycl/test/basic_tests/marray/marray.cpp

@@ -94,5 +94,31 @@ int main() {
   b___ = !mint3{0, 1, 2};
   assert(b___[0] == true && b___[1] == false && b___[2] == false);
 
+  // check copyability
+  constexpr sycl::marray<double, 5> ma;
+  constexpr sycl::marray<double, 5> mb(ma);
+  constexpr sycl::marray<double, 5> mc = ma;
+
+  // check trivially copyability
+  struct Copyable {
+    int a;
+    double b;
+    const char *name;
+  };
+
+  static_assert(std::is_trivially_copyable<sycl::marray<Copyable, 5>>::value,
+                "sycl::marray<Copyable, 5> is not trivially copyable type");
+  static_assert(
+      !std::is_trivially_copyable<sycl::marray<std::string, 5>>::value,
+      "sycl::marray<std::string, 5> is trivially copyable type");
+
+  // check device copyability
+  static_assert(sycl::is_device_copyable<sycl::marray<std::tuple<>, 5>>::value,
+                "sycl::marray<std::tuple<>, 5> is not device copyable type");
+  static_assert(!sycl::is_device_copyable<sycl::marray<std::string, 5>>::value,
+                "sycl::marray<std::string, 5> is device copyable type");
+
+  return 0;
+
   return 0;
 }

sycl/test/basic_tests/valid_kernel_args.cpp

@@ -23,6 +23,10 @@ struct SomeStructure {
   } v;
 };
 
+struct SomeMarrayStructure {
+  cl::sycl::marray<double, 5> points;
+};
+
 #define CHECK_PASSING_TO_KERNEL_BY_VALUE(Type)                                 \
   static_assert(std::is_standard_layout<Type>::value,                          \
                 "Is not standard layouti type.");                     \
@@ -34,3 +38,4 @@ CHECK_PASSING_TO_KERNEL_BY_VALUE(int)
 CHECK_PASSING_TO_KERNEL_BY_VALUE(cl::sycl::cl_uchar4)
 CHECK_PASSING_TO_KERNEL_BY_VALUE(SomeStructure)
 #endif
+CHECK_PASSING_TO_KERNEL_BY_VALUE(SomeMarrayStructure)