Add SSCP implementation for atomic_fence

include/hipSYCL/sycl/libkernel/atomic_fence.hpp

@@ -1,7 +1,7 @@
 /*
  * This file is part of hipSYCL, a SYCL implementation based on CUDA/HIP
  *
- * Copyright (c) 2023 Aksel Alpay
+ * Copyright (c) 2023 Aksel Alpay and contributors
  * All rights reserved.
  *
  * Redistribution and use in source and binary forms, with or without
@@ -31,6 +31,10 @@
 #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 {
@@ -57,6 +61,9 @@ struct atomic_fence_impl<memory_order::seq_cst, memory_scope::work_group>
                           __spv::MemorySemanticsMask::SequentiallyConsistent |
                           __spv::MemorySemanticsMask::WorkgroupMemory)
     );
+    __hipsycl_if_target_sscp(
+      __hipsycl_sscp_fence(memory_order::seq_cst, memory_scope::work_group);
+    );
   }
 };
 
@@ -75,6 +82,9 @@ struct atomic_fence_impl<memory_order::seq_cst, memory_scope::device>
                           __spv::MemorySemanticsMask::CrossWorkgroupMemory |
                           __spv::MemorySemanticsMask::WorkgroupMemory);
     );
+    __hipsycl_if_target_sscp(
+      __hipsycl_sscp_fence(memory_order::seq_cst, memory_scope::device);
+    );
     // 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.

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

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,41 @@
+/*
+ * 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,
+                     __hipsycl_sscp_memory_scope scope) {
+
+  if(scope == hipsycl::sycl::memory_scope::system) {
+    __nvvm_membar_sys();
+  } else if(scope == hipsycl::sycl::memory_scope::device) {
+    __nvvm_membar_gl();
+  } else if(scope <= hipsycl::sycl::memory_scope::work_group) {
+    __nvvm_membar_cta();
+  }
+}

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) {
+
+  __spv::ScopeFlag mem_fence_scope = get_spirv_scope(scope);
+
+  __hipsycl_uint32 flags = get_spirv_memory_semantics(order);
+
+  if(scope <= __hipsycl_sscp_memory_scope::sub_group)
+    flags |= __spv::MemorySemanticsMaskFlag::SubgroupMemory;
+  else if(scope == __hipsycl_sscp_memory_scope::work_group)
+    flags |= __spv::MemorySemanticsMaskFlag::WorkgroupMemory;
+  else if(scope >= __hipsycl_sscp_memory_scope::device)
+    flags |= __spv::MemorySemanticsMaskFlag::CrossWorkgroupMemory;
+
+  __spirv_MemoryBarrier(mem_fence_scope, flags);
+}