[SYCL2020] Implement sycl::atomic_fence

include/hipSYCL/sycl/libkernel/atomic_fence.hpp

@@ -0,0 +1,112 @@
+/*
+ * This file is part of hipSYCL, a SYCL implementation based on CUDA/HIP
+ *
+ * Copyright (c) 2023 Aksel Alpay and contributors
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ *    list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ *    this list of conditions and the following disclaimer in the documentation
+ *    and/or other materials provided with the distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR
+ * ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#ifndef HIPSYCL_ATOMIC_FENCE_HPP
+#define HIPSYCL_ATOMIC_FENCE_HPP
+
+#include "hipSYCL/sycl/libkernel/backend.hpp"
+#include "memory.hpp"
+
+#if HIPSYCL_LIBKERNEL_IS_DEVICE_PASS_SSCP
+#include "hipSYCL/sycl/libkernel/sscp/builtins/fence.hpp"
+#endif
+
+namespace hipsycl {
+namespace sycl {
+namespace detail {
+
+HIPSYCL_KERNEL_TARGET
+inline void atomic_fence(memory_order order, memory_scope scope)
+{
+  __hipsycl_if_target_hiplike(
+    if (order == memory_order::relaxed ||
+        scope == memory_scope::work_item)
+      ;
+    else if(scope <= memory_scope::work_group)
+      __threadfence_block();
+    else
+      __threadfence();
+  );
+
+  __hipsycl_if_target_spirv(
+    __hipsycl_uint32 flags;
+    __spv::Scope fence_scope;
+
+    if(order == memory_order::relaxed)
+      flags = __spv::MemorySemanticsMask::None;
+    else if(order == memory_order::acquire)
+      flags = __spv::MemorySemanticsMask::Acquire;
+    else if(order == memory_order::release)
+      flags = __spv::MemorySemanticsMask::Release;
+    else if(order == memory_order::acq_rel)
+      flags = __spv::MemorySemanticsMask::AcquireRelease;
+    else if(order == memory_order::seq_cst)
+      flags = __spv::MemorySemanticsMask::SequentiallyConsistent;
+
+    if(scope >= memory_scope::sub_group)
+      flags |= __spv::MemorySemanticsMask::SubgroupMemory;
+    if(scope >= memory_scope::work_group)
+      flags |= __spv::MemorySemanticsMask::WorkgroupMemory;
+    if(scope >= memory_scope::device)
+      flags |= __spv::MemorySemanticsMask::CrossWorkgroupMemory;
+
+    if(scope == memory_scope::work_item)
+      fence_scope = __spv::ScopeFlag::Invocation;
+    else if(scope == memory_scope::sub_group)
+      fence_scope = __spv::ScopeFlag::Subgroup;
+    else if(scope == memory_scope::work_group)
+      fence_scope = __spv::ScopeFlag::Workgroup;
+    else if(scope == memory_scope::device)
+      fence_scope = __spv::ScopeFlag::Device;
+    else if(scope == memory_scope::system)
+      fence_scope = __spv::ScopeFlag::CrossDevice;
+
+    __spirv_MemoryBarrier(fence_scope, flags);
+  );
+
+  __hipsycl_if_target_sscp(
+    __hipsycl_sscp_fence(order, scope);
+  );
+
+  // TODO What about CPU?
+  // Empty __hipsycl_if_target_* breaks at compile time w/ nvc++ 22.7 or
+  // older, so comment out that statement for now.
+  //__hipsycl_if_target_host(/* todo */);
+}
+
+} // namespace detail
+
+HIPSYCL_KERNEL_TARGET
+static inline void atomic_fence(memory_order order, memory_scope scope)
+{
+  detail::atomic_fence(order, scope);
+}
+
+} // namespace sycl
+} // namespace hipsycl
+
+#endif

include/hipSYCL/sycl/libkernel/detail/mem_fence.hpp

@@ -28,6 +28,7 @@
 #ifndef HIPSYCL_MEM_FENCE_HPP
 #define HIPSYCL_MEM_FENCE_HPP
 
+#include "hipSYCL/sycl/libkernel/atomic_fence.hpp"
 #include "hipSYCL/sycl/libkernel/backend.hpp"
 #include "hipSYCL/sycl/access.hpp"
 
