From 599c5fb28f94034afd3e9baa6bfdfb63cd0d4a12 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Tue, 25 Feb 2025 09:11:12 +0100 Subject: [PATCH] Replace direct uses of iterator_traits --- thrust/thrust/advance.h | 4 ++-- .../thrust/detail/allocator/copy_construct_range.inl | 8 ++++---- thrust/thrust/detail/memory_algorithms.h | 12 ++++++------ thrust/thrust/iterator/transform_iterator.h | 2 +- .../thrust/system/cuda/detail/adjacent_difference.h | 2 +- thrust/thrust/system/cuda/detail/copy_if.h | 2 +- thrust/thrust/system/cuda/detail/core/util.h | 2 +- thrust/thrust/system/cuda/detail/count.h | 6 +++--- thrust/thrust/system/cuda/detail/extrema.h | 7 +++---- thrust/thrust/system/cuda/detail/find.h | 2 +- thrust/thrust/system/cuda/detail/for_each.h | 2 +- .../system/cuda/detail/internal/copy_cross_system.h | 2 +- thrust/thrust/system/cuda/detail/merge.h | 4 ++-- thrust/thrust/system/cuda/detail/mismatch.h | 2 +- thrust/thrust/system/cuda/detail/partition.h | 4 ++-- thrust/thrust/system/cuda/detail/reduce.h | 6 +++--- thrust/thrust/system/cuda/detail/reduce_by_key.h | 6 +++--- thrust/thrust/system/cuda/detail/replace.h | 4 ++-- thrust/thrust/system/cuda/detail/set_operations.h | 12 ++++++------ thrust/thrust/system/cuda/detail/sort.h | 2 +- thrust/thrust/system/cuda/detail/swap_ranges.h | 6 +++--- thrust/thrust/system/cuda/detail/tabulate.h | 2 +- thrust/thrust/system/cuda/detail/transform.h | 8 ++++---- thrust/thrust/system/cuda/detail/transform_reduce.h | 2 +- thrust/thrust/system/cuda/detail/transform_scan.h | 6 +++--- .../thrust/system/cuda/detail/uninitialized_copy.h | 4 ++-- .../thrust/system/cuda/detail/uninitialized_fill.h | 2 +- thrust/thrust/system/cuda/detail/unique.h | 8 ++++---- thrust/thrust/system/cuda/detail/unique_by_key.h | 6 +++--- .../system/detail/generic/uninitialized_copy.inl | 12 ++++++------ .../system/detail/generic/uninitialized_fill.inl | 8 ++++---- 31 files changed, 77 insertions(+), 78 deletions(-) diff --git a/thrust/thrust/advance.h b/thrust/thrust/advance.h index 801568b6e00..404e8a31c88 100644 --- a/thrust/thrust/advance.h +++ b/thrust/thrust/advance.h @@ -101,7 +101,7 @@ template _CCCL_HOST_DEVICE InputIterator next( InputIterator i -, typename iterator_traits::difference_type n = 1 +, typename ::cuda::std::iterator_traits::difference_type n = 1 ); #endif @@ -135,7 +135,7 @@ template _CCCL_HOST_DEVICE BidirectionalIterator prev( BidirectionalIterator i -, typename iterator_traits::difference_type n = 1 +, typename ::cuda::std::iterator_traits::difference_type n = 1 ); #endif diff --git a/thrust/thrust/detail/allocator/copy_construct_range.inl b/thrust/thrust/detail/allocator/copy_construct_range.inl index 9ebb74eaabf..c02f373b16e 100644 --- a/thrust/thrust/detail/allocator/copy_construct_range.inl +++ b/thrust/thrust/detail/allocator/copy_construct_range.inl @@ -104,8 +104,8 @@ _CCCL_HOST_DEVICE enable_if_convertible_t uniniti thrust::advance(end, n); // create a functor - using InputType = typename iterator_traits::value_type; - using OutputType = typename iterator_traits::value_type; + using InputType = it_value_t; + using OutputType = it_value_t; // do the for_each // note we use to_system to dispatch the for_each @@ -135,8 +135,8 @@ _CCCL_HOST_DEVICE enable_if_convertible_t uniniti ZipIterator begin = thrust::make_zip_iterator(thrust::make_tuple(first, result)); // create a functor - using InputType = typename iterator_traits::value_type; - using OutputType = typename iterator_traits::value_type; + using InputType = it_value_t; + using OutputType = it_value_t; // do the for_each_n // note we use to_system to dispatch the for_each_n diff --git a/thrust/thrust/detail/memory_algorithms.h b/thrust/thrust/detail/memory_algorithms.h index 1da74c4b56d..7ae020c52c3 100644 --- a/thrust/thrust/detail/memory_algorithms.h +++ b/thrust/thrust/detail/memory_algorithms.h @@ -64,7 +64,7 @@ _CCCL_HOST_DEVICE ForwardIt destroy(ForwardIt first, ForwardIt last) noexcept template _CCCL_HOST_DEVICE ForwardIt destroy(Allocator const& alloc, ForwardIt first, ForwardIt last) noexcept { - using T = typename iterator_traits::value_type; + using T = detail::it_value_t; using traits = typename detail::allocator_traits< ::cuda::std::remove_cv_t<::cuda::std::remove_reference_t>>::template rebind_traits::other; @@ -92,7 +92,7 @@ _CCCL_HOST_DEVICE ForwardIt destroy_n(ForwardIt first, Size n) noexcept template _CCCL_HOST_DEVICE ForwardIt destroy_n(Allocator const& alloc, ForwardIt first, Size n) noexcept { - using T = typename iterator_traits::value_type; + using T = detail::it_value_t; using traits = typename detail::allocator_traits< ::cuda::std::remove_cv_t<::cuda::std::remove_reference_t>>::template rebind_traits::other; @@ -109,7 +109,7 @@ _CCCL_HOST_DEVICE ForwardIt destroy_n(Allocator const& alloc, ForwardIt first, S template _CCCL_HOST_DEVICE void uninitialized_construct(ForwardIt first, ForwardIt last, Args const&... args) { - using T = typename iterator_traits::value_type; + using T = detail::it_value_t; ForwardIt current = first; @@ -132,7 +132,7 @@ _CCCL_HOST_DEVICE void uninitialized_construct(ForwardIt first, ForwardIt last, template void uninitialized_construct_with_allocator(Allocator const& alloc, ForwardIt first, ForwardIt last, Args const&... args) { - using T = typename iterator_traits::value_type; + using T = detail::it_value_t; using traits = typename detail::allocator_traits< typename std::remove_cv::type>::type>::template rebind_traits; @@ -159,7 +159,7 @@ void uninitialized_construct_with_allocator(Allocator const& alloc, ForwardIt fi template void uninitialized_construct_n(ForwardIt first, Size n, Args const&... args) { - using T = typename iterator_traits::value_type; + using T = detail::it_value_t; ForwardIt current = first; @@ -182,7 +182,7 @@ void uninitialized_construct_n(ForwardIt first, Size n, Args const&... args) template void uninitialized_construct_n_with_allocator(Allocator const& alloc, ForwardIt first, Size n, Args const&... args) { - using T = typename iterator_traits::value_type; + using T = detail::it_value_t; using traits = typename detail::allocator_traits< typename std::remove_cv::type>::type>::template rebind_traits; diff --git a/thrust/thrust/iterator/transform_iterator.h b/thrust/thrust/iterator/transform_iterator.h index ee55333d103..7783d39bfa6 100644 --- a/thrust/thrust/iterator/transform_iterator.h +++ b/thrust/thrust/iterator/transform_iterator.h @@ -94,7 +94,7 @@ struct make_transform_iterator_base Iterator, value_type, use_default, - typename iterator_traits::iterator_category, + typename ::cuda::std::iterator_traits::iterator_category, reference>; }; diff --git a/thrust/thrust/system/cuda/detail/adjacent_difference.h b/thrust/thrust/system/cuda/detail/adjacent_difference.h index 8459e4e729c..13061c5b86a 100644 --- a/thrust/thrust/system/cuda/detail/adjacent_difference.h +++ b/thrust/thrust/system/cuda/detail/adjacent_difference.h @@ -207,7 +207,7 @@ template OutputIt _CCCL_HOST_DEVICE adjacent_difference(execution_policy& policy, InputIt first, InputIt last, OutputIt result) { - using input_type = typename iterator_traits::value_type; + using input_type = thrust::detail::it_value_t; return cuda_cub::adjacent_difference(policy, first, last, result, minus()); } diff --git a/thrust/thrust/system/cuda/detail/copy_if.h b/thrust/thrust/system/cuda/detail/copy_if.h index 33ba756eafa..778725cb89d 100644 --- a/thrust/thrust/system/cuda/detail/copy_if.h +++ b/thrust/thrust/system/cuda/detail/copy_if.h @@ -188,7 +188,7 @@ THRUST_RUNTIME_FUNCTION OutputIt copy_if( OutputIt output, Predicate predicate) { - using size_type = typename iterator_traits::difference_type; + using size_type = iterator_difference_t; size_type num_items = static_cast(thrust::distance(first, last)); cudaError_t status = cudaSuccess; diff --git a/thrust/thrust/system/cuda/detail/core/util.h b/thrust/thrust/system/cuda/detail/core/util.h index a5f87026d7d..c860fc56bb1 100644 --- a/thrust/thrust/system/cuda/detail/core/util.h +++ b/thrust/thrust/system/cuda/detail/core/util.h @@ -485,7 +485,7 @@ struct get_arch> // BlockLoad // ----------- // a helper metaprogram that returns type of a block loader -template ::value_type> +template > struct BlockLoad { using type = cub::BlockLoad; diff --git a/thrust/thrust/system/cuda/detail/count.h b/thrust/thrust/system/cuda/detail/count.h index 530682ba717..22b373e18c0 100644 --- a/thrust/thrust/system/cuda/detail/count.h +++ b/thrust/thrust/system/cuda/detail/count.h @@ -49,10 +49,10 @@ namespace cuda_cub { template -typename iterator_traits::difference_type _CCCL_HOST_DEVICE +thrust::detail::it_difference_t _CCCL_HOST_DEVICE count_if(execution_policy& policy, InputIt first, InputIt last, UnaryPred unary_pred) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; using flag_iterator_t = transform_iterator; return cuda_cub::reduce_n( @@ -60,7 +60,7 @@ count_if(execution_policy& policy, InputIt first, InputIt last, UnaryPr } template -typename iterator_traits::difference_type _CCCL_HOST_DEVICE +thrust::detail::it_difference_t _CCCL_HOST_DEVICE count(execution_policy& policy, InputIt first, InputIt last, Value const& value) { return cuda_cub::count_if(policy, first, last, thrust::detail::equal_to_value(value)); diff --git a/thrust/thrust/system/cuda/detail/extrema.h b/thrust/thrust/system/cuda/detail/extrema.h index fec5244be00..eb9a645f00f 100644 --- a/thrust/thrust/system/cuda/detail/extrema.h +++ b/thrust/thrust/system/cuda/detail/extrema.h @@ -369,8 +369,8 @@ element(execution_policy& policy, ItemsIt first, ItemsIt last, BinaryPr return last; } - using InputType = typename iterator_traits::value_type; - using IndexType = typename iterator_traits::difference_type; + using InputType = thrust::detail::it_value_t; + using IndexType = thrust::detail::it_difference_t; IndexType num_items = static_cast(thrust::distance(first, last)); @@ -442,8 +442,7 @@ minmax_element(execution_policy& policy, ItemsIt first, ItemsIt last, B } THRUST_CDP_DISPATCH( - (using InputType = typename iterator_traits::value_type; - using IndexType = typename iterator_traits::difference_type; + (using InputType = thrust::detail::it_value_t; using IndexType = thrust::detail::it_difference_t; const auto num_items = static_cast(thrust::distance(first, last)); diff --git a/thrust/thrust/system/cuda/detail/find.h b/thrust/thrust/system/cuda/detail/find.h index fd35aa930ba..1471df050b3 100644 --- a/thrust/thrust/system/cuda/detail/find.h +++ b/thrust/thrust/system/cuda/detail/find.h @@ -97,7 +97,7 @@ template struct transform_input_iterator_t { using self_t = transform_input_iterator_t; - using difference_type = typename iterator_traits::difference_type; + using difference_type = iterator_difference_t; using value_type = ValueType; using pointer = void; using reference = value_type; diff --git a/thrust/thrust/system/cuda/detail/for_each.h b/thrust/thrust/system/cuda/detail/for_each.h index d0a3a65032a..ef84f31c99e 100644 --- a/thrust/thrust/system/cuda/detail/for_each.h +++ b/thrust/thrust/system/cuda/detail/for_each.h @@ -72,7 +72,7 @@ Input THRUST_FUNCTION for_each_n(execution_policy& policy, Input first, template Input THRUST_FUNCTION for_each(execution_policy& policy, Input first, Input last, UnaryOp op) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type count = static_cast(thrust::distance(first, last)); return THRUST_NS_QUALIFIER::cuda_cub::for_each_n(policy, first, count, op); diff --git a/thrust/thrust/system/cuda/detail/internal/copy_cross_system.h b/thrust/thrust/system/cuda/detail/internal/copy_cross_system.h index 2ecd9d02a50..c2249bcc85e 100644 --- a/thrust/thrust/system/cuda/detail/internal/copy_cross_system.h +++ b/thrust/thrust/system/cuda/detail/internal/copy_cross_system.h @@ -85,7 +85,7 @@ OutputIt _CCCL_HOST cross_system_copy_n( thrust::detail::true_type) // trivial copy { - using InputTy = typename iterator_traits::value_type; + using InputTy = thrust::detail::it_value_t; if (n > 0) { trivial_device_copy( diff --git a/thrust/thrust/system/cuda/detail/merge.h b/thrust/thrust/system/cuda/detail/merge.h index 0c39a15f669..c6b480e6c88 100644 --- a/thrust/thrust/system/cuda/detail/merge.h +++ b/thrust/thrust/system/cuda/detail/merge.h @@ -64,7 +64,7 @@ merge(execution_policy& policy, CompareOp compare_op = {}) { THRUST_CDP_DISPATCH( - (using size_type = typename iterator_traits::difference_type; + (using size_type = thrust::detail::it_difference_t; const auto num_keys1 = static_cast(thrust::distance(keys1_begin, keys1_end)); const auto num_keys2 = static_cast(thrust::distance(keys2_begin, keys2_end)); const auto num_keys_out = num_keys1 + num_keys2; @@ -151,7 +151,7 @@ pair _CCCL_HOST_DEVICE merge_by_key( CompareOp compare_op = {}) { THRUST_CDP_DISPATCH( - (using size_type = typename iterator_traits::difference_type; + (using size_type = thrust::detail::it_difference_t; const auto num_keys1 = static_cast(thrust::distance(keys1_begin, keys1_end)); const auto num_keys2 = static_cast(thrust::distance(keys2_begin, keys2_end)); diff --git a/thrust/thrust/system/cuda/detail/mismatch.h b/thrust/thrust/system/cuda/detail/mismatch.h index 8953e54eeed..6e723bf4d75 100644 --- a/thrust/thrust/system/cuda/detail/mismatch.h +++ b/thrust/thrust/system/cuda/detail/mismatch.h @@ -69,7 +69,7 @@ template struct transform_pair_of_input_iterators_t { using self_t = transform_pair_of_input_iterators_t; - using difference_type = typename iterator_traits::difference_type; + using difference_type = thrust::detail::it_difference_t; using value_type = ValueType; using pointer = void; using reference = value_type; diff --git a/thrust/thrust/system/cuda/detail/partition.h b/thrust/thrust/system/cuda/detail/partition.h index 690e1e4dd6d..d016dae268b 100644 --- a/thrust/thrust/system/cuda/detail/partition.h +++ b/thrust/thrust/system/cuda/detail/partition.h @@ -166,7 +166,7 @@ THRUST_RUNTIME_FUNCTION std::size_t partition( OutputIt output, Predicate predicate) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type num_items = thrust::distance(first, last); std::size_t num_selected{}; @@ -241,7 +241,7 @@ THRUST_RUNTIME_FUNCTION InputIt inplace_partition( } // Element type of the input iterator - using value_t = typename iterator_traits::value_type; + using value_t = thrust::detail::it_value_t; std::size_t num_items = static_cast(thrust::distance(first, last)); // Allocate temporary storage, which will serve as the input to the partition diff --git a/thrust/thrust/system/cuda/detail/reduce.h b/thrust/thrust/system/cuda/detail/reduce.h index 7724d91852c..e5014106045 100644 --- a/thrust/thrust/system/cuda/detail/reduce.h +++ b/thrust/thrust/system/cuda/detail/reduce.h @@ -856,7 +856,7 @@ reduce_n(execution_policy& policy, InputIt first, Size num_items, T ini template _CCCL_HOST_DEVICE T reduce(execution_policy& policy, InputIt first, InputIt last, T init, BinaryOp binary_op) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; // FIXME: Check for RA iterator. size_type num_items = static_cast(thrust::distance(first, last)); return cuda_cub::reduce_n(policy, first, num_items, init, binary_op); @@ -869,10 +869,10 @@ _CCCL_HOST_DEVICE T reduce(execution_policy& policy, InputIt first, Inp } template -_CCCL_HOST_DEVICE typename iterator_traits::value_type +_CCCL_HOST_DEVICE thrust::detail::it_value_t reduce(execution_policy& policy, InputIt first, InputIt last) { - using value_type = typename iterator_traits::value_type; + using value_type = thrust::detail::it_value_t; return cuda_cub::reduce(policy, first, last, value_type(0)); } diff --git a/thrust/thrust/system/cuda/detail/reduce_by_key.h b/thrust/thrust/system/cuda/detail/reduce_by_key.h index 05585bebd80..17f01006851 100644 --- a/thrust/thrust/system/cuda/detail/reduce_by_key.h +++ b/thrust/thrust/system/cuda/detail/reduce_by_key.h @@ -150,8 +150,8 @@ template struct ReduceByKeyAgent { - using key_type = typename iterator_traits::value_type; - using value_type = typename iterator_traits::value_type; + using key_type = thrust::detail::it_value_t; + using value_type = thrust::detail::it_value_t; using size_type = Size; using size_value_pair_t = cub::KeyValuePair; @@ -893,7 +893,7 @@ THRUST_RUNTIME_FUNCTION pair reduce_by_key( EqualityOp equality_op, ReductionOp reduction_op) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type num_items = thrust::distance(keys_first, keys_last); diff --git a/thrust/thrust/system/cuda/detail/replace.h b/thrust/thrust/system/cuda/detail/replace.h index 45682ac52c2..12257f2f224 100644 --- a/thrust/thrust/system/cuda/detail/replace.h +++ b/thrust/thrust/system/cuda/detail/replace.h @@ -127,7 +127,7 @@ OutputIt _CCCL_HOST_DEVICE replace_copy_if( Predicate predicate, T const& new_value) { - using output_type = typename iterator_traits::value_type; + using output_type = thrust::detail::it_value_t; using new_value_if_t = __replace::new_value_if_f; return cuda_cub::transform(policy, first, last, result, new_value_if_t(predicate, new_value)); } @@ -142,7 +142,7 @@ OutputIt _CCCL_HOST_DEVICE replace_copy_if( Predicate predicate, T const& new_value) { - using output_type = typename iterator_traits::value_type; + using output_type = thrust::detail::it_value_t; using new_value_if_t = __replace::new_value_if_f; return cuda_cub::transform(policy, first, last, stencil, result, new_value_if_t(predicate, new_value)); } diff --git a/thrust/thrust/system/cuda/detail/set_operations.h b/thrust/thrust/system/cuda/detail/set_operations.h index 0ada139bd35..28c54146190 100644 --- a/thrust/thrust/system/cuda/detail/set_operations.h +++ b/thrust/thrust/system/cuda/detail/set_operations.h @@ -155,7 +155,7 @@ template THRUST_DEVICE_FUNCTION pair balanced_path(It1 keys1, It2 keys2, Size num_keys1, Size num_keys2, Size diag, Size2 levels, CompareOp compare_op) { - using T = typename iterator_traits::value_type; + using T = thrust::detail::it_value_t; Size index1 = merge_path(keys1, num_keys1, keys2, num_keys2, diag, compare_op); Size index2 = diag - index1; @@ -275,10 +275,10 @@ template struct SetOpAgent { - using key1_type = typename iterator_traits::value_type; - using key2_type = typename iterator_traits::value_type; - using value1_type = typename iterator_traits::value_type; - using value2_type = typename iterator_traits::value_type; + using key1_type = thrust::detail::it_value_t; + using key2_type = thrust::detail::it_value_t; + using value1_type = thrust::detail::it_value_t; + using value2_type = thrust::detail::it_value_t; using key_type = key1_type; using value_type = value1_type; @@ -1149,7 +1149,7 @@ THRUST_RUNTIME_FUNCTION pair set_operations( CompareOp compare_op, SetOp set_op) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type num_keys1 = static_cast(thrust::distance(keys1_first, keys1_last)); size_type num_keys2 = static_cast(thrust::distance(keys2_first, keys2_last)); diff --git a/thrust/thrust/system/cuda/detail/sort.h b/thrust/thrust/system/cuda/detail/sort.h index 2b04ee779f7..a32e5bd0b75 100644 --- a/thrust/thrust/system/cuda/detail/sort.h +++ b/thrust/thrust/system/cuda/detail/sort.h @@ -160,7 +160,7 @@ THRUST_RUNTIME_FUNCTION void merge_sort( execution_policy& policy, KeysIt keys_first, KeysIt keys_last, ItemsIt items_first, CompareOp compare_op) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type count = static_cast(thrust::distance(keys_first, keys_last)); diff --git a/thrust/thrust/system/cuda/detail/swap_ranges.h b/thrust/thrust/system/cuda/detail/swap_ranges.h index 73a73d63775..23cd499ad2a 100644 --- a/thrust/thrust/system/cuda/detail/swap_ranges.h +++ b/thrust/thrust/system/cuda/detail/swap_ranges.h @@ -61,8 +61,8 @@ struct swap_f ItemsIt1 items1; ItemsIt2 items2; - using value1_type = typename iterator_traits::value_type; - using value2_type = typename iterator_traits::value_type; + using value1_type = thrust::detail::it_value_t; + using value2_type = thrust::detail::it_value_t; THRUST_FUNCTION swap_f(ItemsIt1 items1_, ItemsIt2 items2_) @@ -88,7 +88,7 @@ template ItemsIt2 _CCCL_HOST_DEVICE swap_ranges(execution_policy& policy, ItemsIt1 first1, ItemsIt1 last1, ItemsIt2 first2) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type num_items = static_cast(thrust::distance(first1, last1)); diff --git a/thrust/thrust/system/cuda/detail/tabulate.h b/thrust/thrust/system/cuda/detail/tabulate.h index 6caec210120..dee20887a77 100644 --- a/thrust/thrust/system/cuda/detail/tabulate.h +++ b/thrust/thrust/system/cuda/detail/tabulate.h @@ -72,7 +72,7 @@ struct functor template void _CCCL_HOST_DEVICE tabulate(execution_policy& policy, Iterator first, Iterator last, TabulateOp tabulate_op) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type count = thrust::distance(first, last); diff --git a/thrust/thrust/system/cuda/detail/transform.h b/thrust/thrust/system/cuda/detail/transform.h index 2baf4810550..6863fba4490 100644 --- a/thrust/thrust/system/cuda/detail/transform.h +++ b/thrust/thrust/system/cuda/detail/transform.h @@ -319,7 +319,7 @@ OutputIt THRUST_FUNCTION transform_if( TransformOp transform_op, Predicate predicate) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type num_items = static_cast(thrust::distance(first, last)); return __transform::unary(policy, first, result, num_items, stencil, transform_op, predicate); } // func transform_if @@ -341,7 +341,7 @@ OutputIt THRUST_FUNCTION transform(execution_policy& policy, InputIt first, InputIt last, OutputIt result, TransformOp transform_op) { THRUST_CDP_DISPATCH( - (using size_type = typename iterator_traits::difference_type; + (using size_type = thrust::detail::it_difference_t; const auto num_items = static_cast(thrust::distance(first, last)); return __transform::cub_transform_many(policy, ::cuda::std::make_tuple(first), result, num_items, transform_op);), (while (first != last) { @@ -372,7 +372,7 @@ OutputIt THRUST_FUNCTION transform_if( TransformOp transform_op, Predicate predicate) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type num_items = static_cast(thrust::distance(first1, last1)); return __transform::binary(policy, first1, first2, result, num_items, stencil, transform_op, predicate); } // func transform_if @@ -387,7 +387,7 @@ OutputIt THRUST_FUNCTION transform( TransformOp transform_op) { THRUST_CDP_DISPATCH( - (using size_type = typename iterator_traits::difference_type; + (using size_type = thrust::detail::it_difference_t; const auto num_items = static_cast(thrust::distance(first1, last1)); return __transform::cub_transform_many( policy, ::cuda::std::make_tuple(first1, first2), result, num_items, transform_op);), diff --git a/thrust/thrust/system/cuda/detail/transform_reduce.h b/thrust/thrust/system/cuda/detail/transform_reduce.h index 95e7e2460fa..ed7d26942f7 100644 --- a/thrust/thrust/system/cuda/detail/transform_reduce.h +++ b/thrust/thrust/system/cuda/detail/transform_reduce.h @@ -129,7 +129,7 @@ template & policy, InputIt first, InputIt last, TransformOp transform_op, T init, ReduceOp reduce_op) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; const size_type num_items = static_cast(thrust::distance(first, last)); THRUST_CDP_DISPATCH( diff --git a/thrust/thrust/system/cuda/detail/transform_scan.h b/thrust/thrust/system/cuda/detail/transform_scan.h index 2791cf0c1d1..335e4ea97c3 100644 --- a/thrust/thrust/system/cuda/detail/transform_scan.h +++ b/thrust/thrust/system/cuda/detail/transform_scan.h @@ -65,7 +65,7 @@ OutputIt _CCCL_HOST_DEVICE transform_inclusive_scan( using result_type = thrust::detail::invoke_result_t; using value_type = ::cuda::std::remove_cvref_t; - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type num_items = static_cast(thrust::distance(first, last)); using transformed_iterator_t = transform_iterator; @@ -86,7 +86,7 @@ OutputIt _CCCL_HOST_DEVICE transform_inclusive_scan( using result_type = thrust::detail::invoke_result_t; using value_type = ::cuda::std::remove_cvref_t; - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type num_items = static_cast(thrust::distance(first, last)); using transformed_iterator_t = transform_iterator; @@ -107,7 +107,7 @@ OutputIt _CCCL_HOST_DEVICE transform_exclusive_scan( // Use the initial value type per https://wg21.link/P0571 using result_type = ::cuda::std::remove_cvref_t; - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type num_items = static_cast(thrust::distance(first, last)); using transformed_iterator_t = transform_iterator; diff --git a/thrust/thrust/system/cuda/detail/uninitialized_copy.h b/thrust/thrust/system/cuda/detail/uninitialized_copy.h index d24cf87e6ed..5be6813ebf7 100644 --- a/thrust/thrust/system/cuda/detail/uninitialized_copy.h +++ b/thrust/thrust/system/cuda/detail/uninitialized_copy.h @@ -58,8 +58,8 @@ struct functor InputIt input; OutputIt output; - using InputType = typename iterator_traits::value_type; - using OutputType = typename iterator_traits::value_type; + using InputType = thrust::detail::it_value_t; + using OutputType = thrust::detail::it_value_t; THRUST_FUNCTION functor(InputIt input_, OutputIt output_) diff --git a/thrust/thrust/system/cuda/detail/uninitialized_fill.h b/thrust/thrust/system/cuda/detail/uninitialized_fill.h index 233f600e2bb..601e0b9f0b1 100644 --- a/thrust/thrust/system/cuda/detail/uninitialized_fill.h +++ b/thrust/thrust/system/cuda/detail/uninitialized_fill.h @@ -58,7 +58,7 @@ struct functor Iterator items; T value; - using value_type = typename iterator_traits::value_type; + using value_type = thrust::detail::it_value_t; THRUST_FUNCTION functor(Iterator items_, T const& value_) diff --git a/thrust/thrust/system/cuda/detail/unique.h b/thrust/thrust/system/cuda/detail/unique.h index 770f7dd631a..0d95ffda5fb 100644 --- a/thrust/thrust/system/cuda/detail/unique.h +++ b/thrust/thrust/system/cuda/detail/unique.h @@ -140,7 +140,7 @@ struct Tuning template struct UniqueAgent { - using item_type = typename iterator_traits::value_type; + using item_type = thrust::detail::it_value_t; using ScanTileState = cub::ScanTileState; @@ -524,7 +524,7 @@ THRUST_RUNTIME_FUNCTION ItemsOutputIt unique( ItemsOutputIt items_result, BinaryPred binary_pred) { - // using size_type = typename iterator_traits::difference_type; + // using size_type = thrust::detail::it_difference_t; using size_type = int; size_type num_items = static_cast(thrust::distance(items_first, items_last)); @@ -590,7 +590,7 @@ unique_copy(execution_policy& policy, InputIt first, InputIt last, Outp template OutputIt _CCCL_HOST_DEVICE unique_copy(execution_policy& policy, InputIt first, InputIt last, OutputIt result) { - using input_type = typename iterator_traits::value_type; + using input_type = thrust::detail::it_value_t; return cuda_cub::unique_copy(policy, first, last, result, equal_to()); } @@ -608,7 +608,7 @@ unique(execution_policy& policy, ForwardIt first, ForwardIt last, Binar template ForwardIt _CCCL_HOST_DEVICE unique(execution_policy& policy, ForwardIt first, ForwardIt last) { - using input_type = typename iterator_traits::value_type; + using input_type = thrust::detail::it_value_t; return cuda_cub::unique(policy, first, last, equal_to()); } diff --git a/thrust/thrust/system/cuda/detail/unique_by_key.h b/thrust/thrust/system/cuda/detail/unique_by_key.h index 468c1bdae6b..fa76ff9ca77 100644 --- a/thrust/thrust/system/cuda/detail/unique_by_key.h +++ b/thrust/thrust/system/cuda/detail/unique_by_key.h @@ -182,7 +182,7 @@ THRUST_RUNTIME_FUNCTION pair unique_by_key( ValOutputIt values_result, BinaryPred binary_pred) { - using size_type = typename iterator_traits::difference_type; + using size_type = thrust::detail::it_difference_t; size_type num_items = static_cast(thrust::distance(keys_first, keys_last)); pair result_end{}; @@ -273,7 +273,7 @@ pair _CCCL_HOST_DEVICE unique_by_key_copy( KeyOutputIt keys_result, ValOutputIt values_result) { - using key_type = typename iterator_traits::value_type; + using key_type = thrust::detail::it_value_t; return cuda_cub::unique_by_key_copy( policy, keys_first, keys_last, values_first, keys_result, values_result, equal_to()); } @@ -298,7 +298,7 @@ template pair _CCCL_HOST_DEVICE unique_by_key(execution_policy& policy, KeyInputIt keys_first, KeyInputIt keys_last, ValInputIt values_first) { - using key_type = typename iterator_traits::value_type; + using key_type = thrust::detail::it_value_t; return cuda_cub::unique_by_key(policy, keys_first, keys_last, values_first, equal_to()); } diff --git a/thrust/thrust/system/detail/generic/uninitialized_copy.inl b/thrust/thrust/system/detail/generic/uninitialized_copy.inl index 470ba5147de..377f529168c 100644 --- a/thrust/thrust/system/detail/generic/uninitialized_copy.inl +++ b/thrust/thrust/system/detail/generic/uninitialized_copy.inl @@ -76,8 +76,8 @@ _CCCL_HOST_DEVICE ForwardIterator uninitialized_copy( thrust::advance(end, n); // create a functor - using InputType = typename iterator_traits::value_type; - using OutputType = typename iterator_traits::value_type; + using InputType = thrust::detail::it_value_t; + using OutputType = thrust::detail::it_value_t; detail::uninitialized_copy_functor f; @@ -116,8 +116,8 @@ _CCCL_HOST_DEVICE ForwardIterator uninitialized_copy_n( ZipIterator zipped_first = thrust::make_zip_iterator(thrust::make_tuple(first, result)); // create a functor - using InputType = typename iterator_traits::value_type; - using OutputType = typename iterator_traits::value_type; + using InputType = thrust::detail::it_value_t; + using OutputType = thrust::detail::it_value_t; detail::uninitialized_copy_functor f; @@ -146,7 +146,7 @@ template & exec, InputIterator first, InputIterator last, ForwardIterator result) { - using ResultType = typename iterator_traits::value_type; + using ResultType = thrust::detail::it_value_t; using ResultTypeHasTrivialCopyConstructor = typename ::cuda::std::is_trivially_copy_constructible::type; @@ -158,7 +158,7 @@ template & exec, InputIterator first, Size n, ForwardIterator result) { - using ResultType = typename iterator_traits::value_type; + using ResultType = thrust::detail::it_value_t; using ResultTypeHasTrivialCopyConstructor = typename ::cuda::std::is_trivially_copy_constructible::type; diff --git a/thrust/thrust/system/detail/generic/uninitialized_fill.inl b/thrust/thrust/system/detail/generic/uninitialized_fill.inl index e8717c23966..02871da5dd6 100644 --- a/thrust/thrust/system/detail/generic/uninitialized_fill.inl +++ b/thrust/thrust/system/detail/generic/uninitialized_fill.inl @@ -60,7 +60,7 @@ _CCCL_HOST_DEVICE void uninitialized_fill( const T& x, thrust::detail::false_type) // ::cuda::std::is_trivially_copy_constructible { - using ValueType = typename iterator_traits::value_type; + using ValueType = thrust::detail::it_value_t; thrust::for_each(exec, first, last, thrust::detail::uninitialized_fill_functor(x)); } // end uninitialized_fill() @@ -84,7 +84,7 @@ _CCCL_HOST_DEVICE ForwardIterator uninitialized_fill_n( const T& x, thrust::detail::false_type) // ::cuda::std::is_trivially_copy_constructible { - using ValueType = typename iterator_traits::value_type; + using ValueType = thrust::detail::it_value_t; return thrust::for_each_n(exec, first, n, thrust::detail::uninitialized_fill_functor(x)); } // end uninitialized_fill() @@ -95,7 +95,7 @@ template _CCCL_HOST_DEVICE void uninitialized_fill( thrust::execution_policy& exec, ForwardIterator first, ForwardIterator last, const T& x) { - using ValueType = typename iterator_traits::value_type; + using ValueType = thrust::detail::it_value_t; using ValueTypeHasTrivialCopyConstructor = ::cuda::std::is_trivially_copy_constructible; @@ -107,7 +107,7 @@ template & exec, ForwardIterator first, Size n, const T& x) { - using ValueType = typename iterator_traits::value_type; + using ValueType = thrust::detail::it_value_t; using ValueTypeHasTrivialCopyConstructor = ::cuda::std::is_trivially_copy_constructible;