enabled integer intrinsics tests

Change-Id: I5d28d556f228240eda2fc0098121ed3b29b041e7


[ROCm/hip commit: 13ce9ece77]

projects/hip/include/hip/hcc_detail/device_functions.h

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
 in the Software without restriction, including without limitation the rights
@@ -28,6 +28,7 @@ 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);
@@ -61,10 +62,11 @@ __device__ static inline int __rhadd(int x, int y)
   int value = z & 0x7FFFFFFF;
   return ((value) >> 1 || sign);
 }
-//__device__ unsigned int __sad(int x, int y, int z);
-/*
-Implement signed version of sad
-*/
+__device__ static inline unsigned int __sad(int x, int y, int z)
+{
+  return x > y ? x - y + z : y - x + z;
+}
+
 __device__ static inline unsigned int __uhadd(unsigned int x, unsigned int y)
 {
   return (x + y) >> 1;

projects/hip/src/device_functions.cpp

@@ -363,7 +363,9 @@ __device__ float __ull2float_rz(unsigned long long int x)
   return (float)x;
 }
 
-
+/*
+Integer Intrinsics
+*/
 
 // integer intrinsic function __poc __clz __ffs __brev
 __device__ unsigned int __popc( unsigned int input)
@@ -486,17 +488,12 @@ struct uintHolder {
   };
 }__attribute__((aligned(8)));
 
-struct uchar2Holder cHoldVal;
-struct ucharHolder cHoldKey;
-struct ucharHolder cHoldOut;
-
-struct intHolder iHold1;
-struct intHolder iHold2;
-struct uintHolder uHold1;
-struct uintHolder uHold2;
 
 __device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s)
 {
+  struct uchar2Holder cHoldVal;
+  struct ucharHolder cHoldKey;
+  struct ucharHolder cHoldOut;
   cHoldKey.ui = s;
   cHoldVal.ui[0] = x;
   cHoldVal.ui[1] = y;
@@ -509,6 +506,8 @@ __device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int
 
 __device__ long long __mul64hi(long long int x, long long int y)
 {
+  struct intHolder iHold1;
+  struct intHolder iHold2;
   iHold1.sl = x;
   iHold2.sl = y;
   iHold1.sl = iHold1.si[1] * iHold2.si[1];
@@ -517,6 +516,8 @@ __device__ long long __mul64hi(long long int x, long long int y)
 
 __device__ unsigned long long __umul64hi(unsigned long long int x, unsigned long long int y)
 {
+  struct uintHolder uHold1;
+  struct uintHolder uHold2;
   uHold1.ul = x;
   uHold2.ul = y;
   uHold1.ul = uHold1.ui[1] * uHold2.ui[1];

projects/hip/tests/src/deviceLib/hipIntegerIntrinsics.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -30,27 +30,27 @@ __device__ void integer_intrinsics()
 {
     __brev((unsigned int)10);
     __brevll((unsigned long long)10);
-    //__byte_perm((unsigned int)0, (unsigned int)0, 0);
+    __byte_perm((unsigned int)0, (unsigned int)0, 0);
     __clz((int)10);
     __clzll((long long)10);
     __ffs((int)10);
     __ffsll((long long)10);
-    //__hadd((int)1, (int)3);
+    __hadd((int)1, (int)3);
     __mul24((int)1, (int)2);
-    //__mul64hi((long long)1, (long long)2);
-    //__mulhi((int)1, (int)2);
+    __mul64hi((long long)1, (long long)2);
+    __mulhi((int)1, (int)2);
     __popc((unsigned int)4);
     __popcll((unsigned long long)4);
     int a = min((int)4, (int)5);
     int b = max((int)4, (int)5);
-    //__rhadd((int)1, (int)2);
-    //__sad((int)1, (int)2, 0);
-    //__uhadd((unsigned int)1, (unsigned int)3);
+    __rhadd((int)1, (int)2);
+    __sad((int)1, (int)2, 0);
+    __uhadd((unsigned int)1, (unsigned int)3);
     __umul24((unsigned int)1, (unsigned int)2);
-    //__umul64hi((unsigned long long)1, (unsigned long long)2);
-    //__umulhi((unsigned int)1, (unsigned int)2);
-    //__urhadd((unsigned int)1, (unsigned int)2);
-    //__usad((unsigned int)1, (unsigned int)2, 0);
+    __umul64hi((unsigned long long)1, (unsigned long long)2);
+    __umulhi((unsigned int)1, (unsigned int)2);
+    __urhadd((unsigned int)1, (unsigned int)2);
+    __usad((unsigned int)1, (unsigned int)2, 0);
 }
 
 __global__ void compileIntegerIntrinsics(hipLaunchParm lp, int ignored)