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

SWDEV-282289 - Use preprocessor directive for atomicAdd to support backward compatibility.

Browse Source 
Change-Id: I73f058040afaa782bf6a491b8761a040fbd5b571
This commit is contained in:
kjayapra-amd
2021-07-12 22:08:56 -04:00
committed by Karthik Jayaprakash
parent 11c28b9fe3
commit a8e0ab621e
2 changed files with 42 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/include/hip/amd_detail/amd_hip_atomic.h
+37
View File
@@ -716,7 +716,25 @@ __device__
inline
float atomicAdd(float* address, float val)
{
#ifndef __HIP_USE_CMPXCHG_FOR_FP_ATOMICS
return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
#else
unsigned int* uaddr{reinterpret_cast<unsigned int*>(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);
#endif
}
#if !defined(__HIPCC_RTC__)
@@ -733,7 +751,26 @@ __device__
inline
double atomicAdd(double* address, double val)
{
#ifndef __HIP_USE_CMPXCHG_FOR_FP_ATOMICS
return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
#else
unsigned long long* uaddr{reinterpret_cast<unsigned long long*>(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);
#endif
}
__device__