From 5f0838300ef7b575ffed290d46b7b0a36f81b62f Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 11 Sep 2018 17:55:28 +0000 Subject: [PATCH 1/4] Use templates for min to prevent ambiguity --- hipamd/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/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 41710e4165..f5294dfc56 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/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/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h index 10967605d4..5457a99cfd 100644 --- a/hipamd/include/hip/hcc_detail/math_functions.h +++ b/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 +__DEVICE__ __host__ inline static const T min(T arg1, T arg2) { + return std::min(arg1, arg2); +} + +template +__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") From 7aa282897f0bc02faf7134e315e835d57622c48d Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 11 Sep 2018 18:48:31 +0000 Subject: [PATCH 2/4] Avoid host min func conflict with gcc min --- .../include/hip/hcc_detail/math_functions.h | 25 ++++++++++--------- 1 file changed, 13 insertions(+), 12 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h index 5457a99cfd..b44d3f00bd 100644 --- a/hipamd/include/hip/hcc_detail/math_functions.h +++ b/hipamd/include/hip/hcc_detail/math_functions.h @@ -1306,30 +1306,22 @@ __DEF_FLOAT_FUN2I(scalbn) #if __HCC__ template -__DEVICE__ __host__ inline static const T min(T arg1, T arg2) { +__DEVICE__ inline static T min(T arg1, T arg2) { return std::min(arg1, arg2); } template -__DEVICE__ __host__ inline static const T max(T arg1, T arg2) { +__DEVICE__ inline static T max(T arg1, T arg2) { return std::max(arg1, arg2); } #else -__device__ inline static int min(int arg1, int arg2) { +__DEVICE__ inline static int min(int arg1, int arg2) { return (arg1 < arg2) ? arg1 : arg2; } -__device__ inline static int max(int arg1, int 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) { @@ -1359,6 +1351,15 @@ __HIP_OVERLOAD2(double, min) #endif +__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); +} + + #pragma pop_macro("__DEF_FLOAT_FUN") #pragma pop_macro("__DEF_FLOAT_FUN2") #pragma pop_macro("__DEF_FLOAT_FUN2I") From ebf565819e7f9fa3443bc681b8c7e75a28760a33 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 12 Sep 2018 14:54:31 +0000 Subject: [PATCH 3/4] Avoid AMP-retrict call to CPU-restrict --- hipamd/include/hip/hcc_detail/math_functions.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h index b44d3f00bd..059aabcc53 100644 --- a/hipamd/include/hip/hcc_detail/math_functions.h +++ b/hipamd/include/hip/hcc_detail/math_functions.h @@ -1307,12 +1307,12 @@ __DEF_FLOAT_FUN2I(scalbn) #if __HCC__ template __DEVICE__ inline static T min(T arg1, T arg2) { - return std::min(arg1, arg2); + return (arg1 < arg2) ? arg1 : arg2; } template __DEVICE__ inline static T max(T arg1, T arg2) { - return std::max(arg1, arg2); + return (arg1 > arg2) ? arg1 : arg2; } #else __DEVICE__ inline static int min(int arg1, int arg2) { From 6fa3ba3e12cd679ead1f3eb8561b8649f2a3dc2e Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Thu, 13 Sep 2018 23:16:20 +0000 Subject: [PATCH 4/4] Fix Tensorflow ambiguous min issue --- .../include/hip/hcc_detail/math_functions.h | 42 +++++++++++++++++++ 1 file changed, 42 insertions(+) diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h index 059aabcc53..6c8510fcbb 100644 --- a/hipamd/include/hip/hcc_detail/math_functions.h +++ b/hipamd/include/hip/hcc_detail/math_functions.h @@ -1310,10 +1310,52 @@ __DEVICE__ inline static T min(T arg1, T arg2) { return (arg1 < arg2) ? arg1 : arg2; } +__DEVICE__ inline static uint32_t min(uint32_t arg1, int32_t arg2) { + return min(arg1, (uint32_t) arg2); +} +/*__DEVICE__ inline static uint32_t min(int32_t arg1, uint32_t arg2) { + return min((uint32_t) arg1, arg2); +} + +__DEVICE__ inline static uint64_t min(uint64_t arg1, int64_t arg2) { + return min(arg1, (uint64_t) arg2); +} +__DEVICE__ inline static uint64_t min(int64_t arg1, uint64_t arg2) { + return min((uint64_t) arg1, arg2); +} + +__DEVICE__ inline static unsigned long long min(unsigned long long arg1, long long arg2) { + return min(arg1, (unsigned long long) arg2); +} +__DEVICE__ inline static unsigned long long min(long long arg1, unsigned long long arg2) { + return min((unsigned long long) arg1, arg2); +}*/ + template __DEVICE__ inline static T max(T arg1, T arg2) { return (arg1 > arg2) ? arg1 : arg2; } + +__DEVICE__ inline static uint32_t max(uint32_t arg1, int32_t arg2) { + return max(arg1, (uint32_t) arg2); +} +__DEVICE__ inline static uint32_t max(int32_t arg1, uint32_t arg2) { + return max((uint32_t) arg1, arg2); +} + +/*__DEVICE__ inline static uint64_t max(uint64_t arg1, int64_t arg2) { + return max(arg1, (uint64_t) arg2); +} +__DEVICE__ inline static uint64_t max(int64_t arg1, uint64_t arg2) { + return max((uint64_t) arg1, arg2); +} + +__DEVICE__ inline static unsigned long long max(unsigned long long arg1, long long arg2) { + return max(arg1, (unsigned long long) arg2); +} +__DEVICE__ inline static unsigned long long max(long long arg1, unsigned long long arg2) { + return max((unsigned long long) arg1, arg2); +}*/ #else __DEVICE__ inline static int min(int arg1, int arg2) { return (arg1 < arg2) ? arg1 : arg2;