Replace internal uses of thrust::tuple with cuda::std::tuple

Author: miscco

cub/test/catch2_test_device_merge.cu

@@ -250,7 +250,7 @@ void test_pairs(
                zip(keys2_h.end(), values2_h.end()),
                zip(reference_keys_h.begin(), reference_values_h.begin()),
                [&](const value_t& a, const value_t& b) {
-                 return compare_op(thrust::get<0>(a), thrust::get<0>(b));
+                 return compare_op(cuda::std::get<0>(a), cuda::std::get<0>(b));
                });
   }
 

cub/test/catch2_test_device_merge_sort.cu

@@ -93,11 +93,11 @@ template <typename CustomT>
 struct tuple_to_custom_op_t
 {
   template <typename KeyT, typename ValueT>
-  __device__ __host__ CustomT operator()(const thrust::tuple<KeyT, ValueT>& val)
+  __device__ __host__ CustomT operator()(const cuda::std::tuple<KeyT, ValueT>& val)
   {
     CustomT custom_val{};
-    custom_val.key = static_cast<std::size_t>(thrust::get<0>(val));
-    custom_val.val = static_cast<std::size_t>(thrust::get<1>(val));
+    custom_val.key = static_cast<std::size_t>(cuda::std::get<0>(val));
+    custom_val.val = static_cast<std::size_t>(cuda::std::get<1>(val));
     return custom_val;
   }
 };

cub/test/catch2_test_device_merge_sort_common.cuh

@@ -22,12 +22,12 @@ struct compare_first_lt_op_t
 {
   /**
    * We need to be able to have two different types for lhs and rhs, as the call to std::stable_sort with a
-   * zip-iterator, will pass a thrust::tuple for lhs and a tuple_of_iterator_references for rhs.
+   * zip-iterator, will pass a cuda::std::tuple for lhs and a tuple_of_iterator_references for rhs.
    */
   template <typename LhsT, typename RhsT>
   __host__ __device__ bool operator()(const LhsT& lhs, const RhsT& rhs) const
   {
-    return thrust::get<0>(lhs) < thrust::get<0>(rhs);
+    return cuda::std::get<0>(lhs) < cuda::std::get<0>(rhs);
   }
 };
 

cub/test/catch2_test_device_merge_sort_iterators.cu

@@ -62,12 +62,12 @@ C2H_TEST("DeviceMergeSort::StableSortKeysCopy works with iterators and is stable
   auto keys_in_it          = cuda::make_zip_iterator(sort_key_it, key_idx_it);
 
   // Perform sort
-  c2h::device_vector<thrust::tuple<key_t, offset_t>> keys_out(
-    num_items, thrust::tuple<key_t, offset_t>{static_cast<key_t>(42), static_cast<offset_t>(42)});
+  c2h::device_vector<cuda::std::tuple<key_t, offset_t>> keys_out(
+    num_items, cuda::std::tuple<key_t, offset_t>{static_cast<key_t>(42), static_cast<offset_t>(42)});
   stable_sort_keys_copy(keys_in_it, keys_out.begin(), num_items, compare_first_lt_op_t{});
 
   // Verify results
-  c2h::host_vector<thrust::tuple<key_t, offset_t>> keys_expected(num_items);
+  c2h::host_vector<cuda::std::tuple<key_t, offset_t>> keys_expected(num_items);
   thrust::copy(keys_in_it, keys_in_it + num_items, keys_expected.begin());
   std::stable_sort(keys_expected.begin(), keys_expected.end(), compare_first_lt_op_t{});
 

cub/test/catch2_test_device_select_flagged_if.cu

@@ -19,9 +19,9 @@ struct predicate_op_wrapper_t
 {
   PredOpT if_pred;
   template <typename FlagT, typename ItemT>
-  __host__ __device__ bool operator()(thrust::tuple<FlagT, ItemT> tuple) const
+  __host__ __device__ bool operator()(cuda::std::tuple<FlagT, ItemT> tuple) const
   {
-    const auto flag = thrust::get<0>(tuple);
+    const auto flag = cuda::std::get<0>(tuple);
     return static_cast<bool>(if_pred(flag));
   }
 };

cub/test/catch2_test_device_select_unique_by_key.cu

@@ -228,7 +228,7 @@ struct project_first
   template <typename Tuple>
   __host__ __device__ bool operator()(const Tuple& lhs, const Tuple& rhs) const
   {
-    return equality_op(thrust::get<0>(lhs), thrust::get<0>(rhs));
+    return equality_op(cuda::std::get<0>(lhs), cuda::std::get<0>(rhs));
   }
 };
 

cub/test/catch2_test_warp_merge_sort.cu

@@ -515,7 +515,7 @@ C2H_TEST("Warp sort on key-value pairs of a partial warp-tile works",
     cpu_kv_pairs,
     segment_sizes,
     params::total_warps,
-    thrust::make_tuple(oob_default, value_type{}),
+    cuda::std::tuple(oob_default, value_type{}),
     params::logical_warp_items);
 
   // Verify results

cub/test/test_device_batch_copy.cu

@@ -6,9 +6,9 @@
 
 #include <thrust/logical.h>
 #include <thrust/sequence.h>
-#include <thrust/tuple.h>
 
 #include <cuda/iterator>
+#include <cuda/std/tuple>
 
 #include <algorithm>
 #include <cstdint>
@@ -90,19 +90,19 @@ GetShuffledRangeOffsets(const c2h::host_vector<RangeSizeT>& range_sizes, const s
 }
 
 template <size_t n, typename... T>
