Merge pull request #630 from senior-zero/fix-main/github/int128

Cleanup CTK version checks
NVIDIA · Mar 6, 2023 · b4f7e23 · b4f7e23
2 parents 5d12837 + 5dce195
commit b4f7e23
Show file tree

Hide file tree

Showing 8 changed files with 45 additions and 27 deletions.
diff --git a/cmake/CubHeaderTesting.cmake b/cmake/CubHeaderTesting.cmake
@@ -30,6 +30,10 @@ function(cub_add_header_test label definitions)
     target_compile_definitions(${headertest_target} PRIVATE ${definitions})
     cub_clone_target_properties(${headertest_target} ${cub_target})
 
+    if (CUB_IN_THRUST)
+      thrust_fix_clang_nvcc_build_for(${headertest_target})
+    endif()
+
     add_dependencies(cub.all.headers ${headertest_target})
     add_dependencies(${config_prefix}.all ${headertest_target})
   endforeach()

diff --git a/cub/util_type.cuh b/cub/util_type.cuh
@@ -39,12 +39,13 @@
 #include <limits>
 #include <type_traits>
 
-#if (__CUDACC_VER_MAJOR__ >= 9 || CUDA_VERSION >= 9000) && !_NVHPC_CUDA
+#include <cuda.h>
+
+#if !_NVHPC_CUDA
     #include <cuda_fp16.h>
 #endif
-#if (__CUDACC_VER_MAJOR__ >= 11 || CUDA_VERSION >= 11000) && !_NVHPC_CUDA &&   \
-  !defined(CUB_DISABLE_BF16_SUPPORT)
-#include <cuda_bf16.h>
+#if !_NVHPC_CUDA && !defined(CUB_DISABLE_BF16_SUPPORT)
+    #include <cuda_bf16.h>
 #endif
 
 #include <cub/detail/uninitialized_copy.cuh>
@@ -62,7 +63,7 @@ CUB_NAMESPACE_BEGIN
 #define CUB_IS_INT128_ENABLED 1
 #endif // !defined(__CUDACC_RTC_INT128__)
 #else  // !defined(__CUDACC_RTC__)
-#if (__CUDACC_VER_MAJOR__ >= 11 || CUDA_VERSION >= 11050) 
+#if CUDA_VERSION >= 11050
 #if (CUB_HOST_COMPILER == CUB_HOST_COMPILER_GCC) || \
     (CUB_HOST_COMPILER == CUB_HOST_COMPILER_CLANG) || \
     defined(__ICC) || defined(_NVHPC_CUDA)
@@ -1107,7 +1108,7 @@ struct FpLimits<double>
     }
 };
 
-#if (__CUDACC_VER_MAJOR__ >= 9 || CUDA_VERSION >= 9000) && !_NVHPC_CUDA
+#if !_NVHPC_CUDA
 template <>
 struct FpLimits<__half>
 {
@@ -1123,8 +1124,7 @@ struct FpLimits<__half>
 };
 #endif
 
