diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index 934de64c93..ebb1194a6c 100755 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/CMakeLists.txt @@ -41,6 +41,7 @@ list(APPEND CMAKE_MODULE_PATH ${HIP_COMMON_DIR}/cmake) option(BUILD_HIPIFY_CLANG "Enable building the CUDA->HIP converter" OFF) option(__HIP_ENABLE_PCH "Enable/Disable pre-compiled hip headers" ON) option(__HIP_ENABLE_RTC "Enable/Disable pre-processed hiprtc shared lib" ON) +option(__HIP_USE_CMPXCHG_FOR_FP_ATOMICS "Enable/Disable FP Atomics version" OFF) if(__HIP_ENABLE_PCH) set(_pchStatus 1) @@ -48,6 +49,10 @@ else() set(_pchStatus 0) endif() +if(__HIP_USE_CMPXCHG_FOR_FP_ATOMICS MATCHES "ON") + add_definitions(-D__HIP_USE_CMPXCHG_FOR_FP_ATOMICS) +endif() + message (STATUS "HIP_COMMON_DIR found at ${HIP_COMMON_DIR}" ) set(HIP_COMMON_INCLUDE_DIR ${HIP_COMMON_DIR}/include) set(HIP_COMMON_BIN_DIR ${HIP_COMMON_DIR}/bin) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h index bba9988035..44dac4aabe 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h @@ -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(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(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__