Fix min/max, icmp asm and add comment for conversion functions

[ROCm/hip commit: 4a7ad93655]

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; }

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); }

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");