From 17d6266150f577d320fcf0fec1cacd55c537fa8f Mon Sep 17 00:00:00 2001
From: Aaron Enye Shi <enye.shi@gmail.com>
Date: Tue, 11 Sep 2018 17:55:28 +0000
Subject: [PATCH] Use templates for min to prevent ambiguity

[ROCm/clr commit: 5f0838300ef7b575ffed290d46b7b0a36f81b62f]
---
 .../include/hip/hcc_detail/hip_runtime.h      | 10 -------
 .../include/hip/hcc_detail/math_functions.h   | 28 +++++++++++++++++++
 2 files changed, 28 insertions(+), 10 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
index 41710e4165..f5294dfc56 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
@@ -199,16 +199,6 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask
 
 #endif  //__HIP_ARCH_GFX803__ == 1
 
-__device__ inline static int min(int arg1, int arg2) {
-    return (arg1 < arg2) ? arg1 : arg2;
-}
-__device__ inline static int max(int arg1, int arg2) {
-    return (arg1 > arg2) ? arg1 : 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__
 
 #if defined __HCC__
diff --git a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h
index 10967605d4..5457a99cfd 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h
@@ -1304,6 +1304,32 @@ float func(float x, int y) \
 }
 __DEF_FLOAT_FUN2I(scalbn)
 
+#if __HCC__
+template<class T>
+__DEVICE__ __host__ inline static const T min(T arg1, T arg2) {
+  return std::min(arg1, arg2);
+}
+
+template<class T>
+__DEVICE__ __host__ inline static const T max(T arg1, T arg2) {
+  return std::max(arg1, arg2);
+}
+#else
+__device__ inline static int min(int arg1, int arg2) {
+  return (arg1 < arg2) ? arg1 : arg2;
+}
+__device__ inline static int max(int arg1, int arg2) {
+  return (arg1 > arg2) ? arg1 : 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);
+}
+
 __DEVICE__
 inline
 float max(float x, float y) {
@@ -1331,6 +1357,8 @@ double min(double x, double y) {
 __HIP_OVERLOAD2(double, max)
 __HIP_OVERLOAD2(double, min)
 
+#endif
+
 #pragma pop_macro("__DEF_FLOAT_FUN")
 #pragma pop_macro("__DEF_FLOAT_FUN2")
 #pragma pop_macro("__DEF_FLOAT_FUN2I")