-#if (__CUDACC_VER_MAJOR__ >= 11 || CUDA_VERSION >= 11000) && !_NVHPC_CUDA &&   \
-  !defined(CUB_DISABLE_BF16_SUPPORT)
+#if !_NVHPC_CUDA && !defined(CUB_DISABLE_BF16_SUPPORT)
 template <>
 struct FpLimits<__nv_bfloat16>
 {
@@ -1278,11 +1278,10 @@ struct NumericTraits<__int128_t>
 
 template <> struct NumericTraits<float> :               BaseTraits<FLOATING_POINT, true, false, unsigned int, float> {};
 template <> struct NumericTraits<double> :              BaseTraits<FLOATING_POINT, true, false, unsigned long long, double> {};
-#if (__CUDACC_VER_MAJOR__ >= 9 || CUDA_VERSION >= 9000) && !_NVHPC_CUDA
+#if !_NVHPC_CUDA
     template <> struct NumericTraits<__half> :          BaseTraits<FLOATING_POINT, true, false, unsigned short, __half> {};
 #endif
-#if (__CUDACC_VER_MAJOR__ >= 11 || CUDA_VERSION >= 11000) && !_NVHPC_CUDA &&   \
-  !defined(CUB_DISABLE_BF16_SUPPORT)
+#if !_NVHPC_CUDA && !defined(CUB_DISABLE_BF16_SUPPORT)
     template <> struct NumericTraits<__nv_bfloat16> :   BaseTraits<FLOATING_POINT, true, false, unsigned short, __nv_bfloat16> {};
 #endif
 

diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
@@ -42,6 +42,10 @@ function(cub_add_example target_name_var example_name example_src cub_target)
   cub_clone_target_properties(${example_target} ${cub_target})
   target_include_directories(${example_target} PRIVATE "${CUB_SOURCE_DIR}/examples")
 
+  if (CUB_IN_THRUST)
+    thrust_fix_clang_nvcc_build_for(${example_target})
+  endif()
+
   # Add to the active configuration's meta target
   add_dependencies(${config_meta_target} ${example_target})
 

diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
@@ -163,6 +163,10 @@ function(cub_add_test target_name_var test_name test_src cub_target)
 
       cub_clone_target_properties(${config_c2h_target} ${cub_target})
       target_link_libraries(${config_c2h_target} PRIVATE CUDA::curand ${cub_target})
+
+      if (CUB_IN_THRUST)
+        thrust_fix_clang_nvcc_build_for(${config_c2h_target})
+      endif()
     endif() # config_c2h_target
 
     if (CUB_SEPARATE_CATCH2)
@@ -191,6 +195,10 @@ function(cub_add_test target_name_var test_name test_src cub_target)
           target_link_options(${config_c2run_target} PRIVATE "-cuda")
         endif()
 
+        if (CUB_IN_THRUST)
+          thrust_fix_clang_nvcc_build_for(${config_c2run_target})
+        endif()
+
         add_test(NAME ${config_c2run_target}
           COMMAND "$<TARGET_FILE:${config_c2run_target}>"
         )
@@ -200,6 +208,10 @@ function(cub_add_test target_name_var test_name test_src cub_target)
       target_link_libraries(${config_c2run_target} PRIVATE ${test_target})
     endif() # CUB_SEPARATE_CATCH2
 
+    if (CUB_IN_THRUST)
+      thrust_fix_clang_nvcc_build_for(${test_target})
+    endif()
+
     target_link_libraries(${test_target} PRIVATE
       ${cub_target}
       ${config_c2h_target}
@@ -220,6 +232,10 @@ function(cub_add_test target_name_var test_name test_src cub_target)
     target_include_directories(${test_target} PRIVATE "${CUB_SOURCE_DIR}/test")
     target_compile_definitions(${test_target} PRIVATE CUB_DEBUG_HOST_ASSERTIONS)
 
+    if (CUB_IN_THRUST)
+      thrust_fix_clang_nvcc_build_for(${test_target})
+    endif()
+
     # Add to the active configuration's meta target
     add_dependencies(${config_meta_target} ${test_target})
 

diff --git a/test/test_device_histogram.cu b/test/test_device_histogram.cu
@@ -44,8 +44,7 @@
 
 #include "test_util.h"
 
-#define TEST_HALF_T \
-  (__CUDACC_VER_MAJOR__ >= 9 || CUDA_VERSION >= 9000) && !_NVHPC_CUDA
+#define TEST_HALF_T !_NVHPC_CUDA
 
 #if TEST_HALF_T 
 #include <cuda_fp16.h>

diff --git a/test/test_device_radix_sort.cu b/test/test_device_radix_sort.cu
@@ -43,11 +43,11 @@
 #include <typeinfo>
 #include <vector>
 
-#if (__CUDACC_VER_MAJOR__ >= 9 || CUDA_VERSION >= 9000) && !_NVHPC_CUDA
+#if !_NVHPC_CUDA
     #include <cuda_fp16.h>
 #endif
 
-#if (__CUDACC_VER_MAJOR__ >= 11 || CUDA_VERSION >= 11000) && !_NVHPC_CUDA
+#if !_NVHPC_CUDA
     #include <cuda_bf16.h>
 #endif
 
@@ -984,14 +984,14 @@ struct UnwrapHalfAndBfloat16 {
     using Type = T;
 };
 
-#if (__CUDACC_VER_MAJOR__ >= 9 || CUDA_VERSION >= 9000) && !_NVHPC_CUDA
+#if !_NVHPC_CUDA
 template <>
 struct UnwrapHalfAndBfloat16<half_t> {
     using Type = __half;
 };
 #endif
 
-#if (__CUDACC_VER_MAJOR__ >= 11 || CUDA_VERSION >= 11000) && !_NVHPC_CUDA
+#if !_NVHPC_CUDA
 template <>
 struct UnwrapHalfAndBfloat16<bfloat16_t> {
     using Type = __nv_bfloat16;
@@ -1962,11 +1962,11 @@ int main(int argc, char** argv)
 #ifdef TEST_EXTENDED_KEY_TYPES
     TestGen<short, false>             (num_items, num_segments);
 
-#if (__CUDACC_VER_MAJOR__ >= 9 || CUDA_VERSION >= 9000) && !_NVHPC_CUDA
+#if !_NVHPC_CUDA
     TestGen<half_t, false>            (num_items, num_segments);
 #endif // CTK >= 9
 
-#if (__CUDACC_VER_MAJOR__ >= 11 || CUDA_VERSION >= 11000) && !_NVHPC_CUDA
+#if !_NVHPC_CUDA
 #if !defined(__ICC)
     // Fails with `-0 != 0` with ICC for unknown reasons. See #333.
     TestGen<bfloat16_t, false>        (num_items, num_segments);

diff --git a/test/test_device_reduce.cu b/test/test_device_reduce.cu
@@ -55,11 +55,9 @@
 #include "test_util.h"
 #include <nv/target>
 
-#define TEST_HALF_T \
-  (__CUDACC_VER_MAJOR__ >= 9 || CUDA_VERSION >= 9000) && !_NVHPC_CUDA
+#define TEST_HALF_T !_NVHPC_CUDA
 
-#define TEST_BF_T \
-  (__CUDACC_VER_MAJOR__ >= 11 || CUDA_VERSION >= 11000) && !_NVHPC_CUDA
+#define TEST_BF_T !_NVHPC_CUDA
 
 #if TEST_HALF_T
 #include <cuda_fp16.h>

diff --git a/test/test_device_segmented_sort.cu b/test/test_device_segmented_sort.cu
@@ -45,11 +45,9 @@
 
 #include <fstream>
 
-#define TEST_HALF_T \
-  (__CUDACC_VER_MAJOR__ >= 9 || CUDA_VERSION >= 9000) && !_NVHPC_CUDA
+#define TEST_HALF_T !_NVHPC_CUDA
 
-#define TEST_BF_T \
-  (__CUDACC_VER_MAJOR__ >= 11 || CUDA_VERSION >= 11000) && !_NVHPC_CUDA
+#define TEST_BF_T !_NVHPC_CUDA
 
 #if TEST_HALF_T
 #include <cuda_fp16.h>