From f0ea51c7869f0bb12bba2b92a9d3bf4693866dbc Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 17 Jan 2017 09:00:09 -0600 Subject: [PATCH] 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 | 13 ++++--------- hipamd/include/hip/hcc_detail/hip_runtime.h | 5 +++++ .../tests/src/deviceLib/hipIntegerIntrinsics.cpp | 1 + hipamd/tests/src/deviceLib/hip_ballot.cpp | 16 +++++++++------- hipamd/tests/src/deviceLib/hip_brev.cpp | 3 +-- hipamd/tests/src/deviceLib/hip_clz.cpp | 2 +- hipamd/tests/src/deviceLib/hip_ffs.cpp | 3 +-- hipamd/tests/src/deviceLib/hip_popc.cpp | 3 +-- 8 files changed, 23 insertions(+), 23 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h index 8eb9d6a46c..bed206c0ed 100644 --- a/hipamd/include/hip/hcc_detail/device_functions.h +++ b/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 */ diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index f6967c3445..e911d17ebb 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/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); diff --git a/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp b/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp index abe0b66e90..f9328e7a2c 100644 --- a/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp +++ b/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" diff --git a/hipamd/tests/src/deviceLib/hip_ballot.cpp b/hipamd/tests/src/deviceLib/hip_ballot.cpp index 236ceb57fe..629e676bc7 100644 --- a/hipamd/tests/src/deviceLib/hip_ballot.cpp +++ b/hipamd/tests/src/deviceLib/hip_ballot.cpp @@ -26,9 +26,11 @@ THE SOFTWARE. #include #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 #include #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; } - diff --git a/hipamd/tests/src/deviceLib/hip_clz.cpp b/hipamd/tests/src/deviceLib/hip_clz.cpp index 5c60b29a2d..869f4406f5 100644 --- a/hipamd/tests/src/deviceLib/hip_clz.cpp +++ b/hipamd/tests/src/deviceLib/hip_clz.cpp @@ -32,6 +32,7 @@ THE SOFTWARE. #include #include #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; } - diff --git a/hipamd/tests/src/deviceLib/hip_ffs.cpp b/hipamd/tests/src/deviceLib/hip_ffs.cpp index dfdc439a21..ba9bd7b9a0 100644 --- a/hipamd/tests/src/deviceLib/hip_ffs.cpp +++ b/hipamd/tests/src/deviceLib/hip_ffs.cpp @@ -32,7 +32,7 @@ THE SOFTWARE. #include #include #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; } - diff --git a/hipamd/tests/src/deviceLib/hip_popc.cpp b/hipamd/tests/src/deviceLib/hip_popc.cpp index b40bbd2000..6fe214c7fa 100644 --- a/hipamd/tests/src/deviceLib/hip_popc.cpp +++ b/hipamd/tests/src/deviceLib/hip_popc.cpp @@ -32,7 +32,7 @@ THE SOFTWARE. #include #include #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; } -