[BACKPORT 3.0] Backport diagnostic suppression machinery (#5281)

* Enhance our deprecation machinery so that we also suppress the right nvcc warnings (#5138)

* Avoid deprecation suppression within a class definition
* Rework deprecation machinery to also use the special NVCC deprecation

* Fix failing warning suppression for nvrtc (#5278)

I was pretty sure that NVRTC sets the cuda compiler to NVCC but it does not.

We need to properly enable the warning suppression for nvrtc too :(

Author: miscco

cub/cub/detail/fast_modulo_division.cuh

@@ -48,7 +48,7 @@
 #include <cuda/std/type_traits> // ::cuda::std::is_integral
 
 #if defined(CCCL_ENABLE_DEVICE_ASSERTIONS)
-_CCCL_NV_DIAG_SUPPRESS(186) // pointless comparison of unsigned integer with zero
+_CCCL_BEGIN_NV_DIAG_SUPPRESS(186) // pointless comparison of unsigned integer with zero
 #endif // CCCL_ENABLE_DEVICE_ASSERTIONS
 
 CUB_NAMESPACE_BEGIN
@@ -242,5 +242,5 @@ _CCCL_DIAG_POP
 CUB_NAMESPACE_END
 
 #if defined(CCCL_ENABLE_DEVICE_ASSERTIONS)
-_CCCL_NV_DIAG_DEFAULT(186)
+_CCCL_END_NV_DIAG_SUPPRESS()
 #endif // CCCL_ENABLE_DEVICE_ASSERTIONS

cub/test/catch2_test_device_adjacent_difference_substract_right.cu

@@ -89,6 +89,7 @@ C2H_TEST("DeviceAdjacentDifference::SubtractRightCopy does not change the input"
   REQUIRE(reference == in);
 }
 
+_CCCL_SUPPRESS_DEPRECATED_PUSH
 template <class T>
 struct ref_diff
 {
@@ -102,19 +103,17 @@ struct ref_diff
   {
     return ulonglong2{rhs.x - lhs.x, rhs.y - lhs.y};
   }
-
-  _CCCL_SUPPRESS_DEPRECATED_PUSH
   __host__ __device__ constexpr ulonglong4 operator()(const ulonglong4& lhs, const ulonglong4& rhs) const noexcept
   {
     return ulonglong4{rhs.x - lhs.x, rhs.y - lhs.y, rhs.z - lhs.z, rhs.w - lhs.w};
   }
-  _CCCL_SUPPRESS_DEPRECATED_POP
 
   __host__ __device__ constexpr long2 operator()(const long2& lhs, const long2& rhs) const noexcept
   {
     return long2{rhs.x - lhs.x, rhs.y - lhs.y};
   }
 };
+_CCCL_SUPPRESS_DEPRECATED_POP
 
 C2H_TEST("DeviceAdjacentDifference::SubtractRight works with iterators", "[device][adjacent_difference]", types)
 {
@@ -196,6 +195,7 @@ C2H_TEST("DeviceAdjacentDifference::SubtractRightCopy works with pointers", "[de
   REQUIRE(reference == out);
 }
 
+_CCCL_SUPPRESS_DEPRECATED_PUSH
 struct cust_diff
 {
   template <class T>
@@ -209,18 +209,17 @@ struct cust_diff
     return ulonglong2{lhs.x - rhs.x, lhs.y - rhs.y};
   }
 
-  _CCCL_SUPPRESS_DEPRECATED_PUSH
   __host__ __device__ constexpr ulonglong4 operator()(const ulonglong4& lhs, const ulonglong4& rhs) const noexcept
   {
     return ulonglong4{lhs.x - rhs.x, lhs.y - rhs.y, lhs.z - rhs.z, lhs.w - rhs.w};
   }
-  _CCCL_SUPPRESS_DEPRECATED_POP
 
   __host__ __device__ constexpr long2 operator()(const long2& lhs, const long2& rhs) const noexcept
   {
     return long2{lhs.x - rhs.x, lhs.y - rhs.y};
   }
 };
+_CCCL_SUPPRESS_DEPRECATED_POP
 
 C2H_TEST("DeviceAdjacentDifference::SubtractRight works with custom difference",
          "[device][adjacent_difference]",

cub/test/catch2_test_device_reduce.cu

@@ -46,8 +46,6 @@ DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMin, device_arg_min);
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::Max, device_max);
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMax, device_arg_max);
 
-// Suppress deprecation warning for the deprecated ArgMin and ArgMax interfaces
-_CCCL_NV_DIAG_SUPPRESS(1444)
 _CCCL_SUPPRESS_DEPRECATED_PUSH
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMin, device_arg_min_old);
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMax, device_arg_max_old);

cub/test/catch2_test_device_reduce_fp_inf.cu

@@ -38,8 +38,6 @@
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMin, device_arg_min);
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMax, device_arg_max);
 
-// Suppress deprecation warning for the deprecated ArgMin and ArgMax interfaces
-_CCCL_NV_DIAG_SUPPRESS(1444)
 _CCCL_SUPPRESS_DEPRECATED_PUSH
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMin, device_arg_min_old);
 DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMax, device_arg_max_old);