@@ -41,20 +42,7 @@ struct mem_fence_impl
   HIPSYCL_KERNEL_TARGET
   static void mem_fence()
   {
-
-    __hipsycl_if_target_hiplike(
-      __threadfence();
-    );
-    __hipsycl_if_target_spirv(
-    __spirv_MemoryBarrier(__spv::Scope::Device,
-                          __spv::MemorySemanticsMask::SequentiallyConsistent |
-                          __spv::MemorySemanticsMask::CrossWorkgroupMemory |
-                          __spv::MemorySemanticsMask::WorkgroupMemory);
-    );
-    // TODO What about CPU?
-    // Empty __hipsycl_if_target_* breaks at compile time w/ nvc++ 22.7 or
-    // older, so comment out that statement for now.
-    //__hipsycl_if_target_host(/* todo */);
+    detail::atomic_fence(memory_order::seq_cst, memory_scope::device);
   }
 
 };
@@ -65,21 +53,10 @@ struct mem_fence_impl<access::fence_space::local_space, M>
   HIPSYCL_KERNEL_TARGET
   static void mem_fence()
   {
-    __hipsycl_if_target_hiplike(
-      __threadfence_block();
-    );
-    __hipsycl_if_target_spirv(
-      __spirv_MemoryBarrier(
-          __spv::Scope::Workgroup,
-          static_cast<uint32_t>(
-              __spv::MemorySemanticsMask::SequentiallyConsistent |
-              __spv::MemorySemanticsMask::WorkgroupMemory));
-    );
+    detail::atomic_fence(memory_order::seq_cst, memory_scope::work_group);
   }
 };
 
-
-
 template <
   access::fence_space Fence_space = access::fence_space::global_and_local,
   access::mode Mode = access::mode::read_write

include/hipSYCL/sycl/libkernel/sscp/builtins/fence.hpp

