Replace uses of thrust::pair with cuda::std::pair

cub/test/catch2_test_c2h_checked_cuda_allocator.cu

@@ -36,7 +36,7 @@ C2H_TEST("c2h::device_vector throws when requested allocations exceed free devic
 C2H_TEST("c2h::device_policy throws when requested allocations exceed free device memory",
          "[c2h][checked_cuda_allocator][device_policy]")
 {
-  thrust::pair<char*, std::ptrdiff_t> buffer{nullptr, 0};
+  cuda::std::pair<char*, std::ptrdiff_t> buffer{nullptr, 0};
   auto policy = thrust::detail::derived_cast(thrust::detail::strip_const(c2h::device_policy));
 
   const std::size_t alloc_bytes = get_alloc_bytes();

libcudacxx/examples/concurrent_hash_table.cu

@@ -112,7 +112,7 @@ public:
   // TODO: Change return type to an enum with three possible values, succeeded,
   // exists, and full.
   template <typename UKey, typename... Args>
-  __host__ __device__ thrust::pair<value_iterator, bool> try_emplace(UKey&& key, Args&&... args)
+  __host__ __device__ cuda::std::pair<value_iterator, bool> try_emplace(UKey&& key, Args&&... args)
   {
     auto index{hash_(key) % capacity_};
     // Linearly probe the storage space up to `capacity_` times; if we haven't
@@ -130,7 +130,7 @@ public:
           new (keys_ + index) key_type(std::forward<UKey>(key));
           new (values_ + index) mapped_type(std::forward<Args>(args)...);
           states_[index].store(state_filled, cuda::std::memory_order_release);
-          return thrust::make_pair(values_ + index, true);
+          return cuda::std::make_pair(values_ + index, true);
         }
       }
       // If we are here, the element we are probing is not empty and we didn't
@@ -142,13 +142,13 @@ public:
       if (key_equal_(keys_[index], key))
       {
         // It is, so the element already exists.
-        return thrust::make_pair(values_ + index, false);
+        return cuda::std::make_pair(values_ + index, false);
       }
       // Otherwise, the element isn't a match, so move on to the next element.
       index = (index + 1) % capacity_;
     }
     // If we are here, the container is full.
-    return thrust::make_pair(value_iterator{}, false);
+    return cuda::std::make_pair(value_iterator{}, false);
   }
 
   __host__ __device__ mapped_type& operator[](key_type const& key)

libcudacxx/include/cuda/std/__memory/construct_at.h

