From c339a526fa8f97331ce3922af7cf75c87c43e0d7 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 13 Jan 2025 17:47:24 +0100 Subject: [PATCH] Use cuda::std::min/max in Thrust (#3364) --- cub/test/catch2_large_array_sort_helper.cuh | 2 +- thrust/examples/set_operations.cu | 1 + .../async/exclusive_scan/large_indices.cu | 6 +- .../async/inclusive_scan/large_indices.cu | 6 +- thrust/testing/cuda/is_partitioned.cu | 2 +- thrust/testing/min_and_max.cu | 56 +++---- thrust/testing/scan.cu | 30 ++-- thrust/thrust/detail/minmax.h | 55 ------- thrust/thrust/detail/vector_base.inl | 29 ++-- thrust/thrust/extrema.h | 153 +----------------- .../system/cuda/detail/adjacent_difference.h | 1 - thrust/thrust/system/cuda/detail/find.h | 5 +- .../thrust/system/cuda/detail/inner_product.h | 1 - thrust/thrust/system/cuda/detail/reduce.h | 5 +- .../thrust/system/cuda/detail/reduce_by_key.h | 1 - .../thrust/system/cuda/detail/scan_by_key.h | 1 - .../system/cuda/detail/set_operations.h | 7 +- .../system/cuda/detail/transform_reduce.h | 1 - thrust/thrust/system/cuda/detail/unique.h | 3 +- .../thrust/system/cuda/detail/unique_by_key.h | 1 - thrust/thrust/system/detail/generic/find.inl | 7 +- .../detail/sequential/stable_merge_sort.inl | 15 +- .../system/tbb/detail/reduce_by_key.inl | 11 +- .../system/tbb/detail/reduce_intervals.h | 5 +- 24 files changed, 96 insertions(+), 308 deletions(-) delete mode 100644 thrust/thrust/detail/minmax.h diff --git a/cub/test/catch2_large_array_sort_helper.cuh b/cub/test/catch2_large_array_sort_helper.cuh index 6c0ed2a48be..e5c095c93dd 100644 --- a/cub/test/catch2_large_array_sort_helper.cuh +++ b/cub/test/catch2_large_array_sort_helper.cuh @@ -140,7 +140,7 @@ public: _CCCL_HOST_DEVICE KeyType operator()(std::size_t idx) const { // The final summary may be padded, so truncate the summary_idx at the last valid idx: - const std::size_t summary_idx = thrust::min(m_num_summaries - 1, idx / m_unpadded_run_size); + const std::size_t summary_idx = cuda::std::min(m_num_summaries - 1, idx / m_unpadded_run_size); const KeyType key = m_is_descending ? static_cast((m_num_summaries - 1 - summary_idx) * m_key_conversion) : static_cast(summary_idx * m_key_conversion); diff --git a/thrust/examples/set_operations.cu b/thrust/examples/set_operations.cu index 2303fb9b69c..8a99b7ca14c 100644 --- a/thrust/examples/set_operations.cu +++ b/thrust/examples/set_operations.cu @@ -1,4 +1,5 @@ #include +#include #include #include #include diff --git a/thrust/testing/async/exclusive_scan/large_indices.cu b/thrust/testing/async/exclusive_scan/large_indices.cu index d03b7fa5662..d3b7e6fabab 100644 --- a/thrust/testing/async/exclusive_scan/large_indices.cu +++ b/thrust/testing/async/exclusive_scan/large_indices.cu @@ -196,17 +196,17 @@ namespace { //------------------------------------------------------------------------------ -// Generate the output sequence using counting iterators and thrust::max<> for +// Generate the output sequence using counting iterators and ::cuda::maximum<> for // custom operator overloads. struct custom_bin_op_overloads { using postfix_args_type = std::tuple< // List any extra arg overloads: - std::tuple> // - initial_value, binop + std::tuple> // - initial_value, binop >; static postfix_args_type generate_postfix_args() { - return postfix_args_type{std::make_tuple(0, thrust::maximum<>{})}; + return postfix_args_type{std::make_tuple(0, ::cuda::maximum<>{})}; } }; diff --git a/thrust/testing/async/inclusive_scan/large_indices.cu b/thrust/testing/async/inclusive_scan/large_indices.cu index 1eb292cc995..aacb19c88ef 100644 --- a/thrust/testing/async/inclusive_scan/large_indices.cu +++ b/thrust/testing/async/inclusive_scan/large_indices.cu @@ -191,17 +191,17 @@ namespace { //------------------------------------------------------------------------------ -// Generate the output sequence using counting iterators and thrust::max<> for +// Generate the output sequence using counting iterators and ::cuda::maximum<> for // custom operator overloads. struct custom_bin_op_overloads { using postfix_args_type = std::tuple< // List any extra arg overloads: - std::tuple> // - custom binary op + std::tuple<::cuda::maximum<>> // - custom binary op >; static postfix_args_type generate_postfix_args() { - return postfix_args_type{std::make_tuple(thrust::maximum<>{})}; + return postfix_args_type{std::make_tuple(::cuda::maximum<>{})}; } }; diff --git a/thrust/testing/cuda/is_partitioned.cu b/thrust/testing/cuda/is_partitioned.cu index 1e02ca38c28..35f59b13ffa 100644 --- a/thrust/testing/cuda/is_partitioned.cu +++ b/thrust/testing/cuda/is_partitioned.cu @@ -26,7 +26,7 @@ void TestIsPartitionedDevice(ExecutionPolicy exec) { size_t n = 1000; - n = thrust::max(n, 2); + n = ::cuda::std::max(n, 2); thrust::device_vector v = unittest::random_integers(n); diff --git a/thrust/testing/min_and_max.cu b/thrust/testing/min_and_max.cu index 7b01fe37450..23355a44157 100644 --- a/thrust/testing/min_and_max.cu +++ b/thrust/testing/min_and_max.cu @@ -9,32 +9,28 @@ struct TestMin { // 2 < 3 T two(2), three(3); - ASSERT_EQUAL(two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two, three)); - ASSERT_EQUAL(two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two, three, thrust::less())); + ASSERT_EQUAL(two, ::cuda::std::min(two, three)); + ASSERT_EQUAL(two, ::cuda::std::min(two, three, thrust::less())); - ASSERT_EQUAL(two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(three, two)); - ASSERT_EQUAL(two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(three, two, thrust::less())); + ASSERT_EQUAL(two, ::cuda::std::min(three, two)); + ASSERT_EQUAL(two, ::cuda::std::min(three, two, thrust::less())); - ASSERT_EQUAL(three, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two, three, thrust::greater())); - ASSERT_EQUAL(three, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(three, two, thrust::greater())); + ASSERT_EQUAL(three, ::cuda::std::min(two, three, thrust::greater())); + ASSERT_EQUAL(three, ::cuda::std::min(three, two, thrust::greater())); using KV = key_value; KV two_and_two(two, two); KV two_and_three(two, three); // the first element breaks ties - ASSERT_EQUAL_QUIET(two_and_two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three)); - ASSERT_EQUAL_QUIET(two_and_three, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two)); + ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::min(two_and_two, two_and_three)); + ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::min(two_and_three, two_and_two)); - ASSERT_EQUAL_QUIET(two_and_two, - thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three, thrust::less())); - ASSERT_EQUAL_QUIET(two_and_three, - thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two, thrust::less())); + ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::min(two_and_two, two_and_three, thrust::less())); + ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::min(two_and_three, two_and_two, thrust::less())); - ASSERT_EQUAL_QUIET( - two_and_two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three, thrust::greater())); - ASSERT_EQUAL_QUIET( - two_and_three, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two, thrust::greater())); + ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::min(two_and_two, two_and_three, thrust::greater())); + ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::min(two_and_three, two_and_two, thrust::greater())); } }; SimpleUnitTest TestMinInstance; @@ -46,32 +42,28 @@ struct TestMax { // 2 < 3 T two(2), three(3); - ASSERT_EQUAL(three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two, three)); - ASSERT_EQUAL(three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two, three, thrust::less())); + ASSERT_EQUAL(three, ::cuda::std::max(two, three)); + ASSERT_EQUAL(three, ::cuda::std::max(two, three, thrust::less())); - ASSERT_EQUAL(three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(three, two)); - ASSERT_EQUAL(three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(three, two, thrust::less())); + ASSERT_EQUAL(three, ::cuda::std::max(three, two)); + ASSERT_EQUAL(three, ::cuda::std::max(three, two, thrust::less())); - ASSERT_EQUAL(two, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two, three, thrust::greater())); - ASSERT_EQUAL(two, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(three, two, thrust::greater())); + ASSERT_EQUAL(two, ::cuda::std::max(two, three, thrust::greater())); + ASSERT_EQUAL(two, ::cuda::std::max(three, two, thrust::greater())); using KV = key_value; KV two_and_two(two, two); KV two_and_three(two, three); // the first element breaks ties - ASSERT_EQUAL_QUIET(two_and_two, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three)); - ASSERT_EQUAL_QUIET(two_and_three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two)); + ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::max(two_and_two, two_and_three)); + ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::max(two_and_three, two_and_two)); - ASSERT_EQUAL_QUIET(two_and_two, - thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three, thrust::less())); - ASSERT_EQUAL_QUIET(two_and_three, - thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two, thrust::less())); + ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::max(two_and_two, two_and_three, thrust::less())); + ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::max(two_and_three, two_and_two, thrust::less())); - ASSERT_EQUAL_QUIET( - two_and_two, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three, thrust::greater())); - ASSERT_EQUAL_QUIET( - two_and_three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two, thrust::greater())); + ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::max(two_and_two, two_and_three, thrust::greater())); + ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::max(two_and_three, two_and_two, thrust::greater())); } }; SimpleUnitTest TestMaxInstance; diff --git a/thrust/testing/scan.cu b/thrust/testing/scan.cu index c30eee013b6..8e8985dde1c 100644 --- a/thrust/testing/scan.cu +++ b/thrust/testing/scan.cu @@ -8,19 +8,11 @@ #include #include +#include #include #include -template -struct max_functor -{ - _CCCL_HOST_DEVICE T operator()(T rhs, T lhs) const - { - return thrust::max(rhs, lhs); - } -}; - template void TestScanSimple() { @@ -289,12 +281,12 @@ struct TestScanWithOperator thrust::host_vector h_output(n); thrust::device_vector d_output(n); - thrust::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), max_functor()); - thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), max_functor()); + thrust::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), cuda::maximum{}); + thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), cuda::maximum{}); ASSERT_EQUAL(d_output, h_output); - thrust::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), T(13), max_functor()); - thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), T(13), max_functor()); + thrust::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), T(13), cuda::maximum{}); + thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), T(13), cuda::maximum{}); ASSERT_EQUAL(d_output, h_output); } }; @@ -311,19 +303,19 @@ struct TestScanWithOperatorToDiscardIterator thrust::discard_iterator<> reference(n); thrust::discard_iterator<> h_result = - thrust::inclusive_scan(h_input.begin(), h_input.end(), thrust::make_discard_iterator(), max_functor()); + thrust::inclusive_scan(h_input.begin(), h_input.end(), thrust::make_discard_iterator(), cuda::maximum{}); thrust::discard_iterator<> d_result = - thrust::inclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), max_functor()); + thrust::inclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), cuda::maximum{}); ASSERT_EQUAL_QUIET(reference, h_result); ASSERT_EQUAL_QUIET(reference, d_result); - h_result = - thrust::exclusive_scan(h_input.begin(), h_input.end(), thrust::make_discard_iterator(), T(13), max_functor()); + h_result = thrust::exclusive_scan( + h_input.begin(), h_input.end(), thrust::make_discard_iterator(), T(13), cuda::maximum{}); - d_result = - thrust::exclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), T(13), max_functor()); + d_result = thrust::exclusive_scan( + d_input.begin(), d_input.end(), thrust::make_discard_iterator(), T(13), cuda::maximum{}); ASSERT_EQUAL_QUIET(reference, h_result); ASSERT_EQUAL_QUIET(reference, d_result); diff --git a/thrust/thrust/detail/minmax.h b/thrust/thrust/detail/minmax.h deleted file mode 100644 index ff467f5d835..00000000000 --- a/thrust/thrust/detail/minmax.h +++ /dev/null @@ -1,55 +0,0 @@ -/* - * Copyright 2008-2013 NVIDIA Corporation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include - -#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) -# pragma GCC system_header -#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) -# pragma clang system_header -#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) -# pragma system_header -#endif // no system header - -THRUST_NAMESPACE_BEGIN - -template -_CCCL_HOST_DEVICE T min THRUST_PREVENT_MACRO_SUBSTITUTION(const T& lhs, const T& rhs, BinaryPredicate comp) -{ - return comp(rhs, lhs) ? rhs : lhs; -} // end min() - -template -_CCCL_HOST_DEVICE T min THRUST_PREVENT_MACRO_SUBSTITUTION(const T& lhs, const T& rhs) -{ - return rhs < lhs ? rhs : lhs; -} // end min() - -template -_CCCL_HOST_DEVICE T max THRUST_PREVENT_MACRO_SUBSTITUTION(const T& lhs, const T& rhs, BinaryPredicate comp) -{ - return comp(lhs, rhs) ? rhs : lhs; -} // end max() - -template -_CCCL_HOST_DEVICE T max THRUST_PREVENT_MACRO_SUBSTITUTION(const T& lhs, const T& rhs) -{ - return lhs < rhs ? rhs : lhs; -} // end max() - -THRUST_NAMESPACE_END diff --git a/thrust/thrust/detail/vector_base.inl b/thrust/thrust/detail/vector_base.inl index 4597f69a047..d66765775ce 100644 --- a/thrust/thrust/detail/vector_base.inl +++ b/thrust/thrust/detail/vector_base.inl @@ -27,7 +27,6 @@ #endif // no system header #include #include -#include #include #include #include @@ -36,6 +35,9 @@ #include #include +#include +#include + #include THRUST_NAMESPACE_BEGIN @@ -348,7 +350,7 @@ void vector_base::reserve(size_type n) size_type new_capacity = n; // do not exceed maximum storage - new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(new_capacity, max_size()); + new_capacity = ::cuda::std::min(new_capacity, max_size()); // create new storage storage_type new_storage(copy_allocator_t(), m_storage, new_capacity); @@ -726,13 +728,14 @@ void vector_base::copy_insert(iterator position, ForwardIterator first const size_type old_size = size(); // compute the new capacity after the allocation - size_type new_capacity = old_size + thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(old_size, num_new_elements); + size_type new_capacity = + old_size + ::cuda::std::max THRUST_PREVENT_MACRO_SUBSTITUTION(old_size, num_new_elements); // allocate exponentially larger new storage - new_capacity = thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(new_capacity, 2 * capacity()); + new_capacity = ::cuda::std::max(new_capacity, 2 * capacity()); // do not exceed maximum storage - new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(new_capacity, max_size()); + new_capacity = ::cuda::std::min(new_capacity, max_size()); if (new_capacity > max_size()) { @@ -797,13 +800,13 @@ void vector_base::append(size_type n) const size_type old_size = size(); // compute the new capacity after the allocation - size_type new_capacity = old_size + thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(old_size, n); + size_type new_capacity = old_size + ::cuda::std::max THRUST_PREVENT_MACRO_SUBSTITUTION(old_size, n); // allocate exponentially larger new storage - new_capacity = thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(new_capacity, 2 * capacity()); + new_capacity = ::cuda::std::max(new_capacity, 2 * capacity()); // do not exceed maximum storage - new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(new_capacity, max_size()); + new_capacity = ::cuda::std::min(new_capacity, max_size()); // create new storage storage_type new_storage(copy_allocator_t(), m_storage, new_capacity); @@ -892,13 +895,13 @@ void vector_base::fill_insert(iterator position, size_type n, const T& const size_type old_size = size(); // compute the new capacity after the allocation - size_type new_capacity = old_size + thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(old_size, n); + size_type new_capacity = old_size + ::cuda::std::max THRUST_PREVENT_MACRO_SUBSTITUTION(old_size, n); // allocate exponentially larger new storage - new_capacity = thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(new_capacity, 2 * capacity()); + new_capacity = ::cuda::std::max(new_capacity, 2 * capacity()); // do not exceed maximum storage - new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(new_capacity, max_size()); + new_capacity = ::cuda::std::min(new_capacity, max_size()); if (new_capacity > max_size()) { @@ -1072,10 +1075,10 @@ void vector_base::allocate_and_copy( } // end if // allocate exponentially larger new storage - size_type allocated_size = thrust::max(requested_size, 2 * capacity()); + size_type allocated_size = ::cuda::std::max(requested_size, 2 * capacity()); // do not exceed maximum storage - allocated_size = thrust::min(allocated_size, max_size()); + allocated_size = ::cuda::std::min(allocated_size, max_size()); if (requested_size > allocated_size) { diff --git a/thrust/thrust/extrema.h b/thrust/thrust/extrema.h index 72b3dcf728d..14c37bb70d3 100644 --- a/thrust/thrust/extrema.h +++ b/thrust/thrust/extrema.h @@ -32,155 +32,13 @@ #include #include -THRUST_NAMESPACE_BEGIN - -/*! This version of \p min returns the smaller of two values, given a comparison operation. - * \param lhs The first value to compare. - * \param rhs The second value to compare. - * \param comp A comparison operation. - * \return The smaller element. - * - * \tparam T is convertible to \p BinaryPredicate's first argument type and to its second argument type. - * \tparam BinaryPredicate is a model of BinaryPredicate. - * - * The following code snippet demonstrates how to use \p min to compute the smaller of two - * key-value objects. - * - * \code - * #include - * ... - * struct key_value - * { - * int key; - * int value; - * }; - * - * struct compare_key_value - * { - * __host__ __device__ - * bool operator()(key_value lhs, key_value rhs) - * { - * return lhs.key < rhs.key; - * } - * }; - * - * ... - * key_value a = {13, 0}; - * key_value b = { 7, 1); - * - * key_value smaller = thrust::min(a, b, compare_key_value()); - * - * // smaller is {7, 1} - * \endcode - * - * \note Returns the first argument when the arguments are equivalent. - * \see max - */ -template -_CCCL_HOST_DEVICE T min THRUST_PREVENT_MACRO_SUBSTITUTION(const T& lhs, const T& rhs, BinaryPredicate comp); - -/*! This version of \p min returns the smaller of two values. - * \param lhs The first value to compare. - * \param rhs The second value to compare. - * \return The smaller element. - * - * \tparam T is a model of LessThan - * Comparable. - * - * The following code snippet demonstrates how to use \p min to compute the smaller of two - * integers. - * - * \code - * #include - * ... - * int a = 13; - * int b = 7; - * - * int smaller = thrust::min(a, b); - * - * // smaller is 7 - * \endcode - * - * \note Returns the first argument when the arguments are equivalent. - * \see max - */ -template -_CCCL_HOST_DEVICE T min THRUST_PREVENT_MACRO_SUBSTITUTION(const T& lhs, const T& rhs); +#include +#include -/*! This version of \p max returns the larger of two values, given a comparison operation. - * \param lhs The first value to compare. - * \param rhs The second value to compare. - * \param comp A comparison operation. - * \return The larger element. - * - * \tparam T is convertible to \p BinaryPredicate's first argument type and to its second argument type. - * \tparam BinaryPredicate is a model of BinaryPredicate. - * - * The following code snippet demonstrates how to use \p max to compute the larger of two - * key-value objects. - * - * \code - * #include - * ... - * struct key_value - * { - * int key; - * int value; - * }; - * - * struct compare_key_value - * { - * __host__ __device__ - * bool operator()(key_value lhs, key_value rhs) - * { - * return lhs.key < rhs.key; - * } - * }; - * - * ... - * key_value a = {13, 0}; - * key_value b = { 7, 1); - * - * key_value larger = thrust::max(a, b, compare_key_value()); - * - * // larger is {13, 0} - * \endcode - * - * \note Returns the first argument when the arguments are equivalent. - * \see min - */ -template -_CCCL_HOST_DEVICE T max THRUST_PREVENT_MACRO_SUBSTITUTION(const T& lhs, const T& rhs, BinaryPredicate comp); +THRUST_NAMESPACE_BEGIN -/*! This version of \p max returns the larger of two values. - * \param lhs The first value to compare. - * \param rhs The second value to compare. - * \return The larger element. - * - * \tparam T is a model of LessThan - * Comparable. - * - * The following code snippet demonstrates how to use \p max to compute the larger of two - * integers. - * - * \code - * #include - * ... - * int a = 13; - * int b = 7; - * - * int larger = thrust::min(a, b); - * - * // larger is 13 - * \endcode - * - * \note Returns the first argument when the arguments are equivalent. - * \see min - */ -template -_CCCL_HOST_DEVICE T max THRUST_PREVENT_MACRO_SUBSTITUTION(const T& lhs, const T& rhs); +using ::cuda::std::max; +using ::cuda::std::min; /*! \addtogroup reductions * \{ @@ -797,4 +655,3 @@ minmax_element(ForwardIterator first, ForwardIterator last, BinaryPredicate comp THRUST_NAMESPACE_END #include -#include diff --git a/thrust/thrust/system/cuda/detail/adjacent_difference.h b/thrust/thrust/system/cuda/detail/adjacent_difference.h index 043e1a571ad..c48a1ddf7e8 100644 --- a/thrust/thrust/system/cuda/detail/adjacent_difference.h +++ b/thrust/thrust/system/cuda/detail/adjacent_difference.h @@ -43,7 +43,6 @@ # include # include -# include # include # include # include diff --git a/thrust/thrust/system/cuda/detail/find.h b/thrust/thrust/system/cuda/detail/find.h index de633c73ebb..bdce49ae640 100644 --- a/thrust/thrust/system/cuda/detail/find.h +++ b/thrust/thrust/system/cuda/detail/find.h @@ -39,7 +39,6 @@ #if _CCCL_HAS_CUDA_COMPILER # include -# include # include # include @@ -79,7 +78,7 @@ struct functor // select the smallest index among true results if (thrust::get<0>(lhs) && thrust::get<0>(rhs)) { - return TupleType(true, (thrust::min)(thrust::get<1>(lhs), thrust::get<1>(rhs))); + return TupleType(true, (::cuda::std::min)(thrust::get<1>(lhs), thrust::get<1>(rhs))); } else if (thrust::get<0>(lhs)) { @@ -113,7 +112,7 @@ find_if_n(execution_policy& policy, InputIt first, Size num_items, Pred // TODO incorporate sizeof(InputType) into interval_threshold and round to multiple of 32 const Size interval_threshold = 1 << 20; - const Size interval_size = (thrust::min)(interval_threshold, num_items); + const Size interval_size = (::cuda::std::min)(interval_threshold, num_items); // force transform_iterator output to bool using XfrmIterator = transform_input_iterator_t; diff --git a/thrust/thrust/system/cuda/detail/inner_product.h b/thrust/thrust/system/cuda/detail/inner_product.h index af41c5ccda8..d30d4f78042 100644 --- a/thrust/thrust/system/cuda/detail/inner_product.h +++ b/thrust/thrust/system/cuda/detail/inner_product.h @@ -37,7 +37,6 @@ #endif // no system header #if _CCCL_HAS_CUDA_COMPILER -# include # include # include diff --git a/thrust/thrust/system/cuda/detail/reduce.h b/thrust/thrust/system/cuda/detail/reduce.h index 9b9401e2891..443063fb3b4 100644 --- a/thrust/thrust/system/cuda/detail/reduce.h +++ b/thrust/thrust/system/cuda/detail/reduce.h @@ -44,7 +44,6 @@ # include # include -# include # include # include # include @@ -647,7 +646,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step( // small, single tile size if (d_temp_storage == nullptr) { - temp_storage_bytes = max(1, vshmem_size); + temp_storage_bytes = ::cuda::std::max(1, vshmem_size); return status; } char* vshmem_ptr = vshmem_size > 0 ? (char*) d_temp_storage : nullptr; @@ -717,7 +716,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step( // if not enough to fill the device with threadblocks // then fill the device with threadblocks - reduce_grid_size = static_cast((min) (num_tiles, static_cast(reduce_device_occupancy))); + reduce_grid_size = static_cast((::cuda::std::min)(num_tiles, static_cast(reduce_device_occupancy))); using drain_agent = AgentLauncher>; AgentPlan drain_plan = drain_agent::get_plan(); diff --git a/thrust/thrust/system/cuda/detail/reduce_by_key.h b/thrust/thrust/system/cuda/detail/reduce_by_key.h index 58aa33fd7e8..cc59c98ab2c 100644 --- a/thrust/thrust/system/cuda/detail/reduce_by_key.h +++ b/thrust/thrust/system/cuda/detail/reduce_by_key.h @@ -44,7 +44,6 @@ # include # include -# include # include # include # include diff --git a/thrust/thrust/system/cuda/detail/scan_by_key.h b/thrust/thrust/system/cuda/detail/scan_by_key.h index 3ffc7898dcb..1240783736c 100644 --- a/thrust/thrust/system/cuda/detail/scan_by_key.h +++ b/thrust/thrust/system/cuda/detail/scan_by_key.h @@ -43,7 +43,6 @@ # include # include -# include # include # include # include diff --git a/thrust/thrust/system/cuda/detail/set_operations.h b/thrust/thrust/system/cuda/detail/set_operations.h index 5f094377421..0ef80c0fb2d 100644 --- a/thrust/thrust/system/cuda/detail/set_operations.h +++ b/thrust/thrust/system/cuda/detail/set_operations.h @@ -52,6 +52,9 @@ # include # include +# include +# include + # include THRUST_NAMESPACE_BEGIN @@ -127,8 +130,8 @@ THRUST_DEVICE_FUNCTION Size merge_path(It1 a, Size aCount, It2 b, Size bCount, S { using T = typename thrust::iterator_traits::value_type; - Size begin = thrust::max(0, diag - bCount); - Size end = thrust::min(diag, aCount); + Size begin = ::cuda::std::max(0, diag - bCount); + Size end = ::cuda::std::min(diag, aCount); while (begin < end) { diff --git a/thrust/thrust/system/cuda/detail/transform_reduce.h b/thrust/thrust/system/cuda/detail/transform_reduce.h index 7306a8ef34a..fbb3054b0bf 100644 --- a/thrust/thrust/system/cuda/detail/transform_reduce.h +++ b/thrust/thrust/system/cuda/detail/transform_reduce.h @@ -43,7 +43,6 @@ # include # include -# include # include # include # include diff --git a/thrust/thrust/system/cuda/detail/unique.h b/thrust/thrust/system/cuda/detail/unique.h index 49d8e2668b6..b8e408254cb 100644 --- a/thrust/thrust/system/cuda/detail/unique.h +++ b/thrust/thrust/system/cuda/detail/unique.h @@ -45,7 +45,6 @@ # include # include -# include # include # include # include @@ -532,7 +531,7 @@ static cudaError_t THRUST_RUNTIME_FUNCTION doit_step( status = tile_status.Init(static_cast(num_tiles), allocations[0], allocation_sizes[0]); CUDA_CUB_RET_IF_FAIL(status); - num_tiles = max(1, num_tiles); + num_tiles = ::cuda::std::max(1, num_tiles); init_agent ia(init_plan, num_tiles, stream, "unique_by_key::init_agent"); ia.launch(tile_status, num_tiles, num_selected_out); CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError()); diff --git a/thrust/thrust/system/cuda/detail/unique_by_key.h b/thrust/thrust/system/cuda/detail/unique_by_key.h index 5b88b9935b3..8ab790933fb 100644 --- a/thrust/thrust/system/cuda/detail/unique_by_key.h +++ b/thrust/thrust/system/cuda/detail/unique_by_key.h @@ -44,7 +44,6 @@ # include # include -# include # include # include # include diff --git a/thrust/thrust/system/detail/generic/find.inl b/thrust/thrust/system/detail/generic/find.inl index 25eeed24e9f..09ec0b8c59e 100644 --- a/thrust/thrust/system/detail/generic/find.inl +++ b/thrust/thrust/system/detail/generic/find.inl @@ -26,7 +26,6 @@ # pragma system_header #endif // no system header #include -#include #include #include #include @@ -34,6 +33,8 @@ #include #include +#include + // Contributed by Erich Elsen THRUST_NAMESPACE_BEGIN @@ -61,7 +62,7 @@ struct find_if_functor // select the smallest index among true results if (thrust::get<0>(lhs) && thrust::get<0>(rhs)) { - return TupleType(true, (thrust::min)(thrust::get<1>(lhs), thrust::get<1>(rhs))); + return TupleType(true, (::cuda::std::min)(thrust::get<1>(lhs), thrust::get<1>(rhs))); } else if (thrust::get<0>(lhs)) { @@ -94,7 +95,7 @@ find_if(thrust::execution_policy& exec, InputIterator first, Inpu // TODO incorporate sizeof(InputType) into interval_threshold and round to multiple of 32 const difference_type interval_threshold = 1 << 20; - const difference_type interval_size = (thrust::min)(interval_threshold, n); + const difference_type interval_size = (::cuda::std::min)(interval_threshold, n); // force transform_iterator output to bool using XfrmIterator = thrust::transform_iterator; diff --git a/thrust/thrust/system/detail/sequential/stable_merge_sort.inl b/thrust/thrust/system/detail/sequential/stable_merge_sort.inl index 415ff2a70a9..e987540c0e2 100644 --- a/thrust/thrust/system/detail/sequential/stable_merge_sort.inl +++ b/thrust/thrust/system/detail/sequential/stable_merge_sort.inl @@ -26,12 +26,13 @@ # pragma system_header #endif // no system header -#include #include #include #include #include +#include + #include THRUST_NAMESPACE_BEGIN @@ -95,7 +96,7 @@ insertion_sort_each(RandomAccessIterator first, RandomAccessIterator last, Size { for (; first < last; first += partition_size) { - RandomAccessIterator partition_last = (thrust::min)(last, first + partition_size); + RandomAccessIterator partition_last = (::cuda::std::min)(last, first + partition_size); thrust::system::detail::sequential::insertion_sort(first, partition_last, comp); } // end for @@ -114,7 +115,7 @@ _CCCL_HOST_DEVICE void insertion_sort_each_by_key( { for (; keys_first < keys_last; keys_first += partition_size, values_first += partition_size) { - RandomAccessIterator1 keys_partition_last = (thrust::min)(keys_last, keys_first + partition_size); + RandomAccessIterator1 keys_partition_last = (::cuda::std::min)(keys_last, keys_first + partition_size); thrust::system::detail::sequential::insertion_sort_by_key(keys_first, keys_partition_last, values_first, comp); } // end for @@ -136,8 +137,8 @@ _CCCL_HOST_DEVICE void merge_adjacent_partitions( { for (; first < last; first += 2 * partition_size, result += 2 * partition_size) { - RandomAccessIterator1 interval_middle = (thrust::min)(last, first + partition_size); - RandomAccessIterator1 interval_last = (thrust::min)(last, interval_middle + partition_size); + RandomAccessIterator1 interval_middle = (::cuda::std::min)(last, first + partition_size); + RandomAccessIterator1 interval_last = (::cuda::std::min)(last, interval_middle + partition_size); thrust::merge(exec, first, interval_middle, interval_middle, interval_last, result, comp); } // end for @@ -165,8 +166,8 @@ _CCCL_HOST_DEVICE void merge_adjacent_partitions_by_key( for (; keys_first < keys_last; keys_first += stride, values_first += stride, keys_result += stride, values_result += stride) { - RandomAccessIterator1 keys_interval_middle = (thrust::min)(keys_last, keys_first + partition_size); - RandomAccessIterator1 keys_interval_last = (thrust::min)(keys_last, keys_interval_middle + partition_size); + RandomAccessIterator1 keys_interval_middle = (::cuda::std::min)(keys_last, keys_first + partition_size); + RandomAccessIterator1 keys_interval_last = (::cuda::std::min)(keys_last, keys_interval_middle + partition_size); RandomAccessIterator2 values_first2 = values_first + (keys_interval_middle - keys_first); diff --git a/thrust/thrust/system/tbb/detail/reduce_by_key.inl b/thrust/thrust/system/tbb/detail/reduce_by_key.inl index f799dfa6087..96c4630ae8f 100644 --- a/thrust/thrust/system/tbb/detail/reduce_by_key.inl +++ b/thrust/thrust/system/tbb/detail/reduce_by_key.inl @@ -25,7 +25,6 @@ #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) # pragma system_header #endif // no system header -#include #include #include #include @@ -36,6 +35,8 @@ #include #include +#include +#include #include #include @@ -197,7 +198,7 @@ struct serial_reduce_by_key_body const size_type interval_idx = r.begin(); const size_type offset_to_first = interval_size * interval_idx; - const size_type offset_to_last = (thrust::min)(n, offset_to_first + interval_size); + const size_type offset_to_last = ::cuda::std::min(n, offset_to_first + interval_size); Iterator1 my_keys_first = keys_first + offset_to_first; Iterator1 my_keys_last = keys_first + offset_to_last; @@ -319,13 +320,13 @@ thrust::pair reduce_by_key( } // count the number of processors - const unsigned int p = thrust::max(1u, std::thread::hardware_concurrency()); + const unsigned int p = ::cuda::std::max(1u, std::thread::hardware_concurrency()); // generate O(P) intervals of sequential work // XXX oversubscribing is a tuning opportunity const unsigned int subscription_rate = 1; - difference_type interval_size = - thrust::min(parallelism_threshold, thrust::max(n, n / (subscription_rate * p))); + difference_type interval_size = ::cuda::std::min( + parallelism_threshold, ::cuda::std::max(n, n / (subscription_rate * p))); difference_type num_intervals = reduce_by_key_detail::divide_ri(n, interval_size); // decompose the input into intervals of size N / num_intervals diff --git a/thrust/thrust/system/tbb/detail/reduce_intervals.h b/thrust/thrust/system/tbb/detail/reduce_intervals.h index 6270650b04d..d29b830823b 100644 --- a/thrust/thrust/system/tbb/detail/reduce_intervals.h +++ b/thrust/thrust/system/tbb/detail/reduce_intervals.h @@ -25,13 +25,14 @@ #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) # pragma system_header #endif // no system header -#include #include #include #include #include #include +#include + #include #include @@ -76,7 +77,7 @@ struct body Size interval_idx = r.begin(); Size offset_to_first = interval_size * interval_idx; - Size offset_to_last = (thrust::min)(n, offset_to_first + interval_size); + Size offset_to_last = (::cuda::std::min)(n, offset_to_first + interval_size); RandomAccessIterator1 my_first = first + offset_to_first; RandomAccessIterator1 my_last = first + offset_to_last;