[SYCL] Implement SYCL2020 sub_group class (#3765)

Author: Alexander Batashev

sycl/include/CL/sycl.hpp

@@ -52,6 +52,7 @@
 #include <CL/sycl/sampler.hpp>
 #include <CL/sycl/specialization_id.hpp>
 #include <CL/sycl/stream.hpp>
+#include <CL/sycl/sub_group.hpp>
 #include <CL/sycl/types.hpp>
 #include <CL/sycl/usm.hpp>
 #include <CL/sycl/version.hpp>

sycl/include/CL/sycl/ONEAPI/sub_group.hpp

@@ -18,6 +18,7 @@
 #include <CL/sycl/detail/spirv.hpp>
 #include <CL/sycl/detail/type_traits.hpp>
 #include <CL/sycl/id.hpp>
+#include <CL/sycl/memory_enums.hpp>
 #include <CL/sycl/range.hpp>
 #include <CL/sycl/types.hpp>
 
@@ -104,6 +105,8 @@ struct sub_group {
   using range_type = range<1>;
   using linear_id_type = uint32_t;
   static constexpr int dimensions = 1;
+  static constexpr sycl::memory_scope fence_scope =
+      sycl::memory_scope::sub_group;
 
   /* --- common interface members --- */
 
@@ -699,6 +702,33 @@ struct sub_group {
 #endif
   }
 
+  linear_id_type get_group_linear_range() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return static_cast<linear_id_type>(get_group_range()[0]);
+#else
+    throw runtime_error("Sub-groups are not supported on host device.",
+                        PI_INVALID_DEVICE);
+#endif
+  }
+
+  linear_id_type get_local_linear_range() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return static_cast<linear_id_type>(get_local_range()[0]);
+#else
+    throw runtime_error("Sub-groups are not supported on host device.",
+                        PI_INVALID_DEVICE);
+#endif
+  }
+
+  bool leader() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return get_local_linear_id() == 0;
+#else
+    throw runtime_error("Sub-groups are not supported on host device.",
+                        PI_INVALID_DEVICE);
+#endif
+  }
+
 protected:
   template <int dimensions> friend class cl::sycl::nd_item;
   friend sub_group this_sub_group();

sycl/include/CL/sycl/group.hpp

@@ -17,6 +17,7 @@
 #include <CL/sycl/device_event.hpp>
 #include <CL/sycl/h_item.hpp>
 #include <CL/sycl/id.hpp>
+#include <CL/sycl/memory_enums.hpp>
 #include <CL/sycl/pointers.hpp>
 #include <CL/sycl/range.hpp>
 #include <stdexcept>
@@ -431,5 +432,8 @@ template <int Dims> group<Dims> this_group() {
 #endif
 }
 
+template <typename Group>
+void group_barrier(Group G, memory_scope FenceScope = Group::fence_scope);
+
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/memory_enums.hpp

@@ -0,0 +1,25 @@
+//==-------------- memory_enums.hpp --- SYCL enums -------------------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/ONEAPI/atomic_enums.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+using ONEAPI::memory_scope;
+
+#if __cplusplus >= 201703L
+inline constexpr auto memory_scope_work_item = memory_scope::work_item;
+inline constexpr auto memory_scope_sub_group = memory_scope::sub_group;
+inline constexpr auto memory_scope_work_group = memory_scope::work_group;
+inline constexpr auto memory_scope_device = memory_scope::device;
+inline constexpr auto memory_scope_system = memory_scope::system;
+#endif
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/nd_item.hpp

@@ -9,7 +9,6 @@
 #pragma once
 
 #include <CL/__spirv/spirv_ops.hpp>
-#include <CL/sycl/ONEAPI/sub_group.hpp>
 #include <CL/sycl/access/access.hpp>
 #include <CL/sycl/detail/defines.hpp>
 #include <CL/sycl/detail/helpers.hpp>
@@ -18,6 +17,7 @@
 #include <CL/sycl/item.hpp>
 #include <CL/sycl/nd_range.hpp>
 #include <CL/sycl/range.hpp>
+#include <CL/sycl/sub_group.hpp>
 
 #include <cstddef>
 #include <stdexcept>
@@ -67,7 +67,7 @@ template <int dimensions = 1> class nd_item {
 
   group<dimensions> get_group() const { return Group; }
 
-  ONEAPI::sub_group get_sub_group() const { return ONEAPI::sub_group(); }
+  sub_group get_sub_group() const { return sub_group(); }
 
   size_t __SYCL_ALWAYS_INLINE get_group(int dimension) const {
     size_t Size = Group[dimension];

sycl/include/CL/sycl/sub_group.hpp

@@ -0,0 +1,37 @@
+//==----------- sub_group.hpp --- SYCL sub-group ---------------------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/ONEAPI/sub_group.hpp>
+#include <CL/sycl/group.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+using ONEAPI::sub_group;
+// TODO move the entire sub_group class implementation to this file once
+// breaking changes are allowed.
+
+template <>
+inline void group_barrier<sub_group>(sub_group Group, memory_scope FenceScope) {
+  (void)Group;
+  (void)FenceScope;
+#ifdef __SYCL_DEVICE_ONLY__
+  __spirv_ControlBarrier(__spv::Scope::Subgroup,
+                         detail::spirv::getScope(FenceScope),
+                         __spv::MemorySemanticsMask::AcquireRelease |
+                             __spv::MemorySemanticsMask::SubgroupMemory |
+                             __spv::MemorySemanticsMask::WorkgroupMemory |
+                             __spv::MemorySemanticsMask::CrossWorkgroupMemory);
+#else
+  throw runtime_error("Sub-groups are not supported on host device.",
+                      PI_INVALID_DEVICE);
+#endif
+}
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)