Use cuda::std::min/max in Thrust (NVIDIA#3364)

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<KeyType>((m_num_summaries - 1 - summary_idx) * m_key_conversion)
                                         : static_cast<KeyType>(summary_idx * m_key_conversion);
 

thrust/examples/set_operations.cu

@@ -1,4 +1,5 @@
 #include <thrust/device_vector.h>
+#include <thrust/extrema.h>
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/merge.h>
 #include <thrust/set_operations.h>

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<uint64_t, thrust::maximum<>> // - initial_value, binop
+    std::tuple<uint64_t, ::cuda::maximum<>> // - 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<>{})};
   }
 };
 

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<thrust::maximum<>> // - 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<>{})};
   }
 };
 

thrust/testing/cuda/is_partitioned.cu

@@ -26,7 +26,7 @@ void TestIsPartitionedDevice(ExecutionPolicy exec)
 {
   size_t n = 1000;
 
-  n = thrust::max<size_t>(n, 2);
+  n = ::cuda::std::max<size_t>(n, 2);
 
   thrust::device_vector<int> v = unittest::random_integers<int>(n);
 

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<T>()));
+    ASSERT_EQUAL(two, ::cuda::std::min(two, three));
+    ASSERT_EQUAL(two, ::cuda::std::min(two, three, thrust::less<T>()));
 
-    ASSERT_EQUAL(two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(three, two));
-    ASSERT_EQUAL(two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(three, two, thrust::less<T>()));
+    ASSERT_EQUAL(two, ::cuda::std::min(three, two));
+    ASSERT_EQUAL(two, ::cuda::std::min(three, two, thrust::less<T>()));
 
-    ASSERT_EQUAL(three, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two, three, thrust::greater<T>()));
-    ASSERT_EQUAL(three, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(three, two, thrust::greater<T>()));
+    ASSERT_EQUAL(three, ::cuda::std::min(two, three, thrust::greater<T>()));
+    ASSERT_EQUAL(three, ::cuda::std::min(three, two, thrust::greater<T>()));
 
     using KV = key_value<T, T>;
     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<KV>()));
-    ASSERT_EQUAL_QUIET(two_and_three,
-                       thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two, thrust::less<KV>()));
+    ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::min(two_and_two, two_and_three, thrust::less<KV>()));
+    ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::min(two_and_three, two_and_two, thrust::less<KV>()));
 
-    ASSERT_EQUAL_QUIET(
-      two_and_two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three, thrust::greater<KV>()));
-    ASSERT_EQUAL_QUIET(
-      two_and_three, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two, thrust::greater<KV>()));
+    ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::min(two_and_two, two_and_three, thrust::greater<KV>()));
+    ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::min(two_and_three, two_and_two, thrust::greater<KV>()));
   }
 };
 SimpleUnitTest<TestMin, NumericTypes> 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<T>()));
+    ASSERT_EQUAL(three, ::cuda::std::max(two, three));
+    ASSERT_EQUAL(three, ::cuda::std::max(two, three, thrust::less<T>()));
 
-    ASSERT_EQUAL(three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(three, two));
-    ASSERT_EQUAL(three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(three, two, thrust::less<T>()));
+    ASSERT_EQUAL(three, ::cuda::std::max(three, two));
+    ASSERT_EQUAL(three, ::cuda::std::max(three, two, thrust::less<T>()));
 
-    ASSERT_EQUAL(two, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two, three, thrust::greater<T>()));
-    ASSERT_EQUAL(two, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(three, two, thrust::greater<T>()));
+    ASSERT_EQUAL(two, ::cuda::std::max(two, three, thrust::greater<T>()));
+    ASSERT_EQUAL(two, ::cuda::std::max(three, two, thrust::greater<T>()));
 
     using KV = key_value<T, T>;
     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<KV>()));
-    ASSERT_EQUAL_QUIET(two_and_three,
-                       thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two, thrust::less<KV>()));
+    ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::max(two_and_two, two_and_three, thrust::less<KV>()));
+    ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::max(two_and_three, two_and_two, thrust::less<KV>()));
 
-    ASSERT_EQUAL_QUIET(
-      two_and_two, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three, thrust::greater<KV>()));
-    ASSERT_EQUAL_QUIET(
-      two_and_three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two, thrust::greater<KV>()));
+    ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::max(two_and_two, two_and_three, thrust::greater<KV>()));
+    ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::max(two_and_three, two_and_two, thrust::greater<KV>()));
   }
 };
 SimpleUnitTest<TestMax, NumericTypes> TestMaxInstance;

thrust/testing/scan.cu

@@ -8,19 +8,11 @@
 #include <thrust/iterator/retag.h>
 #include <thrust/scan.h>
 
