From 542d3219dc0fdb526a0211e280f6f5b8afb3eb7d Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Wed, 25 Nov 2020 10:46:09 -0500 Subject: [PATCH] Modify atomicAdd float to use clang builtin SWDEV-261229: Change atomicAdd float implementation to use clang builtin along with deprecating atomicAddNoRet Change-Id: I5e98034fec54ee1a2b85d5f56273e34945d49ea3 --- hipamd/include/hip/hcc_detail/hip_atomic.h | 17 ++--------------- 1 file changed, 2 insertions(+), 15 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_atomic.h b/hipamd/include/hip/hcc_detail/hip_atomic.h index d00ebcdabb..a1370ce377 100644 --- a/hipamd/include/hip/hcc_detail/hip_atomic.h +++ b/hipamd/include/hip/hcc_detail/hip_atomic.h @@ -57,23 +57,10 @@ __device__ inline float atomicAdd(float* address, float val) { - unsigned int* uaddr{reinterpret_cast(address)}; - unsigned int r{__atomic_load_n(uaddr, __ATOMIC_RELAXED)}; - - unsigned int old; - do { - old = __atomic_load_n(uaddr, __ATOMIC_RELAXED); - - if (r != old) { r = old; continue; } - - r = atomicCAS(uaddr, r, __float_as_uint(val + __uint_as_float(r))); - - if (r == old) break; - } while (true); - - return __uint_as_float(r); + return __atomic_fetch_add(address, val, __ATOMIC_RELAXED); } +DEPRECATED("use atomicAdd instead") __device__ inline void atomicAddNoRet(float* address, float val)