Fix NVRTC 13.3 warning bug in CCCL and CCCL.C (#9171)

Author: NaderAlAwar

c/parallel/src/util/build_utils.h

@@ -24,6 +24,10 @@ namespace cccl::detail
  */
 inline void extend_args_with_build_config(std::vector<const char*>& args, const cccl_build_config* config)
 {
+  // Work around an NVRTC 13.3 diagnostic pragma bug in CUDA Toolkit headers that leaves deprecated vector type warnings
+  // unsuppressed when including headers such as cuda_fp8.h, cuda_fp6.h, and cuda_fp4.h.
+  args.push_back("-D__NV_NO_VECTOR_DEPRECATION_DIAG");
+
   if (config)
   {
     // Add extra compile flags

c/parallel/test/test_util.h

@@ -89,9 +89,9 @@ inline std::string compile(const std::string& source)
   REQUIRE(NVRTC_SUCCESS == nvrtcCreateProgram(&prog, source.c_str(), "op.cu", 0, nullptr, nullptr));
 
   // TEST_CTK_PATH needed to include cuda_fp16.h
-  const char* options[] = {"--std=c++17", "-rdc=true", "-dlto", TEST_CTK_PATH};
+  const char* options[] = {"--std=c++17", "-rdc=true", "-dlto", "-D__NV_NO_VECTOR_DEPRECATION_DIAG", TEST_CTK_PATH};
 
-  if (nvrtcCompileProgram(prog, 4, options) != NVRTC_SUCCESS)
+  if (nvrtcCompileProgram(prog, 5, options) != NVRTC_SUCCESS)
   {
     size_t log_size{};
     REQUIRE(NVRTC_SUCCESS == nvrtcGetProgramLogSize(prog, &log_size));

cub/cub/util_type.cuh

@@ -346,6 +346,7 @@ __CUB_ALIGN_BYTES(int4, 16)
 __CUB_ALIGN_BYTES(uint4, 16)
 __CUB_ALIGN_BYTES(float4, 16)
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 __CUB_ALIGN_BYTES(long4, 16)
 __CUB_ALIGN_BYTES(ulong4, 16)
 _CCCL_SUPPRESS_DEPRECATED_POP
@@ -359,6 +360,7 @@ __CUB_ALIGN_BYTES(longlong2, 16)
 __CUB_ALIGN_BYTES(ulonglong2, 16)
 __CUB_ALIGN_BYTES(double2, 16)
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 __CUB_ALIGN_BYTES(longlong4, 16)
 __CUB_ALIGN_BYTES(ulonglong4, 16)
 __CUB_ALIGN_BYTES(double4, 16)
@@ -645,18 +647,21 @@ CUB_DEFINE_VECTOR_TYPE(signed char,        char)
 CUB_DEFINE_VECTOR_TYPE(short,              short)
 CUB_DEFINE_VECTOR_TYPE(int,                int)
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 CUB_DEFINE_VECTOR_TYPE(long,               long)
 CUB_DEFINE_VECTOR_TYPE(long long,          longlong)
 _CCCL_SUPPRESS_DEPRECATED_POP
 CUB_DEFINE_VECTOR_TYPE(unsigned char,      uchar)
 CUB_DEFINE_VECTOR_TYPE(unsigned short,     ushort)
 CUB_DEFINE_VECTOR_TYPE(unsigned int,       uint)
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 CUB_DEFINE_VECTOR_TYPE(unsigned long,      ulong)
 CUB_DEFINE_VECTOR_TYPE(unsigned long long, ulonglong)
 _CCCL_SUPPRESS_DEPRECATED_POP
 CUB_DEFINE_VECTOR_TYPE(float,              float)
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 CUB_DEFINE_VECTOR_TYPE(double,             double)
 _CCCL_SUPPRESS_DEPRECATED_POP
 CUB_DEFINE_VECTOR_TYPE(bool,               uchar)

cub/test/catch2_test_device_adjacent_difference_substract_right.cu

@@ -65,6 +65,7 @@ C2H_TEST("DeviceAdjacentDifference::SubtractRightCopy does not change the input"
 }
 
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 template <class T>
 struct ref_diff
 {
@@ -171,6 +172,7 @@ C2H_TEST("DeviceAdjacentDifference::SubtractRightCopy works with pointers", "[de
 }
 
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 struct cust_diff
 {
   template <class T>

cub/test/catch2_test_device_histogram.cu

@@ -31,6 +31,7 @@ DECLARE_LAUNCH_WRAPPER(cub::DeviceHistogram::HistogramEven, histogram_even);
 DECLARE_LAUNCH_WRAPPER(cub::DeviceHistogram::HistogramRange, histogram_range);
 
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 DECLARE_TMPL_LAUNCH_WRAPPER(cub::DeviceHistogram::MultiHistogramEven,
                             multi_histogram_even,
                             ESCAPE_LIST(int Channels, int ActiveChannels),

cub/test/catch2_test_device_reduce.cu

@@ -26,6 +26,7 @@ DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::Max, device_max);
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMax, device_arg_max);
 
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMin, device_arg_min_old);
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMax, device_arg_max_old);
 _CCCL_SUPPRESS_DEPRECATED_POP

cub/test/catch2_test_device_reduce.cuh

@@ -204,6 +204,7 @@ inline void init_default_constant(uchar3& val, int element_val = 2)
 }
 
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 inline void init_default_constant(ulonglong4& val, int element_val = 2)
 {
   const auto element_init = static_cast<unsigned long long>(element_val);

cub/test/catch2_test_device_reduce_fp_inf.cu

@@ -14,6 +14,7 @@ DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMin, device_arg_min);
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMax, device_arg_max);
 
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMin, device_arg_min_old);
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMax, device_arg_max_old);
 _CCCL_SUPPRESS_DEPRECATED_POP

cub/test/catch2_test_device_select_unique.cu

@@ -38,6 +38,7 @@ inline ulonglong2 to_bound(const unsigned long long bound)
 }
 
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 template <>
 inline ulonglong4 to_bound(const unsigned long long bound)
 {

cub/test/catch2_test_device_select_unique_by_key.cu

@@ -27,6 +27,7 @@ inline ulonglong2 to_bound(const unsigned long long bound)
 }
 
 _CCCL_SUPPRESS_DEPRECATED_PUSH
+_CCCL_SUPPRESS_DEPRECATED_NVRTC_DIAG
 template <>
 inline ulonglong4 to_bound(const unsigned long long bound)
 {