From 71f96388ede083efce1cc43e2704b05ad8626136 Mon Sep 17 00:00:00 2001 From: Brian Sumner Date: Fri, 27 Jul 2018 09:29:29 -0700 Subject: [PATCH] Drop use of NVCC_COMPAT since it is always set --- include/hip/hcc_detail/device_functions.h | 41 ----------------------- include/hip/hcc_detail/hip_runtime.h | 2 -- tests/src/deviceLib/hip_anyall.cpp | 13 ------- tests/src/deviceLib/hip_clz.cpp | 9 ----- tests/src/deviceLib/hip_ffs.cpp | 8 ----- 5 files changed, 73 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 6455fa6cd1..5e1c8a4e6b 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -48,69 +48,35 @@ __device__ static inline unsigned int __popcll(unsigned long long int input) { } __device__ static inline unsigned int __clz(unsigned int input) { -#ifdef NVCC_COMPAT return input == 0 ? 32 : __builtin_clz(input); -#else - return input == 0 ? -1 : __builtin_clz(input); -#endif } __device__ static inline unsigned int __clzll(unsigned long long int input) { -#ifdef NVCC_COMPAT return input == 0 ? 64 : ( input == 0 ? -1 : __builtin_clzl(input) ); -#else - return input == 0 ? -1 : __builtin_clzl(input); -#endif } __device__ static inline unsigned int __clz(int input) { -#ifdef NVCC_COMPAT return input == 0 ? 32 : ( input > 0 ? __builtin_clz(input) : __builtin_clz(~input) ); -#else - if (input == 0) return -1; - return input > 0 ? __builtin_clz(input) : __builtin_clz(~input); -#endif } __device__ static inline unsigned int __clzll(long long int input) { -#ifdef NVCC_COMPAT return input == 0 ? 64 : input > 0 ? __builtin_clzl(input) : __builtin_clzl(~input); -#else - if (input == 0) return -1; - return input > 0 ? __builtin_clzl(input) : __builtin_clzl(~input); -#endif } __device__ static inline unsigned int __ffs(unsigned int input) { -#ifdef NVCC_COMPAT return ( input == 0 ? -1 : __builtin_ctz(input) ) + 1; -#else - return input == 0 ? -1 : __builtin_ctz(input); -#endif } __device__ static inline unsigned int __ffsll(unsigned long long int input) { -#ifdef NVCC_COMPAT return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1; -#else - return input == 0 ? -1 : __builtin_ctzl(input); -#endif } __device__ static inline unsigned int __ffs(int input) { -#ifdef NVCC_COMPAT return ( input == 0 ? -1 : __builtin_ctz(input) ) + 1; -#else - return input == 0 ? -1 : __builtin_ctz(input); -#endif } __device__ static inline unsigned int __ffsll(long long int input) { -#ifdef NVCC_COMPAT return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1; -#else - return input == 0 ? -1 : __builtin_ctzl(input); -#endif } __device__ static inline unsigned int __brev(unsigned int input) { return __llvm_bitrev_b32(input); } @@ -673,14 +639,7 @@ int __all(int predicate) { __device__ inline int __any(int predicate) { -#ifdef NVCC_COMPAT - if (__ockl_wfany_i32(predicate) != 0) - return 1; - else - return 0; -#else return __ockl_wfany_i32(predicate); -#endif } // XXX from llvm/include/llvm/IR/InstrTypes.h diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index c2ae6e8e4f..17f987a745 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -57,8 +57,6 @@ THE SOFTWARE. #if __HCC_OR_HIP_CLANG__ -// Define NVCC_COMPAT for CUDA compatibility -#define NVCC_COMPAT #define CUDA_SUCCESS hipSuccess #include diff --git a/tests/src/deviceLib/hip_anyall.cpp b/tests/src/deviceLib/hip_anyall.cpp index 21ad9f871c..934c4dbcfd 100644 --- a/tests/src/deviceLib/hip_anyall.cpp +++ b/tests/src/deviceLib/hip_anyall.cpp @@ -86,28 +86,15 @@ int main(int argc, char* argv[]) { printf("warp no. %d __all = %d \n", i, host_all[i]); if (host_all[i] != 1) ++allcount; -#if defined(__HIP_PLATFORM_HCC__) && !defined(NVCC_COMPAT) - if (host_any[i] != 64) ++anycount; -#else if (host_any[i] != 1) ++anycount; -#endif } -#if defined(__HIP_PLATFORM_HCC__) && !defined(NVCC_COMPAT) - if (anycount == 1 && allcount == 1) - printf("PASSED\n"); - else { - printf("FAILED\n"); - return EXIT_FAILURE; - } -#else if (anycount == 0 && allcount == 1) printf("PASSED\n"); else { printf("FAILED\n"); return EXIT_FAILURE; } -#endif return EXIT_SUCCESS; } diff --git a/tests/src/deviceLib/hip_clz.cpp b/tests/src/deviceLib/hip_clz.cpp index 6f3ed33a39..985da04967 100644 --- a/tests/src/deviceLib/hip_clz.cpp +++ b/tests/src/deviceLib/hip_clz.cpp @@ -45,12 +45,7 @@ THE SOFTWARE. unsigned int firstbit_u32(unsigned int a) { if (a == 0) { -#if defined(__HIP_PLATFORM_HCC__) && !defined(NVCC_COMPAT) - - return -1; -#else return 32; -#endif } unsigned int pos = 0; while ((int)a > 0) { @@ -62,11 +57,7 @@ unsigned int firstbit_u32(unsigned int a) { unsigned int firstbit_u64(unsigned long long int a) { if (a == 0) { -#if defined(__HIP_PLATFORM_HCC__) && !defined(NVCC_COMPAT) - return -1; -#else return 64; -#endif } unsigned int pos = 0; while ((long long int)a > 0) { diff --git a/tests/src/deviceLib/hip_ffs.cpp b/tests/src/deviceLib/hip_ffs.cpp index 220c742464..04cf38caa8 100644 --- a/tests/src/deviceLib/hip_ffs.cpp +++ b/tests/src/deviceLib/hip_ffs.cpp @@ -49,21 +49,13 @@ THE SOFTWARE. template int lastbit(T a) { if (a == 0) -#if defined(__HIP_PLATFORM_HCC__) && !defined(NVCC_COMPAT) - return -1; -#else return 0; -#endif int pos = 1; while ((a & 1) != 1) { a >>= 1; pos++; } -#if defined(__HIP_PLATFORM_HCC__) && !defined(NVCC_COMPAT) - return pos - 1; -#else return pos; -#endif }