SWDEV-333033 - add safe and unsafe atomic min and max including gfx940 and add missing nvidia support

Change-Id: I829a0a5fd49c510e77eabbcb92d1a415ef6b5a4c


[ROCm/clr commit: 4375b9f5b9]

projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h

@@ -427,20 +427,30 @@ unsigned long long atomicMin_system(unsigned long long* address, unsigned long l
 
 __device__
 inline
-float atomicMin(float* address, float val) {
-  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);
+float atomicMin(float* addr, float val) {
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+  return unsafeAtomicMin(address, val);
+#else
+  #if __has_builtin(__hip_atomic_load) && \
+      __has_builtin(__hip_atomic_compare_exchange_strong)
+  float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  bool done = false;
+  while (!done && value > val) {
+    done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
   }
-
   return value;
+  #else
+  unsigned int *uaddr = (unsigned int *)addr;
+  unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+  bool done = false;
+  while (!done && __uint_as_float(value) > val) {
+    done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  }
+  return __uint_as_float(value);
+  #endif
+#endif
 }
 
 __device__
@@ -463,20 +473,30 @@ float atomicMin_system(float* address, float val) {
 
 __device__
 inline
-double atomicMin(double* address, double val) {
-  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);
+double atomicMin(double* addr, double val) {
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+  return unsafeAtomicMin(address, val);
+#else
+  #if __has_builtin(__hip_atomic_load) && \
+      __has_builtin(__hip_atomic_compare_exchange_strong)
+  double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  bool done = false;
+  while (!done && value > val) {
+    done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
   }
-
   return value;
+  #else
+  unsigned long long *uaddr = (unsigned long long *)addr;
+  unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+  bool done = false;
+  while (!done && __longlong_as_double(value) > val) {
+    done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  }
+  return __longlong_as_double(value);
+  #endif
+#endif
 }
 
 __device__
@@ -547,20 +567,30 @@ unsigned long long atomicMax_system(unsigned long long* address, unsigned long l
 
 __device__
 inline
-float atomicMax(float* address, float val) {
-  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);
+float atomicMax(float* addr, float val) {
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+  return unsafeAtomicMax(addr, val);
+#else
+  #if __has_builtin(__hip_atomic_load) && \
+      __has_builtin(__hip_atomic_compare_exchange_strong)
+  float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  bool done = false;
+  while (!done && value < val) {
+    done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
   }
-
   return value;
+  #else
+  unsigned int *uaddr = (unsigned int *)addr;
+  unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+  bool done = false;
+  while (!done && __uint_as_float(value) < val) {
+    done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  }
+  return __uint_as_float(value);
+  #endif
+#endif
 }
 
 __device__
@@ -583,20 +613,30 @@ float atomicMax_system(float* address, float val) {
 
 __device__
 inline
-double atomicMax(double* address, double val) {
-  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);
+double atomicMax(double* addr, double val) {
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+  return unsafeAtomicMax(addr, val);
+#else
+  #if __has_builtin(__hip_atomic_load) && \
+      __has_builtin(__hip_atomic_compare_exchange_strong)
+  double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  bool done = false;
+  while (!done && value < val) {
+    done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
   }
-
   return value;
+  #else
+  unsigned long long *uaddr = (unsigned long long *)addr;
+  unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+  bool done = false;
+  while (!done && __longlong_as_double(value) < val) {
+    done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  }
+  return __longlong_as_double(value);
+  #endif
+#endif
 }
 
 __device__

projects/clr/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h

@@ -50,7 +50,10 @@ THE SOFTWARE.
  * @return Original value contained in \p addr.
  */
 __device__ inline float unsafeAtomicAdd(float* addr, float value) {
-#if defined(__gfx90a__) &&                                                     \
+#if defined(__gfx940__) &&                                                     \
+    __has_builtin(__builtin_amdgcn_flat_atomic_fadd_f32)
+  return __builtin_amdgcn_flat_atomic_fadd_f32(addr, value);
+#elif defined(__gfx90a__) &&                                                   \
     __has_builtin(__builtin_amdgcn_is_shared) &&                               \
     __has_builtin(__builtin_amdgcn_is_private) &&                              \
     __has_builtin(__builtin_amdgcn_ds_atomic_fadd_f32) &&                      \
@@ -73,6 +76,78 @@ __device__ inline float unsafeAtomicAdd(float* addr, float value) {
 #endif
 }
 
+/**
+ * @brief Unsafe floating point rmw atomic max.
+ *
+ * Performs a relaxed read-modify-write floating point atomic max with
+ * device memory scope. The original value at \p addr is returned and
+ * the value at \p addr is replaced by \p val if greater.
+ *
+ * @note This operation is currently identical to that performed by
+ * atomicMax and is included for completeness.
+ *
+ * @param [in,out] addr Pointer to value to be updated
+ * @param [in] val Value used to update the value at \p addr.
+ * @return Original value contained in \p addr.
+ */
+__device__ inline float unsafeAtomicMax(float* addr, float val) {
+  #if __has_builtin(__hip_atomic_load) && \
+      __has_builtin(__hip_atomic_compare_exchange_strong)
+  float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  bool done = false;
+  while (!done && value < val) {
+    done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  }
+  return value;
+  #else
+  unsigned int *uaddr = (unsigned int *)addr;
+  unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+  bool done = false;
+  while (!done && __uint_as_float(value) < val) {
+    done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  }
+  return __uint_as_float(value);
+  #endif
+}
+
+/**
+ * @brief Unsafe floating point rmw atomic min.
+ *
+ * Performs a relaxed read-modify-write floating point atomic min with
+ * device memory scope. The original value at \p addr is returned and
+ * the value at \p addr is replaced by \p val if lesser.
+ *
+ * @note This operation is currently identical to that performed by
+ * atomicMin and is included for completeness.
+ *
+ * @param [in,out] addr Pointer to value to be updated
+ * @param [in] val Value used to update the value at \p addr.
+ * @return Original value contained in \p addr.
+ */
+__device__ inline float unsafeAtomicMin(float* addr, float val) {
+  #if __has_builtin(__hip_atomic_load) && \
+      __has_builtin(__hip_atomic_compare_exchange_strong)
+  float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  bool done = false;
+  while (!done && value > val) {
+    done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  }
+  return value;
+  #else
+  unsigned int *uaddr = (unsigned int *)addr;
+  unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+  bool done = false;
+  while (!done && __uint_as_float(value) > val) {
+    done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  }
+  return __uint_as_float(value);
+  #endif
+}
+
 /**
  * @brief Unsafe double precision rmw atomic add.
  *
@@ -95,12 +170,12 @@ __device__ inline float unsafeAtomicAdd(float* addr, float value) {
  * Passing in global segment addresses in fine grain allocations will result in
  * undefined behavior and are not supported.
  *
- * @param [in,out] addr Pointer to value to be increment by \p value.
+ * @param [in,out] addr Pointer to value to be updated.
  * @param [in] value Value by \p addr is to be incremented.
  * @return Original value contained in \p addr.
  */
 __device__ inline double unsafeAtomicAdd(double* addr, double value) {
-#if defined(__gfx90a__) &&                                                     \
+#if (defined(__gfx90a__) || defined(__gfx940_)) &&                              \
     __has_builtin(__builtin_amdgcn_flat_atomic_fadd_f64)
   return __builtin_amdgcn_flat_atomic_fadd_f64(addr, value);
 #elif defined (__hip_atomic_fetch_add)
@@ -110,6 +185,112 @@ __device__ inline double unsafeAtomicAdd(double* addr, double value) {
 #endif
 }
 
+/**
+ * @brief Unsafe double precision rmw atomic max.
+ *
+ * Performs a relaxed read-modify-write double precision atomic max with
+ * device memory scope. Original value at \p addr is returned and
+ * the value of \p addr is updated with \p val if greater.
+ *
+ * @note This operation currently only performs different operations for
+ * the gfx90a target. Other devices continue to use safe atomics.
+ *
+ * It can be used to generate code that uses fast hardware floating point atomic
+ * operations which may handle rounding and subnormal values differently than
+ * non-atomic floating point operations.
+ *
+ * The operation is not always safe and can have undefined behavior unless
+ * following condition are met:
+ *
+ * - \p addr is at least 8 byte aligned
+ * - If \p addr is a global segment address, it is in a coarse grain allocation.
+ * Passing in global segment addresses in fine grain allocations will result in
+ * undefined behavior and are not supported.
+ *
+ * @param [in,out] addr Pointer to value to be updated.
+ * @param [in] val Value used to updated the contents at \p addr
+ * @return Original value contained at \p addr.
+ */
+__device__ inline double unsafeAtomicMax(double* addr, double val) {
+#if (defined(__gfx90a__) || defined(__gfx940__)) && \
+    __has_builtin(__builtin_amdgcn_flat_atomic_fmax_f64)
+  return __builtin_amdgcn_flat_atomic_fmax_f64(addr, val);
+#else
+  #if __has_builtin(__hip_atomic_load) && \
+      __has_builtin(__hip_atomic_compare_exchange_strong)
+  double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  bool done = false;
+  while (!done && value < val) {
+    done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  }
+  return value;
+  #else
+  unsigned long long *uaddr = (unsigned long long *)addr;
+  unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+  bool done = false;
+  while (!done && __longlong_as_double(value) < val) {
+    done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  }
+  return __longlong_as_double(value);
+  #endif
+#endif
+}
+
+/**
+ * @brief Unsafe double precision rmw atomic min.
+ *
+ * Performs a relaxed read-modify-write double precision atomic min with
+ * device memory scope. Original value at \p addr is returned and
+ * the value of \p addr is updated with \p val if lesser.
+ *
+ * @note This operation currently only performs different operations for
+ * the gfx90a target. Other devices continue to use safe atomics.
+ *
+ * It can be used to generate code that uses fast hardware floating point atomic
+ * operations which may handle rounding and subnormal values differently than
+ * non-atomic floating point operations.
+ *
+ * The operation is not always safe and can have undefined behavior unless
+ * following condition are met:
+ *
+ * - \p addr is at least 8 byte aligned
+ * - If \p addr is a global segment address, it is in a coarse grain allocation.
+ * Passing in global segment addresses in fine grain allocations will result in
+ * undefined behavior and are not supported.
+ *
+ * @param [in,out] addr Pointer to value to be updated.
+ * @param [in] val Value used to updated the contents at \p addr
+ * @return Original value contained at \p addr.
+ */
+__device__ inline double unsafeAtomicMin(double* addr, double val) {
+#if (defined(__gfx90a__) || defined(__gfx940__)) && \
+    __has_builtin(__builtin_amdgcn_flat_atomic_fmin_f64)
+  return __builtin_amdgcn_flat_atomic_fmin_f64(addr, val);
+#else
+  #if __has_builtin(__hip_atomic_load) && \
+      __has_builtin(__hip_atomic_compare_exchange_strong)
+  double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  bool done = false;
+  while (!done && value > val) {
+    done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  }
+  return value;
+  #else
+  unsigned long long *uaddr = (unsigned long long *)addr;
+  unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+  bool done = false;
+  while (!done && __longlong_as_double(value) > val) {
+    done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  }
+  return __longlong_as_double(value);
+  #endif
+#endif
+}
+
 /**
  * @brief Safe floating point rmw atomic add.
  *
@@ -163,6 +344,78 @@ __device__ inline float safeAtomicAdd(float* addr, float value) {
 #endif
 }
 
+/**
+ * @brief Safe floating point rmw atomic max.
+ *
+ * Performs a relaxed read-modify-write floating point atomic max with
+ * device memory scope. The original value at \p addr is returned and
+ * the value at \p addr is replaced by \p val if greater.
+ *
+ * @note This operation ensures that, on all targets, we produce safe atomics.
+ * This will be the case even when -munsafe-fp-atomics is passed into the compiler.
+ *
+ * @param [in,out] addr Pointer to value to be updated
+ * @param [in] val Value used to update the value at \p addr.
+ * @return Original value contained in \p addr.
+ */
+__device__ inline float safeAtomicMax(float* addr, float val) {
+  #if __has_builtin(__hip_atomic_load) && \
+      __has_builtin(__hip_atomic_compare_exchange_strong)
+  float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  bool done = false;
+  while (!done && value < val) {
+    done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  }
+  return value;
+  #else
+  unsigned int *uaddr = (unsigned int *)addr;
+  unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+  bool done = false;
+  while (!done && __uint_as_float(value) < val) {
+    done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  }
+  return __uint_as_float(value);
+  #endif
+}
+
+/**
+ * @brief Safe floating point rmw atomic min.
+ *
+ * Performs a relaxed read-modify-write floating point atomic min with
+ * device memory scope. The original value at \p addr is returned and
+ * the value at \p addr is replaced by \p val if lesser.
+ *
+ * @note This operation ensures that, on all targets, we produce safe atomics.
+ * This will be the case even when -munsafe-fp-atomics is passed into the compiler.
+ *
+ * @param [in,out] addr Pointer to value to be updated
+ * @param [in] val Value used to update the value at \p addr.
+ * @return Original value contained in \p addr.
+ */
+__device__ inline float safeAtomicMin(float* addr, float val) {
+  #if __has_builtin(__hip_atomic_load) && \
+      __has_builtin(__hip_atomic_compare_exchange_strong)
+  float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  bool done = false;
+  while (!done && value > val) {
+    done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  }
+  return value;
+  #else
+  unsigned int *uaddr = (unsigned int *)addr;
+  unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+  bool done = false;
+  while (!done && __uint_as_float(value) > val) {
+    done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  }
+  return __uint_as_float(value);
+  #endif
+}
+
 /**
  * @brief Safe double precision rmw atomic add.
  *
@@ -178,7 +431,7 @@ __device__ inline float safeAtomicAdd(float* addr, float value) {
  * @return Original value contained in \p addr.
  */
 __device__ inline double safeAtomicAdd(double* addr, double value) {
-#if defined(__gfx90a__) &&                                                    \
+#if (defined(__gfx90a__) || defined(__gfx940__)) &&                                                    \
     __has_builtin(__hip_atomic_fetch_add)
   // On gfx90a, with the __hip_atomic_fetch_add builtin, relaxed system-scope
   // atomics will produce safe CAS loops, but are otherwise not different than
@@ -215,4 +468,99 @@ __device__ inline double safeAtomicAdd(double* addr, double value) {
 #endif // __has_builtin(__hip_atomic_fetch_add)
 #endif
 }
+
+/**
+ * @brief Safe double precision rmw atomic max.
+ *
+ * Performs a relaxed read-modify-write double precision atomic max with
+ * device memory scope. Original value at \p addr is returned and
+ * the value of \p addr is updated with \p val if greater.
+ *
+ * @note This operation ensures that, on all targets, we produce safe atomics.
+ * This will be the case even when -munsafe-fp-atomics is passed into the compiler.
+ *
+ * @param [in,out] addr Pointer to value to be updated.
+ * @param [in] val Value used to updated the contents at \p addr
+ * @return Original value contained at \p addr.
+ */
+__device__ inline double safeAtomicMax(double* addr, double val) {
+  #if __has_builtin(__builtin_amdgcn_is_private)
+  if (__builtin_amdgcn_is_private(
+          (const __attribute__((address_space(0))) void*)addr)) {
+    double old = *addr;
+    *addr = __builtin_fmax(old, val);
+    return old;
+  } else {
+  #endif
+  #if __has_builtin(__hip_atomic_load) && \
+      __has_builtin(__hip_atomic_compare_exchange_strong)
+  double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  bool done = false;
+  while (!done && value < val) {
+    done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  }
+  return value;
+  #else
+  unsigned long long *uaddr = (unsigned long long *)addr;
+  unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+  bool done = false;
+  while (!done && __longlong_as_double(value) < val) {
+    done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  }
+  return __longlong_as_double(value);
+  #endif
+  #if __has_builtin(__builtin_amdgcn_is_private)
+  }
+  #endif
+}
+
+/**
+ * @brief Safe double precision rmw atomic min.
+ *
+ * Performs a relaxed read-modify-write double precision atomic min with
+ * device memory scope. Original value at \p addr is returned and
+ * the value of \p addr is updated with \p val if lesser.
+ *
+ * @note This operation ensures that, on all targets, we produce safe atomics.
+ * This will be the case even when -munsafe-fp-atomics is passed into the compiler.
+ *
+ * @param [in,out] addr Pointer to value to be updated.
+ * @param [in] val Value used to updated the contents at \p addr
+ * @return Original value contained at \p addr.
+ */
+__device__ inline double safeAtomicMin(double* addr, double val) {
+  #if __has_builtin(__builtin_amdgcn_is_private)
+  if (__builtin_amdgcn_is_private(
+           (const __attribute__((address_space(0))) void*)addr)) {
+    double old = *addr;
+    *addr = __builtin_fmin(old, val);
+    return old;
+  } else {
+  #endif
+  #if __has_builtin(__hip_atomic_load) && \
+      __has_builtin(__hip_atomic_compare_exchange_strong)
+  double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  bool done = false;
+  while (!done && value > val) {
+    done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+  }
+  return value;
+  #else
+  unsigned long long *uaddr = (unsigned long long *)addr;
+  unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+  bool done = false;
+  while (!done && __longlong_as_double(value) > val) {
+    done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false,
+               __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  }
+  return __longlong_as_double(value);
+  #endif
+  #if __has_builtin(__builtin_amdgcn_is_private)
+  }
+  #endif
+}
+
 #endif

projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_atomics.h

@@ -0,0 +1,75 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_ATOMICS_H
+#define HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_ATOMICS_H
+
+
+__device__ inline float atomicMax(float* addr, float val) {
+    unsigned int *uaddr = (unsigned int *)addr;
+    float value = __uint_as_float(*uaddr);
+
+    while (value < val) {
+        value = __uint_as_float(atomicCAS(uaddr, __float_as_uint(value),
+                                                 __float_as_uint(val)));
+    }
+    return value;
+}
+
+__device__ inline double atomicMax(double* addr, double val) {
+    unsigned long long* uaddr  = (unsigned long long *)addr;
+    double value = __longlong_as_double(*uaddr);
+
+    while (value < val) {
+        value = __longlong_as_double(atomicCAS(uaddr,
+                                        __double_as_longlong(value),
+                                        __double_as_longlong(val)));
+    }
+
+    return value;
+}
+
+__device__ inline float atomicMin(float* addr, float val) {
+    unsigned int *uaddr = (unsigned int *)addr;
+    float value = __uint_as_float(*uaddr);
+
+    while (value > val) {
+        value = __uint_as_float(atomicCAS(uaddr, __float_as_uint(value),
+                                                 __float_as_uint(val)));
+    }
+    return value;
+}
+
+__device__ inline double atomicMin(double* addr, double val) {
+    unsigned long long* uaddr  = (unsigned long long *)addr;
+    double value = __longlong_as_double(*uaddr);
+
+    while (value > val) {
+        value = __longlong_as_double(atomicCAS(uaddr,
+                                         __double_as_longlong(value),
+                                         __double_as_longlong(val)));
+    }
+
+    return value;
+}
+
+#endif

projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime.h

@@ -76,6 +76,7 @@ typedef int hipLaunchParm;
 
 #ifdef __CUDACC__
 
+#include "nvidia_hip_atomics.h"
 #include "nvidia_hip_unsafe_atomics.h"
 
 #define hipThreadIdx_x threadIdx.x

projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_unsafe_atomics.h

@@ -44,6 +44,22 @@ __device__ inline double unsafeAtomicAdd(double* addr, double value) {
 #endif
 }
 
+__device__ inline float unsafeAtomicMax(float* addr, float value) {
+    return atomicMax(addr, value);
+}
+
+__device__ inline double unsafeAtomicMax(double* addr, double val) {
+    return atomicMax(addr, val);
+}
+
+__device__ inline float unsafeAtomicMin(float* addr, float value) {
+    return atomicMin(addr, value);
+}
+
+__device__ inline double unsafeAtomicMin(double* addr, double val) {
+    return atomicMin(addr, val);
+}
+
 __device__ inline float safeAtomicAdd(float* addr, float value) {
     return atomicAdd(addr, value);
 }
@@ -65,4 +81,20 @@ __device__ inline double safeAtomicAdd(double* addr, double value) {
 #endif
 }
 
+__device__ inline float safeAtomicMax(float* addr, float value) {
+    return atomicMax(addr, value);
+}
+
+__device__ inline double safeAtomicMax(double* addr, double val) {
+    return atomicMax(addr, val);
+}
+
+__device__ inline float safeAtomicMin(float* addr, float value) {
+    return atomicMin(addr, value);
+}
+
+__device__ inline double safeAtomicMin(double* addr, double val) {
+    return atomicMin(addr, val);
+}
+
 #endif