@@ -0,0 +1,38 @@
+/*
+ * This file is part of hipSYCL, a SYCL implementation based on CUDA/HIP
+ *
+ * Copyright (c) 2023 Aksel Alpay and contributors
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ *    list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ *    this list of conditions and the following disclaimer in the documentation
+ *    and/or other materials provided with the distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR
+ * ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#include "builtin_config.hpp"
+
+#ifndef HIPSYCL_SSCP_FENCE_BUILTINS_HPP
+#define HIPSYCL_SSCP_FENCE_BUILTINS_HPP
+
+#define HIPSYCL_SSCP_CONVERGENT_BUILTIN HIPSYCL_SSCP_BUILTIN __attribute__((convergent))
+
+HIPSYCL_SSCP_CONVERGENT_BUILTIN void
+__hipsycl_sscp_fence(__hipsycl_sscp_memory_order, __hipsycl_sscp_memory_scope);
+
+#endif

include/hipSYCL/sycl/sycl.hpp

@@ -72,6 +72,7 @@
 #include "libkernel/builtins.hpp"
 #include "libkernel/atomic.hpp"
 #include "libkernel/atomic_ref.hpp"
+#include "libkernel/atomic_fence.hpp"
 #include "libkernel/stream.hpp"
 #include "libkernel/sub_group.hpp"
 #include "libkernel/group_traits.hpp"

src/libkernel/sscp/amdgpu/fence.cpp

@@ -0,0 +1,87 @@
+/*
+ * This file is part of hipSYCL, a SYCL implementation based on CUDA/HIP
+ *
+ * Copyright (c) 2023 Aksel Alpay and contributors
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ *    list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ *    this list of conditions and the following disclaimer in the documentation
+ *    and/or other materials provided with the distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR
+ * ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#include "hipSYCL/sycl/libkernel/sscp/builtins/fence.hpp"
+
+enum amdgpu_memory_order {
+  relaxed = __ATOMIC_RELAXED,
+  acquire = __ATOMIC_ACQUIRE,
+  release = __ATOMIC_RELEASE,
+  acq_rel = __ATOMIC_ACQ_REL,
+  seq_cst = __ATOMIC_SEQ_CST
+};
+
+enum amdgpu_memory_scope {
+  work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
+  work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
+  device = __OPENCL_MEMORY_SCOPE_DEVICE,
+  all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
+  sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
+};
+
+extern "C" void
+__atomic_work_item_fence(unsigned mem_fence_flags, amdgpu_memory_order, amdgpu_memory_scope);
+
+__attribute__((always_inline)) amdgpu_memory_order
+__hipsycl_amdgpu_get_mem_order(__hipsycl_sscp_memory_order order) {
+
+  if(order == __hipsycl_sscp_memory_order::relaxed)
+    return relaxed;
+  else if(order == __hipsycl_sscp_memory_order::acquire)
+    return acquire;
+  else if(order == __hipsycl_sscp_memory_order::release)
+    return release;
+  else if(order == __hipsycl_sscp_memory_order::acq_rel)
+    return acq_rel;
+  else
+    return seq_cst;
+}
+
+__attribute__((always_inline)) amdgpu_memory_scope
+__hipsycl_amdgpu_get_mem_scope(__hipsycl_sscp_memory_scope scope) {
+
+  if(scope == __hipsycl_sscp_memory_scope::work_item)
+    return work_item;
+  else if(scope == __hipsycl_sscp_memory_scope::sub_group)
+    return sub_group;
+  else if(scope == __hipsycl_sscp_memory_scope::work_group)
+    return work_group;
+  else if(scope == __hipsycl_sscp_memory_scope::device)
+    return device;
+  else
+    return all_svm_devices;
+}
+
+HIPSYCL_SSCP_CONVERGENT_BUILTIN void
+__hipsycl_sscp_fence(__hipsycl_sscp_memory_order order,
+                     __hipsycl_sscp_memory_scope scope) {
+
+  auto mem_order = __hipsycl_amdgpu_get_mem_order(order);
+  auto mem_scope = __hipsycl_amdgpu_get_mem_scope(scope);
+
+  __atomic_work_item_fence(0, mem_order, mem_scope);
+}

src/libkernel/sscp/ptx/fence.cpp

@@ -0,0 +1,43 @@
+/*
+ * This file is part of hipSYCL, a SYCL implementation based on CUDA/HIP
+ *
+ * Copyright (c) 2023 Aksel Alpay and contributors
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ *    list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ *    this list of conditions and the following disclaimer in the documentation
+ *    and/or other materials provided with the distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR
+ * ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#include "hipSYCL/sycl/libkernel/sscp/builtins/fence.hpp"
+
+HIPSYCL_SSCP_CONVERGENT_BUILTIN void
+__hipsycl_sscp_fence(__hipsycl_sscp_memory_order order,
+                     __hipsycl_sscp_memory_scope scope) {
+
+  if (order == hipsycl::sycl::memory_order::relaxed ||
+      scope == hipsycl::sycl::memory_scope::work_item)
+    ;
+  else if(scope <= hipsycl::sycl::memory_scope::work_group)
+    __nvvm_membar_cta();
+  else if(scope == hipsycl::sycl::memory_scope::device)
+    __nvvm_membar_gl();
+  else if(scope == hipsycl::sycl::memory_scope::system)
+    __nvvm_membar_sys();
+}

src/libkernel/sscp/spirv/fence.cpp

@@ -0,0 +1,50 @@
+/*
+ * This file is part of hipSYCL, a SYCL implementation based on CUDA/HIP
+ *
+ * Copyright (c) 2023 Aksel Alpay and contributors
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ *    list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ *    this list of conditions and the following disclaimer in the documentation
+ *    and/or other materials provided with the distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR
+ * ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#include "hipSYCL/sycl/libkernel/sscp/builtins/fence.hpp"
+#include "hipSYCL/sycl/libkernel/sscp/builtins/spirv/spirv_common.hpp"
+
+__attribute__((convergent)) extern "C" void
+__spirv_MemoryBarrier(__spv::ScopeFlag Memory, __hipsycl_uint32 Semantics);
+
+HIPSYCL_SSCP_CONVERGENT_BUILTIN void
+__hipsycl_sscp_fence(__hipsycl_sscp_memory_order order,
+                     __hipsycl_sscp_memory_scope scope) {
+
+  __hipsycl_uint32 flags = get_spirv_memory_semantics(order);
+
+  if(scope >= __hipsycl_sscp_memory_scope::sub_group)
+    flags |= __spv::MemorySemanticsMaskFlag::SubgroupMemory;
+  if(scope >= __hipsycl_sscp_memory_scope::work_group)
+    flags |= __spv::MemorySemanticsMaskFlag::WorkgroupMemory;
+  if(scope >= __hipsycl_sscp_memory_scope::device)
+    flags |= __spv::MemorySemanticsMaskFlag::CrossWorkgroupMemory;
+
+  __spv::ScopeFlag fence_scope = get_spirv_scope(scope);
+
+  __spirv_MemoryBarrier(fence_scope, flags);
+}