From 8e1a7fdd09a17beb23d2ca03ffc8d4a21cd70ea6 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Thu, 5 Jul 2018 20:15:41 +0000 Subject: [PATCH] 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: 47d78e372e03e3c6d52081bee080f8712597ad82] --- .../hip/include/hip/hcc_detail/device_library_decls.h | 3 +++ projects/hip/include/hip/hcc_detail/hip_runtime.h | 11 +++++++++-- projects/hip/src/device_util.cpp | 7 ------- 3 files changed, 12 insertions(+), 9 deletions(-) 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; }