@@ -158,7 +158,7 @@ _CCCL_API constexpr _ForwardIterator __destroy(_ForwardIterator, _ForwardIterato
 
 _CCCL_EXEC_CHECK_DISABLE
 template <class _Tp>
-_CCCL_API constexpr void __destroy_at(_Tp* __loc)
+_CCCL_API constexpr void __destroy_at([[maybe_unused]] _Tp* __loc)
 {
   _CCCL_ASSERT(__loc != nullptr, "null pointer given to __destroy_at");
   if constexpr (is_trivially_destructible_v<_Tp>)
@@ -201,7 +201,7 @@ __reverse_destroy(_BidirectionalIterator __first, _BidirectionalIterator __last)
 
 _CCCL_EXEC_CHECK_DISABLE
 template <class _Tp>
-_CCCL_API inline _CCCL_CONSTEXPR_CXX20 void destroy_at(_Tp* __loc)
+_CCCL_API inline _CCCL_CONSTEXPR_CXX20 void destroy_at([[maybe_unused]] _Tp* __loc)
 {
   _CCCL_ASSERT(__loc != nullptr, "null pointer given to __destroy_at");
   if constexpr (is_trivially_destructible_v<_Tp>)

thrust/benchmarks/bench/set_operations/difference_by_key.cu

@@ -14,7 +14,7 @@ struct op_t
             class InputIterator4,
             class OutputIterator1,
             class OutputIterator2>
-  __host__ thrust::pair<OutputIterator1, OutputIterator2> operator()(
+  __host__ cuda::std::pair<OutputIterator1, OutputIterator2> operator()(
     const PolicyT& policy,
     InputIterator1 keys_first1,
     InputIterator1 keys_last1,

thrust/benchmarks/bench/set_operations/intersection_by_key.cu

@@ -14,7 +14,7 @@ struct op_t
             class InputIterator4,
             class OutputIterator1,
             class OutputIterator2>
-  __host__ thrust::pair<OutputIterator1, OutputIterator2> operator()(
+  __host__ cuda::std::pair<OutputIterator1, OutputIterator2> operator()(
     const PolicyT& policy,
     InputIterator1 keys_first1,
     InputIterator1 keys_last1,

thrust/benchmarks/bench/set_operations/symmetric_difference_by_key.cu

@@ -14,7 +14,7 @@ struct op_t
             class InputIterator4,
             class OutputIterator1,
             class OutputIterator2>
-  __host__ thrust::pair<OutputIterator1, OutputIterator2> operator()(
+  __host__ cuda::std::pair<OutputIterator1, OutputIterator2> operator()(
     const PolicyT& policy,
     InputIterator1 keys_first1,
     InputIterator1 keys_last1,

thrust/benchmarks/bench/set_operations/union_by_key.cu

@@ -14,7 +14,7 @@ struct op_t
             class InputIterator4,
             class OutputIterator1,
             class OutputIterator2>
-  __host__ thrust::pair<OutputIterator1, OutputIterator2> operator()(
+  __host__ cuda::std::pair<OutputIterator1, OutputIterator2> operator()(
     const PolicyT& policy,
     InputIterator1 keys_first1,
     InputIterator1 keys_last1,

thrust/examples/bounding_box.cu

@@ -1,9 +1,10 @@
 #include <thrust/device_vector.h>
 #include <thrust/extrema.h>
-#include <thrust/pair.h>
 #include <thrust/random.h>
 #include <thrust/transform_reduce.h>
 
+#include <cuda/std/utility>
+
 // This example shows how to compute a bounding box
 // for a set of points in two dimensions.
 

thrust/examples/cuda/custom_temporary_allocation.cu

@@ -1,10 +1,11 @@
 #include <thrust/generate.h>
 #include <thrust/host_vector.h>
-#include <thrust/pair.h>
 #include <thrust/sort.h>
 #include <thrust/system/cuda/execution_policy.h>
 #include <thrust/system/cuda/vector.h>
 
+#include <cuda/std/utility>
+
 #include <cassert>
 #include <cstdlib>
 #include <iostream>

thrust/testing/catch2_test_binary_search.cu

@@ -197,11 +197,11 @@ TEMPLATE_LIST_TEST_CASE("ScalarEqualRangeSimple", "[binary_search]", vector_list
 }
 
 template <typename ForwardIterator, typename LessThanComparable>
-thrust::pair<ForwardIterator, ForwardIterator>
+cuda::std::pair<ForwardIterator, ForwardIterator>
 equal_range(my_system& system, ForwardIterator first, ForwardIterator /*last*/, const LessThanComparable& /*value*/)
 {
   system.validate_dispatch();
-  return thrust::make_pair(first, first);
+  return cuda::std::make_pair(first, first);
 }
 
 TEST_CASE("ScalarEqualRangeDispatchExplicit", "[binary_search]")
@@ -215,11 +215,11 @@ TEST_CASE("ScalarEqualRangeDispatchExplicit", "[binary_search]")
 }
 
 template <typename ForwardIterator, typename LessThanComparable>
-thrust::pair<ForwardIterator, ForwardIterator>
+cuda::std::pair<ForwardIterator, ForwardIterator>
 equal_range(my_tag, ForwardIterator first, ForwardIterator /*last*/, const LessThanComparable& /*value*/)
 {
   *first = 13;
-  return thrust::make_pair(first, first);
+  return cuda::std::make_pair(first, first);
 }
 
 TEST_CASE("ScalarEqualRangeDispatchImplicit", "[binary_search]")

thrust/testing/catch2_test_transform.cu

@@ -2,10 +2,11 @@
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/iterator/zip_iterator.h>
-#include <thrust/pair.h>
 #include <thrust/transform.h>
 #include <thrust/tuple.h>
 
+#include <cuda/std/utility>
+
 #include "catch2_test_helper.h"
 #include "unittest/random.h"
 #include "unittest/special_types.h"
@@ -376,9 +377,9 @@ TEMPLATE_LIST_TEST_CASE("UnaryToDiscardIterator", "[transform]", variable_list)
 struct repeat2
 {
   template <typename T>
-  _CCCL_HOST_DEVICE thrust::pair<T, T> operator()(T x)
+  _CCCL_HOST_DEVICE cuda::std::pair<T, T> operator()(T x)
   {
-    return thrust::make_pair(x, x);
+    return cuda::std::make_pair(x, x);
   }
 };
 

thrust/testing/cuda/binary_search.cu

@@ -1,16 +1,17 @@
 #include <thrust/binary_search.h>
 #include <thrust/device_vector.h>
 #include <thrust/distance.h>
-#include <thrust/pair.h>
 #include <thrust/sequence.h>
 
+#include <cuda/std/utility>
+
 #include <unittest/unittest.h>
 
 void TestEqualRangeOnStream()
 { // Regression test for GH issue #921 (nvbug 2173437)
   using vector_t   = typename thrust::device_vector<int>;
   using iterator_t = typename vector_t::iterator;
-  using result_t   = thrust::pair<iterator_t, iterator_t>;
+  using result_t   = cuda::std::pair<iterator_t, iterator_t>;
 
   vector_t input(10);
   thrust::sequence(thrust::device, input.begin(), input.end(), 0);

thrust/testing/cuda/memory.cu

@@ -50,7 +50,7 @@ void TestGetTemporaryBufferDeviceSeq()
   const std::ptrdiff_t n = 9001;
 
   using pointer         = thrust::pointer<int, thrust::detail::seq_t>;
-  using ptr_and_sz_type = thrust::pair<pointer, std::ptrdiff_t>;
+  using ptr_and_sz_type = cuda::std::pair<pointer, std::ptrdiff_t>;
   thrust::device_vector<ptr_and_sz_type> d_result(1);
 
   get_temporary_buffer_kernel<<<1, 1>>>(n, d_result.begin());

thrust/testing/cuda/merge_by_key.cu

@@ -55,7 +55,7 @@ void TestMergeByKeyDevice(ExecutionPolicy exec)
 
   using Iterator = typename thrust::device_vector<int>::iterator;
 
-  thrust::device_vector<thrust::pair<Iterator, Iterator>> result_ends(1);
+  thrust::device_vector<cuda::std::pair<Iterator, Iterator>> result_ends(1);
 
   merge_by_key_kernel<<<1, 1>>>(
     exec,
@@ -71,7 +71,7 @@ void TestMergeByKeyDevice(ExecutionPolicy exec)
   cudaError_t const err = cudaDeviceSynchronize();
   ASSERT_EQUAL(cudaSuccess, err);
 
-  thrust::pair<Iterator, Iterator> ends = result_ends[0];
+  cuda::std::pair<Iterator, Iterator> ends = result_ends[0];
 
   ASSERT_EQUAL_QUIET(result_key.end(), ends.first);
   ASSERT_EQUAL_QUIET(result_val.end(), ends.second);
@@ -120,7 +120,7 @@ void TestMergeByKeyCudaStreams()
   cudaStream_t s;
   cudaStreamCreate(&s);
 
-  thrust::pair<Iterator, Iterator> ends = thrust::merge_by_key(
+  cuda::std::pair<Iterator, Iterator> ends = thrust::merge_by_key(
     thrust::cuda::par.on(s),
     a_key.begin(),
     a_key.end(),

thrust/testing/cuda/minmax_element.cu

@@ -30,7 +30,7 @@ void TestMinMaxElementDevice(ExecutionPolicy exec)
   typename thrust::device_vector<int>::iterator d_max;
 
   using pair_type =
-    thrust::pair<typename thrust::device_vector<int>::iterator, typename thrust::device_vector<int>::iterator>;
+    cuda::std::pair<typename thrust::device_vector<int>::iterator, typename thrust::device_vector<int>::iterator>;
 
   thrust::device_vector<pair_type> d_result(1);
 

thrust/testing/cuda/mismatch.cu

@@ -20,7 +20,7 @@ void TestMismatchDevice(ExecutionPolicy exec)
   thrust::device_vector<int> b = {1, 2, 4, 3};
 
   using pair_type =
-    thrust::pair<typename thrust::device_vector<int>::iterator, typename thrust::device_vector<int>::iterator>;
+    cuda::std::pair<typename thrust::device_vector<int>::iterator, typename thrust::device_vector<int>::iterator>;
 
   thrust::device_vector<pair_type> d_result(1);
 

thrust/testing/cuda/pair_sort.cu

@@ -1,7 +1,8 @@
 #include <thrust/execution_policy.h>
-#include <thrust/pair.h>
 #include <thrust/sort.h>
 
+#include <cuda/std/utility>
+
 #include <unittest/unittest.h>
 
 #ifdef THRUST_TEST_DEVICE_SIDE
@@ -14,17 +15,17 @@ __global__ void stable_sort_kernel(ExecutionPolicy exec, Iterator first, Iterato
 struct make_pair_functor
 {
   template <typename T1, typename T2>
-  _CCCL_HOST_DEVICE thrust::pair<T1, T2> operator()(const T1& x, const T2& y)
+  _CCCL_HOST_DEVICE cuda::std::pair<T1, T2> operator()(const T1& x, const T2& y)
   {
-    return thrust::make_pair(x, y);
+    return cuda::std::make_pair(x, y);
   } // end operator()()
 }; // end make_pair_functor
 
 template <typename ExecutionPolicy>
 void TestPairStableSortDevice(ExecutionPolicy exec)
 {
   size_t n = 10000;
-  using P  = thrust::pair<int, int>;
+  using P  = cuda::std::pair<int, int>;
 
   thrust::host_vector<int> h_p1 = unittest::random_integers<int>(n);
   thrust::host_vector<int> h_p2 = unittest::random_integers<int>(n);

thrust/testing/cuda/pair_sort_by_key.cu

@@ -1,9 +1,10 @@
 #include <thrust/execution_policy.h>
-#include <thrust/pair.h>
 #include <thrust/sequence.h>
 #include <thrust/sort.h>
 #include <thrust/transform.h>
 
+#include <cuda/std/utility>
+
 #include <unittest/unittest.h>
 
 #ifdef THRUST_TEST_DEVICE_SIDE
@@ -17,17 +18,17 @@ stable_sort_by_key_kernel(ExecutionPolicy exec, Iterator1 keys_first, Iterator1
 struct make_pair_functor
 {
   template <typename T1, typename T2>
-  _CCCL_HOST_DEVICE thrust::pair<T1, T2> operator()(const T1& x, const T2& y)
+  _CCCL_HOST_DEVICE cuda::std::pair<T1, T2> operator()(const T1& x, const T2& y)
   {
-    return thrust::make_pair(x, y);
+    return cuda::std::make_pair(x, y);
   } // end operator()()
 }; // end make_pair_functor
 
 template <typename ExecutionPolicy>
 void TestPairStableSortByKeyDevice(ExecutionPolicy exec)
 {
   size_t n = 10000;
-  using P  = thrust::pair<int, int>;
+  using P  = cuda::std::pair<int, int>;
 
   // host arrays
   thrust::host_vector<int> h_p1 = unittest::random_integers<int>(n);

thrust/testing/cuda/partition.cu

@@ -187,7 +187,7 @@ void TestPartitionCopyDevice(ExecutionPolicy exec)
   thrust::device_vector<int> true_results(2);
   thrust::device_vector<int> false_results(3);
 
-  using pair_type = thrust::pair<iterator, iterator>;
+  using pair_type = cuda::std::pair<iterator, iterator>;
   thrust::device_vector<pair_type> iterators(1);
 
   partition_copy_kernel<<<1, 1>>>(
@@ -273,7 +273,7 @@ void TestPartitionCopyStencilDevice(ExecutionPolicy exec)
   thrust::device_vector<int> false_results(3);
 
   using iterator  = typename thrust::device_vector<int>::iterator;
-  using pair_type = thrust::pair<iterator, iterator>;
+  using pair_type = cuda::std::pair<iterator, iterator>;
   thrust::device_vector<pair_type> iterators(1);
 
   partition_copy_kernel<<<1, 1>>>(
@@ -474,7 +474,7 @@ void TestStablePartitionCopyDevice(ExecutionPolicy exec)
   thrust::device_vector<int> true_results(2);
   thrust::device_vector<int> false_results(3);
 
-  using pair_type = thrust::pair<iterator, iterator>;
+  using pair_type = cuda::std::pair<iterator, iterator>;
   thrust::device_vector<pair_type> iterators(1);
 
   stable_partition_copy_kernel<<<1, 1>>>(
@@ -560,7 +560,7 @@ void TestStablePartitionCopyStencilDevice(ExecutionPolicy exec)
   thrust::device_vector<int> false_results(3);
 
   using iterator  = typename thrust::device_vector<int>::iterator;
-  using pair_type = thrust::pair<iterator, iterator>;
+  using pair_type = cuda::std::pair<iterator, iterator>;
   thrust::device_vector<pair_type> iterators(1);
 
   stable_partition_copy_kernel<<<1, 1>>>(

thrust/testing/cuda/reduce_by_key.cu

@@ -164,7 +164,7 @@ void TestReduceByKeyDevice(ExecutionPolicy exec)
   thrust::device_vector<T> values;
 
   using iterator_pair =
-    typename thrust::pair<typename thrust::device_vector<T>::iterator, typename thrust::device_vector<T>::iterator>;
+    typename cuda::std::pair<typename thrust::device_vector<T>::iterator, typename thrust::device_vector<T>::iterator>;
 
   thrust::device_vector<iterator_pair> new_last_vec(1);
   iterator_pair new_last;
@@ -293,7 +293,7 @@ void TestReduceByKeyCudaStreams(ExecutionPolicy policy)
   Vector keys;
   Vector values;
 
-  thrust::pair<Vector::iterator, Vector::iterator> new_last;
+  cuda::std::pair<Vector::iterator, Vector::iterator> new_last;
 
   // basic test
   initialize_keys(keys);

thrust/testing/cuda/set_difference_by_key.cu

@@ -41,7 +41,7 @@ void TestSetDifferenceByKeyDevice(ExecutionPolicy exec)
   Vector ref_key{2, 5}, ref_val{0, 0};
   Vector result_key(2), result_val(2);
 
-  using iter_pair = thrust::pair<Iterator, Iterator>;
+  using iter_pair = cuda::std::pair<Iterator, Iterator>;
 
   thrust::device_vector<iter_pair> end_vec(1);
 
@@ -94,7 +94,7 @@ void TestSetDifferenceByKeyCudaStreams()
   cudaStream_t s;
   cudaStreamCreate(&s);
 
-  thrust::pair<Iterator, Iterator> end = thrust::set_difference_by_key(
+  cuda::std::pair<Iterator, Iterator> end = thrust::set_difference_by_key(
     thrust::cuda::par.on(s),
     a_key.begin(),
     a_key.end(),