From 13ce9ece77f6109efc09dd7fdc1ba9181dc56894 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 17 Jan 2017 09:59:08 -0600 Subject: [PATCH] enabled integer intrinsics tests Change-Id: I5d28d556f228240eda2fc0098121ed3b29b041e7 --- include/hip/hcc_detail/device_functions.h | 12 ++++++---- src/device_functions.cpp | 19 ++++++++-------- tests/src/deviceLib/hipIntegerIntrinsics.cpp | 24 ++++++++++---------- 3 files changed, 29 insertions(+), 26 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index fb5a1a6c18..06beeb23f8 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/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; diff --git a/src/device_functions.cpp b/src/device_functions.cpp index 7fb67b787a..4b0eb9a5ff 100644 --- a/src/device_functions.cpp +++ b/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]; diff --git a/tests/src/deviceLib/hipIntegerIntrinsics.cpp b/tests/src/deviceLib/hipIntegerIntrinsics.cpp index f9328e7a2c..63530574d8 100644 --- a/tests/src/deviceLib/hipIntegerIntrinsics.cpp +++ b/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)