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

Change-Id: I73f058040afaa782bf6a491b8761a040fbd5b571


[ROCm/clr commit: a8e0ab621e]

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)

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<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__