Merge pull request #603 from b-sumner/master

Drop use of NVCC_COMPAT since it is always set

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

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 <hip/hip_runtime_api.h>

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;
 }

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) {

tests/src/deviceLib/hip_ffs.cpp

@@ -49,21 +49,13 @@ THE SOFTWARE.
 template <typename T>
 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
 }