fixed broken tests and device code for integer intrinsics

1. Fixed build issues with new Integer intrinsics
2. Changed tests to work exactly as CUDA code
3. Still some integer intrinsics need to be supported

Change-Id: Ie6f4171259cf4da517436895d4f6f01e01f59b11

hipamd/include/hip/hcc_detail/device_functions.h

@@ -44,21 +44,21 @@ __device__ static inline int __mul24(int x, int y)
   return __hip_hc_ir_mul24_int(x, y);
 }
 __device__ long long int __mul64hi(long long int x, long long int y);
-__device__ int __mulhi(int x, int y)
+__device__ static inline int __mulhi(int x, int y)
 {
   return __hip_hc_ir_mulhi_int(x, y);
 }
 __device__ unsigned int __popc( unsigned int x);
 __device__ unsigned int __popcll( unsigned long long int x);
-__device__ int __rhadd(int x, int y)
+__device__ static inline int __rhadd(int x, int y)
 {
   return (x + y + 1) >> 1;
 }
 //__device__ unsigned int __sad(int x, int y, int z);
 /*
-Implemented signed version of sad
+Implement signed version of sad
 */
-__device__ unsigned int __uhadd(unsigned int x, unsigned int y)
+__device__ static inline unsigned int __uhadd(unsigned int x, unsigned int y)
 {
   return (x + y) >> 1;
 }
@@ -71,11 +71,6 @@ __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);
 
-// warp vote function __all __any __ballot
-__device__ int __all(  int input);
-__device__ int __any( int input);
-__device__  unsigned long long int __ballot( int input);
-
 /*
 Rounding modes are not yet supported in HIP
 */

hipamd/include/hip/hcc_detail/hip_runtime.h

@@ -420,6 +420,11 @@ __device__ unsigned int atomicInc(unsigned int* address,
 __device__ unsigned int atomicDec(unsigned int* address,
                        unsigned int val);
 
+                       // warp vote function __all __any __ballot
+__device__ int __all(  int input);
+__device__ int __any( int input);
+__device__  unsigned long long int __ballot( int input);
+
 // warp shuffle functions
 #ifdef __cplusplus
 __device__ int __shfl(int input, int lane, int width=warpSize);

hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp

@@ -20,6 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 #include "hip/hip_runtime.h"
+#include "hip/device_functions.h"
 #include "test_common.h"
 
 #pragma GCC diagnostic ignored "-Wall"

hipamd/tests/src/deviceLib/hip_ballot.cpp

@@ -26,9 +26,11 @@ THE SOFTWARE.
 #include <iostream>
 
 #include "hip/hip_runtime.h"
+#include "hip/device_functions.h"
+
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 
-__global__ void 
+__global__ void
 	gpu_ballot(hipLaunchParm lp, unsigned int* device_ballot, int Num_Warps_per_Block,int pshift)
 {
 
@@ -39,7 +41,7 @@ __global__ void
 #else
 	atomicAdd(&device_ballot[warp_num+hipBlockIdx_x*Num_Warps_per_Block],__popc(__ballot(tid - 245)));
 #endif
- 
+
 }
 
 
@@ -47,24 +49,24 @@ int main(int argc, char *argv[])
 { int warpSize, pshift;
   hipDeviceProp_t devProp;
   hipGetDeviceProperties(&devProp, 0);
-  
+
   warpSize = devProp.warpSize;
 
   int w = warpSize;
-  pshift = 0; 
+  pshift = 0;
   while (w >>= 1) ++pshift;
-  
+
   unsigned int Num_Threads_per_Block      = 512;
   unsigned int Num_Blocks_per_Grid        = 1;
   unsigned int Num_Warps_per_Block        = Num_Threads_per_Block/warpSize;
   unsigned int Num_Warps_per_Grid         = (Num_Threads_per_Block*Num_Blocks_per_Grid)/warpSize;
   unsigned int* host_ballot = (unsigned int*)malloc(Num_Warps_per_Grid*sizeof(unsigned int));
-  unsigned int* device_ballot; 
+  unsigned int* device_ballot;
   HIP_ASSERT(hipMalloc((void**)&device_ballot, Num_Warps_per_Grid*sizeof(unsigned int)));
   int divergent_count =0;
   for (int i=0; i<Num_Warps_per_Grid; i++) host_ballot[i] = 0;
 
- 
+
   HIP_ASSERT(hipMemcpy(device_ballot, host_ballot, Num_Warps_per_Grid*sizeof(unsigned int), hipMemcpyHostToDevice));
 
   hipLaunchKernel(gpu_ballot, dim3(Num_Blocks_per_Grid),dim3(Num_Threads_per_Block),0,0, device_ballot,Num_Warps_per_Block,pshift);

hipamd/tests/src/deviceLib/hip_brev.cpp

@@ -32,7 +32,7 @@ THE SOFTWARE.
 #include <stdlib.h>
 #include <iostream>
 #include "hip/hip_runtime.h"
-
+#include "hip/device_functions.h"
 
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 
@@ -181,4 +181,3 @@ int main() {
 
   return errors;
 }
-

hipamd/tests/src/deviceLib/hip_clz.cpp

@@ -32,6 +32,7 @@ THE SOFTWARE.
 #include <stdlib.h>
 #include <iostream>
 #include "hip/hip_runtime.h"
+#include "hip/device_functions.h"
 
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 #define WIDTH     8
@@ -188,4 +189,3 @@ int main() {
 
   return errors;
 }
-

hipamd/tests/src/deviceLib/hip_ffs.cpp

@@ -32,7 +32,7 @@ THE SOFTWARE.
 #include <stdlib.h>
 #include <iostream>
 #include "hip/hip_runtime.h"
-
+#include "hip/device_functions.h"
 
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 
@@ -184,4 +184,3 @@ int main() {
 
   return errors;
 }
-

hipamd/tests/src/deviceLib/hip_popc.cpp

@@ -32,7 +32,7 @@ THE SOFTWARE.
 #include <stdlib.h>
 #include <iostream>
 #include "hip/hip_runtime.h"
-
+#include "hip/device_functions.h"
 
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 
@@ -172,4 +172,3 @@ int main() {
 
   return errors;
 }
-