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

Change-Id: I829a0a5fd49c510e77eabbcb92d1a415ef6b5a4c

hipnv/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

hipnv/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

hipnv/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