[SYClomatic-test] Add help function for Block Load header

help_function/help_function.xml

@@ -133,17 +133,17 @@
     <test testName="onedpl_test_tagged_pointer" configFile="config/TEMPLATE_help_function_usm.xml" />
     <test testName="onedpl_test_temporary_allocation" configFile="config/TEMPLATE_help_function_usm.xml" />
     <test testName="onedpl_test_transform_if" configFile="config/TEMPLATE_help_function_skip_cuda_backend.xml" />
-    <!-- <test testName="onedpl_test_transform" configFile="config/TEMPLATE_help_function_skip_cuda_backend.xml" /> -->
+    <test testName="onedpl_test_transform" configFile="config/TEMPLATE_help_function_skip_cuda_backend.xml" />
     <test testName="onedpl_test_transform_output_iterator" configFile="config/TEMPLATE_help_function_skip_cuda_backend.xml" />
     <test testName="onedpl_test_group_sort" configFile="config/TEMPLATE_help_function.xml" />
+    <test testName="onedpl_test_group_load" configFile="config/TEMPLATE_help_function.xml" />
     <test testName="onedpl_test_transform_reduce" configFile="config/TEMPLATE_help_function_skip_double.xml"  splitGroup="double" />
     <!-- <test testName="onedpl_test_translate_key" configFile="config/TEMPLATE_help_function_skip_double.xml"  splitGroup="double" /> -->
     <!-- <test testName="onedpl_test_uninitialized_fill" configFile="config/TEMPLATE_help_function_skip_cuda_backend.xml" /> -->
     <test testName="onedpl_test_unique_by_key_copy" configFile="config/TEMPLATE_help_function_skip_cuda_backend.xml" />
     <test testName="onedpl_test_unique_by_key" configFile="config/TEMPLATE_help_function_skip_cuda_backend.xml" />
     <test testName="onedpl_test_unique" configFile="config/TEMPLATE_help_function_skip_cuda_backend.xml" />
     <test testName="onedpl_test_vector" configFile="config/TEMPLATE_help_function_skip_cuda_backend.xml" />
-    <test testName="onedpl_test_group_exchange" configFile="config/TEMPLATE_help_function.xml" />
     <test testName="util_select_from_sub_group" configFile="config/TEMPLATE_help_function_usm.xml" />
     <test testName="util_shift_sub_group_left" configFile="config/TEMPLATE_help_function_usm.xml" />
     <test testName="util_shift_sub_group_right" configFile="config/TEMPLATE_help_function_usm.xml" />

help_function/src/onedpl_test_group_load.cpp

@@ -0,0 +1,106 @@
+// ====------ onedpl_test_group_load.cpp---------- -*- C++ -* ----===////
+//
+// 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 <sycl/sycl.hpp>
+#include <dpct/dpct.hpp>
+#include <dpct/dpl_utils.hpp>
+#include <iostream>
+
+
+bool helper_function(const int* ptr,const char* func_name){
+  int expected[512];
+  for (int i = 0; i < 512; ++i) {
+      expected[i] = i;
+  }
+  for (int i = 0; i < 512; ++i) {
+    if (ptr[i] != expected[i]) {
+      std::cout << func_name <<" failed\n";
+      std::ostream_iterator<int> Iter(std::cout, ", ");
+      std::copy(ptr, ptr + 512, Iter);
+      std::cout << std::endl;
+      return false;
+    }
+  }
+
+  std::cout << func_name <<" pass\n";
+  return true;
+}
+
+bool test_load_blocked() {
+  sycl::queue q;
+  int data[512];
+  for (int i = 0; i < 512; i++) data[i] = i;
+
+  sycl::buffer<int, 1> buffer(data, 512);
+  q.submit([&](sycl::handler &h) {
+    sycl::accessor dacc(buffer, h, sycl::read_write);
+    h.parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)),
+        [=](sycl::nd_item<3> item) {
+          int thread_data[4];
+          auto *d = dacc.get_multi_ptr<sycl::access::decorated::yes>().get();
+          dpct::group::load_blocked<128>(item.get_local_linear_id(), d, thread_data);
+        });
+  });
+  q.wait_and_throw();
+
+  sycl::host_accessor data_accessor(buffer, sycl::read_only);
+  const int *ptr = data_accessor.get_multi_ptr<sycl::access::decorated::yes>();
+  return helper_function(ptr,"test_load_blocked");
+}
+
+bool test_load_striped() {
+  sycl::queue q;
+  int data[512];
+  for (int i = 0; i < 512; i++) data[i] = i;
+
+  sycl::buffer<int, 1> buffer(data, 512);
+  q.submit([&](sycl::handler &h) {
+    sycl::accessor dacc(buffer, h, sycl::read_write);
+    h.parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)),
+        [=](sycl::nd_item<3> item) {
+          int thread_data[4];
+          auto *d = dacc.get_multi_ptr<sycl::access::decorated::yes>().get();
+          dpct::group::load_striped<128>(item.get_local_linear_id(), d, thread_data);
+        });
+  });
+  q.wait_and_throw();
+
+  sycl::host_accessor data_accessor(buffer, sycl::read_only);
+  const int *ptr = data_accessor.get_multi_ptr<sycl::access::decorated::yes>();
+  return helper_function(ptr,"test_load_blocked");
+}
+
+bool test_load_subgroup_striped() {
+  sycl::queue q;
+  int data[512];
+  for (int i = 0; i < 512; i++) data[i] = i;
+
+  sycl::buffer<int, 1> buffer(data, 512);
+  q.submit([&](sycl::handler &h) {
+    sycl::accessor dacc(buffer, h, sycl::read_write);
+    h.parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)),
+        [=](sycl::nd_item<3> item) {
+          int thread_data[4];
+          auto *d = dacc.get_multi_ptr<sycl::access::decorated::yes>().get();
+          dpct::group::load_subgroup_striped<128, item>(item, item.get_local_linear_id(), d, thread_data);
+        });
+  });
+  q.wait_and_throw();
+
+  sycl::host_accessor data_accessor(buffer, sycl::read_only);
+  const int *ptr = data_accessor.get_multi_ptr<sycl::access::decorated::yes>();
+  return helper_function(ptr,"test_load_blocked");
+}
+
+
+int main() {
+  return !(test_load_blocked() && test_load_striped() && test_load_warp_striped());
+}