+#include <cuda/functional>
 #include <cuda/std/array>
 
 #include <unittest/unittest.h>
 
-template <typename T>
-struct max_functor
-{
-  _CCCL_HOST_DEVICE T operator()(T rhs, T lhs) const
-  {
-    return thrust::max(rhs, lhs);
-  }
-};
-
 template <class Vector>
 void TestScanSimple()
 {
@@ -289,12 +281,12 @@ struct TestScanWithOperator
     thrust::host_vector<T> h_output(n);
     thrust::device_vector<T> d_output(n);
 
-    thrust::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), max_functor<T>());
-    thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), max_functor<T>());
+    thrust::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), cuda::maximum<T>{});
+    thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), cuda::maximum<T>{});
     ASSERT_EQUAL(d_output, h_output);
 
-    thrust::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), T(13), max_functor<T>());
-    thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), T(13), max_functor<T>());
+    thrust::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), T(13), cuda::maximum<T>{});
+    thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), T(13), cuda::maximum<T>{});
     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<T>());
+      thrust::inclusive_scan(h_input.begin(), h_input.end(), thrust::make_discard_iterator(), cuda::maximum<T>{});
 
     thrust::discard_iterator<> d_result =
-      thrust::inclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), max_functor<T>());
+      thrust::inclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), cuda::maximum<T>{});
 
     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<T>());
+    h_result = thrust::exclusive_scan(
+      h_input.begin(), h_input.end(), thrust::make_discard_iterator(), T(13), cuda::maximum<T>{});
 
-    d_result =
-      thrust::exclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), T(13), max_functor<T>());
+    d_result = thrust::exclusive_scan(
+      d_input.begin(), d_input.end(), thrust::make_discard_iterator(), T(13), cuda::maximum<T>{});
 
     ASSERT_EQUAL_QUIET(reference, h_result);
     ASSERT_EQUAL_QUIET(reference, d_result);

thrust/thrust/detail/vector_base.inl

@@ -27,7 +27,6 @@
 #endif // no system header
 #include <thrust/advance.h>
 #include <thrust/detail/copy.h>
-#include <thrust/detail/minmax.h>
 #include <thrust/detail/overlapped_copy.h>
 #include <thrust/detail/temporary_array.h>
 #include <thrust/detail/type_traits.h>
@@ -36,6 +35,9 @@
 #include <thrust/equal.h>
 #include <thrust/iterator/iterator_traits.h>
 
+#include <cuda/std/__algorithm/max.h>
+#include <cuda/std/__algorithm/min.h>
+
 #include <stdexcept>
 
 THRUST_NAMESPACE_BEGIN
@@ -348,7 +350,7 @@ void vector_base<T, Alloc>::reserve(size_type n)
     size_type new_capacity = n;
 
     // do not exceed maximum storage
-    new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION<size_type>(new_capacity, max_size());
+    new_capacity = ::cuda::std::min<size_type>(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<T, Alloc>::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<size_type>(new_capacity, 2 * capacity());
+      new_capacity = ::cuda::std::max<size_type>(new_capacity, 2 * capacity());
 
       // do not exceed maximum storage
-      new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION<size_type>(new_capacity, max_size());
+      new_capacity = ::cuda::std::min<size_type>(new_capacity, max_size());
 
       if (new_capacity > max_size())
       {
@@ -797,13 +800,13 @@ void vector_base<T, Alloc>::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<size_type>(new_capacity, 2 * capacity());
+      new_capacity = ::cuda::std::max<size_type>(new_capacity, 2 * capacity());
 
       // do not exceed maximum storage
-      new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION<size_type>(new_capacity, max_size());
+      new_capacity = ::cuda::std::min<size_type>(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<T, Alloc>::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<size_type>(new_capacity, 2 * capacity());
+      new_capacity = ::cuda::std::max<size_type>(new_capacity, 2 * capacity());
 
       // do not exceed maximum storage
-      new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION<size_type>(new_capacity, max_size());
+      new_capacity = ::cuda::std::min<size_type>(new_capacity, max_size());
 
       if (new_capacity > max_size())
       {
@@ -1072,10 +1075,10 @@ void vector_base<T, Alloc>::allocate_and_copy(
   } // end if
 
   // allocate exponentially larger new storage
-  size_type allocated_size = thrust::max<size_type>(requested_size, 2 * capacity());
+  size_type allocated_size = ::cuda::std::max<size_type>(requested_size, 2 * capacity());
 
   // do not exceed maximum storage
-  allocated_size = thrust::min<size_type>(allocated_size, max_size());
+  allocated_size = ::cuda::std::min<size_type>(allocated_size, max_size());
 
   if (requested_size > allocated_size)
   {