[SYCL] Enable async_work_group_copy for scalar and vector bool types

sycl/include/CL/sycl/detail/type_traits.hpp

@@ -207,6 +207,19 @@ template <typename T>
 struct is_vector_arithmetic
     : bool_constant<is_vec<T>::value && is_arithmetic<T>::value> {};
 
+// is_bool
+template <typename T>
+struct is_scalar_bool
+    : bool_constant<std::is_same<remove_cv_t<T>, bool>::value> {};
+
+template <typename T>
+struct is_vector_bool
+    : bool_constant<is_vec<T>::value &&
+                    is_scalar_bool<vector_element_t<T>>::value> {};
+
+template <typename T>
+struct is_bool : bool_constant<is_scalar_bool<vector_element_t<T>>::value> {};
+
 // is_pointer
 template <typename T> struct is_pointer_impl : std::false_type {};
 

sycl/include/CL/sycl/group.hpp

@@ -274,58 +274,99 @@ template <int Dimensions = 1> class group {
     __spirv_MemoryBarrier(__spv::Scope::Workgroup, flags);
   }
 
+  /// Asynchronously copies a number of elements specified by \p numElements
+  /// from the source pointed by \p src to destination pointed by \p dest
+  /// with a source stride specified by \p srcStride, and returns a SYCL
+  /// device_event which can be used to wait on the completion of the copy.
+  /// Permitted types for dataT are all scalar and vector types, except boolean.
   template <typename dataT>
