diff --git a/thrust/thrust/system/cuda/detail/count.h b/thrust/thrust/system/cuda/detail/count.h index cb9b7017902..530682ba717 100644 --- a/thrust/thrust/system/cuda/detail/count.h +++ b/thrust/thrust/system/cuda/detail/count.h @@ -40,6 +40,7 @@ # include # include +# include # include # include @@ -52,7 +53,7 @@ typename iterator_traits::difference_type _CCCL_HOST_DEVICE count_if(execution_policy& policy, InputIt first, InputIt last, UnaryPred unary_pred) { using size_type = typename iterator_traits::difference_type; - using flag_iterator_t = transform_input_iterator_t; + using flag_iterator_t = transform_iterator; return cuda_cub::reduce_n( policy, flag_iterator_t(first, unary_pred), thrust::distance(first, last), size_type(0), plus()); diff --git a/thrust/thrust/system/cuda/detail/extrema.h b/thrust/thrust/system/cuda/detail/extrema.h index f10c3578173..617eb8bbc79 100644 --- a/thrust/thrust/system/cuda/detail/extrema.h +++ b/thrust/thrust/system/cuda/detail/extrema.h @@ -45,6 +45,8 @@ # include # include # include +# include +# include # include # include # include @@ -370,10 +372,10 @@ element(execution_policy& policy, ItemsIt first, ItemsIt last, BinaryPr IndexType num_items = static_cast(thrust::distance(first, last)); - using iterator_tuple = tuple>; + using iterator_tuple = tuple>; using zip_iterator = zip_iterator; - iterator_tuple iter_tuple = thrust::make_tuple(first, counting_iterator_t(0)); + iterator_tuple iter_tuple = thrust::make_tuple(first, counting_iterator(0)); using arg_min_t = ArgFunctor; using T = tuple; @@ -443,15 +445,15 @@ minmax_element(execution_policy& policy, ItemsIt first, ItemsIt last, B const auto num_items = static_cast(thrust::distance(first, last)); - using iterator_tuple = tuple>; + using iterator_tuple = tuple>; using zip_iterator = zip_iterator; - iterator_tuple iter_tuple = thrust::make_tuple(first, counting_iterator_t(0)); + iterator_tuple iter_tuple = thrust::make_tuple(first, counting_iterator(0)); using arg_minmax_t = __extrema::arg_minmax_f; using two_pairs_type = typename arg_minmax_t::two_pairs_type; using duplicate_t = typename arg_minmax_t::duplicate_tuple; - using transform_t = transform_input_iterator_t; + using transform_t = transform_iterator; zip_iterator begin = make_zip_iterator(iter_tuple); two_pairs_type result = __extrema::extrema( diff --git a/thrust/thrust/system/cuda/detail/find.h b/thrust/thrust/system/cuda/detail/find.h index de633c73ebb..4b4e11c8ed9 100644 --- a/thrust/thrust/system/cuda/detail/find.h +++ b/thrust/thrust/system/cuda/detail/find.h @@ -41,6 +41,8 @@ # include # include +# include +# include # include THRUST_NAMESPACE_BEGIN @@ -116,11 +118,11 @@ find_if_n(execution_policy& policy, InputIt first, Size num_items, Pred const Size interval_size = (thrust::min)(interval_threshold, num_items); // force transform_iterator output to bool - using XfrmIterator = transform_input_iterator_t; - using IteratorTuple = thrust::tuple>; + using XfrmIterator = transform_iterator; + using IteratorTuple = thrust::tuple>; using ZipIterator = thrust::zip_iterator; - IteratorTuple iter_tuple = thrust::make_tuple(XfrmIterator(first, predicate), counting_iterator_t(0)); + IteratorTuple iter_tuple = thrust::make_tuple(XfrmIterator(first, predicate), counting_iterator(0)); ZipIterator begin = thrust::make_zip_iterator(iter_tuple); ZipIterator end = begin + num_items; diff --git a/thrust/thrust/system/cuda/detail/inner_product.h b/thrust/thrust/system/cuda/detail/inner_product.h index af41c5ccda8..1c36d5e256f 100644 --- a/thrust/thrust/system/cuda/detail/inner_product.h +++ b/thrust/thrust/system/cuda/detail/inner_product.h @@ -39,15 +39,15 @@ #if _CCCL_HAS_CUDA_COMPILER # include # include +# include +# include # include - -# include +# include THRUST_NAMESPACE_BEGIN namespace cuda_cub { - template T _CCCL_HOST_DEVICE inner_product( execution_policy& policy, @@ -58,11 +58,9 @@ T _CCCL_HOST_DEVICE inner_product( ReduceOp reduce_op, ProductOp product_op) { - using size_type = typename iterator_traits::difference_type; - size_type num_items = static_cast(thrust::distance(first1, last1)); - using binop_iterator_t = transform_pair_of_input_iterators_t; - - return cuda_cub::reduce_n(policy, binop_iterator_t(first1, first2, product_op), num_items, init, reduce_op); + const auto n = thrust::distance(first1, last1); + const auto first = make_transform_iterator(make_zip_iterator(first1, first2), make_zip_function(product_op)); + return cuda_cub::reduce_n(policy, first, n, init, reduce_op); } template @@ -71,7 +69,6 @@ inner_product(execution_policy& policy, InputIt1 first1, InputIt1 last1 { return cuda_cub::inner_product(policy, first1, last1, first2, init, plus(), multiplies()); } - } // namespace cuda_cub THRUST_NAMESPACE_END diff --git a/thrust/thrust/system/cuda/detail/mismatch.h b/thrust/thrust/system/cuda/detail/mismatch.h index fc7a878b7d5..9987799faca 100644 --- a/thrust/thrust/system/cuda/detail/mismatch.h +++ b/thrust/thrust/system/cuda/detail/mismatch.h @@ -40,15 +40,14 @@ # include # include +# include # include # include - -# include +# include THRUST_NAMESPACE_BEGIN namespace cuda_cub { - template pair _CCCL_HOST_DEVICE mismatch(execution_policy& policy, InputIt1 first1, InputIt1 last1, InputIt2 first2, BinaryPred binary_pred); @@ -69,15 +68,12 @@ template pair _CCCL_HOST_DEVICE mismatch(execution_policy& policy, InputIt1 first1, InputIt1 last1, InputIt2 first2, BinaryPred binary_pred) { - using transform_t = transform_pair_of_input_iterators_t; - - transform_t transform_first = transform_t(first1, first2, binary_pred); - - transform_t result = cuda_cub::find_if_not( - policy, transform_first, transform_first + thrust::distance(first1, last1), ::cuda::std::__identity{}); - - return thrust::make_pair(first1 + thrust::distance(transform_first, result), - first2 + thrust::distance(transform_first, result)); + const auto n = thrust::distance(first1, last1); + const auto first = make_zip_iterator(first1, first2); + const auto last = make_zip_iterator(last1, first2 + n); + const auto mismatch_pos = cuda_cub::find_if_not(policy, first, last, make_zip_function(binary_pred)); + const auto dist = thrust::distance(first, mismatch_pos); + return thrust::make_pair(first1 + dist, first2 + dist); } template diff --git a/thrust/thrust/system/cuda/detail/transform_scan.h b/thrust/thrust/system/cuda/detail/transform_scan.h index ed94edc7d47..2227249ab16 100644 --- a/thrust/thrust/system/cuda/detail/transform_scan.h +++ b/thrust/thrust/system/cuda/detail/transform_scan.h @@ -39,6 +39,7 @@ #if _CCCL_HAS_CUDA_COMPILER # include # include +# include # include # include @@ -66,7 +67,7 @@ OutputIt _CCCL_HOST_DEVICE transform_inclusive_scan( using size_type = typename iterator_traits::difference_type; size_type num_items = static_cast(thrust::distance(first, last)); - using transformed_iterator_t = transform_input_iterator_t; + using transformed_iterator_t = transform_iterator; return cuda_cub::inclusive_scan_n(policy, transformed_iterator_t(first, transform_op), num_items, result, scan_op); } @@ -87,7 +88,7 @@ OutputIt _CCCL_HOST_DEVICE transform_inclusive_scan( using size_type = typename iterator_traits::difference_type; size_type num_items = static_cast(thrust::distance(first, last)); - using transformed_iterator_t = transform_input_iterator_t; + using transformed_iterator_t = transform_iterator; return cuda_cub::inclusive_scan_n( policy, transformed_iterator_t(first, transform_op), num_items, result, init, scan_op); @@ -108,7 +109,7 @@ OutputIt _CCCL_HOST_DEVICE transform_exclusive_scan( using size_type = typename iterator_traits::difference_type; size_type num_items = static_cast(thrust::distance(first, last)); - using transformed_iterator_t = transform_input_iterator_t; + using transformed_iterator_t = transform_iterator; return cuda_cub::exclusive_scan_n( policy, transformed_iterator_t(first, transform_op), num_items, result, init, scan_op); diff --git a/thrust/thrust/system/cuda/detail/util.h b/thrust/thrust/system/cuda/detail/util.h index 4cdde4508d0..27f0e1af288 100644 --- a/thrust/thrust/system/cuda/detail/util.h +++ b/thrust/thrust/system/cuda/detail/util.h @@ -249,12 +249,13 @@ _CCCL_HOST_DEVICE inline void throw_on_error(cudaError_t status, char const* msg } } -// FIXME: Move the iterators elsewhere. - +// deprecated [Since 2.8] template -struct transform_input_iterator_t +struct CCCL_DEPRECATED_BECAUSE("Use thrust::transform_iterator") transform_input_iterator_t { - using self_t = transform_input_iterator_t; + _CCCL_SUPPRESS_DEPRECATED_PUSH + using self_t = transform_input_iterator_t; + _CCCL_SUPPRESS_DEPRECATED_POP using difference_type = typename iterator_traits::difference_type; using value_type = ValueType; using pointer = void; @@ -358,10 +359,14 @@ struct transform_input_iterator_t } }; // struct transform_input_iterarot_t +// deprecated [Since 2.8] template -struct transform_pair_of_input_iterators_t +struct CCCL_DEPRECATED_BECAUSE("Use thrust::transform_iterator of a thrust::zip_iterator") + transform_pair_of_input_iterators_t { - using self_t = transform_pair_of_input_iterators_t; + _CCCL_SUPPRESS_DEPRECATED_PUSH + using self_t = transform_pair_of_input_iterators_t; + _CCCL_SUPPRESS_DEPRECATED_POP using difference_type = typename iterator_traits::difference_type; using value_type = ValueType; using pointer = void; @@ -488,10 +493,13 @@ struct CCCL_DEPRECATED_BECAUSE("Use cuda::std::identity") identity } }; +// deprecated [Since 2.8] template -struct counting_iterator_t +struct CCCL_DEPRECATED_BECAUSE("Use thrust::counting_iterator") counting_iterator_t { - using self_t = counting_iterator_t; + _CCCL_SUPPRESS_DEPRECATED_PUSH + using self_t = counting_iterator_t; + _CCCL_SUPPRESS_DEPRECATED_POP using difference_type = T; using value_type = T; using pointer = void;