Address some feedback, implement barrier

Author: Alexander Batashev

sycl/include/CL/__spirv/spirv_ops.hpp

@@ -567,9 +567,6 @@ extern SYCL_EXTERNAL void
 __spirv_ocl_prefetch(const __attribute__((opencl_global)) char *Ptr,
                      size_t NumBytes) noexcept;
 
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT bool
-    __spirv_GroupNonUniformElect(__spv::Scope::Flag) noexcept;
-
 #else // if !__SYCL_DEVICE_ONLY__
 
 template <typename dataT>

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

@@ -20,6 +20,7 @@
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/range.hpp>
 #include <CL/sycl/types.hpp>
+#include <CL/sycl/enums.hpp>
 
 #include <type_traits>
 
@@ -104,7 +105,7 @@ struct sub_group {
   using range_type = range<1>;
   using linear_id_type = uint32_t;
   static constexpr int dimensions = 1;
-  static constexpr memory_scope fence_scope = memory_scope::sub_group;
+  static constexpr sycl::memory_scope fence_scope = sycl::memory_scope::sub_group;
 
   /* --- common interface members --- */
 
@@ -720,7 +721,7 @@ struct sub_group {
 
   bool leader() const {
 #ifdef __SYCL_DEVICE_ONLY__
-    return detail::spirv::GroupNonUniformElect<sub_group>();
+    return get_local_linear_id() == 0;
 #else
     throw runtime_error("Sub-groups are not supported on host device.",
                         PI_INVALID_DEVICE);

sycl/include/CL/sycl/detail/spirv.hpp

@@ -15,6 +15,7 @@
 #include <CL/sycl/detail/helpers.hpp>
 #include <CL/sycl/detail/type_traits.hpp>
 #include <CL/sycl/id.hpp>
+#include <CL/sycl/enums.hpp>
 
 #ifdef __SYCL_DEVICE_ONLY__
 __SYCL_INLINE_NAMESPACE(cl) {
@@ -259,6 +260,22 @@ getScope(ONEAPI::memory_scope Scope) {
   }
 }
 
+constexpr __spv::Scope::Flag
+getScope(memory_scope Scope) {
+  switch (Scope) {
+  case memory_scope::work_item:
+    return __spv::Scope::Invocation;
+  case memory_scope::sub_group:
+    return __spv::Scope::Subgroup;
+  case memory_scope::work_group:
+    return __spv::Scope::Workgroup;
+  case memory_scope::device:
+    return __spv::Scope::Device;
+  case memory_scope::system:
+    return __spv::Scope::CrossDevice;
+  }
+}
+
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_integral<T>::value, T>
 AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr,
@@ -735,10 +752,6 @@ EnableIfGenericShuffle<T> SubgroupShuffleUp(T x, id<1> local_id) {
   return Result;
 }
 
-template <typename Group> bool GroupNonUniformElect() {
-  return __spirv_GroupNonUniformElect(group_scope<Group>::value);
-}
-
 } // namespace spirv
 } // namespace detail
 } // namespace sycl

sycl/include/CL/sycl/enums.hpp

@@ -0,0 +1,30 @@
+//==-------------- 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
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+enum class memory_scope : unsigned char {
+      work_item = 0,
+      sub_group = 1,
+      work_group = 2,
+      device = 3,
+      system = 4
+};
+
+#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
+}
+}
+

sycl/include/CL/sycl/group.hpp

@@ -19,6 +19,7 @@
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/pointers.hpp>
 #include <CL/sycl/range.hpp>
+#include <CL/sycl/enums.hpp>
 #include <stdexcept>
 #include <type_traits>
 
@@ -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/sub_group.hpp

@@ -8,12 +8,29 @@
 
 #pragma once
 
+#include <CL/sycl/group.hpp>
 #include <CL/sycl/ONEAPI/sub_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)