diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 89a48147ef..450eb778ce 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..6c8510fcbb 100644 --- a/hipamd/include/hip/hcc_detail/math_functions.h +++ b/hipamd/include/hip/hcc_detail/math_functions.h @@ -1304,6 +1304,66 @@ float func(float x, int y) \ } __DEF_FLOAT_FUN2I(scalbn) +#if __HCC__ +template +__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; +} +__DEVICE__ inline static int max(int arg1, int arg2) { + return (arg1 > arg2) ? arg1 : arg2; +} + __DEVICE__ inline float max(float x, float y) { @@ -1331,6 +1391,17 @@ double min(double x, double y) { __HIP_OVERLOAD2(double, max) __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")