[libomptarget] Implement atomic inc and fence functions for AMDGCN using clang builtins

[saiislam]

Summary

These functions use __builtin_amdgcn_atomic_inc32():

uint32_t atomicInc(uint32_t *address);
uint32_t atomicInc(uint32_t *address, uint32_t max);

These functions use builtin_amdgcn_fence():
kmpc_impl_threadfence()
kmpc_impl_threadfence_block()
kmpc_impl_threadfence_system()

They will take place of current mechanism of directly calling IR functions.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_threadfence.hip

11 ↗

(On Diff #275393)

Since these are one line wrappers around intrinsics, it's probably better to implement them as INLINE annotated functions in target_impl.h. Less noise in the filesystem, can do some optimisation on the single-tu level before calling opt.

openmp/libomptarget/deviceRTLs/amdgcn/src/hip_atomics.h

37

Please delete the unused one. Also, should need to mark it INLINE instead of DEVICE to avoid duplicate symbol errors. I'm surprised this linked.

[saiislam]

1. Removed unused functions
2. Changed to INLINE
3. Moved fence functions to target_impl.h

[JonChesterfield]

This looks right, thank you. We need to build with trunk clang at present anyway. I think we'll be able to build with 11.0 once that forks.

openmp/libomptarget/deviceRTLs/amdgcn/src/hip_atomics.h

 //===---- hip_atomics.h - Declarations of hip atomic functions ---- C++ -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef OMPTARGET_AMDGCN_HIP_ATOMICS_H
 #define OMPTARGET_AMDGCN_HIP_ATOMICS_H
 
 #include "target_impl.h"
 
-// inc requires an amdgcn specific intrinsic which is not yet available
-DEVICE unsigned atomicInc(unsigned *address);
-DEVICE unsigned atomicInc(unsigned *address, unsigned max);
-DEVICE int atomicInc(int *address);
-
 namespace {
 
 template <typename T> DEVICE T atomicAdd(T *address, T val) {
   return __atomic_fetch_add(address, val, __ATOMIC_SEQ_CST);
 }
 
 template <typename T> DEVICE T atomicMax(T *address, T val) {
   return __atomic_fetch_max(address, val, __ATOMIC_SEQ_CST);
 }
 
 template <typename T> DEVICE T atomicExch(T *address, T val) {
   T r;
   __atomic_exchange(address, &val, &r, __ATOMIC_SEQ_CST);
   return r;
 }
 
 template <typename T> DEVICE T atomicCAS(T *address, T compare, T val) {
   (void)__atomic_compare_exchange(address, &compare, &val, false,
                                   __ATOMIC_SEQ_CST, __ATOMIC_RELAXED);
   return compare;
 }
 
+INLINE uint32_t atomicInc(uint32_t *address, uint32_t max) {
+  return __builtin_amdgcn_atomic_inc32(address, max, __ATOMIC_SEQ_CST, "");
+}
+
 } // namespace
 #endif

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h

 //===------- target_impl.h - AMDGCN OpenMP GPU implementation ----- HIP -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Declarations and definitions of target specific functions and constants
 //
 //===----------------------------------------------------------------------===//
 #ifndef OMPTARGET_AMDGCN_TARGET_IMPL_H
 #define OMPTARGET_AMDGCN_TARGET_IMPL_H
 
 #ifndef __AMDGCN__
 #error "amdgcn target_impl.h expects to be compiled under __AMDGCN__"
 #endif
 
 #include "amdgcn_interface.h"
 
 #include <assert.h>
 #include <inttypes.h>
 #include <stddef.h>
 #include <stdint.h>

 DEVICE __kmpc_impl_lanemask_t __kmpc_impl_activemask();
 
 DEVICE int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t, int32_t Var,
                                      int32_t SrcLane);
 
 DEVICE int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t, int32_t Var,
                                           uint32_t Delta, int32_t Width);
 
 INLINE void __kmpc_impl_syncthreads() { __builtin_amdgcn_s_barrier(); }
 
 INLINE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t) {
   // AMDGCN doesn't need to sync threads in a warp
 }
 
 INLINE void __kmpc_impl_named_sync(int barrier, uint32_t num_threads) {
   // we have protected the master warp from releasing from its barrier
   // due to a full workgroup barrier in the middle of a work function.
   // So it is ok to issue a full workgroup barrier here.
   __builtin_amdgcn_s_barrier();
 }
 
-DEVICE void __kmpc_impl_threadfence(void);
-DEVICE void __kmpc_impl_threadfence_block(void);
-DEVICE void __kmpc_impl_threadfence_system(void);
+INLINE void __kmpc_impl_threadfence() {
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
+}
+
+INLINE void __kmpc_impl_threadfence_block() {
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
+}
+
+INLINE void __kmpc_impl_threadfence_system() {
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
+}
 
 // Calls to the AMDGCN layer (assuming 1D layout)
 INLINE int GetThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
 INLINE int GetBlockIdInKernel() { return __builtin_amdgcn_workgroup_id_x(); }
 DEVICE int GetNumberOfBlocksInKernel();
 DEVICE int GetNumberOfThreadsInBlock();
 DEVICE unsigned GetWarpId();
 DEVICE unsigned GetLaneId();
 
 DEVICE bool __kmpc_impl_is_first_active_thread();
 
 // Locks