From 2b89bb572aa07140f7b60e576fdaae4c25cf75ec Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 10 Jul 2018 18:56:48 +0000 Subject: [PATCH] Fix min/max, icmp asm and add comment for conversion functions [ROCm/hip commit: 4a7ad93655954a81e1b9e863f9ae7fa2cc92e25d] --- projects/hip/include/hip/hcc_detail/device_functions.h | 3 ++- projects/hip/include/hip/hcc_detail/hip_runtime.h | 4 ++-- projects/hip/include/hip/hcc_detail/llvm_intrinsics.h | 4 +++- 3 files changed, 7 insertions(+), 4 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h index 5466982878..47c25c87d7 100644 --- a/projects/hip/include/hip/hcc_detail/device_functions.h +++ b/projects/hip/include/hip/hcc_detail/device_functions.h @@ -440,7 +440,8 @@ __device__ static inline char4 __hip_hc_mul8pk(char4 in1, char4 in2) { } /* -Rounding modes are not yet supported in HIP + * Rounding modes are not yet supported in HIP + * TODO: Conversion functions are not correct, need to fix when BE is ready */ __device__ static inline float __double2float_rd(double x) { return (double)x; } diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index 267f970377..13eaf4fda4 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -202,10 +202,10 @@ __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 (int)(__ocml_fmin_f32((float) arg1, (float) arg2)); + return (arg1 < arg2) ? arg1 : arg2; } __device__ inline static int max(int arg1, int arg2) { - return (int)(__ocml_fmax_f32((float) arg1, (float) arg2)); + return (arg1 > arg2) ? arg1 : arg2; } __host__ inline static int min(int arg1, int arg2) { return std::min(arg1, arg2); } diff --git a/projects/hip/include/hip/hcc_detail/llvm_intrinsics.h b/projects/hip/include/hip/hcc_detail/llvm_intrinsics.h index 6f2fc45626..dc6fd05c52 100644 --- a/projects/hip/include/hip/hcc_detail/llvm_intrinsics.h +++ b/projects/hip/include/hip/hcc_detail/llvm_intrinsics.h @@ -31,7 +31,9 @@ THE SOFTWARE. #include "hip/hcc_detail/host_defines.h" -__device__ ulong __llvm_amdgcn_icmp_i32(uint x, uint y, uint z) __asm("llvm.amdgcn.icmp.i32"); +__device__ +__attribute__((convergent)) +ulong __llvm_amdgcn_icmp_i32(uint x, uint y, uint z) __asm("llvm.amdgcn.icmp.i32"); __device__ unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize");