From c49ef9143747e3ff21713b4439529eb2a694d036 Mon Sep 17 00:00:00 2001
From: "Maronas, Marcos" <marcos.maronas@intel.com>
Date: Wed, 8 Feb 2023 08:07:47 -0800
Subject: [PATCH 1/5] Adds extensive testing for SYCL2020 implicitly device
 copyable types.

Signed-off-by: Maronas, Marcos <marcos.maronas@intel.com>
---
 SYCL/Basic/device_implicitly_copyable.cpp | 509 ++++++++++++++++++++++
 1 file changed, 509 insertions(+)
 create mode 100644 SYCL/Basic/device_implicitly_copyable.cpp

diff --git a/SYCL/Basic/device_implicitly_copyable.cpp b/SYCL/Basic/device_implicitly_copyable.cpp
new file mode 100644
index 0000000000..1e77e0ec79
--- /dev/null
+++ b/SYCL/Basic/device_implicitly_copyable.cpp
@@ -0,0 +1,509 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+//==--device_implicitly_copyable.cpp - SYCL implicit device copyable test --==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#include <cassert>
+#include <iostream>
+
+#include <sycl/sycl.hpp>
+
+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() {
+  constexpr size_t arr_size = 5;
+  constexpr int ref_val = 14;
+  sycl::queue q;
+  {
+    std::pair<int, float> pair_arr[arr_size];
+    std::pair<int, float> pair{ref_val, ref_val};
+    std::pair<int, float> result_pair_arr[arr_size];
+    std::pair<int, float> result_pair;
+
+    for (auto i = 0; i < arr_size; i++) {
+      pair_arr[i].first = i;
+      pair_arr[i].second = i;
+    }
+
+    {
+      sycl::buffer<std::pair<int, float>, 1> buf_pair_arr{
+          result_pair_arr, sycl::range<1>(arr_size)};
+      sycl::buffer<std::pair<int, float>, 1> buf_pair{&result_pair,
+                                                      sycl::range<1>(1)};
+
+      q.submit([&](sycl::handler &cgh) {
+         auto acc_pair_arr =
+             sycl::accessor{buf_pair_arr, cgh, sycl::read_write};
+         auto acc_pair = sycl::accessor{buf_pair, cgh, sycl::read_write};
+         cgh.single_task([=]() {
+           for (auto i = 0; i < arr_size; i++) {
+             acc_pair_arr[i] = pair_arr[i];
+           }
+           acc_pair[0] = pair;
+         });
+       }).wait_and_throw();
+    }
+
+    for (auto i = 0; i < arr_size; i++) {
+      assert(result_pair_arr[i].first == i);
+      assert(result_pair_arr[i].second == i);
+    }
+    assert(result_pair.first == ref_val && result_pair.second == ref_val);
+  }
+
+  {
+    std::pair<ACopyable, float> pair_arr[arr_size];
+    std::pair<ACopyable, float> pair{ACopyable(ref_val), ref_val};
+    std::pair<ACopyable, float> result_pair_arr[arr_size];
+    std::pair<ACopyable, float> result_pair;
+
+    for (auto i = 0; i < arr_size; i++) {
+      pair_arr[i].first = ACopyable(i);
+      pair_arr[i].second = i;
+    }
+
+    {
+      sycl::buffer<std::pair<ACopyable, float>, 1> buf_pair_arr{
+          result_pair_arr, sycl::range<1>(arr_size)};
+      sycl::buffer<std::pair<ACopyable, float>, 1> buf_pair{&result_pair,
+                                                            sycl::range<1>(1)};
+
+      q.submit([&](sycl::handler &cgh) {
+         auto acc_pair_arr =
+             sycl::accessor{buf_pair_arr, cgh, sycl::read_write};
+         auto acc_pair = sycl::accessor{buf_pair, cgh, sycl::read_write};
+         cgh.single_task([=]() {
+           for (auto i = 0; i < arr_size; i++) {
+             acc_pair_arr[i] = pair_arr[i];
+           }
+           acc_pair[0] = pair;
+         });
+       }).wait_and_throw();
+    }
+
+    for (auto i = 0; i < arr_size; i++) {
+      assert(result_pair_arr[i].first.i == i);
+      assert(result_pair_arr[i].second == i);
+    }
+    assert(result_pair.first.i == ref_val && result_pair.second == ref_val);
+  }
+
+  {
+    std::tuple<int, float, bool> tuple_arr[arr_size];
+    std::tuple<int, float, bool> tuple{ref_val, ref_val, true};
+    std::tuple<int, float, bool> result_tuple_arr[arr_size];
+    std::tuple<int, float, bool> result_tuple;
+
+    for (auto i = 0; i < arr_size; i++) {
+      auto &t = tuple_arr[i];
+      std::get<0>(t) = i;
+      std::get<1>(t) = i;
+      std::get<2>(t) = true;
+    }
+
+    {
+      sycl::buffer<std::tuple<int, float, bool>, 1> buf_tuple_arr{
+          result_tuple_arr, sycl::range<1>(arr_size)};
+      sycl::buffer<std::tuple<int, float, bool>, 1> buf_tuple{
+          &result_tuple, sycl::range<1>(1)};
+
+      q.submit([&](sycl::handler &cgh) {
+         auto acc_tuple_arr =
+             sycl::accessor{buf_tuple_arr, cgh, sycl::read_write};
+         auto acc_tuple = sycl::accessor{buf_tuple, cgh, sycl::read_write};
+         cgh.single_task([=]() {
+           for (auto i = 0; i < arr_size; i++) {
+             acc_tuple_arr[i] = tuple_arr[i];
+           }
+           acc_tuple[0] = tuple;
+         });
+       }).wait_and_throw();
+    }
+
+    for (auto i = 0; i < arr_size; i++) {
+      auto t = result_tuple_arr[i];
+      assert(std::get<0>(t) == i);
+      assert(std::get<1>(t) == i);
+      assert(std::get<2>(t) == true);
+    }
+    assert(std::get<0>(result_tuple) == ref_val);
+    assert(std::get<1>(result_tuple) == ref_val);
+    assert(std::get<2>(result_tuple) == true);
+  }
+
+  {
+    std::tuple<ACopyable, float, bool> tuple_arr[arr_size];
+    std::tuple<ACopyable, float, bool> tuple{ACopyable(ref_val), ref_val, true};
+    std::tuple<ACopyable, float, bool> result_tuple_arr[arr_size];
+    std::tuple<ACopyable, float, bool> result_tuple;
+
+    for (auto i = 0; i < arr_size; i++) {
+      auto &t = tuple_arr[i];
+      std::get<0>(t) = ACopyable(i);
+      std::get<1>(t) = i;
+      std::get<2>(t) = true;
+    }
+
+    {
+      sycl::buffer<std::tuple<ACopyable, float, bool>, 1> buf_tuple_arr{
+          result_tuple_arr, sycl::range<1>(arr_size)};
+      sycl::buffer<std::tuple<ACopyable, float, bool>, 1> buf_tuple{
+          &result_tuple, sycl::range<1>(1)};
+
+      q.submit([&](sycl::handler &cgh) {
+         auto acc_tuple_arr =
+             sycl::accessor{buf_tuple_arr, cgh, sycl::read_write};
+         auto acc_tuple = sycl::accessor{buf_tuple, cgh, sycl::read_write};
+         cgh.single_task([=]() {
+           for (auto i = 0; i < arr_size; i++) {
+             acc_tuple_arr[i] = tuple_arr[i];
+           }
+           acc_tuple[0] = tuple;
+         });
+       }).wait_and_throw();
+    }
+
+    for (auto i = 0; i < arr_size; i++) {
+      auto t = result_tuple_arr[i];
+      assert(std::get<0>(t).i == i);
+      assert(std::get<1>(t) == i);
+      assert(std::get<2>(t) == true);
+    }
+    assert(std::get<0>(result_tuple).i == ref_val);
+    assert(std::get<1>(result_tuple) == ref_val);
+    assert(std::get<2>(result_tuple) == true);
+  }
+
+  {
+    std::variant<int, float, bool> variant_arr[arr_size];
+    std::variant<int, float, bool> variant{14};
+    std::variant<int, float, bool> result_variant_arr[arr_size];
+    std::variant<int, float, bool> result_variant;
+
+    constexpr int variant_size = 3;
+    for (auto i = 0; i < arr_size; i++) {
+      auto &v = variant_arr[i];
+      auto index = i % variant_size;
+      if (index == 0) {
+        v = i;
+      } else if (index == 1) {
+        v = (float)i;
+      } else {
+        v = true;
+      }
+    }
+
+    {
+      sycl::buffer<std::variant<int, float, bool>, 1> buf_variant_arr{
+          result_variant_arr, sycl::range<1>(arr_size)};
+      sycl::buffer<std::variant<int, float, bool>, 1> buf_variant{
+          &result_variant, sycl::range<1>(1)};
+
+      q.submit([&](sycl::handler &cgh) {
+         auto acc_variant_arr =
+             sycl::accessor{buf_variant_arr, cgh, sycl::read_write};
+         auto acc_variant = sycl::accessor{buf_variant, cgh, sycl::read_write};
+         cgh.single_task([=]() {
+           for (auto i = 0; i < arr_size; i++) {
+             acc_variant_arr[i] = variant_arr[i];
+           }
+           acc_variant[0] = variant;
+         });
+       }).wait_and_throw();
+    }
+
+    for (auto i = 0; i < arr_size; i++) {
+      auto v = result_variant_arr[i];
+      auto index = i % variant_size;
+      if (index == 0) {
+        assert(std::get<0>(v) == i);
+      } else if (index == 1) {
+        assert(std::get<1>(v) == i);
+      } else {
+        assert(std::get<2>(v) == true);
+      }
+    }
+    assert(std::get<0>(result_variant) == ref_val);
+  }
+
+  {
+    std::variant<ACopyable> variant_arr[arr_size];
+    std::variant<ACopyable> variant;
+    std::variant<ACopyable> result_variant_arr[arr_size];
+    std::variant<ACopyable> result_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, arr_size> arr_arr[arr_size];
+    std::array<int, arr_size> arr;
+    std::array<int, arr_size> result_arr_arr[arr_size];
+    std::array<int, arr_size> result_arr;
+
+    for (auto i = 0; i < arr_size; i++) {
+      auto &a = arr_arr[i];
+      for (auto j = 0; j < arr_size; j++) {
+        a[j] = j;
+      }
+      arr[i] = i;
+    }
+
+    {
+      sycl::buffer<std::array<int, arr_size>, 1> buf_arr_arr{
+          result_arr_arr, sycl::range<1>(arr_size)};
+      sycl::buffer<std::array<int, arr_size>, 1> buf_arr{&result_arr,
+                                                         sycl::range<1>(1)};
+
+      q.submit([&](sycl::handler &cgh) {
+         auto acc_arr_arr = sycl::accessor{buf_arr_arr, cgh, sycl::read_write};
+         auto acc_arr = sycl::accessor{buf_arr, cgh, sycl::read_write};
+         cgh.single_task([=]() {
+           for (auto i = 0; i < arr_size; i++) {
+             acc_arr_arr[i] = arr_arr[i];
+           }
+           acc_arr[0] = arr;
+         });
+       }).wait_and_throw();
+    }
+
+    for (auto i = 0; i < arr_size; i++) {
+      auto a = result_arr_arr[i];
+      for (auto j = 0; j < arr_size; j++) {
+        assert(a[j] == j);
+      }
+      assert(result_arr[i] == i);
+    }
+  }
+
+  {
+    std::array<ACopyable, arr_size> arr_arr[arr_size];
+    std::array<ACopyable, arr_size> arr;
+    std::array<ACopyable, arr_size> result_arr_arr[arr_size];
+    std::array<ACopyable, arr_size> result_arr;
+
+    for (auto i = 0; i < arr_size; i++) {
+      auto &a = arr_arr[i];
+      for (auto j = 0; j < arr_size; j++) {
+        a[j] = ACopyable(j);
+      }
+      arr[i] = ACopyable(i);
+    }
+
+    {
+      sycl::buffer<std::array<ACopyable, arr_size>, 1> buf_arr_arr{
+          result_arr_arr, sycl::range<1>(arr_size)};
+      sycl::buffer<std::array<ACopyable, arr_size>, 1> buf_arr{
+          &result_arr, sycl::range<1>(1)};
+
+      q.submit([&](sycl::handler &cgh) {
+         auto acc_arr_arr = sycl::accessor{buf_arr_arr, cgh, sycl::read_write};
+         auto acc_arr = sycl::accessor{buf_arr, cgh, sycl::read_write};
+         cgh.single_task([=]() {
+           for (auto i = 0; i < arr_size; i++) {
+             acc_arr_arr[i] = arr_arr[i];
+           }
+           acc_arr[0] = arr;
+         });
+       }).wait_and_throw();
+    }
+
+    for (auto i = 0; i < arr_size; i++) {
+      auto a = result_arr_arr[i];
+      for (auto j = 0; j < arr_size; j++) {
+        assert(a[j].i == j);
+      }
+      assert(result_arr[i].i == i);
+    }
+  }
+
+  {
+    std::optional<int> opt_arr[arr_size];
+    std::optional<int> opt;
+    std::optional<int> result_opt_arr[arr_size];
+    std::optional<int> result_opt;
+
+    for (auto i = 0; i < arr_size; i++) {
+      opt_arr[i] = i;
+    }
+    opt = ref_val;
+
+    {
+      sycl::buffer<std::optional<int>, 1> buf_opt_arr{result_opt_arr,
+                                                      sycl::range<1>(arr_size)};
+      sycl::buffer<std::optional<int>, 1> buf_opt{&result_opt,
+                                                  sycl::range<1>(1)};
+
+      q.submit([&](sycl::handler &cgh) {
+         auto acc_opt_arr = sycl::accessor{buf_opt_arr, cgh, sycl::read_write};
+         auto acc_opt = sycl::accessor{buf_opt, cgh, sycl::read_write};
+         cgh.single_task([=]() {
+           for (auto i = 0; i < arr_size; i++) {
+             acc_opt_arr[i] = opt_arr[i];
+           }
+           acc_opt[0] = opt;
+         });
+       }).wait_and_throw();
+    }
+
+    for (auto i = 0; i < arr_size; i++) {
+      assert(result_opt_arr[i] == i);
+    }
+
+      assert(result_opt == ref_val);
+  }
+
+  {
+    std::optional<ACopyable> opt_arr[arr_size];
+    std::optional<ACopyable> opt;
+    std::optional<ACopyable> result_opt_arr[arr_size];
+    std::optional<ACopyable> result_opt;
+
+    for (auto i = 0; i < arr_size; i++) {
+      opt_arr[i] = ACopyable(i);
+    }
+    opt = ACopyable(ref_val);
+
+    {
+      sycl::buffer<std::optional<ACopyable>, 1> buf_opt_arr{
+          result_opt_arr, sycl::range<1>(arr_size)};
+      sycl::buffer<std::optional<ACopyable>, 1> buf_opt{&result_opt,
+                                                        sycl::range<1>(1)};
+
+      q.submit([&](sycl::handler &cgh) {
+         auto acc_opt_arr = sycl::accessor{buf_opt_arr, cgh, sycl::read_write};
+         auto acc_opt = sycl::accessor{buf_opt, cgh, sycl::read_write};
+         cgh.single_task([=]() {
+           for (auto i = 0; i < arr_size; i++) {
+             acc_opt_arr[i] = opt_arr[i];
+           }
+           acc_opt[0] = opt;
+         });
+       }).wait_and_throw();
+    }
+
+    for (auto i = 0; i < arr_size; i++) {
+      assert(result_opt_arr[i]->i == i);
+    }
+
+    assert(result_opt->i == ref_val);
+  }
+
+  {
+    std::string strv_arr_val[arr_size];
+    std::string strv_val{std::to_string(ref_val)};
+    std::string_view strv_arr[arr_size];
+    std::string_view strv{strv_val};
+    std::string_view result_strv_arr[arr_size];
+    std::string_view result_strv;
+
+    for (auto i = 0; i < arr_size; i++) {
+      strv_arr_val[i] = std::to_string(i);
+      strv_arr[i] = std::string_view{strv_arr_val[i]};
+    }
+
+    {
+      sycl::buffer<std::string_view, 1> buf_string_view_arr{
+          result_strv_arr, sycl::range<1>(arr_size)};
+      sycl::buffer<std::string_view, 1> buf_string_view{&result_strv,
+                                                        sycl::range<1>(1)};
+
+      q.submit([&](sycl::handler &cgh) {
+         auto acc_string_view_arr =
+             sycl::accessor{buf_string_view_arr, cgh, sycl::read_write};
+         auto acc_string_view =
+             sycl::accessor{buf_string_view, cgh, sycl::read_write};
+         cgh.single_task([=]() {
+           for (auto i = 0; i < arr_size; i++) {
+             acc_string_view_arr[i] = strv_arr[i];
+           }
+           acc_string_view[0] = strv;
+         });
+       }).wait_and_throw();
+    }
+
+    for (auto i = 0; i < arr_size; i++) {
+      assert(result_strv_arr[i] == std::to_string(i));
+    }
+
+    assert(result_strv == std::to_string(ref_val));
+  }
+
+#if __cpp_lib_span >= 202002
+  {
+    std::vector<int> v(arr_size);
+    std::span<int> s{v.data(), arr_size};
+    std::span<int> result_s{v.data(), arr_size};
+
+    for (auto i = 0; i < arr_size; i++) {
+      s[i] = i;
+    }
+
+    {
+      std::buffer<std::span<int>, 1> buf_span{&result_s, std::range<1>(1)};
+
+      q.submit([&](std::handler &cgh) {
+         auto acc_span_arr = std::accessor{buf_span, cgh, std::read_write};
+         cgh.single_task([=]() {
+           for (auto i = 0; i < arr_size; i++) {
+             acc_span_arr[0][i] = s[i];
+           }
+         });
+       }).wait_and_throw();
+    }
+
+    for (auto i = 0; i < arr_size; i++) {
+      assert(result_s[i] == i);
+    }
+  }
+#endif
+
+  {
+    std::vector<int> v(arr_size);
+    sycl::span<int> s{v.data(), arr_size};
+    sycl::span<int> result_s{v.data(), arr_size};
+
+    for (auto i = 0; i < arr_size; i++) {
+      s[i] = i;
+    }
+
+    {
+      sycl::buffer<sycl::span<int>, 1> buf_span{&result_s, sycl::range<1>(1)};
+
+      q.submit([&](sycl::handler &cgh) {
+         auto acc_span_arr = sycl::accessor{buf_span, cgh, sycl::read_write};
+         cgh.single_task([=]() {
+           for (auto i = 0; i < arr_size; i++) {
+             acc_span_arr[0][i] = s[i];
+           }
+         });
+       }).wait_and_throw();
+    }
+
+    for (auto i = 0; i < arr_size; i++) {
+      assert(result_s[i] == i);
+    }
+  }
+
+  std::cout << "Test passed" << std::endl;
+}

