From 2e9868d597e880145f7638624917364caae0fa0a Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Fri, 25 Oct 2019 11:14:17 +0100
Subject: [PATCH] Fix deadlock, remove old __sync_* use. (#1584)

This fixes a deadlock introduced by the switch to TTAS loops, and is therefore mildly urgent (to prevent the CI from hoovering in the broken code).

[ROCm/hip commit: a855a13c22b3b77cd7f7f731dc98041399afe130]
---
 .../hip/include/hip/hcc_detail/hip_atomic.h   | 24 +++++++++----------
 1 file changed, 12 insertions(+), 12 deletions(-)

diff --git a/projects/hip/include/hip/hcc_detail/hip_atomic.h b/projects/hip/include/hip/hcc_detail/hip_atomic.h
index 2c13411319..263f639e96 100644
--- a/projects/hip/include/hip/hcc_detail/hip_atomic.h
+++ b/projects/hip/include/hip/hcc_detail/hip_atomic.h
@@ -58,15 +58,15 @@ inline
 float atomicAdd(float* address, float val)
 {
     unsigned int* uaddr{reinterpret_cast<unsigned int*>(address)};
-    unsigned int old{__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+    unsigned int r{__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
 
-    unsigned int r;
+    unsigned int old;
     do {
-        r = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+        old = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
 
         if (r != old) { r = old; continue; }
 
-        old = atomicCAS(uaddr, r, __float_as_uint(val + __uint_as_float(r)));
+        r = atomicCAS(uaddr, r, __float_as_uint(val + __uint_as_float(r)));
 
         if (r == old) break;
     } while (true);
@@ -78,15 +78,15 @@ inline
 double atomicAdd(double* address, double val)
 {
     unsigned long long* uaddr{reinterpret_cast<unsigned long long*>(address)};
-    unsigned long long old{__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+    unsigned long long r{__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
 
-    unsigned long long r;
+    unsigned long long old;
     do {
-        r = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+        old = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
 
         if (r != old) { r = old; continue; }
 
-        old = atomicCAS(
+        r = atomicCAS(
             uaddr, r, __double_as_longlong(val + __longlong_as_double(r)));
 
         if (r == old) break;
@@ -140,13 +140,13 @@ __device__
 inline
 int atomicMin(int* address, int val)
 {
-    return __sync_fetch_and_min(address, val);
+    return __atomic_fetch_min(address, val, __ATOMIC_RELAXED);
 }
 __device__
 inline
 unsigned int atomicMin(unsigned int* address, unsigned int val)
 {
-    return __sync_fetch_and_umin(address, val);
+    return __atomic_fetch_min(address, val, __ATOMIC_RELAXED);
 }
 __device__
 inline
@@ -169,13 +169,13 @@ __device__
 inline
 int atomicMax(int* address, int val)
 {
-    return __sync_fetch_and_max(address, val);
+    return __atomic_fetch_max(address, val, __ATOMIC_RELAXED);
 }
 __device__
 inline
 unsigned int atomicMax(unsigned int* address, unsigned int val)
 {
-    return __sync_fetch_and_umax(address, val);
+    return __atomic_fetch_max(address, val, __ATOMIC_RELAXED);
 }
 __device__
 inline