[SYCL] SYCL 2020 atomic_fence (#4053)

Link to the spec: https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:barriers-fences

Author: Alexander Batashev

sycl/include/CL/sycl.hpp

@@ -19,6 +19,7 @@
 #include <CL/sycl/accessor.hpp>
 #include <CL/sycl/aspects.hpp>
 #include <CL/sycl/atomic.hpp>
+#include <CL/sycl/atomic_fence.hpp>
 #include <CL/sycl/backend.hpp>
 #include <CL/sycl/buffer.hpp>
 #include <CL/sycl/builtins.hpp>

sycl/include/CL/sycl/ONEAPI/atomic_fence.hpp

@@ -23,6 +23,7 @@ namespace detail {
 using namespace cl::sycl::detail;
 }
 
+__SYCL2020_DEPRECATED("use sycl::atomic_fence instead")
 static inline void atomic_fence(memory_order order, memory_scope scope) {
 #ifdef __SYCL_DEVICE_ONLY__
   auto SPIRVOrder = detail::spirv::getMemorySemanticsMask(order);

sycl/include/CL/sycl/atomic.hpp

@@ -11,6 +11,7 @@
 #include <CL/__spirv/spirv_ops.hpp>
 #include <CL/sycl/access/access.hpp>
 #include <CL/sycl/detail/helpers.hpp>
+#include <CL/sycl/memory_enums.hpp>
 
 #ifndef __SYCL_DEVICE_ONLY__
 #include <atomic>
@@ -26,8 +27,6 @@
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 
-enum class memory_order : int { relaxed = 0 };
-
 // Forward declaration
 template <typename pointerT, access::address_space addressSpace>
 class multi_ptr;
@@ -82,9 +81,6 @@ static inline std::memory_order
 getStdMemoryOrder(__spv::MemorySemanticsMask::Flag) {
   return std::memory_order_relaxed;
 }
-static inline std::memory_order getStdMemoryOrder(::cl::sycl::memory_order) {
-  return std::memory_order_relaxed;
-}
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/atomic_fence.hpp

@@ -0,0 +1,35 @@
+//==--------- atomic_fence.hpp - SYCL 2020 atomic_fence --------------------==//
+//
+// 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/__spirv/spirv_ops.hpp>
+#include <CL/sycl/detail/spirv.hpp>
+#include <CL/sycl/memory_enums.hpp>
+
+#ifndef __SYCL_DEVICE_ONLY__
+#include <atomic>
+#endif
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+
+static inline void atomic_fence(memory_order order, memory_scope scope) {
+#ifdef __SYCL_DEVICE_ONLY__
+  auto SPIRVOrder = detail::spirv::getMemorySemanticsMask(order);
+  auto SPIRVScope = detail::spirv::getScope(scope);
+  __spirv_MemoryBarrier(SPIRVScope, static_cast<uint32_t>(SPIRVOrder));
+#else
+  (void)scope;
+  auto StdOrder = detail::getStdMemoryOrder(order);
+  atomic_thread_fence(StdOrder);
+#endif
+}
+
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

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/memory_enums.hpp>
 
 #ifdef __SYCL_DEVICE_ONLY__
 __SYCL_INLINE_NAMESPACE(cl) {
@@ -216,24 +217,28 @@ EnableIfGenericBroadcast<T> GroupBroadcast(T x, id<Dimensions> local_id) {
 
 // Single happens-before means semantics should always apply to all spaces
 // Although consume is unsupported, forwarding to acquire is valid
-static inline constexpr __spv::MemorySemanticsMask::Flag
-getMemorySemanticsMask(ONEAPI::memory_order Order) {
+template <typename T>
+static inline constexpr typename std::enable_if<
+    std::is_same<T, sycl::ONEAPI::memory_order>::value ||
+        std::is_same<T, sycl::memory_order>::value,
+    __spv::MemorySemanticsMask::Flag>::type
+getMemorySemanticsMask(T Order) {
   __spv::MemorySemanticsMask::Flag SpvOrder = __spv::MemorySemanticsMask::None;
   switch (Order) {
-  case ONEAPI::memory_order::relaxed:
+  case T::relaxed:
     SpvOrder = __spv::MemorySemanticsMask::None;
     break;
-  case ONEAPI::memory_order::__consume_unsupported:
-  case ONEAPI::memory_order::acquire:
+  case T::__consume_unsupported:
+  case T::acquire:
     SpvOrder = __spv::MemorySemanticsMask::Acquire;
     break;
-  case ONEAPI::memory_order::release:
+  case T::release:
     SpvOrder = __spv::MemorySemanticsMask::Release;
     break;
-  case ONEAPI::memory_order::acq_rel:
+  case T::acq_rel:
     SpvOrder = __spv::MemorySemanticsMask::AcquireRelease;
     break;
-  case ONEAPI::memory_order::seq_cst:
+  case T::seq_cst:
     SpvOrder = __spv::MemorySemanticsMask::SequentiallyConsistent;
     break;
   }

sycl/include/CL/sycl/memory_enums.hpp

@@ -14,12 +14,51 @@ __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 using ONEAPI::memory_scope;
 
+enum class memory_order : int {
+  relaxed = 0,
+  acquire = 1,
+  __consume_unsupported =
+      2, // helps optimizer when mapping to std::memory_order
+  release = 3,
+  acq_rel = 4,
+  seq_cst = 5
+};
+
 #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;
+
+inline constexpr auto memory_order_relaxed = memory_order::relaxed;
+inline constexpr auto memory_order_acquire = memory_order::acquire;
+inline constexpr auto memory_order_release = memory_order::release;
+inline constexpr auto memory_order_acq_rel = memory_order::acq_rel;
+inline constexpr auto memory_order_seq_cst = memory_order::seq_cst;
 #endif
+
+#ifndef __SYCL_DEVICE_ONLY__
+namespace detail {
+
+static constexpr std::memory_order getStdMemoryOrder(sycl::memory_order order) {
+  switch (order) {
+  case memory_order::relaxed:
+    return std::memory_order_relaxed;
+  case memory_order::__consume_unsupported:
+    return std::memory_order_consume;
+  case memory_order::acquire:
+    return std::memory_order_acquire;
+  case memory_order::release:
+    return std::memory_order_release;
+  case memory_order::acq_rel:
+    return std::memory_order_acq_rel;
+  case memory_order::seq_cst:
+    return std::memory_order_seq_cst;
+  }
+}
+
+} // namespace detail
+#endif // __SYCL_DEVICE_ONLY__
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/test/check_device_code/atomic_fence.cpp

@@ -0,0 +1,28 @@
+// RUN: %clangxx -fsycl-device-only -fsycl-unnamed-lambda -S -Xclang -emit-llvm %s -o - | FileCheck %s
+
+#include <sycl/sycl.hpp>
+
+int main() {
+  sycl::queue Q;
+
+  Q.single_task([] {
+    // CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 2, i32 896) #2
+    sycl::atomic_fence(sycl::memory_order::relaxed,
+                       sycl::memory_scope::work_group);
+    // CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 2, i32 898) #2
+    sycl::atomic_fence(sycl::memory_order::acquire,
+                       sycl::memory_scope::work_group);
+    // CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 2, i32 900) #2
+    sycl::atomic_fence(sycl::memory_order::release,
+                       sycl::memory_scope::work_group);
+    // CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 2, i32 904) #2
+    sycl::atomic_fence(sycl::memory_order::acq_rel,
+                       sycl::memory_scope::work_group);
+    // CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 2, i32 912) #2
+    sycl::atomic_fence(sycl::memory_order::seq_cst,
+                       sycl::memory_scope::work_group);
+  });
+  Q.wait();
+
+  return 0;
+}

sycl/test/warnings/sycl_2020_deprecations.cpp

@@ -4,6 +4,7 @@
 // RUN: %clangxx %fsycl-host-only -fsyntax-only -sycl-std=1.2.1 -Xclang -verify -Xclang -verify-ignore-unexpected=note %s -o %t.out
 
 #include <CL/sycl.hpp>
+#include <CL/sycl/ONEAPI/atomic_fence.hpp>
 
 int main() {
   cl_context ClCtx;
@@ -100,5 +101,9 @@ int main() {
   auto MCA = sycl::info::device::max_constant_args;
   (void)MCA;
 
+  // expected-warning@+1{{'atomic_fence' is deprecated: use sycl::atomic_fence instead}}
+  sycl::ONEAPI::atomic_fence(sycl::ONEAPI::memory_order::relaxed,
+                             sycl::ONEAPI::memory_scope::work_group);
+
   return 0;
 }