-std::enable_if_t<n >= thrust::tuple_size<thrust::tuple<T...>>::value>
-print_tuple(std::ostream&, const thrust::tuple<T...>&)
+std::enable_if_t<n >= cuda::std::tuple_size<cuda::std::tuple<T...>>::value>
+print_tuple(std::ostream&, const cuda::std::tuple<T...>&)
 {}
 
 template <size_t n, typename... T>
-std::enable_if_t<n + 1 <= thrust::tuple_size<thrust::tuple<T...>>::value>
-print_tuple(std::ostream& os, const thrust::tuple<T...>& tup)
+std::enable_if_t<n + 1 <= cuda::std::tuple_size<cuda::std::tuple<T...>>::value>
+print_tuple(std::ostream& os, const cuda::std::tuple<T...>& tup)
 {
   if constexpr (n != 0)
   {
     os << ", ";
   }
-  os << thrust::get<n>(tup);
+  os << cuda::std::get<n>(tup);
   print_tuple<n + 1>(os, tup);
 }
 
@@ -399,7 +399,7 @@ int main(int argc, char** argv)
   for (const auto& size_range : size_ranges)
   {
     // The most granular type being copied.
-    using AtomicCopyT         = thrust::tuple<int64_t, int32_t, int16_t, char, char>;
+    using AtomicCopyT         = cuda::std::tuple<int64_t, int32_t, int16_t, char, char>;
     RangeSizeT min_range_size = static_cast<RangeSizeT>(cuda::round_up(size_range.first, sizeof(AtomicCopyT)));
     RangeSizeT max_range_size =
       static_cast<RangeSizeT>(cuda::round_up(size_range.second, static_cast<RangeSizeT>(sizeof(AtomicCopyT))));

cudax/examples/stf/thrust_zip_iterator.cu

@@ -30,10 +30,10 @@ using namespace cuda::experimental::stf;
 // Functor to apply the transformation
 struct my_transform_functor
 {
-  __host__ __device__ int operator()(const thrust::tuple<int, char>& t) const
+  __host__ __device__ int operator()(const cuda::std::tuple<int, char>& t) const
   {
-    int a  = thrust::get<0>(t);
-    char b = thrust::get<1>(t);
+    int a  = cuda::std::get<0>(t);
+    char b = cuda::std::get<1>(t);
     return a + static_cast<int>(b); // Example operation
   }
 };
@@ -50,10 +50,10 @@ void thrust_algorithm(context& ctx, ZippedIt& first, ZippedIt& last, OutIt& outp
   size_t num_elements = cuda::std::distance(first, last);
 
   // Extract underlying iterators from the zip iterator
-  auto itA = thrust::get<0>(first.get_iterator_tuple());
+  auto itA = cuda::std::get<0>(first.get_iterator_tuple());
   int* A   = thrust::raw_pointer_cast(&(*itA));
 
-  auto itB = thrust::get<1>(first.get_iterator_tuple());
+  auto itB = cuda::std::get<1>(first.get_iterator_tuple());
   char* B  = thrust::raw_pointer_cast(&(*itB));
 
   int* C = thrust::raw_pointer_cast(output.data());
@@ -66,7 +66,7 @@ void thrust_algorithm(context& ctx, ZippedIt& first, ZippedIt& last, OutIt& outp
   ctx.task(lA.read(), lB.read(), lC.write())->*[](cudaStream_t stream, auto dA, auto dB, auto dC) {
     // Reconstruct a zipped iterator from the data instances passed to the lambda function
     size_t num_elements = dA.size();
-    auto dfirst         = thrust::make_zip_iterator(thrust::make_tuple(dA.data_handle(), dB.data_handle()));
+    auto dfirst         = thrust::make_zip_iterator(cuda::std::tuple(dA.data_handle(), dB.data_handle()));
     auto dlast          = dfirst + num_elements;
 
     // Create a device pointer from the raw pointer
@@ -94,8 +94,8 @@ int main()
   B[1] = 'y';
   B[2] = 'z';
 
-  auto first = thrust::make_zip_iterator(thrust::make_tuple(A.begin(), B.begin()));
-  auto last  = thrust::make_zip_iterator(thrust::make_tuple(A.end(), B.end()));
+  auto first = thrust::make_zip_iterator(cuda::std::tuple(A.begin(), B.begin()));
+  auto last  = thrust::make_zip_iterator(cuda::std::tuple(A.end(), B.end()));
 
   thrust_algorithm(ctx, first, last, C, data_place::current_device());
 
@@ -114,8 +114,8 @@ int main()
   hB[1] = 'y';
   hB[2] = 'z';
 
-  auto hfirst = thrust::make_zip_iterator(thrust::make_tuple(hA.begin(), hB.begin()));
-  auto hlast  = thrust::make_zip_iterator(thrust::make_tuple(hA.end(), hB.end()));
+  auto hfirst = thrust::make_zip_iterator(cuda::std::tuple(hA.begin(), hB.begin()));
+  auto hlast  = thrust::make_zip_iterator(cuda::std::tuple(hA.end(), hB.end()));
 
   thrust_algorithm(ctx, hfirst, hlast, hC, data_place::host());
 

thrust/examples/arbitrary_transformation.cu

@@ -45,7 +45,7 @@ struct arbitrary_functor1
   __host__ __device__ void operator()(Tuple t)
   {
     // D[i] = A[i] + B[i] * C[i];
-    thrust::get<3>(t) = thrust::get<0>(t) + thrust::get<1>(t) * thrust::get<2>(t);
+    cuda::std::get<3>(t) = cuda::std::get<0>(t) + cuda::std::get<1>(t) * cuda::std::get<2>(t);
   }
 };