From 4882ad2ea37f1e206ef37e869e5a56b7ab904755 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 22 Jan 2025 17:22:25 +0100 Subject: [PATCH 1/2] Deprecate and replace CUB iterators existing in Thrust Fixes: #3261 --- cub/cub/device/device_run_length_encode.cuh | 9 +++------ .../dispatch/dispatch_streaming_reduce.cuh | 17 ++--------------- cub/test/catch2_test_device_reduce.cu | 4 ---- cub/test/catch2_test_device_reduce_fp_inf.cu | 3 --- .../catch2_test_device_reduce_large_offsets.cu | 3 --- .../catch2_test_device_run_length_encode.cu | 14 -------------- cub/test/catch2_test_util_type.cu | 11 +++++------ 7 files changed, 10 insertions(+), 51 deletions(-) diff --git a/cub/cub/device/device_run_length_encode.cuh b/cub/cub/device/device_run_length_encode.cuh index 72d134c8c04..21ed7bd77ec 100644 --- a/cub/cub/device/device_run_length_encode.cuh +++ b/cub/cub/device/device_run_length_encode.cuh @@ -48,7 +48,8 @@ #include #include #include -#include + +#include #include @@ -200,9 +201,7 @@ struct DeviceRunLengthEncode using length_t = cub::detail::non_void_value_t; // Generator type for providing 1s values for run-length reduction - _CCCL_SUPPRESS_DEPRECATED_PUSH - using lengths_input_iterator_t = ConstantInputIterator; - _CCCL_SUPPRESS_DEPRECATED_POP + using lengths_input_iterator_t = THRUST_NS_QUALIFIER::constant_iterator; using accum_t = ::cuda::std::__accumulator_t; @@ -210,7 +209,6 @@ struct DeviceRunLengthEncode using policy_t = detail::rle::encode::policy_hub; - _CCCL_SUPPRESS_DEPRECATED_PUSH return DispatchReduceByKey< InputIteratorT, UniqueOutputIteratorT, @@ -232,7 +230,6 @@ struct DeviceRunLengthEncode reduction_op(), num_items, stream); - _CCCL_SUPPRESS_DEPRECATED_POP } //! @rst diff --git a/cub/cub/device/dispatch/dispatch_streaming_reduce.cuh b/cub/cub/device/dispatch/dispatch_streaming_reduce.cuh index 2c41bfd267f..a4679453b04 100644 --- a/cub/cub/device/dispatch/dispatch_streaming_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_streaming_reduce.cuh @@ -15,8 +15,8 @@ #include #include -#include +#include #include #include @@ -25,8 +25,6 @@ #ifndef _CCCL_DOXYGEN_INVOKED // Do not document -// suppress deprecation warnings for ConstantInputIterator -_CCCL_SUPPRESS_DEPRECATED_PUSH CUB_NAMESPACE_BEGIN namespace detail::reduce @@ -189,12 +187,6 @@ template , PerPartitionOffsetT, ReductionOpT>> struct dispatch_streaming_arg_reduce_t { -# if _CCCL_COMPILER(NVHPC) - // NVHPC fails to suppress a deprecation when the alias is inside the function below, so we put it here and span a - // deprecation suppression region across the entire file as well - using constant_offset_it_t = ConstantInputIterator; -# endif // _CCCL_COMPILER(NVHPC) - // Internal dispatch routine for computing a device-wide argument extremum, like `ArgMin` and `ArgMax` // // @param[in] d_temp_storage @@ -234,11 +226,7 @@ struct dispatch_streaming_arg_reduce_t cudaStream_t stream) { // Constant iterator to provide the offset of the current partition for the user-provided input iterator -# if !_CCCL_COMPILER(NVHPC) - _CCCL_SUPPRESS_DEPRECATED_PUSH - using constant_offset_it_t = ConstantInputIterator; - _CCCL_SUPPRESS_DEPRECATED_POP -# endif + using constant_offset_it_t = THRUST_NS_QUALIFIER::constant_iterator; // Wrapped input iterator to produce index-value tuples, i.e., -tuples // We make sure to offset the user-provided input iterator by the current partition's offset @@ -382,7 +370,6 @@ struct dispatch_streaming_arg_reduce_t }; } // namespace detail::reduce -_CCCL_SUPPRESS_DEPRECATED_POP CUB_NAMESPACE_END #endif // !_CCCL_DOXYGEN_INVOKED diff --git a/cub/test/catch2_test_device_reduce.cu b/cub/test/catch2_test_device_reduce.cu index 2142259cdc6..da739a9dc35 100644 --- a/cub/test/catch2_test_device_reduce.cu +++ b/cub/test/catch2_test_device_reduce.cu @@ -39,10 +39,6 @@ #include #include -// need to suppress deprecation warnings for ConstantInputIterator in the cudafe1.stub.c file, so there is no matching -// _CCCL_SUPPRESS_DEPRECATED_POP at the end of this file -_CCCL_SUPPRESS_DEPRECATED_PUSH - DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::Reduce, device_reduce); DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::Sum, device_sum); DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::Min, device_min); diff --git a/cub/test/catch2_test_device_reduce_fp_inf.cu b/cub/test/catch2_test_device_reduce_fp_inf.cu index 17b49f1c651..101a30f8b65 100644 --- a/cub/test/catch2_test_device_reduce_fp_inf.cu +++ b/cub/test/catch2_test_device_reduce_fp_inf.cu @@ -45,9 +45,6 @@ DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMin, device_arg_min_old); DECLARE_LAUNCH_WRAPPER(cub::DeviceReduce::ArgMax, device_arg_max_old); _CCCL_SUPPRESS_DEPRECATED_POP -// suppress deprecation of ConstantInputIterator in cudafe1.stub.c file -_CCCL_SUPPRESS_DEPRECATED_PUSH - // %PARAM% TEST_LAUNCH lid 0:1 C2H_TEST("Device reduce arg{min,max} works with inf items", "[reduce][device]") diff --git a/cub/test/catch2_test_device_reduce_large_offsets.cu b/cub/test/catch2_test_device_reduce_large_offsets.cu index 18b7ba5d205..2f9123d8658 100644 --- a/cub/test/catch2_test_device_reduce_large_offsets.cu +++ b/cub/test/catch2_test_device_reduce_large_offsets.cu @@ -24,9 +24,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 of ConstantInputIterator in cudafe1.stub.c file -_CCCL_SUPPRESS_DEPRECATED_PUSH - // %PARAM% TEST_LAUNCH lid 0:1:2 // List of offset types to test diff --git a/cub/test/catch2_test_device_run_length_encode.cu b/cub/test/catch2_test_device_run_length_encode.cu index b001f619958..4fee40e99ef 100644 --- a/cub/test/catch2_test_device_run_length_encode.cu +++ b/cub/test/catch2_test_device_run_length_encode.cu @@ -25,13 +25,6 @@ * ******************************************************************************/ -#include - -#if _CCCL_COMPILER(NVHPC) -// to suppress warnings for CountingInputIterator -_CCCL_SUPPRESS_DEPRECATED_PUSH -#endif // _CCCL_COMPILER(NVHPC) - #include "insert_nested_NVTX_range_guard.h" // above header needs to be included first @@ -50,9 +43,6 @@ _CCCL_SUPPRESS_DEPRECATED_PUSH DECLARE_LAUNCH_WRAPPER(cub::DeviceRunLengthEncode::Encode, run_length_encode); -// suppress deprecation of ConstantInputIterator in cudafe1.stub.c file -_CCCL_SUPPRESS_DEPRECATED_PUSH - // %PARAM% TEST_LAUNCH lid 0:1:2 using all_types = @@ -274,7 +264,3 @@ C2H_TEST("DeviceRunLengthEncode::Encode can handle leading NaN", "[device][run_l REQUIRE(out_counts == reference_counts); REQUIRE(out_num_runs == reference_num_runs); } - -#if _CCCL_COMPILER(NVHPC) -_CCCL_SUPPRESS_DEPRECATED_POP -#endif // _CCCL_COMPILER(NVHPC) diff --git a/cub/test/catch2_test_util_type.cu b/cub/test/catch2_test_util_type.cu index d09392b1fb8..31a4dbbb50b 100644 --- a/cub/test/catch2_test_util_type.cu +++ b/cub/test/catch2_test_util_type.cu @@ -25,10 +25,11 @@ * ******************************************************************************/ -#include -#include #include +#include +#include + #include #include @@ -36,10 +37,9 @@ C2H_TEST("Tests non_void_value_t", "[util][type]") { - _CCCL_SUPPRESS_DEPRECATED_PUSH using fallback_t = float; - using void_fancy_it = cub::DiscardOutputIterator; - using non_void_fancy_it = cub::CountingInputIterator; + using void_fancy_it = thrust::discard_iterator; + using non_void_fancy_it = thrust::counting_iterator; // falls back for const void* STATIC_REQUIRE(::cuda::std::is_same>::value); - _CCCL_SUPPRESS_DEPRECATED_POP } CUB_DEFINE_DETECT_NESTED_TYPE(cat_detect, cat); From 5779408c91a029644821ac70c8c68e6c2ef31bb2 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jan 2025 23:19:32 +0100 Subject: [PATCH 2/2] Consider thrust::discard_iterator's value_type void --- cub/cub/util_type.cuh | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/cub/cub/util_type.cuh b/cub/cub/util_type.cuh index c185bd27327..15b9a327859 100644 --- a/cub/cub/util_type.cuh +++ b/cub/cub/util_type.cuh @@ -45,6 +45,8 @@ #include +#include + #include #include #include @@ -107,7 +109,13 @@ struct non_void_value_impl template struct non_void_value_impl { - using type = ::cuda::std::_If<::cuda::std::is_void>::value, FallbackT, value_t>; + // we consider thrust::discard_iterator's value_type as `void` as well, so users can switch from + // cub::DiscardInputIterator to thrust::discard_iterator. + using type = + ::cuda::std::_If<::cuda::std::is_void>::value + || ::cuda::std::is_same, THRUST_NS_QUALIFIER::discard_iterator<>::value_type>::value, + FallbackT, + value_t>; }; /**