Add tests for accessor version of copy_to

Author: fineg74

SYCL/ESIMD/regression/copyto_char_test.cpp

@@ -26,7 +26,7 @@ using shared_vector = std::vector<DataT, shared_allocator<DataT>>;
 
 using DataT = char;
 
-template <int NumElems> int test_to_copy() {
+template <int NumElems, bool IsAcc> int test_to_copy() {
   sycl::queue queue;
   constexpr int NumSelectedElems = NumElems / 3;
   constexpr int Stride = 2;
@@ -47,21 +47,41 @@ template <int NumElems> int test_to_copy() {
     ref_data_for_fill[i] =
         initial_ref_data[i] + value_for_increase_ref_data_for_fill;
   }
-  queue.submit([&](sycl::handler &cgh) {
-    DataT *init_ref_ptr = initial_ref_data.data();
-    DataT *ref_data_for_fill_ptr = ref_data_for_fill.data();
-    DataT *const out_ptr = result.data();
-
-    cgh.single_task([=]() SYCL_ESIMD_KERNEL {
-      simd<DataT, NumElems> src_simd_obj;
-      src_simd_obj.copy_from(init_ref_ptr);
-      simd<DataT, NumSelectedElems> dst_simd_obj;
-      dst_simd_obj.copy_from(ref_data_for_fill_ptr);
-      src_simd_obj.template select<NumSelectedElems, Stride>(Offset) =
-          dst_simd_obj;
-      src_simd_obj.copy_to(out_ptr);
+  if constexpr (IsAcc) {
+    sycl::buffer<DataT> output_buf(result.data(), result.size());
+    queue.submit([&](sycl::handler &cgh) {
+      DataT *init_ref_ptr = initial_ref_data.data();
+      DataT *ref_data_for_fill_ptr = ref_data_for_fill.data();
+      auto acc =
+          output_buf.template get_access<sycl::access::mode::read_write>(cgh);
+
+      cgh.single_task([=]() SYCL_ESIMD_KERNEL {
+        simd<DataT, NumElems> src_simd_obj;
+        src_simd_obj.copy_from(init_ref_ptr);
+        simd<DataT, NumSelectedElems> dst_simd_obj;
+        dst_simd_obj.copy_from(ref_data_for_fill_ptr);
+        src_simd_obj.template select<NumSelectedElems, Stride>(Offset) =
+            dst_simd_obj;
+        src_simd_obj.copy_to(acc, 0);
+      });
     });
-  });
+  } else {
+    queue.submit([&](sycl::handler &cgh) {
+      DataT *init_ref_ptr = initial_ref_data.data();
+      DataT *ref_data_for_fill_ptr = ref_data_for_fill.data();
+      DataT *const out_ptr = result.data();
+
+      cgh.single_task([=]() SYCL_ESIMD_KERNEL {
+        simd<DataT, NumElems> src_simd_obj;
+        src_simd_obj.copy_from(init_ref_ptr);
+        simd<DataT, NumSelectedElems> dst_simd_obj;
+        dst_simd_obj.copy_from(ref_data_for_fill_ptr);
+        src_simd_obj.template select<NumSelectedElems, Stride>(Offset) =
+            dst_simd_obj;
+        src_simd_obj.copy_to(out_ptr);
+      });
+    });
+  }
   queue.wait_and_throw();
 
   std::vector<size_t> selected_indices;
@@ -114,13 +134,23 @@ template <int NumElems> int test_to_copy() {
 int main() {
   int test_result = 0;
 
-  test_result |= test_to_copy<16>();
-  test_result |= test_to_copy<32 + 8>();
-  test_result |= test_to_copy<32 + 10>();
-  test_result |= test_to_copy<10>();
-  test_result |= test_to_copy<8>();
-  test_result |= test_to_copy<7>();
-  test_result |= test_to_copy<15>();
+  test_result |= test_to_copy<16, false>();
+  test_result |= test_to_copy<32 + 8, false>();
+  test_result |= test_to_copy<32 + 10, false>();
+  test_result |= test_to_copy<10, false>();
+  test_result |= test_to_copy<8, false>();
+  test_result |= test_to_copy<7, false>();
+  test_result |= test_to_copy<15, false>();
+  test_result |= test_to_copy<32, false>();
+
+  test_result |= test_to_copy<16, true>();
+  test_result |= test_to_copy<32 + 8, true>();
+  test_result |= test_to_copy<32 + 10, true>();
+  test_result |= test_to_copy<10, true>();
+  test_result |= test_to_copy<8, true>();
+  test_result |= test_to_copy<7, true>();
+  test_result |= test_to_copy<15, true>();
+  test_result |= test_to_copy<32, true>();
 
   if (!test_result) {
     std::cout << "Test passed" << std::endl;