From 9783590b0a0fb3c8ab7f626e080c6db552f5971e Mon Sep 17 00:00:00 2001
From: "Maronas, Marcos" <marcos.maronas@intel.com>
Date: Wed, 8 Feb 2023 08:28:57 -0800
Subject: [PATCH 2/5] Fixes clang-format issues.

Signed-off-by: Maronas, Marcos <marcos.maronas@intel.com>
---
 SYCL/Basic/device_implicitly_copyable.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/SYCL/Basic/device_implicitly_copyable.cpp b/SYCL/Basic/device_implicitly_copyable.cpp
index 1e77e0ec79..14e2d074d6 100644
--- a/SYCL/Basic/device_implicitly_copyable.cpp
+++ b/SYCL/Basic/device_implicitly_copyable.cpp
@@ -370,7 +370,7 @@ int main() {
       assert(result_opt_arr[i] == i);
     }
 
-      assert(result_opt == ref_val);
+    assert(result_opt == ref_val);
   }
 
   {

From 6fc1e34fb801881d16bbdc131d86b969b73ff1aa Mon Sep 17 00:00:00 2001
From: "Maronas, Marcos" <marcos.maronas@intel.com>
Date: Thu, 9 Feb 2023 05:54:49 -0800
Subject: [PATCH 3/5] Moving common code to a function.

Signed-off-by: Maronas, Marcos <marcos.maronas@intel.com>
---
 SYCL/Basic/device_implicitly_copyable.cpp | 362 +++++++---------------
 1 file changed, 113 insertions(+), 249 deletions(-)

diff --git a/SYCL/Basic/device_implicitly_copyable.cpp b/SYCL/Basic/device_implicitly_copyable.cpp
index 14e2d074d6..ae6362be0d 100644
--- a/SYCL/Basic/device_implicitly_copyable.cpp
+++ b/SYCL/Basic/device_implicitly_copyable.cpp
@@ -24,39 +24,42 @@ struct ACopyable {
 
 template <> struct sycl::is_device_copyable<ACopyable> : std::true_type {};
 
+template <typename DataT, size_t ArrSize>
+void CaptureAndCopy(const DataT *data_arr, const DataT &data_scalar,
+                    DataT *result_arr, DataT *result_scalar, sycl::queue &q) {
+  sycl::buffer<DataT, 1> buf_arr{result_arr, sycl::range<1>(ArrSize)};
+  sycl::buffer<DataT, 1> buf_scalar{result_scalar, sycl::range<1>(1)};
+
+  q.submit([&](sycl::handler &cgh) {
+    auto acc_arr = sycl::accessor{buf_arr, cgh, sycl::read_write};
+    auto acc_scalar = sycl::accessor{buf_scalar, cgh, sycl::read_write};
+    cgh.single_task([=]() {
+      for (auto i = 0; i < ArrSize; i++) {
+        acc_arr[i] = data_arr[i];
+      }
+      acc_scalar[0] = data_scalar;
+    });
+  });
+}
+
 int main() {
   constexpr size_t arr_size = 5;
   constexpr int ref_val = 14;
   sycl::queue q;
   {
-    std::pair<int, float> pair_arr[arr_size];
-    std::pair<int, float> pair{ref_val, ref_val};
-    std::pair<int, float> result_pair_arr[arr_size];
-    std::pair<int, float> result_pair;
+    using DataT = std::pair<int, float>;
+    DataT pair_arr[arr_size];
+    DataT pair{ref_val, ref_val};
+    DataT result_pair_arr[arr_size];
+    DataT result_pair;
 
     for (auto i = 0; i < arr_size; i++) {
       pair_arr[i].first = i;
       pair_arr[i].second = i;
     }
 
-    {
-      sycl::buffer<std::pair<int, float>, 1> buf_pair_arr{
-          result_pair_arr, sycl::range<1>(arr_size)};
-      sycl::buffer<std::pair<int, float>, 1> buf_pair{&result_pair,
-                                                      sycl::range<1>(1)};
-
-      q.submit([&](sycl::handler &cgh) {
-         auto acc_pair_arr =
-             sycl::accessor{buf_pair_arr, cgh, sycl::read_write};
-         auto acc_pair = sycl::accessor{buf_pair, cgh, sycl::read_write};
-         cgh.single_task([=]() {
-           for (auto i = 0; i < arr_size; i++) {
-             acc_pair_arr[i] = pair_arr[i];
-           }
-           acc_pair[0] = pair;
-         });
-       }).wait_and_throw();
-    }
+    CaptureAndCopy<DataT, arr_size>(pair_arr, pair, result_pair_arr,
+                                    &result_pair, q);
 
     for (auto i = 0; i < arr_size; i++) {
       assert(result_pair_arr[i].first == i);
@@ -66,34 +69,19 @@ int main() {
   }
 
   {
-    std::pair<ACopyable, float> pair_arr[arr_size];
-    std::pair<ACopyable, float> pair{ACopyable(ref_val), ref_val};
-    std::pair<ACopyable, float> result_pair_arr[arr_size];
-    std::pair<ACopyable, float> result_pair;
+    using DataT = std::pair<ACopyable, float>;
+    DataT pair_arr[arr_size];
+    DataT pair{ACopyable(ref_val), ref_val};
+    DataT result_pair_arr[arr_size];
+    DataT result_pair;
 
     for (auto i = 0; i < arr_size; i++) {
       pair_arr[i].first = ACopyable(i);
       pair_arr[i].second = i;
     }
 
-    {
-      sycl::buffer<std::pair<ACopyable, float>, 1> buf_pair_arr{
-          result_pair_arr, sycl::range<1>(arr_size)};
-      sycl::buffer<std::pair<ACopyable, float>, 1> buf_pair{&result_pair,
-                                                            sycl::range<1>(1)};
-
-      q.submit([&](sycl::handler &cgh) {
-         auto acc_pair_arr =
-             sycl::accessor{buf_pair_arr, cgh, sycl::read_write};
-         auto acc_pair = sycl::accessor{buf_pair, cgh, sycl::read_write};
-         cgh.single_task([=]() {
-           for (auto i = 0; i < arr_size; i++) {
-             acc_pair_arr[i] = pair_arr[i];
-           }
-           acc_pair[0] = pair;
-         });
-       }).wait_and_throw();
-    }
+    CaptureAndCopy<DataT, arr_size>(pair_arr, pair, result_pair_arr,
+                                    &result_pair, q);
 
     for (auto i = 0; i < arr_size; i++) {
       assert(result_pair_arr[i].first.i == i);
@@ -103,10 +91,11 @@ int main() {
   }
 
   {
-    std::tuple<int, float, bool> tuple_arr[arr_size];
-    std::tuple<int, float, bool> tuple{ref_val, ref_val, true};
-    std::tuple<int, float, bool> result_tuple_arr[arr_size];
-    std::tuple<int, float, bool> result_tuple;
+    using DataT = std::tuple<int, float, bool>;
+    DataT tuple_arr[arr_size];
+    DataT tuple{ref_val, ref_val, true};
+    DataT result_tuple_arr[arr_size];
+    DataT result_tuple;
 
     for (auto i = 0; i < arr_size; i++) {
       auto &t = tuple_arr[i];
@@ -115,24 +104,8 @@ int main() {
       std::get<2>(t) = true;
     }
 
-    {
-      sycl::buffer<std::tuple<int, float, bool>, 1> buf_tuple_arr{
-          result_tuple_arr, sycl::range<1>(arr_size)};
-      sycl::buffer<std::tuple<int, float, bool>, 1> buf_tuple{
-          &result_tuple, sycl::range<1>(1)};
-
-      q.submit([&](sycl::handler &cgh) {
-         auto acc_tuple_arr =
-             sycl::accessor{buf_tuple_arr, cgh, sycl::read_write};
-         auto acc_tuple = sycl::accessor{buf_tuple, cgh, sycl::read_write};
-         cgh.single_task([=]() {
-           for (auto i = 0; i < arr_size; i++) {
-             acc_tuple_arr[i] = tuple_arr[i];
-           }
-           acc_tuple[0] = tuple;
-         });
-       }).wait_and_throw();
-    }
+    CaptureAndCopy<DataT, arr_size>(tuple_arr, tuple, result_tuple_arr,
+                                    &result_tuple, q);
 
     for (auto i = 0; i < arr_size; i++) {
       auto t = result_tuple_arr[i];
@@ -146,10 +119,11 @@ int main() {
   }
 
   {
-    std::tuple<ACopyable, float, bool> tuple_arr[arr_size];
-    std::tuple<ACopyable, float, bool> tuple{ACopyable(ref_val), ref_val, true};
-    std::tuple<ACopyable, float, bool> result_tuple_arr[arr_size];
-    std::tuple<ACopyable, float, bool> result_tuple;
+    using DataT = std::tuple<ACopyable, float, bool>;
+    DataT tuple_arr[arr_size];
+    DataT tuple{ACopyable(ref_val), ref_val, true};
+    DataT result_tuple_arr[arr_size];
+    DataT result_tuple;
 
     for (auto i = 0; i < arr_size; i++) {
       auto &t = tuple_arr[i];
@@ -158,24 +132,8 @@ int main() {
       std::get<2>(t) = true;
     }
 
-    {
-      sycl::buffer<std::tuple<ACopyable, float, bool>, 1> buf_tuple_arr{
-          result_tuple_arr, sycl::range<1>(arr_size)};
-      sycl::buffer<std::tuple<ACopyable, float, bool>, 1> buf_tuple{
-          &result_tuple, sycl::range<1>(1)};
-
-      q.submit([&](sycl::handler &cgh) {
-         auto acc_tuple_arr =
-             sycl::accessor{buf_tuple_arr, cgh, sycl::read_write};
-         auto acc_tuple = sycl::accessor{buf_tuple, cgh, sycl::read_write};
-         cgh.single_task([=]() {
-           for (auto i = 0; i < arr_size; i++) {
-             acc_tuple_arr[i] = tuple_arr[i];
-           }
-           acc_tuple[0] = tuple;
-         });
-       }).wait_and_throw();
-    }
+    CaptureAndCopy<DataT, arr_size>(tuple_arr, tuple, result_tuple_arr,
+                                    &result_tuple, q);
 
     for (auto i = 0; i < arr_size; i++) {
       auto t = result_tuple_arr[i];
@@ -189,10 +147,11 @@ int main() {
   }
 
   {
-    std::variant<int, float, bool> variant_arr[arr_size];
-    std::variant<int, float, bool> variant{14};
-    std::variant<int, float, bool> result_variant_arr[arr_size];
-    std::variant<int, float, bool> result_variant;
+    using DataT = std::variant<int, float, bool>;
+    DataT variant_arr[arr_size];
+    DataT variant{14};
+    DataT result_variant_arr[arr_size];
+    DataT result_variant;
 
     constexpr int variant_size = 3;
     for (auto i = 0; i < arr_size; i++) {
@@ -207,24 +166,8 @@ int main() {
       }
     }
 
-    {
-      sycl::buffer<std::variant<int, float, bool>, 1> buf_variant_arr{
-          result_variant_arr, sycl::range<1>(arr_size)};
-      sycl::buffer<std::variant<int, float, bool>, 1> buf_variant{
-          &result_variant, sycl::range<1>(1)};
-
-      q.submit([&](sycl::handler &cgh) {
-         auto acc_variant_arr =
-             sycl::accessor{buf_variant_arr, cgh, sycl::read_write};
-         auto acc_variant = sycl::accessor{buf_variant, cgh, sycl::read_write};
-         cgh.single_task([=]() {
-           for (auto i = 0; i < arr_size; i++) {
-             acc_variant_arr[i] = variant_arr[i];
-           }
-           acc_variant[0] = variant;
-         });
-       }).wait_and_throw();
-    }
+    CaptureAndCopy<DataT, arr_size>(variant_arr, variant, result_variant_arr,
+                                    &result_variant, q);
 
     for (auto i = 0; i < arr_size; i++) {
       auto v = result_variant_arr[i];
@@ -241,10 +184,11 @@ int main() {
   }
 
   {
-    std::variant<ACopyable> variant_arr[arr_size];
-    std::variant<ACopyable> variant;
-    std::variant<ACopyable> result_variant_arr[arr_size];
-    std::variant<ACopyable> result_variant;
+    using DataT = std::variant<ACopyable>;
+    DataT variant_arr[arr_size];
+    DataT variant;
+    DataT result_variant_arr[arr_size];
+    DataT result_variant;
     q.submit([&](sycl::handler &cgh) {
        cgh.single_task([=]() {
          // std::variant with complex types relies on virtual functions, so
@@ -256,10 +200,11 @@ int main() {
   }
 
   {
-    std::array<int, arr_size> arr_arr[arr_size];
-    std::array<int, arr_size> arr;
-    std::array<int, arr_size> result_arr_arr[arr_size];
-    std::array<int, arr_size> result_arr;
+    using DataT = std::array<int, arr_size>;
+    DataT arr_arr[arr_size];
+    DataT arr;
+    DataT result_arr_arr[arr_size];
+    DataT result_arr;
 
     for (auto i = 0; i < arr_size; i++) {
       auto &a = arr_arr[i];
@@ -269,23 +214,8 @@ int main() {
       arr[i] = i;
     }
 
-    {
-      sycl::buffer<std::array<int, arr_size>, 1> buf_arr_arr{
-          result_arr_arr, sycl::range<1>(arr_size)};
-      sycl::buffer<std::array<int, arr_size>, 1> buf_arr{&result_arr,
-                                                         sycl::range<1>(1)};
-
-      q.submit([&](sycl::handler &cgh) {
-         auto acc_arr_arr = sycl::accessor{buf_arr_arr, cgh, sycl::read_write};
-         auto acc_arr = sycl::accessor{buf_arr, cgh, sycl::read_write};
-         cgh.single_task([=]() {
-           for (auto i = 0; i < arr_size; i++) {
-             acc_arr_arr[i] = arr_arr[i];
-           }
-           acc_arr[0] = arr;
-         });
-       }).wait_and_throw();
-    }
+    CaptureAndCopy<DataT, arr_size>(arr_arr, arr, result_arr_arr, &result_arr,
+                                    q);
 
     for (auto i = 0; i < arr_size; i++) {
       auto a = result_arr_arr[i];
@@ -297,10 +227,11 @@ int main() {
   }
 
   {
-    std::array<ACopyable, arr_size> arr_arr[arr_size];
-    std::array<ACopyable, arr_size> arr;
-    std::array<ACopyable, arr_size> result_arr_arr[arr_size];
-    std::array<ACopyable, arr_size> result_arr;
+    using DataT = std::array<ACopyable, arr_size>;
+    DataT arr_arr[arr_size];
+    DataT arr;
+    DataT result_arr_arr[arr_size];
+    DataT result_arr;
 
     for (auto i = 0; i < arr_size; i++) {
       auto &a = arr_arr[i];
@@ -310,23 +241,8 @@ int main() {
       arr[i] = ACopyable(i);
     }
 
-    {
-      sycl::buffer<std::array<ACopyable, arr_size>, 1> buf_arr_arr{
-          result_arr_arr, sycl::range<1>(arr_size)};
-      sycl::buffer<std::array<ACopyable, arr_size>, 1> buf_arr{
-          &result_arr, sycl::range<1>(1)};
-
-      q.submit([&](sycl::handler &cgh) {
-         auto acc_arr_arr = sycl::accessor{buf_arr_arr, cgh, sycl::read_write};
-         auto acc_arr = sycl::accessor{buf_arr, cgh, sycl::read_write};
-         cgh.single_task([=]() {
-           for (auto i = 0; i < arr_size; i++) {
-             acc_arr_arr[i] = arr_arr[i];
-           }
-           acc_arr[0] = arr;
-         });
-       }).wait_and_throw();
-    }
+    CaptureAndCopy<DataT, arr_size>(arr_arr, arr, result_arr_arr, &result_arr,
+                                    q);
 
     for (auto i = 0; i < arr_size; i++) {
       auto a = result_arr_arr[i];
@@ -338,33 +254,19 @@ int main() {
   }
 
   {
-    std::optional<int> opt_arr[arr_size];
-    std::optional<int> opt;
-    std::optional<int> result_opt_arr[arr_size];
-    std::optional<int> result_opt;
+    using DataT = std::optional<int>;
+    DataT opt_arr[arr_size];
+    DataT opt;
+    DataT result_opt_arr[arr_size];
+    DataT result_opt;
 
     for (auto i = 0; i < arr_size; i++) {
       opt_arr[i] = i;
     }
     opt = ref_val;
 
-    {
-      sycl::buffer<std::optional<int>, 1> buf_opt_arr{result_opt_arr,
-                                                      sycl::range<1>(arr_size)};
-      sycl::buffer<std::optional<int>, 1> buf_opt{&result_opt,
-                                                  sycl::range<1>(1)};
-
-      q.submit([&](sycl::handler &cgh) {
-         auto acc_opt_arr = sycl::accessor{buf_opt_arr, cgh, sycl::read_write};
-         auto acc_opt = sycl::accessor{buf_opt, cgh, sycl::read_write};
-         cgh.single_task([=]() {
-           for (auto i = 0; i < arr_size; i++) {
-             acc_opt_arr[i] = opt_arr[i];
-           }
-           acc_opt[0] = opt;
-         });
-       }).wait_and_throw();
-    }
+    CaptureAndCopy<DataT, arr_size>(opt_arr, opt, result_opt_arr, &result_opt,
+                                    q);
 
     for (auto i = 0; i < arr_size; i++) {
       assert(result_opt_arr[i] == i);
@@ -374,33 +276,19 @@ int main() {
   }
 
   {
-    std::optional<ACopyable> opt_arr[arr_size];
-    std::optional<ACopyable> opt;
-    std::optional<ACopyable> result_opt_arr[arr_size];
-    std::optional<ACopyable> result_opt;
+    using DataT = std::optional<ACopyable>;
+    DataT opt_arr[arr_size];
+    DataT opt;
+    DataT result_opt_arr[arr_size];
+    DataT result_opt;
 
     for (auto i = 0; i < arr_size; i++) {
       opt_arr[i] = ACopyable(i);
     }
     opt = ACopyable(ref_val);
 
-    {
-      sycl::buffer<std::optional<ACopyable>, 1> buf_opt_arr{
-          result_opt_arr, sycl::range<1>(arr_size)};
-      sycl::buffer<std::optional<ACopyable>, 1> buf_opt{&result_opt,
-                                                        sycl::range<1>(1)};
-
-      q.submit([&](sycl::handler &cgh) {
-         auto acc_opt_arr = sycl::accessor{buf_opt_arr, cgh, sycl::read_write};
-         auto acc_opt = sycl::accessor{buf_opt, cgh, sycl::read_write};
-         cgh.single_task([=]() {
-           for (auto i = 0; i < arr_size; i++) {
-             acc_opt_arr[i] = opt_arr[i];
-           }
-           acc_opt[0] = opt;
-         });
-       }).wait_and_throw();
-    }
+    CaptureAndCopy<DataT, arr_size>(opt_arr, opt, result_opt_arr, &result_opt,
+                                    q);
 
     for (auto i = 0; i < arr_size; i++) {
       assert(result_opt_arr[i]->i == i);
@@ -410,37 +298,21 @@ int main() {
   }
 
   {
+    using DataT = std ::string_view;
     std::string strv_arr_val[arr_size];
     std::string strv_val{std::to_string(ref_val)};
-    std::string_view strv_arr[arr_size];
-    std::string_view strv{strv_val};
-    std::string_view result_strv_arr[arr_size];
-    std::string_view result_strv;
+    DataT strv_arr[arr_size];
+    DataT strv{strv_val};
+    DataT result_strv_arr[arr_size];
+    DataT result_strv;
 
     for (auto i = 0; i < arr_size; i++) {
       strv_arr_val[i] = std::to_string(i);
       strv_arr[i] = std::string_view{strv_arr_val[i]};
     }
 
-    {
-      sycl::buffer<std::string_view, 1> buf_string_view_arr{
-          result_strv_arr, sycl::range<1>(arr_size)};
-      sycl::buffer<std::string_view, 1> buf_string_view{&result_strv,
-                                                        sycl::range<1>(1)};
-
-      q.submit([&](sycl::handler &cgh) {
-         auto acc_string_view_arr =
-             sycl::accessor{buf_string_view_arr, cgh, sycl::read_write};
-         auto acc_string_view =
-             sycl::accessor{buf_string_view, cgh, sycl::read_write};
-         cgh.single_task([=]() {
-           for (auto i = 0; i < arr_size; i++) {
-             acc_string_view_arr[i] = strv_arr[i];
-           }
-           acc_string_view[0] = strv;
-         });
-       }).wait_and_throw();
-    }
+    CaptureAndCopy<DataT, arr_size>(strv_arr, strv, result_strv_arr,
+                                    &result_strv, q);
 
     for (auto i = 0; i < arr_size; i++) {
       assert(result_strv_arr[i] == std::to_string(i));
@@ -451,57 +323,49 @@ int main() {
 
 #if __cpp_lib_span >= 202002
   {
+    using DataT = std::span<int>;
     std::vector<int> v(arr_size);
-    std::span<int> s{v.data(), arr_size};
-    std::span<int> result_s{v.data(), arr_size};
+    DataT s_arr[arr_size];
+    DataT s{v.data(), arr_size};
+    DataT result_s_arr[arr_size];
+    DataT result_s{v.data(), arr_size};
 
     for (auto i = 0; i < arr_size; i++) {
       s[i] = i;
+      s_arr[i] = s;
     }
 
-    {
-      std::buffer<std::span<int>, 1> buf_span{&result_s, std::range<1>(1)};
-
-      q.submit([&](std::handler &cgh) {
-         auto acc_span_arr = std::accessor{buf_span, cgh, std::read_write};
-         cgh.single_task([=]() {
-           for (auto i = 0; i < arr_size; i++) {
-             acc_span_arr[0][i] = s[i];
-           }
-         });
-       }).wait_and_throw();
-    }
+    CaptureAndCopy<DataT, arr_size>(s_arr, s, result_s_arr, &result_s, q);
 
     for (auto i = 0; i < arr_size; i++) {
       assert(result_s[i] == i);
+      for (auto j = 0; j < arr_size; j++) {
+        assert(result_s_arr[i][j] == j);
+      }
     }
   }
 #endif
 
   {
+    using DataT = sycl::span<int>;
     std::vector<int> v(arr_size);
-    sycl::span<int> s{v.data(), arr_size};
-    sycl::span<int> result_s{v.data(), arr_size};
+    DataT s_arr[arr_size];
+    DataT s{v.data(), arr_size};
+    DataT result_s_arr[arr_size];
+    DataT result_s{v.data(), arr_size};
 
     for (auto i = 0; i < arr_size; i++) {
       s[i] = i;
+      s_arr[i] = s;
     }
 
-    {
-      sycl::buffer<sycl::span<int>, 1> buf_span{&result_s, sycl::range<1>(1)};
-
-      q.submit([&](sycl::handler &cgh) {
-         auto acc_span_arr = sycl::accessor{buf_span, cgh, sycl::read_write};
-         cgh.single_task([=]() {
-           for (auto i = 0; i < arr_size; i++) {
-             acc_span_arr[0][i] = s[i];
-           }
-         });
-       }).wait_and_throw();
-    }
+    CaptureAndCopy<DataT, arr_size>(s_arr, s, result_s_arr, &result_s, q);
 
     for (auto i = 0; i < arr_size; i++) {
       assert(result_s[i] == i);
+      for (auto j = 0; j < arr_size; j++) {
+        assert(result_s_arr[i][j] == j);
+      }
     }
   }
 

From 433aee7a731f327b0b10dea69c597a6b75254be4 Mon Sep 17 00:00:00 2001
From: "Maronas, Marcos" <marcos.maronas@intel.com>
Date: Fri, 10 Feb 2023 07:12:03 -0800
Subject: [PATCH 4/5] Clarifies comment.

Signed-off-by: Maronas, Marcos <marcos.maronas@intel.com>
---
 SYCL/Basic/device_implicitly_copyable.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/SYCL/Basic/device_implicitly_copyable.cpp b/SYCL/Basic/device_implicitly_copyable.cpp
index ae6362be0d..4895cea5ae 100644
--- a/SYCL/Basic/device_implicitly_copyable.cpp
+++ b/SYCL/Basic/device_implicitly_copyable.cpp
@@ -191,8 +191,8 @@ int main() {
     DataT result_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
+         // Some implementations of 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);
        });

From 356355a0a4ce69dc19fb492643d6f369e907652d Mon Sep 17 00:00:00 2001
From: "Maronas, Marcos" <marcos.maronas@intel.com>
Date: Tue, 14 Feb 2023 07:46:48 -0800
Subject: [PATCH 5/5] Fixes failing test in devices.

Signed-off-by: Maronas, Marcos <marcos.maronas@intel.com>
---
 SYCL/Basic/device_implicitly_copyable.cpp | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/SYCL/Basic/device_implicitly_copyable.cpp b/SYCL/Basic/device_implicitly_copyable.cpp
index 4895cea5ae..203c1ff190 100644
--- a/SYCL/Basic/device_implicitly_copyable.cpp
+++ b/SYCL/Basic/device_implicitly_copyable.cpp
@@ -27,6 +27,9 @@ template <> struct sycl::is_device_copyable<ACopyable> : std::true_type {};
 template <typename DataT, size_t ArrSize>
 void CaptureAndCopy(const DataT *data_arr, const DataT &data_scalar,
                     DataT *result_arr, DataT *result_scalar, sycl::queue &q) {
+  // We need to copy data_arr, otherwise when using a device it tries to use the host memory
+  DataT cpy_data_arr[ArrSize];
+  std::memcpy(cpy_data_arr, data_arr, sizeof(cpy_data_arr));
   sycl::buffer<DataT, 1> buf_arr{result_arr, sycl::range<1>(ArrSize)};
   sycl::buffer<DataT, 1> buf_scalar{result_scalar, sycl::range<1>(1)};
 
@@ -35,7 +38,7 @@ void CaptureAndCopy(const DataT *data_arr, const DataT &data_scalar,
     auto acc_scalar = sycl::accessor{buf_scalar, cgh, sycl::read_write};
     cgh.single_task([=]() {
       for (auto i = 0; i < ArrSize; i++) {
-        acc_arr[i] = data_arr[i];
+        acc_arr[i] = cpy_data_arr[i];
       }
       acc_scalar[0] = data_scalar;
     });