From a8fde39f74afa181112994b906e4c06087d07850 Mon Sep 17 00:00:00 2001
From: Satyanvesh Dittakavi <Satyanvesh.Dittakavi@amd.com>
Date: Fri, 3 Dec 2021 19:34:18 +0000
Subject: [PATCH] SWDEV-311744 - Update floating point max/min atomic
 implementation

Clang builtins doesnt yet support max/min atomics with float/double
parameters, Hence modify atomicMax/Min to use compare and swap
implementation

Change-Id: Ia351f9b136129fe25f0eddcbbb17549553247182


[ROCm/clr commit: f3c5f12d267561a2d04a8139db8e9d4aecbb0590]
---
 .../include/hip/amd_detail/amd_hip_atomic.h   | 112 ++++++++++++++++--
 1 file changed, 104 insertions(+), 8 deletions(-)

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 43b8293e47..43fbcb1fe9 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
@@ -412,25 +412,73 @@ unsigned long long atomicMin_system(unsigned long long* address, unsigned long l
 __device__
 inline
 float atomicMin(float* address, float val) {
-  return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  unsigned int* uaddr { reinterpret_cast<unsigned int*>(address) };
+  #if __has_builtin(__hip_atomic_load)
+    unsigned int tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT)};
+  #else
+    unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+  #endif
+  float value = __uint_as_float(tmp);
+
+  while (val < value) {
+    value = atomicCAS(address, value, val);
+  }
+
+  return value;
 }
 
 __device__
 inline
 float atomicMin_system(float* address, float val) {
-  return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+  unsigned int* uaddr { reinterpret_cast<unsigned int*>(address) };
+  #if __has_builtin(__hip_atomic_load)
+    unsigned int tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM)};
+  #else
+    unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+  #endif
+  float value = __uint_as_float(tmp);
+
+  while (val < value) {
+    value = atomicCAS(address, value, val);
+  }
+
+  return value;
 }
 
 __device__
 inline
 double atomicMin(double* address, double val) {
-  return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  unsigned long long* uaddr { reinterpret_cast<unsigned long long*>(address) };
+  #if __has_builtin(__hip_atomic_load)
+    unsigned long long tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT)};
+  #else
+    unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+  #endif
+  double value = __longlong_as_double(tmp);
+
+  while (val < value) {
+    value = atomicCAS(address, value, val);
+  }
+
+  return value;
 }
 
 __device__
 inline
 double atomicMin_system(double* address, double val) {
-  return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+  unsigned long long* uaddr { reinterpret_cast<unsigned long long*>(address) };
+  #if __has_builtin(__hip_atomic_load)
+    unsigned long long tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM)};
+  #else
+    unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+  #endif
+  double value = __longlong_as_double(tmp);
+
+  while (val < value) {
+    value = atomicCAS(address, value, val);
+  }
+
+  return value;
 }
 
 __device__
@@ -484,25 +532,73 @@ unsigned long long atomicMax_system(unsigned long long* address, unsigned long l
 __device__
 inline
 float atomicMax(float* address, float val) {
-  return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  unsigned int* uaddr { reinterpret_cast<unsigned int*>(address) };
+  #if __has_builtin(__hip_atomic_load)
+    unsigned int tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT)};
+  #else
+    unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+  #endif
+  float value = __uint_as_float(tmp);
+
+  while (value < val) {
+    value = atomicCAS(address, value, val);
+  }
+
+  return value;
 }
 
 __device__
 inline
 float atomicMax_system(float* address, float val) {
-  return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+  unsigned int* uaddr { reinterpret_cast<unsigned int*>(address) };
+  #if __has_builtin(__hip_atomic_load)
+    unsigned int tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM)};
+  #else
+    unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+  #endif
+  float value = __uint_as_float(tmp);
+
+  while (value < val) {
+    value = atomicCAS(address, value, val);
+  }
+
+  return value;
 }
 
 __device__
 inline
 double atomicMax(double* address, double val) {
-  return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  unsigned long long* uaddr { reinterpret_cast<unsigned long long*>(address) };
+  #if __has_builtin(__hip_atomic_load)
+    unsigned long long tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT)};
+  #else
+    unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+  #endif
+  double value = __longlong_as_double(tmp);
+
+  while (value < val) {
+    value = atomicCAS(address, value, val);
+  }
+
+  return value;
 }
 
 __device__
 inline
 double atomicMax_system(double* address, double val) {
-  return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+  unsigned long long* uaddr { reinterpret_cast<unsigned long long*>(address) };
+  #if __has_builtin(__hip_atomic_load)
+    unsigned long long tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM)};
+  #else
+    unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+  #endif
+  double value = __longlong_as_double(tmp);
+
+  while (value < val) {
+      value = atomicCAS(address, value, val);
+  }
+
+  return value;
 }
 
 __device__