-  device_event async_work_group_copy(local_ptr<dataT> dest,
-                                     global_ptr<dataT> src,
-                                     size_t numElements) const {
+  detail::enable_if_t<!detail::is_bool<dataT>::value, device_event>
+  async_work_group_copy(local_ptr<dataT> dest, global_ptr<dataT> src,
+                        size_t numElements, size_t srcStride) const {
     using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
     using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
 
-    __ocl_event_t e = OpGroupAsyncCopyGlobalToLocal(
+    __ocl_event_t E = OpGroupAsyncCopyGlobalToLocal(
         __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
-        numElements, 1, 0);
-    return device_event(&e);
+        numElements, srcStride, 0);
+    return device_event(&E);
   }
 
+  /// Asynchronously copies a number of elements specified by \p numElements
+  /// from the source pointed by \p src to destination pointed by \p dest with
+  /// the destination stride specified by \p destStride, and returns a SYCL
+  /// device_event which can be used to wait on the completion of the copy.
+  /// Permitted types for dataT are all scalar and vector types, except boolean.
   template <typename dataT>
-  device_event async_work_group_copy(global_ptr<dataT> dest,
-                                     local_ptr<dataT> src,
-                                     size_t numElements) const {
+  detail::enable_if_t<!detail::is_bool<dataT>::value, device_event>
+  async_work_group_copy(global_ptr<dataT> dest, local_ptr<dataT> src,
+                        size_t numElements, size_t destStride) const {
     using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
     using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
 
-    __ocl_event_t e = OpGroupAsyncCopyLocalToGlobal(
+    __ocl_event_t E = OpGroupAsyncCopyLocalToGlobal(
         __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
-        numElements, 1, 0);
-    return device_event(&e);
+        numElements, destStride, 0);
+    return device_event(&E);
+  }
+
+  /// Specialization for scalar bool type.
+  /// Asynchronously copies a number of elements specified by \p NumElements
+  /// from the source pointed by \p Src to destination pointed by \p Dest
+  /// with a stride specified by \p Stride, and returns a SYCL device_event
+  /// which can be used to wait on the completion of the copy.
+  template <typename T, access::address_space DestS, access::address_space SrcS>
+  detail::enable_if_t<detail::is_scalar_bool<T>::value, device_event>
+  async_work_group_copy(multi_ptr<T, DestS> Dest, multi_ptr<T, SrcS> Src,
+                        size_t NumElements, size_t Stride) const {
+    static_assert(sizeof(bool) == sizeof(uint8_t),
+                  "Async copy to/from bool memory is not supported.");
+    auto DestP =
+        multi_ptr<uint8_t, DestS>(reinterpret_cast<uint8_t *>(Dest.get()));
+    auto SrcP =
+        multi_ptr<uint8_t, SrcS>(reinterpret_cast<uint8_t *>(Src.get()));
+    return async_work_group_copy(DestP, SrcP, NumElements, Stride);
+  }
+
+  /// Specialization for vector bool type.
+  /// Asynchronously copies a number of elements specified by \p NumElements
+  /// from the source pointed by \p Src to destination pointed by \p Dest
+  /// with a stride specified by \p Stride, and returns a SYCL device_event
+  /// which can be used to wait on the completion of the copy.
+  template <typename T, access::address_space DestS, access::address_space SrcS>
+  detail::enable_if_t<detail::is_vector_bool<T>::value, device_event>
+  async_work_group_copy(multi_ptr<T, DestS> Dest, multi_ptr<T, SrcS> Src,
+                        size_t NumElements, size_t Stride) const {
+    static_assert(sizeof(bool) == sizeof(uint8_t),
+                  "Async copy to/from bool memory is not supported.");
+    using VecT = detail::change_base_type_t<T, uint8_t>;
+    auto DestP = multi_ptr<VecT, DestS>(reinterpret_cast<VecT *>(Dest.get()));
+    auto SrcP = multi_ptr<VecT, SrcS>(reinterpret_cast<VecT *>(Src.get()));
+    return async_work_group_copy(DestP, SrcP, NumElements, Stride);
   }
 
+  /// Asynchronously copies a number of elements specified by \p numElements
+  /// from the source pointed by \p src to destination pointed by \p dest and
+  /// returns a SYCL device_event which can be used to wait on the completion
+  /// of the copy.
+  /// Permitted types for dataT are all scalar and vector types.
   template <typename dataT>
   device_event async_work_group_copy(local_ptr<dataT> dest,
                                      global_ptr<dataT> src,
-                                     size_t numElements,
-                                     size_t srcStride) const {
-    using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
-    using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
-
-    __ocl_event_t e = OpGroupAsyncCopyGlobalToLocal(
-        __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
-        numElements, srcStride, 0);
-    return device_event(&e);
+                                     size_t numElements) const {
+    return async_work_group_copy(dest, src, numElements, 1);
   }
 
+  /// Asynchronously copies a number of elements specified by \p numElements
+  /// from the source pointed by \p src to destination pointed by \p dest and
+  /// returns a SYCL device_event which can be used to wait on the completion
+  /// of the copy.
+  /// Permitted types for dataT are all scalar and vector types.
   template <typename dataT>
   device_event async_work_group_copy(global_ptr<dataT> dest,
                                      local_ptr<dataT> src,
-                                     size_t numElements,
-                                     size_t destStride) const {
-    using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
-    using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
-
-    __ocl_event_t e = OpGroupAsyncCopyLocalToGlobal(
-        __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
-        numElements, destStride, 0);
-    return device_event(&e);
+                                     size_t numElements) const {
+    return async_work_group_copy(dest, src, numElements, 1);
   }
 
   template <typename... eventTN>

sycl/test/basic_tests/group_async_copy.cpp

@@ -0,0 +1,160 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.run
+// RUN: %GPU_RUN_PLACEHOLDER %t.run
+// RUN: %CPU_RUN_PLACEHOLDER %t.run
+// RUN: %ACC_RUN_PLACEHOLDER %t.run
+// RUN: env SYCL_DEVICE_FILTER=host %t.run
+
+#include <CL/sycl.hpp>
+#include <iostream>
+#include <typeinfo>
+
+using namespace cl::sycl;
+
+template <typename T> class KernelName;
+
+// Define the number of work items to enqueue.
+const size_t NElems = 32;
+const size_t WorkGroupSize = 8;
+const size_t NWorkGroups = NElems / WorkGroupSize;
+
+template <typename T> void initInputBuffer(buffer<T, 1> &Buf, size_t Stride) {
+  auto Acc = Buf.template get_access<access::mode::write>();
+  for (size_t I = 0; I < Buf.get_count(); I += WorkGroupSize) {
+    for (size_t J = 0; J < WorkGroupSize; J++)
+      Acc[I + J] = I + J + ((J % Stride == 0) ? 100 : 0);
+  }
+}
+
+template <typename T> void initOutputBuffer(buffer<T, 1> &Buf) {
+  auto Acc = Buf.template get_access<access::mode::write>();
+  for (size_t I = 0; I < Buf.get_count(); I++)
+    Acc[I] = 0;
+}
+
+template <typename T> struct is_vec : std::false_type {};
+template <typename T, size_t N> struct is_vec<vec<T, N>> : std::true_type {};
+
+template <typename T> bool checkEqual(vec<T, 1> A, size_t B) {
+  T TB = B;
+  return A.s0() == TB;
+}
+
+template <typename T> bool checkEqual(vec<T, 4> A, size_t B) {
+  T TB = B;
+  return A.x() == TB && A.y() == TB && A.z() == TB && A.w() == TB;
+}
+
+template <typename T>
+typename std::enable_if<!is_vec<T>::value, bool>::type checkEqual(T A,
+                                                                  size_t B) {
+  T TB = B;
+  return A == TB;
+}
+
+template <typename T> std::string toString(vec<T, 1> A) {
+  std::string R("(");
+  return R + std::to_string(A.s0()) + ")";
+}
+
+template <typename T> std::string toString(vec<T, 4> A) {
+  std::string R("(");
+  R += std::to_string(A.x()) + "," + std::to_string(A.y()) + "," +
+       std::to_string(A.z()) + "," + std::to_string(A.w()) + ")";
+  return R;
+}
+
+template <typename T = void>
+typename std::enable_if<!is_vec<T>::value, std::string>::type toString(T A) {
+  return std::to_string(A);
+}
+
+template <typename T> int checkResults(buffer<T, 1> &OutBuf, size_t Stride) {
+  auto Out = OutBuf.template get_access<access::mode::read>();
+  int EarlyFailout = 20;
+
+  for (size_t I = 0; I < OutBuf.get_count(); I += WorkGroupSize) {
+    for (size_t J = 0; J < WorkGroupSize; J++) {
+      size_t ExpectedVal = (J % Stride == 0) ? (100 + I + J) : 0;
+      if (!checkEqual(Out[I + J], ExpectedVal)) {
+        std::cerr << std::string(typeid(T).name()) + ": Stride=" << Stride
+                  << " : Incorrect value at index " << I + J
+                  << " : Expected: " << toString(ExpectedVal)
+                  << ", Computed: " << toString(Out[I + J]) << "\n";
+        if (--EarlyFailout == 0)
+          return 1;
+      }
+    }
+  }
+  return EarlyFailout - 20;
+}
+
+template <typename T> int test(size_t Stride) {
+  queue Q;
+
+  buffer<T, 1> InBuf(NElems);
+  buffer<T, 1> OutBuf(NElems);
+
+  initInputBuffer(InBuf, Stride);
+  initOutputBuffer(OutBuf);
+
+  Q.submit([&](handler &CGH) {
+     auto In = InBuf.template get_access<access::mode::read>(CGH);
+     auto Out = OutBuf.template get_access<access::mode::write>(CGH);
+     accessor<T, 1, access::mode::read_write, access::target::local> Local(
+         range<1>{WorkGroupSize}, CGH);
+
+     nd_range<1> NDR{range<1>(NElems), range<1>(WorkGroupSize)};
+     CGH.parallel_for<KernelName<T>>(NDR, [=](nd_item<1> NDId) {
+       auto GrId = NDId.get_group_linear_id();
+       auto Group = NDId.get_group();
+       size_t NElemsToCopy =
+           WorkGroupSize / Stride + ((WorkGroupSize % Stride) ? 1 : 0);
+       size_t Offset = GrId * WorkGroupSize;
+       if (Stride == 1) { // Check the version without stride arg.
+         auto E = NDId.async_work_group_copy(
+             Local.get_pointer(), In.get_pointer() + Offset, NElemsToCopy);
+         E.wait();
+       } else {
+         auto E = NDId.async_work_group_copy(Local.get_pointer(),
+                                             In.get_pointer() + Offset,
+                                             NElemsToCopy, Stride);
+         E.wait();
+       }
+
+       if (Stride == 1) { // Check the version without stride arg.
+         auto E = Group.async_work_group_copy(
+             Out.get_pointer() + Offset, Local.get_pointer(), NElemsToCopy);
+         Group.wait_for(E);
+       } else {
+         auto E = Group.async_work_group_copy(Out.get_pointer() + Offset,
+                                              Local.get_pointer(), NElemsToCopy,
+                                              Stride);
+         Group.wait_for(E);
+       }
+     });
+   }).wait();
+
+  return checkResults(OutBuf, Stride);
+}
+
+int main() {
+  for (int Stride = 1; Stride < WorkGroupSize; Stride++) {
+    if (test<int>(Stride))
+      return 1;
+    if (test<vec<int, 1>>(Stride))
+      return 1;
+    if (test<int4>(Stride))
+      return 1;
+    if (test<bool>(Stride))
+      return 1;
+    if (test<vec<bool, 1>>(Stride))
+      return 1;
+    if (test<vec<bool, 4>>(Stride))
+      return 1;
+    if (test<cl::sycl::cl_bool>(Stride))
+      return 1;
+  }
+
+  std::cout << "Test passed.\n";
+  return 0;
+}