[SYCL] Extends support for SYCL 2020 implicitly device copyable types

sycl/include/sycl/types.hpp

@@ -58,6 +58,7 @@
 #include <array>
 #include <cmath>
 #include <cstring>
+#include <variant>
 #ifndef __SYCL_DEVICE_ONLY__
 #include <cfenv>
 #endif
@@ -2401,6 +2402,35 @@ struct is_device_copyable<
 template <typename T>
 inline constexpr bool is_device_copyable_v = is_device_copyable<T>::value;
 
+// std::array<T, 0> is implicitly device copyable type.
+template <typename T>
+struct is_device_copyable<std::array<T, 0>> : std::true_type {};
+
+// std::array<T, N> is implicitly device copyable type if T is device copyable
+// and it is not trivially copyable (if it is trivially copyable it is device
+// copyable by default)
+template <typename T, std::size_t N>
+struct is_device_copyable<
+    std::array<T, N>,
+    std::enable_if_t<!std::is_trivially_copyable<std::array<T, N>>::value>>
+    : is_device_copyable<T> {};
+
+// std::optional<T> is implicitly device copyable type if T is device copyable
+// and it is not trivially copyable (if it is trivially copyable it is device
+// copyable by default)
+template <typename T>
+struct is_device_copyable<
+    std::optional<T>,
+    std::enable_if_t<!std::is_trivially_copyable<std::optional<T>>::value>>
+    : is_device_copyable<T> {};
+
+// std::pair<T1, T2> is implicitly device copyable type if T1 and T2 are device
+// copyable
+template <typename T1, typename T2>
+struct is_device_copyable<std::pair<T1, T2>>
+    : detail::bool_constant<is_device_copyable<T1>::value &&
+                            is_device_copyable<T2>::value> {};
+
 // std::tuple<> is implicitly device copyable type.
 template <> struct is_device_copyable<std::tuple<>> : std::true_type {};
 
@@ -2411,6 +2441,19 @@ struct is_device_copyable<std::tuple<T, Ts...>>
     : detail::bool_constant<is_device_copyable<T>::value &&
                             is_device_copyable<std::tuple<Ts...>>::value> {};
 
+// std::variant<> is implicitly device copyable type
+template <> struct is_device_copyable<std::variant<>> : std::true_type {};
+
+// std::variant<Ts...> is implicitly device copyable type if each type T of
+// Ts... is device copyable, and it is not trivially copyable (if it is
+// trivially copyable it is device copyable by default)
+template <typename T, typename... Ts>
+struct is_device_copyable<std::variant<T, Ts...>,
+                          std::enable_if_t<!std::is_trivially_copyable<
+                              std::variant<T, Ts...>>::value>>
+    : detail::bool_constant<is_device_copyable<T>::value &&
+                            is_device_copyable<std::variant<Ts...>>::value> {};
+
 // 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).
@@ -2420,6 +2463,14 @@ struct is_device_copyable<
                                          !std::is_trivially_copyable<T>::value>>
     : std::true_type {};
 
