diff --git a/projects/hip/include/hip/hcc_detail/device_library_decls.h b/projects/hip/include/hip/hcc_detail/device_library_decls.h
index 82c39b24f0..a7e81a1968 100644
--- a/projects/hip/include/hip/hcc_detail/device_library_decls.h
+++ b/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)))
 
diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h
index 3e0f2e27c5..48818fb15d 100644
--- a/projects/hip/include/hip/hcc_detail/hip_runtime.h
+++ b/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__
 
diff --git a/projects/hip/src/device_util.cpp b/projects/hip/src/device_util.cpp
index 5107acd8c6..34e198c61d 100644
--- a/projects/hip/src/device_util.cpp
+++ b/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; }