AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-268989 - Modify atomicAdd double to use clang builtin

Browse Source 
Change-Id: If9d40046df2f3f5af7cc1bd6c935fbe7d686e7d8


[ROCm/hip commit: ecdbaf73cf]
This commit is contained in:
Satyanvesh Dittakavi
2021-01-20 02:33:05 -05:00
parent 601b7f6d7b
commit 8a803a1592
1 changed files with 1 additions and 16 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/include/hip/amd_detail/hip_atomic.h
+1 -16
View File
@@ -72,22 +72,7 @@ __device__
inline inline
double atomicAdd(double* address, double val) double atomicAdd(double* address, double val)
{ {
unsigned long long* uaddr{reinterpret_cast<unsigned long long*>(address)}; return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
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);
} }
__device__ __device__