Implement min/max functions in HIP header

Remove using hc::precise_math min and max. Instead we can use ocml directly for device and std:: for host.


[ROCm/hip commit: 47d78e372e]

projects/hip/include/hip/hcc_detail/device_library_decls.h

@@ -43,6 +43,9 @@ extern "C" __device__ float __ocml_rint_f32(float);
 extern "C" __device__ float __ocml_ceil_f32(float);
 extern "C" __device__ float __ocml_trunc_f32(float);
 
+extern "C" __device__ float __ocml_fmin_f32(float, float);
+extern "C" __device__ float __ocml_fmax_f32(float, float);
+
 // Introduce local address space
 #define __local __attribute__((address_space(3)))
 

projects/hip/include/hip/hcc_detail/hip_runtime.h

@@ -201,8 +201,15 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask
 
 #endif  //__HIP_ARCH_GFX803__ == 1
 
-__host__ __device__ int min(int arg1, int arg2);
-__host__ __device__ int max(int arg1, int arg2);
+__device__ inline static int min(int arg1, int arg2) {
+    return (int)(__ocml_fmin_f32((float) arg1, (float) arg2));
+}
+__device__ inline static int max(int arg1, int arg2) {
+    return (int)(__ocml_fmax_f32((float) arg1, (float) arg2));
+}
+
+__host__ inline static int min(int arg1, int arg2) { return std::min(arg1, arg2); }
+__host__ inline static int max(int arg1, int arg2) { return std::max(arg1, arg2); }
 
 #endif  // __HCC_OR_HIP_CLANG__
 

projects/hip/src/device_util.cpp

@@ -147,11 +147,4 @@ __device__ void* __hip_hc_memset(void* dst, uint8_t val, size_t size) {
 // abort
 __device__ void abort() { return hc::abort(); }
 
-__host__ __device__ int min(int arg1, int arg2) {
-    return (int)(hc::precise_math::fmin((float)arg1, (float)arg2));
-}
-__host__ __device__ int max(int arg1, int arg2) {
-    return (int)(hc::precise_math::fmax((float)arg1, (float)arg2));
-}
-
 __host__ void* __get_dynamicgroupbaseptr() { return nullptr; }