+// array is device copyable if element type is device copyable and it is also
+// not trivially copyable (if the element type is trivially copyable, the array
+// is device copyable by default).
+template <typename T, std::size_t N>
+struct is_device_copyable<
+    T[N], std::enable_if_t<!std::is_trivially_copyable<T>::value>>
+    : is_device_copyable<T> {};
+
 namespace detail {
 template <typename T, typename = void>
 struct IsDeprecatedDeviceCopyable : std::false_type {};
@@ -2433,6 +2484,10 @@ struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020")
                             std::is_trivially_destructible<T>::value &&
                             !is_device_copyable<T>::value>> : std::true_type {};
 
+template <typename T, int N>
+struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020")
+    IsDeprecatedDeviceCopyable<T[N]> : IsDeprecatedDeviceCopyable<T> {};
+
 #ifdef __SYCL_DEVICE_ONLY__
 // Checks that the fields of the type T with indices 0 to (NumFieldsToCheck - 1)
 // are device copyable.

sycl/test/basic_tests/implicit_device_copyable_types.cpp

@@ -0,0 +1,161 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+
+#include <sycl/sycl.hpp>
+#include <variant>
+
+struct ACopyable {
+  int i;
+  ACopyable() = default;
+  ACopyable(int _i) : i(_i) {}
+  ACopyable(const ACopyable &x) : i(x.i) {}
+};
+
+template <> struct sycl::is_device_copyable<ACopyable> : std::true_type {};
+
+int main() {
+  sycl::queue q;
+  {
+    std::pair<int, float> pair_arr[5];
+    std::pair<int, float> pair;
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() {
+         std::pair<int, float> p0 = pair_arr[0];
+         std::pair<int, float> p = pair;
+       });
+     }).wait_and_throw();
+  }
+
+  {
+    std::pair<ACopyable, float> pair_arr[5];
+    std::pair<ACopyable, float> pair;
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() {
+         std::pair<ACopyable, float> p0 = pair_arr[0];
+         std::pair<ACopyable, float> p = pair;
+       });
+     }).wait_and_throw();
+  }
+
+  {
+    std::tuple<int, float, bool> tuple_arr[5];
+    std::tuple<int, float, bool> tuple;
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() {
+         std::tuple<int, float, bool> t0 = tuple_arr[0];
+         std::tuple<int, float, bool> t = tuple;
+       });
+     }).wait_and_throw();
+  }
+
+  {
+    std::tuple<ACopyable, float, bool> tuple_arr[5];
+    std::tuple<ACopyable, float, bool> tuple;
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() {
+         std::tuple<ACopyable, float, bool> t0 = tuple_arr[0];
+         std::tuple<ACopyable, float, bool> t = tuple;
+       });
+     }).wait_and_throw();
+  }
+
+  {
+    std::variant<int, float, bool> variant_arr[5];
+    std::variant<int, float, bool> variant;
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() {
+         std::variant<int, float, bool> v0 = variant_arr[0];
+         std::variant<int, float, bool> v = variant;
+       });
+     }).wait_and_throw();
+  }
+
+  {
+    std::variant<ACopyable> variant_arr[5];
+    std::variant<ACopyable> variant;
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() {
+         // std::variant with complex types relies on virtual functions, so
+         // they cannot be used within sycl kernels
+         auto size = sizeof(variant_arr[0]);
+         size = sizeof(variant);
+       });
+     }).wait_and_throw();
+  }
+
+  {
+    std::array<int, 513> arr_arr[5];
+    std::array<int, 513> arr;
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() {
+         std::array<int, 513> arr0 = arr_arr[0];
+         std::array<int, 513> a = arr;
+       });
+     }).wait_and_throw();
+  }
+
+  {
+    std::array<ACopyable, 513> arr_arr[5];
+    std::array<ACopyable, 513> arr;
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() {
+         std::array<ACopyable, 513> arr0 = arr_arr[0];
+         std::array<ACopyable, 513> a = arr;
+       });
+     }).wait_and_throw();
+  }
+
+  {
+    sycl::queue q{};
+    std::optional<int> opt_arr[5];
+    std::optional<int> opt;
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() {
+         std::optional<int> o0 = opt_arr[0];
+         std::optional<int> o = opt;
+       });
+     }).wait_and_throw();
+  }
+
+  {
+    sycl::queue q{};
+    std::optional<ACopyable> opt_arr[5];
+    std::optional<ACopyable> opt;
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() {
+         std::optional<ACopyable> o0 = opt_arr[0];
+         std::optional<ACopyable> o = opt;
+       });
+     }).wait_and_throw();
+  }
+
+  {
+    std::string_view strv_arr[5];
+    std::string_view strv;
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() {
+         std::string_view str0 = strv_arr[0];
+         std::string_view str = strv;
+       });
+     }).wait_and_throw();
+  }
+
+#if __cpp_lib_span >= 202002
+  {
+    std::vector<int> v(5);
+    std::span<int> s{v.begin(), 4};
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() { int x = s[0]; });
+     }).wait_and_throw();
+  }
+#endif
+
+  {
+    std::vector<int> v(5);
+    sycl::span<int> s{v.data(), 4};
+    q.submit([&](sycl::handler &cgh) {
+       cgh.single_task([=]() { int x = s[0]; });
+     }).wait_and_throw();
+  }
+
+  return 0;
+}