Replace thrust::tuple with cuda::std::tuple

cpp/benchmarks/common/generate_input.cu

@@ -33,6 +33,7 @@
 #include <rmm/device_uvector.hpp>
 #include <rmm/mr/device/per_device_resource.hpp>
 
+#include <cuda/std/tuple>
 #include <thrust/binary_search.h>
 #include <thrust/copy.h>
 #include <thrust/device_ptr.h>
@@ -51,7 +52,6 @@
 #include <thrust/scan.h>
 #include <thrust/tabulate.h>
 #include <thrust/transform.h>
-#include <thrust/tuple.h>
 
 #include <cuda/functional>
 
@@ -464,9 +464,9 @@ std::unique_ptr<cudf::column> create_random_column(data_profile const& profile,
 
 struct valid_or_zero {
   template <typename T>
-  __device__ T operator()(thrust::tuple<T, bool> len_valid) const
+  __device__ T operator()(cuda::std::tuple<T, bool> len_valid) const
   {
-    return thrust::get<1>(len_valid) ? thrust::get<0>(len_valid) : T{0};
+    return cuda::std::get<1>(len_valid) ? cuda::std::get<0>(len_valid) : T{0};
   }
 };
 
@@ -481,10 +481,10 @@ struct string_generator {
   // range 32-127 is ASCII; 127-136 will be multi-byte UTF-8
   {
   }
-  __device__ void operator()(thrust::tuple<cudf::size_type, cudf::size_type> str_begin_end)
+  __device__ void operator()(cuda::std::tuple<cudf::size_type, cudf::size_type> str_begin_end)
   {
-    auto begin = thrust::get<0>(str_begin_end);
-    auto end   = thrust::get<1>(str_begin_end);
+    auto begin = cuda::std::get<0>(str_begin_end);
+    auto end   = cuda::std::get<1>(str_begin_end);
     engine.discard(begin);
     for (auto i = begin; i < end; ++i) {
       auto ch = char_dist(engine);
@@ -519,7 +519,7 @@ std::unique_ptr<cudf::column> create_random_utf8_string_column(data_profile cons
     cuda::proclaim_return_type<cudf::size_type>([] __device__(auto) { return 0; }),
     thrust::logical_not<bool>{});
   auto valid_lengths = thrust::make_transform_iterator(
-    thrust::make_zip_iterator(thrust::make_tuple(lengths.begin(), null_mask.begin())),
+    thrust::make_zip_iterator(cuda::std::make_tuple(lengths.begin(), null_mask.begin())),
     valid_or_zero{});
   rmm::device_uvector<cudf::size_type> offsets(num_rows + 1, cudf::get_default_stream());
   thrust::exclusive_scan(

cpp/benchmarks/string/url_decode.cu

@@ -26,12 +26,12 @@
 #include <cudf/types.hpp>
 #include <cudf/utilities/default_stream.hpp>
 
+#include <cuda/std/tuple>
 #include <thrust/execution_policy.h>
 #include <thrust/for_each.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/random.h>
-#include <thrust/tuple.h>
 
 struct url_string_generator {
   char* chars;
@@ -43,10 +43,10 @@ struct url_string_generator {
   {
   }
 
-  __device__ void operator()(thrust::tuple<cudf::size_type, cudf::size_type> str_begin_end)
+  __device__ void operator()(cuda::std::tuple<cudf::size_type, cudf::size_type> str_begin_end)
   {
-    auto begin = thrust::get<0>(str_begin_end);
-    auto end   = thrust::get<1>(str_begin_end);
+    auto begin = cuda::std::get<0>(str_begin_end);
+    auto end   = cuda::std::get<1>(str_begin_end);
     engine.discard(begin);
     for (auto i = begin; i < end; ++i) {
       if (esc_seq_dist(engine) < esc_seq_chance and i < end - 3) {

cpp/include/cudf/detail/merge.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2018-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -29,8 +29,8 @@ namespace detail {
 enum class side : bool { LEFT, RIGHT };
 
 /**
- * @brief Tagged index type: `thrust::get<0>` indicates left/right side,
- * `thrust::get<1>` indicates the row index
+ * @brief Tagged index type: `cuda::std::get<0>` indicates left/right side,
+ * `cuda::std::get<1>` indicates the row index
  */
 using index_type = thrust::pair<side, cudf::size_type>;
 

cpp/include/cudf/detail/null_mask.cuh

@@ -30,12 +30,12 @@
 #include <cub/block/block_reduce.cuh>
 #include <cub/device/device_segmented_reduce.cuh>
 
+#include <cuda/std/tuple>
 #include <thrust/for_each.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/transform.h>
-#include <thrust/tuple.h>
 
 #include <cuda/functional>
 
@@ -334,8 +334,8 @@ rmm::device_uvector<size_type> segmented_count_bits(bitmask_type const* bitmask,
       thrust::make_zip_iterator(first_bit_indices_begin, last_bit_indices_begin);
     auto segment_length_iterator = thrust::transform_iterator(
       segments_begin, cuda::proclaim_return_type<size_type>([] __device__(auto const& segment) {
-        auto const begin = thrust::get<0>(segment);
-        auto const end   = thrust::get<1>(segment);
+        auto const begin = cuda::std::get<0>(segment);
+        auto const end   = cuda::std::get<1>(segment);
         return end - begin;
       }));
     thrust::transform(rmm::exec_policy(stream),
@@ -546,8 +546,8 @@ std::pair<rmm::device_buffer, size_type> segmented_null_mask_reduction(
     thrust::make_zip_iterator(first_bit_indices_begin, last_bit_indices_begin);
   auto const segment_length_iterator = thrust::make_transform_iterator(
     segments_begin, cuda::proclaim_return_type<size_type>([] __device__(auto const& segment) {
-      auto const begin = thrust::get<0>(segment);
-      auto const end   = thrust::get<1>(segment);
+      auto const begin = cuda::std::get<0>(segment);
+      auto const end   = cuda::std::get<1>(segment);
       return end - begin;
     }));
 
@@ -579,8 +579,8 @@ std::pair<rmm::device_buffer, size_type> segmented_null_mask_reduction(
     length_and_valid_count,
     length_and_valid_count + num_segments,
     [null_handling, valid_initial_value] __device__(auto const& length_and_valid_count) {
-      auto const length      = thrust::get<0>(length_and_valid_count);
-      auto const valid_count = thrust::get<1>(length_and_valid_count);
+      auto const length      = cuda::std::get<0>(length_and_valid_count);
+      auto const valid_count = cuda::std::get<1>(length_and_valid_count);
       return (null_handling == null_policy::EXCLUDE)
                ? (valid_initial_value.value_or(false) || valid_count > 0)
                : (valid_initial_value.value_or(length > 0) && valid_count == length);

cpp/include/cudf/detail/replace/nulls.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -17,13 +17,13 @@
 
 #include <cudf/types.hpp>
 
+#include <cuda/std/tuple>
 #include <thrust/functional.h>
-#include <thrust/tuple.h>
 
 namespace cudf {
 namespace detail {
 
-using idx_valid_pair_t = thrust::tuple<cudf::size_type, bool>;
+using idx_valid_pair_t = cuda::std::tuple<cudf::size_type, bool>;
 
 /**
  * @brief Functor used by `replace_nulls(replace_policy)` to determine the index to gather from in
@@ -36,7 +36,7 @@ using idx_valid_pair_t = thrust::tuple<cudf::size_type, bool>;
 struct replace_policy_functor {
   __device__ idx_valid_pair_t operator()(idx_valid_pair_t const& lhs, idx_valid_pair_t const& rhs)
   {
-    return thrust::get<1>(rhs) ? rhs : lhs;
+    return cuda::std::get<1>(rhs) ? rhs : lhs;
   }
 };
 

cpp/include/cudf/strings/detail/merge.cuh

@@ -27,10 +27,10 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/std/tuple>
 #include <thrust/for_each.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
-#include <thrust/tuple.h>
 
 #include <cuda/functional>
 
@@ -42,7 +42,7 @@ namespace detail {
  *
  * Caller must set the validity mask in the output column.
  *
- * @tparam row_order_iterator This must be an iterator for type thrust::tuple<side,size_type>.
+ * @tparam row_order_iterator This must be an iterator for type cuda::std::tuple<side,size_type>.
  *
  * @param lhs First column.
  * @param rhs Second column.

cpp/include/cudf/strings/detail/strings_column_factories.cuh

@@ -26,6 +26,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/distance.h>
 #include <thrust/for_each.h>
@@ -35,7 +36,6 @@
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/pair.h>
 #include <thrust/transform.h>
-#include <thrust/tuple.h>
 
 #include <cuda/functional>
 
@@ -120,13 +120,13 @@ std::unique_ptr<column> make_strings_column(IndexPairIterator begin,
       auto chars_data = rmm::device_uvector<char>(bytes, stream, mr);
       auto d_chars    = chars_data.data();
       auto copy_chars = [d_chars] __device__(auto item) {
-        string_index_pair const str = thrust::get<0>(item);
-        size_type const offset      = thrust::get<1>(item);
+        string_index_pair const str = cuda::std::get<0>(item);
+        size_type const offset      = cuda::std::get<1>(item);
         if (str.first != nullptr) memcpy(d_chars + offset, str.first, str.second);
       };
       thrust::for_each_n(rmm::exec_policy(stream),
                          thrust::make_zip_iterator(
-                           thrust::make_tuple(begin, offsets_view.template begin<size_type>())),
+                           cuda::std::make_tuple(begin, offsets_view.template begin<size_type>())),
                          strings_count,
                          copy_chars);
       return chars_data;

cpp/include/cudf_test/tdigest_utilities.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -27,20 +27,20 @@
 #include <rmm/exec_policy.hpp>
 #include <rmm/mr/device/per_device_resource.hpp>
 
+#include <cuda/std/tuple>
 #include <thrust/device_vector.h>
 #include <thrust/execution_policy.h>
 #include <thrust/extrema.h>
 #include <thrust/host_vector.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/transform.h>
-#include <thrust/tuple.h>
 
 // for use with groupby and reduction aggregation tests.
 
 namespace cudf {
 namespace test {
 
-using expected_value = thrust::tuple<size_type, double, double>;
+using expected_value = cuda::std::tuple<size_type, double, double>;
 
 /**
  * @brief Device functor to compute min of a sequence of values serially.

cpp/src/copying/contiguous_split.cu

@@ -32,6 +32,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/std/tuple>
 #include <thrust/binary_search.h>
 #include <thrust/execution_policy.h>
 #include <thrust/for_each.h>
@@ -43,7 +44,6 @@
 #include <thrust/reduce.h>
 #include <thrust/scan.h>
 #include <thrust/transform.h>
-#include <thrust/tuple.h>
 
 #include <cuda/functional>
 
@@ -947,7 +947,7 @@ struct size_of_helper {
  */
 struct num_batches_func {
   thrust::pair<std::size_t, std::size_t> const* const batches;
-  __device__ std::size_t operator()(size_type i) const { return thrust::get<0>(batches[i]); }
+  __device__ std::size_t operator()(size_type i) const { return cuda::std::get<0>(batches[i]); }
 };
 
 /**
@@ -1466,7 +1466,7 @@ std::unique_ptr<chunk_iteration_state> chunk_iteration_state::create(
      out_to_in_index] __device__(size_type i) {
       size_type const in_buf_index = out_to_in_index(i);
       size_type const batch_index  = i - d_batch_offsets[in_buf_index];
-      auto const batch_size        = thrust::get<1>(batches[in_buf_index]);
+      auto const batch_size        = cuda::std::get<1>(batches[in_buf_index]);
       dst_buf_info const& in       = d_orig_dst_buf_info[in_buf_index];
 
       // adjust info

cpp/src/groupby/sort/group_correlation.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -27,13 +27,13 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/std/tuple>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/reduce.h>
 #include <thrust/transform.h>
-#include <thrust/tuple.h>
 
 #include <type_traits>
 
@@ -187,7 +187,7 @@ std::unique_ptr<column> group_correlation(column_view const& covariance,
   CUDF_EXPECTS(covariance.type().id() == type_id::FLOAT64, "Covariance result must be FLOAT64");
   auto stddev0_ptr = stddev_0.begin<result_type>();
   auto stddev1_ptr = stddev_1.begin<result_type>();
-  auto stddev_iter = thrust::make_zip_iterator(thrust::make_tuple(stddev0_ptr, stddev1_ptr));
+  auto stddev_iter = thrust::make_zip_iterator(cuda::std::make_tuple(stddev0_ptr, stddev1_ptr));
   auto result      = make_numeric_column(covariance.type(),
                                     covariance.size(),
                                     cudf::detail::copy_bitmask(covariance, stream, mr),
@@ -201,7 +201,7 @@ std::unique_ptr<column> group_correlation(column_view const& covariance,
                     stddev_iter,
                     d_result,
                     [] __device__(auto const covariance, auto const stddev) {
-                      return covariance / thrust::get<0>(stddev) / thrust::get<1>(stddev);
+                      return covariance / cuda::std::get<0>(stddev) / cuda::std::get<1>(stddev);
                     });
 
   result->set_null_count(covariance.null_count());

cpp/src/groupby/sort/group_merge_m2.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -26,13 +26,13 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/std/tuple>
 #include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/reduce.h>
 #include <thrust/transform.h>
-#include <thrust/tuple.h>
 
 namespace cudf {
 namespace groupby {
@@ -90,7 +90,7 @@ struct merge_fn {
     // This case should never happen, because all groups are non-empty as the results of
     // aggregation. Here we just to make sure we cover this case.
     if (start_idx == end_idx) {
-      return thrust::make_tuple(size_type{0}, result_type{0}, result_type{0}, int8_t{0});
+      return cuda::std::make_tuple(size_type{0}, result_type{0}, result_type{0}, int8_t{0});
     }
 
     // If `(n = d_counts[idx]) > 0` then `d_means[idx] != null` and `d_M2s[idx] != null`.
@@ -121,7 +121,7 @@ struct merge_fn {
     // zero), then the output is a null.
     auto const is_valid = int8_t{merge_vals.count > 0};
 
-    return thrust::make_tuple(merge_vals.count, merge_vals.mean, merge_vals.M2, is_valid);
+    return cuda::std::make_tuple(merge_vals.count, merge_vals.mean, merge_vals.M2, is_valid);
   }
 };
 
@@ -157,13 +157,13 @@ std::unique_ptr<column> group_merge_m2(column_view const& values,
 
   // Perform merging for all the aggregations. Their output (and their validity data) are written
   // out concurrently through an output zip iterator.
-  using iterator_tuple  = thrust::tuple<size_type*, result_type*, result_type*, int8_t*>;
+  using iterator_tuple  = cuda::std::tuple<size_type*, result_type*, result_type*, int8_t*>;
   using output_iterator = thrust::zip_iterator<iterator_tuple>;
   auto const out_iter =
-    output_iterator{thrust::make_tuple(result_counts->mutable_view().template data<size_type>(),
-                                       result_means->mutable_view().template data<result_type>(),
-                                       result_M2s->mutable_view().template data<result_type>(),
-                                       validities.begin())};
+    output_iterator{cuda::std::make_tuple(result_counts->mutable_view().template data<size_type>(),
+                                          result_means->mutable_view().template data<result_type>(),
+                                          result_M2s->mutable_view().template data<result_type>(),
+                                          validities.begin())};
 
   auto const count_valid = values.child(0);
   auto const mean_values = values.child(1);