[SYCL] Add tests for SYCL2020 sub_groups features (#283)

Author: alexbatashev

SYCL/SubGroup/barrier.cpp

@@ -16,9 +16,11 @@
 #include <CL/sycl.hpp>
 #include <limits>
 #include <numeric>
-template <typename T> class sycl_subgr;
+
+template <typename T, bool UseNewSyntax> class sycl_subgr;
 using namespace cl::sycl;
-template <typename T> void check(queue &Queue, size_t G = 240, size_t L = 60) {
+template <typename T, bool UseNewSyntax = false>
+void check(queue &Queue, size_t G = 240, size_t L = 60) {
   try {
     nd_range<1> NdRange(G, L);
     std::vector<T> data(G);
@@ -29,21 +31,26 @@ template <typename T> void check(queue &Queue, size_t G = 240, size_t L = 60) {
       auto addacc = addbuf.template get_access<access::mode::read_write>(cgh);
       auto sgsizeacc = sgsizebuf.get_access<access::mode::read_write>(cgh);
 
-      cgh.parallel_for<sycl_subgr<T>>(NdRange, [=](nd_item<1> NdItem) {
-        ONEAPI::sub_group SG = NdItem.get_sub_group();
-        size_t lid = SG.get_local_id().get(0);
-        size_t gid = NdItem.get_global_id(0);
-        size_t SGoff = gid - lid;
+      cgh.parallel_for<sycl_subgr<T, UseNewSyntax>>(
+          NdRange, [=](nd_item<1> NdItem) {
+            ONEAPI::sub_group SG = NdItem.get_sub_group();
+            size_t lid = SG.get_local_id().get(0);
+            size_t gid = NdItem.get_global_id(0);
+            size_t SGoff = gid - lid;
 
-        T res = 0;
-        for (size_t i = 0; i <= lid; i++) {
-          res += addacc[SGoff + i];
-        }
-        SG.barrier(access::fence_space::global_space);
-        addacc[gid] = res;
-        if (NdItem.get_global_id(0) == 0)
-          sgsizeacc[0] = SG.get_max_local_range()[0];
-      });
+            T res = 0;
+            for (size_t i = 0; i <= lid; i++) {
+              res += addacc[SGoff + i];
+            }
+            if constexpr (UseNewSyntax) {
+              group_barrier(SG);
+            } else {
+              SG.barrier(access::fence_space::global_space);
+            }
+            addacc[gid] = res;
+            if (NdItem.get_global_id(0) == 0)
+              sgsizeacc[0] = SG.get_max_local_range()[0];
+          });
     });
     auto addacc = addbuf.template get_access<access::mode::read_write>();
     auto sgsizeacc = sgsizebuf.get_access<access::mode::read_write>();
@@ -79,8 +86,14 @@ int main() {
   check<long>(Queue);
   check<unsigned long>(Queue);
   check<float>(Queue);
+  check<int, true>(Queue);
+  check<unsigned int, true>(Queue);
+  check<long, true>(Queue);
+  check<unsigned long, true>(Queue);
+  check<float, true>(Queue);
   if (Queue.get_device().has_extension("cl_khr_fp64")) {
     check<double>(Queue);
+    check<double, true>(Queue);
   }
   std::cout << "Test passed." << std::endl;
   return 0;

SYCL/SubGroup/sub_groups_sycl2020.cpp

@@ -0,0 +1,57 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+//
+// TODO enable test on CUDA once kernel_bundle is supported
+// UNSUPPORTED: cuda
+
+#include <sycl/sycl.hpp>
+
+class TestKernel;
+
+int main() {
+  sycl::queue Q;
+  sycl::buffer<int, 3> Buf{sycl::range{3, 32, 32}};
+
+  sycl::kernel_id TestKernelID = sycl::get_kernel_id<TestKernel>();
+
+  sycl::kernel_bundle KernelBundle =
+      sycl::get_kernel_bundle<sycl::bundle_state::executable>(Q.get_context(),
+                                                              {TestKernelID});
+
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.use_kernel_bundle(KernelBundle);
+    sycl::accessor Acc{Buf, CGH, sycl::write_only};
+    CGH.parallel_for<TestKernel>(
+        sycl::nd_range<1>(sycl::range{32}, sycl::range{32}),
+        [=](sycl::nd_item<1> item) {
+          auto SG = item.get_sub_group();
+          Acc[0][SG.get_group_linear_id()][SG.get_local_linear_id()] =
+              SG.leader();
+          Acc[1][SG.get_group_linear_id()][SG.get_local_linear_id()] =
+              SG.get_group_linear_range();
+          Acc[2][SG.get_group_linear_id()][SG.get_local_linear_id()] =
+              SG.get_local_linear_range();
+        });
+  });
+
+  sycl::host_accessor Acc{Buf, sycl::read_only};
+
+  sycl::kernel Kernel = KernelBundle.get_kernel(TestKernelID);
+
+  const size_t SubgroupSize =
+      Kernel.get_info<sycl::info::kernel_device_specific::max_sub_group_size>(
+          Q.get_device(), sycl::range{32, 1, 1});
+  const size_t MaxNumSubgroups = 32 / SubgroupSize;
+
+  for (size_t SGNo = 0; SGNo < MaxNumSubgroups; SGNo++) {
+    for (size_t WINo = 0; WINo < SubgroupSize; WINo++) {
+      const int Leader = WINo == 0 ? 1 : 0;
+      assert(Acc[0][SGNo][WINo] == Leader);
+      assert(Acc[1][SGNo][WINo] == MaxNumSubgroups);
+      assert(Acc[2][SGNo][WINo] == SubgroupSize);
+    }
+  }
+
+  return 0;
+}