[NFC] Remove unsafeAtomicAdd from runtime

Change-Id: I66802f13ff61fcdd07fb2b4886e14687face5113
ROCm · Sep 12, 2024 · 6995684 · 6995684
1 parent 8ed841c
commit 6995684
Show file tree

Hide file tree

Showing 2 changed files with 0 additions and 32 deletions.
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -537,8 +537,6 @@ __OMP_RTL(__kmpc_syncwarp, false, Void, Int64)
 
 __OMP_RTL(__kmpc_is_generic_main_thread_id, false, Int8, Int32)
 
-__OMP_RTL(__kmpc_unsafeAtomicAdd, false, Float, FloatPtr, Float)
-
 __OMP_RTL(__kmpc_atomicCASLoopAdd_float, false, Void, FloatPtr, Float)
 
 __OMP_RTL(__kmpc_atomicCASLoopAdd_double, false, Void, DoublePtr, Double)

diff --git a/offload/DeviceRTL/src/Synchronization.cpp b/offload/DeviceRTL/src/Synchronization.cpp
@@ -150,8 +150,6 @@ uint64_t atomicAdd(uint64_t *Address, uint64_t Val,
   return __atomic_fetch_add(Address, Val, Ordering);
 }
 
-float unsafeAtomicAdd(float *addr, float value);
-
 #if defined(__gfx941__)
 #define ATOMIC_CAS_LOOP_MIN(TY) void atomicCASLoopMin(TY *addr, TY val);
 
@@ -377,28 +375,6 @@ void setLock(omp_lock_t *Lock) {
   // test_lock will now return true for any thread in the warp
 }
 
-#if defined(__gfx90a__) && __has_builtin(__builtin_amdgcn_is_shared) &&        \
-    __has_builtin(__builtin_amdgcn_is_private) &&                              \
-    __has_builtin(__builtin_amdgcn_ds_atomic_fadd_f32) &&                      \
-    __has_builtin(__builtin_amdgcn_global_atomic_fadd_f32)
-// This function is called for gfx90a only and single precision
-// floating point type
-float unsafeAtomicAdd(float *addr, float value) {
-  if (__builtin_amdgcn_is_shared(
-          (const __attribute__((address_space(0))) void *)addr))
-    return __builtin_amdgcn_ds_atomic_fadd_f32(
-        (const __attribute__((address_space(3))) float *)addr, value);
-  else if (__builtin_amdgcn_is_private(
-               (const __attribute__((address_space(0))) void *)addr)) {
-    float temp = *addr;
-    *addr = temp + value;
-    return temp;
-  }
-  return __builtin_amdgcn_global_atomic_fadd_f32(
-      (const __attribute__((address_space(1))) float *)addr, value);
-}
-#endif // if defined(gfx90a) &&
-
 void unsetCriticalLock(omp_lock_t *Lock) {
   (void)atomicExchange((uint32_t *)Lock, UNSET, atomic::acq_rel);
 }
@@ -557,8 +533,6 @@ void setLock(omp_lock_t *Lock) {
   } // wait for 0 to be the read value
 }
 
-float unsafeAtomicAdd(float *addr, float value) { return 0.0; }
-
 #pragma omp end declare variant
 ///}
 
@@ -770,10 +744,6 @@ void omp_unset_lock(omp_lock_t *Lock) { impl::unsetLock(Lock); }
 
 int omp_test_lock(omp_lock_t *Lock) { return impl::testLock(Lock); }
 
-float __kmpc_unsafeAtomicAdd(float *addr, float value) {
-  return impl::unsafeAtomicAdd(addr, value);
-}
-
 #if defined(__gfx941__)
 #define KMPC_ATOMIC_CAS_LOOP_MIN(TY)                                           \
   void __kmpc_atomicCASLoopMin_##TY(TY *addr, TY val) {                        \