Revert changes for atomic FADD support when address is in LDS (#1701)

This reverts PR #1591 and follow-on PR #1695

include/hip/hcc_detail/hip_atomic.h

@@ -55,7 +55,7 @@ unsigned long long atomicAdd(
 }
 __device__
 inline
-float atomicAdd_impl(float* address, float val)
+float atomicAdd(float* address, float val)
 {
     unsigned int* uaddr{reinterpret_cast<unsigned int*>(address)};
     unsigned int r{__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
@@ -75,34 +75,6 @@ float atomicAdd_impl(float* address, float val)
 }
 __device__
 inline
-bool __hip_is_shared(const __attribute__((address_space(0))) void* ptr) noexcept
-{   // TODO: this is ersatz for __builtin_amdgcn_is_shared.
-    #if defined(__HIP_DEVICE_COMPILE__)
-        const unsigned int gp = reinterpret_cast<unsigned long long>(ptr);
-
-        return
-            gp ==(__builtin_amdgcn_s_getreg((15 << 11) | (16 << 6) | 15) << 16);
-    #else
-        return false;
-    #endif
-}
-__device__
-inline
-float atomicAdd(float* address, float val)
-{
-    using GP = const __attribute__((address_space(0))) void*;
-    using LP = __attribute__((address_space(3))) float*;
-
-    #if __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ || __HIP_ARCH_GFX908__
-        if (__hip_is_shared((GP) address)) {
-            return __builtin_amdgcn_ds_faddf((LP) address, val, 0, 0, false);
-        }
-    #endif
-
-    return atomicAdd_impl(address, val);
-}
-__device__
-inline
 double atomicAdd(double* address, double val)
 {
     unsigned long long* uaddr{reinterpret_cast<unsigned long long*>(address)};