Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Backport to 2.8: Deprecate and replace thrust::cuda_cub iterators (#3422) #3510

Merged
merged 1 commit into from
Jan 30, 2025
Merged

Backport to 2.8: Deprecate and replace thrust::cuda_cub iterators (#3422) #3510

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion thrust/thrust/system/cuda/detail/count.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,6 +40,7 @@
# include <thrust/system/cuda/config.h>

# include <thrust/distance.h>
# include <thrust/iterator/transform_iterator.h>
# include <thrust/system/cuda/detail/reduce.h>
# include <thrust/system/cuda/detail/util.h>

Expand All @@ -52,7 +53,7 @@ typename iterator_traits<InputIt>::difference_type _CCCL_HOST_DEVICE
count_if(execution_policy<Derived>& policy, InputIt first, InputIt last, UnaryPred unary_pred)
{
using size_type = typename iterator_traits<InputIt>::difference_type;
using flag_iterator_t = transform_input_iterator_t<size_type, InputIt, UnaryPred>;
using flag_iterator_t = transform_iterator<UnaryPred, InputIt, size_type, size_type>;

return cuda_cub::reduce_n(
policy, flag_iterator_t(first, unary_pred), thrust::distance(first, last), size_type(0), plus<size_type>());
Expand Down
12 changes: 7 additions & 5 deletions thrust/thrust/system/cuda/detail/extrema.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,6 +45,8 @@
# include <thrust/detail/temporary_array.h>
# include <thrust/distance.h>
# include <thrust/extrema.h>
# include <thrust/iterator/counting_iterator.h>
# include <thrust/iterator/transform_iterator.h>
# include <thrust/pair.h>
# include <thrust/system/cuda/detail/cdp_dispatch.h>
# include <thrust/system/cuda/detail/reduce.h>
Expand Down Expand Up @@ -370,10 +372,10 @@ element(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last, BinaryPr

IndexType num_items = static_cast<IndexType>(thrust::distance(first, last));

using iterator_tuple = tuple<ItemsIt, counting_iterator_t<IndexType>>;
using iterator_tuple = tuple<ItemsIt, counting_iterator<IndexType>>;
using zip_iterator = zip_iterator<iterator_tuple>;

iterator_tuple iter_tuple = thrust::make_tuple(first, counting_iterator_t<IndexType>(0));
iterator_tuple iter_tuple = thrust::make_tuple(first, counting_iterator<IndexType>(0));

using arg_min_t = ArgFunctor<InputType, IndexType, BinaryPred>;
using T = tuple<InputType, IndexType>;
Expand Down Expand Up @@ -443,15 +445,15 @@ minmax_element(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last, B

const auto num_items = static_cast<IndexType>(thrust::distance(first, last));

using iterator_tuple = tuple<ItemsIt, counting_iterator_t<IndexType>>;
using iterator_tuple = tuple<ItemsIt, counting_iterator<IndexType>>;
using zip_iterator = zip_iterator<iterator_tuple>;

iterator_tuple iter_tuple = thrust::make_tuple(first, counting_iterator_t<IndexType>(0));
iterator_tuple iter_tuple = thrust::make_tuple(first, counting_iterator<IndexType>(0));

using arg_minmax_t = __extrema::arg_minmax_f<InputType, IndexType, BinaryPred>;
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<two_pairs_type, zip_iterator, duplicate_t>;
using transform_t = transform_iterator<duplicate_t, zip_iterator, two_pairs_type, two_pairs_type>;

zip_iterator begin = make_zip_iterator(iter_tuple);
two_pairs_type result = __extrema::extrema(
Expand Down
8 changes: 5 additions & 3 deletions thrust/thrust/system/cuda/detail/find.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,6 +41,8 @@

# include <thrust/detail/minmax.h>
# include <thrust/distance.h>
# include <thrust/iterator/counting_iterator.h>
# include <thrust/iterator/transform_iterator.h>
# include <thrust/system/cuda/detail/execution_policy.h>

THRUST_NAMESPACE_BEGIN
Expand Down Expand Up @@ -116,11 +118,11 @@ find_if_n(execution_policy<Derived>& 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<bool, InputIt, Predicate>;
using IteratorTuple = thrust::tuple<XfrmIterator, counting_iterator_t<Size>>;
using XfrmIterator = transform_iterator<Predicate, InputIt, bool, bool>;
using IteratorTuple = thrust::tuple<XfrmIterator, counting_iterator<Size>>;
using ZipIterator = thrust::zip_iterator<IteratorTuple>;

IteratorTuple iter_tuple = thrust::make_tuple(XfrmIterator(first, predicate), counting_iterator_t<Size>(0));
IteratorTuple iter_tuple = thrust::make_tuple(XfrmIterator(first, predicate), counting_iterator<Size>(0));

ZipIterator begin = thrust::make_zip_iterator(iter_tuple);
ZipIterator end = begin + num_items;
Expand Down
15 changes: 6 additions & 9 deletions thrust/thrust/system/cuda/detail/inner_product.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,15 +39,15 @@
#if _CCCL_HAS_CUDA_COMPILER
# include <thrust/detail/minmax.h>
# include <thrust/distance.h>
# include <thrust/iterator/transform_iterator.h>
# include <thrust/iterator/zip_iterator.h>
# include <thrust/system/cuda/detail/reduce.h>

# include <iterator>
# include <thrust/zip_function.h>

THRUST_NAMESPACE_BEGIN

namespace cuda_cub
{

template <class Derived, class InputIt1, class InputIt2, class T, class ReduceOp, class ProductOp>
T _CCCL_HOST_DEVICE inner_product(
execution_policy<Derived>& policy,
Expand All @@ -58,11 +58,9 @@ T _CCCL_HOST_DEVICE inner_product(
ReduceOp reduce_op,
ProductOp product_op)
{
using size_type = typename iterator_traits<InputIt1>::difference_type;
size_type num_items = static_cast<size_type>(thrust::distance(first1, last1));
using binop_iterator_t = transform_pair_of_input_iterators_t<T, InputIt1, InputIt2, ProductOp>;

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 <class Derived, class InputIt1, class InputIt2, class T>
Expand All @@ -71,7 +69,6 @@ inner_product(execution_policy<Derived>& policy, InputIt1 first1, InputIt1 last1
{
return cuda_cub::inner_product(policy, first1, last1, first2, init, plus<T>(), multiplies<T>());
}

} // namespace cuda_cub

THRUST_NAMESPACE_END
Expand Down
20 changes: 8 additions & 12 deletions thrust/thrust/system/cuda/detail/mismatch.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,15 +40,14 @@
# include <thrust/system/cuda/config.h>

# include <thrust/distance.h>
# include <thrust/iterator/zip_iterator.h>
# include <thrust/pair.h>
# include <thrust/system/cuda/detail/execution_policy.h>

# include <cuda/std/__functional/identity.h>
# include <thrust/zip_function.h>

THRUST_NAMESPACE_BEGIN
namespace cuda_cub
{

template <class Derived, class InputIt1, class InputIt2, class BinaryPred>
pair<InputIt1, InputIt2> _CCCL_HOST_DEVICE
mismatch(execution_policy<Derived>& policy, InputIt1 first1, InputIt1 last1, InputIt2 first2, BinaryPred binary_pred);
Expand All @@ -69,15 +68,12 @@ template <class Derived, class InputIt1, class InputIt2, class BinaryPred>
pair<InputIt1, InputIt2> _CCCL_HOST_DEVICE
mismatch(execution_policy<Derived>& policy, InputIt1 first1, InputIt1 last1, InputIt2 first2, BinaryPred binary_pred)
{
using transform_t = transform_pair_of_input_iterators_t<bool, InputIt1, InputIt2, BinaryPred>;

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 <class Derived, class InputIt1, class InputIt2>
Expand Down
7 changes: 4 additions & 3 deletions thrust/thrust/system/cuda/detail/transform_scan.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,6 +39,7 @@
#if _CCCL_HAS_CUDA_COMPILER
# include <thrust/detail/type_traits.h>
# include <thrust/distance.h>
# include <thrust/iterator/transform_iterator.h>
# include <thrust/system/cuda/detail/scan.h>

# include <cuda/std/type_traits>
Expand Down Expand Up @@ -66,7 +67,7 @@ OutputIt _CCCL_HOST_DEVICE transform_inclusive_scan(

using size_type = typename iterator_traits<InputIt>::difference_type;
size_type num_items = static_cast<size_type>(thrust::distance(first, last));
using transformed_iterator_t = transform_input_iterator_t<value_type, InputIt, TransformOp>;
using transformed_iterator_t = transform_iterator<TransformOp, InputIt, value_type, value_type>;

return cuda_cub::inclusive_scan_n(policy, transformed_iterator_t(first, transform_op), num_items, result, scan_op);
}
Expand All @@ -87,7 +88,7 @@ OutputIt _CCCL_HOST_DEVICE transform_inclusive_scan(

using size_type = typename iterator_traits<InputIt>::difference_type;
size_type num_items = static_cast<size_type>(thrust::distance(first, last));
using transformed_iterator_t = transform_input_iterator_t<value_type, InputIt, TransformOp>;
using transformed_iterator_t = transform_iterator<TransformOp, InputIt, value_type, value_type>;

return cuda_cub::inclusive_scan_n(
policy, transformed_iterator_t(first, transform_op), num_items, result, init, scan_op);
Expand All @@ -108,7 +109,7 @@ OutputIt _CCCL_HOST_DEVICE transform_exclusive_scan(

using size_type = typename iterator_traits<InputIt>::difference_type;
size_type num_items = static_cast<size_type>(thrust::distance(first, last));
using transformed_iterator_t = transform_input_iterator_t<result_type, InputIt, TransformOp>;
using transformed_iterator_t = transform_iterator<TransformOp, InputIt, result_type, result_type>;

return cuda_cub::exclusive_scan_n(
policy, transformed_iterator_t(first, transform_op), num_items, result, init, scan_op);
Expand Down
24 changes: 16 additions & 8 deletions thrust/thrust/system/cuda/detail/util.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <class ValueType, class InputIt, class UnaryOp>
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<InputIt>::difference_type;
using value_type = ValueType;
using pointer = void;
Expand Down Expand Up @@ -358,10 +359,14 @@ struct transform_input_iterator_t
}
}; // struct transform_input_iterarot_t

// deprecated [Since 2.8]
template <class ValueType, class InputIt1, class InputIt2, class BinaryOp>
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<InputIt1>::difference_type;
using value_type = ValueType;
using pointer = void;
Expand Down Expand Up @@ -488,10 +493,13 @@ struct CCCL_DEPRECATED_BECAUSE("Use cuda::std::identity") identity
}
};

// deprecated [Since 2.8]
template <class T>
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;
Expand Down
Loading