intel · bader · Aug 5, 2020 · Aug 5, 2020
@@ -51,9 +51,7 @@ John Pennycook, Intel (john 'dot' pennycook 'at' intel 'dot' com)
 
 == Dependencies
 
-This extension is written against the SYCL 1.2.1 specification, Revision v1.2.1-6 and the following extensions:
-
-- +SYCL_INTEL_extended_atomics+
+This extension is written against the SYCL 1.2.1 specification, Revision v1.2.1-6.
 
 == Overview
 
@@ -69,10 +67,6 @@ The extension introduces the following functions:
 - +reduce+
 - +exclusive_scan+
 - +inclusive_scan+
-- +barrier+
-
-The definitions and behavior of the following functions are based on equivalents in the SYCL 2020 provisional specification:
-- +barrier+
 
 === Alignment with OpenCL vs C++
 
@@ -258,21 +252,6 @@ The return types of the collective functions in {cpp}17 are not deduced from the
 |Perform an inclusive scan over the values in the range [_first_, _last_) using the operator _binary_op_, which must be one of the group algorithms library function objects.  The value written to +result + i+ is the inclusive scan of the first +i+ values in the range and an initial value specified by _init_. Returns a pointer to the end of the output range. _first_, _last_, _result_, _binary_op_ and _init_ must be the same for all work-items in the group. _binary_op(init, *first)_ must return a value of type _T_.
 |===
 
-==== Synchronization
-
-The behavior of memory fences in this section is aligned with the single happens-before relationship defined by the +SYCL_INTEL_extended_atomics+ extension.
-
-|===
-|Function|Description
-
-|+template <typename Group> void barrier(Group g);+
-|Synchronize all work-items in the group, and ensure that all memory accesses to any address space prior to the barrier are visible to all work-items in the group after the barrier.  The scope of the group memory fences implied by this barrier is the narrowest scope including all work-items in the group.
-
-|+template <typename Group> void barrier(Group g, memory_scope scope);+
-|Synchronize all work-items in the group, and ensure that all memory accesses to any address space prior to the barrier are visible to all work-items specified by _scope_ after the barrier.  The scope of the group memory fences implied by this barrier is controlled by _scope_ and must be broader than the narrowest scope including all work-items in the group.  If the specified _scope_ is narrower than the narrowest scope including all work-items in the group, the _scope_ argument is ignored.
-
-|===
-
 == Issues
 
 None.
@@ -291,7 +270,6 @@ None.
 |========================================
 |Rev|Date|Author|Changes
 |1|2020-01-30|John Pennycook|*Initial public working draft*
-|2|2020-07-28|John Pennycook|*Add group barrier*
 |========================================
 
 //************************************************************************

@@ -70,7 +70,6 @@ This extension adds sub-group support to all of the functions from +SYCL_INTEL_g
 - +reduce+
 - +exclusive_scan+
 - +inclusive_scan+
-- +barrier+
 
 It additionally introduces a number of functions that are currently specific to sub-groups:
 
@@ -166,7 +165,6 @@ None.
 |========================================
 |Rev|Date|Author|Changes
 |1|2020-03-16|John Pennycook|*Initial public working draft*
-|2|2020-07-28|John Pennycook|*Add group barrier*
 |========================================
 
 //************************************************************************

@@ -13,7 +13,6 @@
 #include <CL/sycl/detail/spirv.hpp>
 #include <CL/sycl/detail/type_traits.hpp>
 #include <CL/sycl/group.hpp>
-#include <CL/sycl/intel/atomic.hpp>
 #include <CL/sycl/intel/functional.hpp>
 #include <CL/sycl/intel/sub_group.hpp>
 
@@ -78,15 +77,6 @@ template <> inline id<3> linear_id_to_id(range<3> r, size_t linear_id) {
   return result;
 }
 
-// TODO: Replace with Group::fence_scope from SYCL 2020 provisional
-template <typename Group> struct FenceScope {
-  static constexpr intel::memory_scope value = intel::memory_scope::work_group;
-};
-
-template <> struct FenceScope<intel::sub_group> {
-  static constexpr intel::memory_scope value = intel::memory_scope::sub_group;
-};
-
 template <typename T, class BinaryOperation> struct identity {};
 
 template <typename T, typename V> struct identity<T, intel::plus<V>> {
@@ -906,42 +896,6 @@ template <typename Group> bool leader(Group g) {
 #endif
 }
 
-template <typename Group> void barrier(Group, memory_scope scope) {
-  static_assert(sycl::detail::is_generic_group<Group>::value,
-                "Group algorithms only support the sycl::group and "
-                "intel::sub_group class.");
-#ifdef __SYCL_DEVICE_ONLY__
-  // MemoryScope must be broader than Group scope for correctness
-  auto GroupScope = detail::FenceScope<Group>::value;
-  auto BroadestScope = (scope > GroupScope) ? scope : GroupScope;
-  auto MemoryScope = sycl::detail::spirv::getScope(BroadestScope);
-  auto ExecutionScope = sycl::detail::spirv::group_scope<Group>::value;
-  __spirv_ControlBarrier(ExecutionScope, MemoryScope,
-                         __spv::MemorySemanticsMask::AcquireRelease |
-                             __spv::MemorySemanticsMask::SubgroupMemory |
-                             __spv::MemorySemanticsMask::WorkgroupMemory |
-                             __spv::MemorySemanticsMask::CrossWorkgroupMemory);
-#else
-  (void)scope;
-  throw runtime_error("Group algorithms are not supported on host device.",
-                      PI_INVALID_DEVICE);
-#endif
-}
-
-template <typename Group> void barrier(Group g) {
-  static_assert(sycl::detail::is_generic_group<Group>::value,
-                "Group algorithms only support the sycl::group and "
-                "intel::sub_group class.");
-#ifdef __SYCL_DEVICE_ONLY__
-  auto MemoryScope = detail::FenceScope<Group>::value;
-  barrier(g, MemoryScope);
-#else
-  (void)g;
-  throw runtime_error("Group algorithms are not supported on host device.",
-                      PI_INVALID_DEVICE);
-#endif
-}
-
 } // namespace intel
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)