From 02c7f3a70fff8df3bf482e014ad74ea2279cd7a2 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 17 Jan 2017 09:27:51 -0600
Subject: [PATCH] added last few integer intrinsic support 1. Added usad,
 umulhi, urhadd 2. Corrected implementation of __hadd, __hradd 3. TODO:
 __sad(). It gets tricky as ISA sees them as unsigned

Change-Id: Ibd2c2133b462f9393f3990355706386c79256bba
---
 include/hip/hcc_detail/device_functions.h | 26 ++++++++++++++++++-----
 1 file changed, 21 insertions(+), 5 deletions(-)

diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h
index bed206c0ed..fb5a1a6c18 100644
--- a/include/hip/hcc_detail/device_functions.h
+++ b/include/hip/hcc_detail/device_functions.h
@@ -27,6 +27,7 @@ extern "C" unsigned int __hip_hc_ir_umul24_int(unsigned int, unsigned int);
 extern "C" signed int __hip_hc_ir_mul24_int(signed int, signed int);
 extern "C" signed int __hip_hc_ir_mulhi_int(signed int, signed int);
 extern "C" unsigned int __hip_hc_ir_umulhi_int(unsigned int, unsigned int);
+extern "C" unsigned int __hip_hc_ir_usad_int(unsigned int, unsigned int, unsigned int);
 // integer intrinsic function __poc __clz __ffs __brev
 __device__ unsigned int __brev( unsigned int x);
 __device__ unsigned long long int __brevll( unsigned long long int x);
@@ -37,7 +38,10 @@ __device__ unsigned int __ffs(int x);
 __device__ unsigned int __ffsll(long long int x);
 __device__ static inline unsigned int __hadd(int x, int y)
 {
-  return (x + y) >> 1;
+  int z = x + y;
+  int sign = z & 0x8000000;
+  int value = z & 0x7FFFFFFF;
+  return ((value) >> 1 || sign);
 }
 __device__ static inline int __mul24(int x, int y)
 {
@@ -52,7 +56,10 @@ __device__ unsigned int __popc( unsigned int x);
 __device__ unsigned int __popcll( unsigned long long int x);
 __device__ static inline int __rhadd(int x, int y)
 {
-  return (x + y + 1) >> 1;
+  int z = x + y + 1;
+  int sign = z & 0x8000000;
+  int value = z & 0x7FFFFFFF;
+  return ((value) >> 1 || sign);
 }
 //__device__ unsigned int __sad(int x, int y, int z);
 /*
@@ -67,9 +74,18 @@ __device__ static inline int __umul24(unsigned int x, unsigned int y)
   return __hip_hc_ir_umul24_int(x, y);
 }
 __device__ unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y);
-__device__ unsigned int __umulhi(unsigned int x, unsigned int y);
-__device__ unsigned int __urhadd(unsigned int x, unsigned int y);
-__device__ unsigned int __usad(unsigned int x, unsigned int y, unsigned int z);
+__device__ static inline unsigned int __umulhi(unsigned int x, unsigned int y)
+{
+  return __hip_hc_ir_umulhi_int(x, y);
+}
+__device__ static inline unsigned int __urhadd(unsigned int x, unsigned int y)
+{
+  return (x + y + 1) >> 1;
+}
+__device__ static inline unsigned int __usad(unsigned int x, unsigned int y, unsigned int z)
+{
+  return __hip_hc_ir_usad_int(x, y, z);
+}
 
 /*
 Rounding modes are not yet supported in HIP