cudax/include/cuda/experimental/__async/sender/env.cuh

@@ -46,7 +46,7 @@ _CCCL_DIAG_SUPPRESS_CLANG("-Wdeprecated-copy")
 // warning #20012-D: __device__ annotation is ignored on a
 // function("inplace_stop_source") that is explicitly defaulted on its first
 // declaration
-_CCCL_NV_DIAG_SUPPRESS(20012)
+_CCCL_BEGIN_NV_DIAG_SUPPRESS(20012)
 
 namespace cuda::experimental::__async
 {
@@ -204,7 +204,7 @@ template <class _Ty>
 using env_of_t = decltype(get_env(__declval<_Ty>()));
 } // namespace cuda::experimental::__async
 
-_CCCL_NV_DIAG_DEFAULT(20012)
+_CCCL_END_NV_DIAG_SUPPRESS()
 
 _CCCL_DIAG_POP
 

cudax/include/cuda/experimental/__async/sender/epilogue.cuh

@@ -13,4 +13,8 @@
 
 #undef _CUDAX_ASYNC_PROLOGUE_INCLUDED
 
+#if _CCCL_CUDA_COMPILER(NVHPC)
+_CCCL_END_NV_DIAG_SUPPRESS()
+#endif // _CCCL_CUDA_COMPILER(NVHPC)
+
 _CCCL_DIAG_POP

cudax/include/cuda/experimental/__async/sender/prologue.cuh

@@ -22,3 +22,7 @@ _CCCL_DIAG_SUPPRESS_MSVC(4848) // [[no_unique_address]] prior to C++20 as a vend
 _CCCL_DIAG_SUPPRESS_GCC("-Wmissing-braces")
 _CCCL_DIAG_SUPPRESS_CLANG("-Wmissing-braces")
 _CCCL_DIAG_SUPPRESS_MSVC(5246) // missing braces around initializer
+
+#if _CCCL_CUDA_COMPILER(NVHPC)
+_CCCL_BEGIN_NV_DIAG_SUPPRESS(cuda_compile)
+#endif // _CCCL_CUDA_COMPILER(NVHPC)

cudax/include/cuda/experimental/__async/sender/stop_token.cuh

@@ -38,7 +38,7 @@
 // warning #20012-D: __device__ annotation is ignored on a
 // function("inplace_stop_source") that is explicitly defaulted on its first
 // declaration
-_CCCL_NV_DIAG_SUPPRESS(20012)
+_CCCL_BEGIN_NV_DIAG_SUPPRESS(20012)
 
 namespace cuda::experimental::__async
 {
@@ -481,7 +481,7 @@ template <class _Token, class _Callback>
 using stop_callback_for_t = typename _Token::template callback_type<_Callback>;
 } // namespace cuda::experimental::__async
 
-_CCCL_NV_DIAG_DEFAULT(20012)
+_CCCL_END_NV_DIAG_SUPPRESS()
 
 #include <cuda/experimental/__async/sender/epilogue.cuh>
 

cudax/include/cuda/experimental/__async/sender/utility.cuh

@@ -122,7 +122,7 @@ _CUDAX_API constexpr _CUDA_VSTD::decay_t<_Ty> __decay_copy(_Ty&& __ty) noexcept(
 _CCCL_DIAG_PUSH
 _CCCL_DIAG_SUPPRESS_GCC("-Wnon-template-friend")
 _CCCL_DIAG_SUPPRESS_NVHPC(probable_guiding_friend)
-_CCCL_NV_DIAG_SUPPRESS(probable_guiding_friend)
+_CCCL_BEGIN_NV_DIAG_SUPPRESS(probable_guiding_friend)
 
 // __zip/__unzip is for keeping type names short. It has the unfortunate side
 // effect of obfuscating the types.
@@ -196,7 +196,7 @@ using __unzip = decltype(__slot_allocated(_Id())());
 using __ignore_this_typedef [[maybe_unused]] = __zip<void>;
 } // namespace
 
-_CCCL_NV_DIAG_DEFAULT(probable_guiding_friend)
+_CCCL_END_NV_DIAG_SUPPRESS()
 _CCCL_DIAG_POP
 
 } // namespace cuda::experimental::__async

cudax/include/cuda/experimental/__async/sender/when_all.cuh

@@ -242,13 +242,13 @@ extern __completion_metadata<_CUDA_VSTD::true_type, // There will be no value co
 
 _CCCL_DIAG_PUSH
 _CCCL_DIAG_SUPPRESS_GCC("-Wunused-value")
-_CCCL_NV_DIAG_SUPPRESS(expr_has_no_effect)
+_CCCL_BEGIN_NV_DIAG_SUPPRESS(expr_has_no_effect)
 _CCCL_DIAG_SUPPRESS_NVHPC(expr_has_no_effect)
 
 template <size_t... _Offsets>
 inline constexpr size_t __last_offset = (0, ..., _Offsets);
 
-_CCCL_NV_DIAG_DEFAULT(expr_has_no_effect)
+_CCCL_END_NV_DIAG_SUPPRESS()
 _CCCL_DIAG_POP
 
 template <size_t _Count, size_t... _Offsets>