From ecdbaf73cfdafc173bc7151559ba7fc103dbd5ef Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Wed, 20 Jan 2021 02:33:05 -0500 Subject: [PATCH] SWDEV-268989 - Modify atomicAdd double to use clang builtin Change-Id: If9d40046df2f3f5af7cc1bd6c935fbe7d686e7d8 --- include/hip/amd_detail/hip_atomic.h | 17 +---------------- 1 file changed, 1 insertion(+), 16 deletions(-) diff --git a/include/hip/amd_detail/hip_atomic.h b/include/hip/amd_detail/hip_atomic.h index a1370ce377..798254f2a1 100644 --- a/include/hip/amd_detail/hip_atomic.h +++ b/include/hip/amd_detail/hip_atomic.h @@ -72,22 +72,7 @@ __device__ inline double atomicAdd(double* address, double val) { - unsigned long long* uaddr{reinterpret_cast(address)}; - unsigned long long r{__atomic_load_n(uaddr, __ATOMIC_RELAXED)}; - - unsigned long long old; - do { - old = __atomic_load_n(uaddr, __ATOMIC_RELAXED); - - if (r != old) { r = old; continue; } - - r = atomicCAS( - uaddr, r, __double_as_longlong(val + __longlong_as_double(r))); - - if (r == old) break; - } while (true); - - return __longlong_as_double(r); + return __atomic_fetch_add(address, val, __ATOMIC_RELAXED); } __device__