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

cpp/src/hdbscan/condensed_hierarchy.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -16,6 +16,7 @@
 
 #include <cub/cub.cuh>
 #include <cuda/functional>
+#include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/device_ptr.h>
 #include <thrust/execution_policy.h>
@@ -26,7 +27,6 @@
 #include <thrust/sort.h>
 #include <thrust/transform.h>
 #include <thrust/transform_reduce.h>
-#include <thrust/tuple.h>
 
 namespace ML {
 namespace HDBSCAN {
@@ -37,15 +37,15 @@ struct TupleComp {
   __host__ __device__ bool operator()(const one& t1, const two& t2)
   {
     // sort first by each parent,
-    if (thrust::get<0>(t1) < thrust::get<0>(t2)) return true;
-    if (thrust::get<0>(t1) > thrust::get<0>(t2)) return false;
+    if (cuda::std::get<0>(t1) < cuda::std::get<0>(t2)) return true;
+    if (cuda::std::get<0>(t1) > cuda::std::get<0>(t2)) return false;
 
     // within each parent, sort by each child,
-    if (thrust::get<1>(t1) < thrust::get<1>(t2)) return true;
-    if (thrust::get<1>(t1) > thrust::get<1>(t2)) return false;
+    if (cuda::std::get<1>(t1) < cuda::std::get<1>(t2)) return true;
+    if (cuda::std::get<1>(t1) > cuda::std::get<1>(t2)) return false;
 
     // then sort by value in descending order
-    return thrust::get<2>(t1) < thrust::get<2>(t2);
+    return cuda::std::get<2>(t1) < cuda::std::get<2>(t2);
   }
 };
 
@@ -91,9 +91,9 @@ CondensedHierarchy<value_idx, value_t>::CondensedHierarchy(const raft::handle_t&
 
   n_clusters = max_cluster - min_cluster + 1;
 
-  auto sort_keys =
-    thrust::make_zip_iterator(thrust::make_tuple(parents.begin(), children.begin(), sizes.begin()));
-  auto sort_values = thrust::make_zip_iterator(thrust::make_tuple(lambdas.begin()));
+  auto sort_keys = thrust::make_zip_iterator(
+    cuda::std::make_tuple(parents.begin(), children.begin(), sizes.begin()));
+  auto sort_values = thrust::make_zip_iterator(cuda::std::make_tuple(lambdas.begin()));
 
   thrust::sort_by_key(thrust::cuda::par.on(handle.get_stream()),
                       sort_keys,
@@ -157,17 +157,17 @@ void CondensedHierarchy<value_idx, value_t>::condense(value_idx* full_parents,
   sizes.resize(n_edges, stream);
 
   auto in = thrust::make_zip_iterator(
-    thrust::make_tuple(full_parents, full_children, full_lambdas, full_sizes));
+    cuda::std::make_tuple(full_parents, full_children, full_lambdas, full_sizes));
 
   auto out = thrust::make_zip_iterator(
-    thrust::make_tuple(parents.data(), children.data(), lambdas.data(), sizes.data()));
+    cuda::std::make_tuple(parents.data(), children.data(), lambdas.data(), sizes.data()));
 
   thrust::copy_if(thrust::cuda::par.on(stream),
                   in,
                   in + size,
                   out,
-                  [=] __device__(thrust::tuple<value_idx, value_idx, value_t, value_idx> tup) {
-                    return thrust::get<3>(tup) != -1;
+                  [=] __device__(cuda::std::tuple<value_idx, value_idx, value_t, value_idx> tup) {
+                    return cuda::std::get<3>(tup) != -1;
                   });
 
   auto parents_ptr = thrust::device_pointer_cast(parents.data());
@@ -179,9 +179,9 @@ void CondensedHierarchy<value_idx, value_t>::condense(value_idx* full_parents,
 
   n_clusters = max_cluster - min_cluster + 1;
 
-  auto sort_keys =
-    thrust::make_zip_iterator(thrust::make_tuple(parents.begin(), children.begin(), sizes.begin()));
-  auto sort_values = thrust::make_zip_iterator(thrust::make_tuple(lambdas.begin()));
+  auto sort_keys = thrust::make_zip_iterator(
+    cuda::std::make_tuple(parents.begin(), children.begin(), sizes.begin()));
+  auto sort_values = thrust::make_zip_iterator(cuda::std::make_tuple(lambdas.begin()));
 
   thrust::sort_by_key(
     thrust::cuda::par.on(stream), sort_keys, sort_keys + n_edges, sort_values, TupleComp());

cpp/src/hdbscan/detail/reachability.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -20,7 +20,6 @@
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/transform.h>
-#include <thrust/tuple.h>
 
 namespace ML {
 namespace HDBSCAN {

cpp/src/hdbscan/detail/select.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -20,6 +20,7 @@
 
 #include <cub/cub.cuh>
 #include <cuda/std/functional>
+#include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/execution_policy.h>
 #include <thrust/fill.h>
@@ -31,7 +32,6 @@
 #include <thrust/sort.h>
 #include <thrust/transform.h>
 #include <thrust/transform_reduce.h>
-#include <thrust/tuple.h>
 
 #include <algorithm>
 
@@ -170,15 +170,16 @@ void excess_of_mass(const raft::handle_t& handle,
 
   value_idx* cluster_sizes_ptr = cluster_sizes.data();
 
-  auto out = thrust::make_zip_iterator(thrust::make_tuple(parents, children, sizes));
+  auto out = thrust::make_zip_iterator(cuda::std::make_tuple(parents, children, sizes));
   thrust::for_each(exec_policy,
                    out,
                    out + cluster_tree_edges,
-                   [=] __device__(const thrust::tuple<value_idx, value_idx, value_idx>& tup) {
+                   [=] __device__(const cuda::std::tuple<value_idx, value_idx, value_idx>& tup) {
                      // if parent is root (0), add to cluster_sizes_ptr
-                     if (thrust::get<0>(tup) == 0) cluster_sizes_ptr[0] += thrust::get<2>(tup);
+                     if (cuda::std::get<0>(tup) == 0)
+                       cluster_sizes_ptr[0] += cuda::std::get<2>(tup);
 
-                     cluster_sizes_ptr[thrust::get<1>(tup)] = thrust::get<2>(tup);
+                     cluster_sizes_ptr[cuda::std::get<1>(tup)] = cuda::std::get<2>(tup);
                    });
 
   /**
@@ -325,7 +326,7 @@ void cluster_epsilon_search(const raft::handle_t& handle,
                   [] __device__(auto cluster) { return cluster; });
 
   // sort lambdas and parents by children for epsilon search
-  auto start = thrust::make_zip_iterator(thrust::make_tuple(parents, lambdas));
+  auto start = thrust::make_zip_iterator(cuda::std::make_tuple(parents, lambdas));
   thrust::sort_by_key(thrust_policy, children, children + cluster_tree_edges, start);
   rmm::device_uvector<value_t> eps(cluster_tree_edges, stream);
   thrust::transform(

cpp/src/hdbscan/detail/stabilities.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -19,6 +19,7 @@
 #include <rmm/exec_policy.hpp>
 
 #include <cub/cub.cuh>
+#include <cuda/std/tuple>
 #include <thrust/execution_policy.h>
 #include <thrust/fill.h>
 #include <thrust/for_each.h>
@@ -27,7 +28,6 @@
 #include <thrust/reduce.h>
 #include <thrust/sort.h>
 #include <thrust/transform.h>
-#include <thrust/tuple.h>
 
 #include <algorithm>
 
@@ -115,10 +115,10 @@ void compute_stabilities(const raft::handle_t& handle,
   // finally, we find minimum between initialized births where parent=child
   // and births of parents for their children
   auto births_zip =
-    thrust::make_zip_iterator(thrust::make_tuple(births.data(), births_parent_min.data()));
-  auto min_op = [] __device__(const thrust::tuple<value_t, value_t>& birth_pair) {
-    auto birth             = thrust::get<0>(birth_pair);
-    auto births_parent_min = thrust::get<1>(birth_pair);
+    thrust::make_zip_iterator(cuda::std::make_tuple(births.data(), births_parent_min.data()));
+  auto min_op = [] __device__(const cuda::std::tuple<value_t, value_t>& birth_pair) {
+    auto birth             = cuda::std::get<0>(birth_pair);
+    auto births_parent_min = cuda::std::get<1>(birth_pair);
 
     return birth < births_parent_min ? birth : births_parent_min;
   };
@@ -179,13 +179,13 @@ void get_stability_scores(const raft::handle_t& handle,
    */
 
   auto enumeration = thrust::make_zip_iterator(
-    thrust::make_tuple(thrust::make_counting_iterator(0), cluster_sizes.data()));
+    cuda::std::make_tuple(thrust::make_counting_iterator(0), cluster_sizes.data()));
   thrust::for_each(exec_policy,
                    enumeration,
                    enumeration + n_condensed_clusters,
-                   [=] __device__(thrust::tuple<value_idx, value_idx> tup) {
-                     value_idx size        = thrust::get<1>(tup);
-                     value_idx c           = thrust::get<0>(tup);
+                   [=] __device__(cuda::std::tuple<value_idx, value_idx> tup) {
+                     value_idx size        = cuda::std::get<1>(tup);
+                     value_idx c           = cuda::std::get<0>(tup);
                      value_idx out_cluster = label_map[c];
 
                      if (out_cluster >= 0) {

cpp/src/hdbscan/detail/utils.h

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -24,6 +24,7 @@
 
 #include <cub/cub.cuh>
 #include <cuda/functional>
+#include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/execution_policy.h>
 #include <thrust/for_each.h>
@@ -33,7 +34,6 @@
 #include <thrust/sort.h>
 #include <thrust/transform.h>
 #include <thrust/transform_reduce.h>
-#include <thrust/tuple.h>
 
 #include <algorithm>
 
@@ -107,9 +107,9 @@ Common::CondensedHierarchy<value_idx, value_t> make_cluster_tree(
   rmm::device_uvector<value_t> cluster_lambdas(cluster_tree_edges, stream);
   rmm::device_uvector<value_idx> cluster_sizes(cluster_tree_edges, stream);
 
-  auto in = thrust::make_zip_iterator(thrust::make_tuple(parents, children, lambdas, sizes));
+  auto in = thrust::make_zip_iterator(cuda::std::make_tuple(parents, children, lambdas, sizes));
 
-  auto out = thrust::make_zip_iterator(thrust::make_tuple(
+  auto out = thrust::make_zip_iterator(cuda::std::make_tuple(
     cluster_parents.data(), cluster_children.data(), cluster_lambdas.data(), cluster_sizes.data()));
 
   thrust::copy_if(thrust_policy,

cpp/src/hdbscan/prediction_data.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -16,6 +16,7 @@
 #include <rmm/exec_policy.hpp>
 
 #include <cub/cub.cuh>
+#include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/count.h>
 #include <thrust/execution_policy.h>
@@ -67,13 +68,13 @@ void build_index_into_children(const raft::handle_t& handle,
   auto counting = thrust::make_counting_iterator<value_idx>(0);
 
   auto index_op = [index_into_children] __device__(auto t) {
-    index_into_children[thrust::get<0>(t)] = thrust::get<1>(t);
+    index_into_children[cuda::std::get<0>(t)] = cuda::std::get<1>(t);
     return;
   };
   thrust::for_each(
     exec_policy,
-    thrust::make_zip_iterator(thrust::make_tuple(children, counting)),
-    thrust::make_zip_iterator(thrust::make_tuple(children + n_edges, counting + n_edges)),
+    thrust::make_zip_iterator(cuda::std::make_tuple(children, counting)),
+    thrust::make_zip_iterator(cuda::std::make_tuple(children + n_edges, counting + n_edges)),
     index_op);
 }
 /**

cpp/src/svm/linear.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -27,12 +27,12 @@
 
 #include <cuda/functional>
 #include <cuda/std/functional>
+#include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/device_ptr.h>
 #include <thrust/execution_policy.h>
 #include <thrust/fill.h>
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/tuple.h>
 
 #include <cublas_v2.h>
 #include <omp.h>
@@ -322,8 +322,8 @@ int fit(const raft::handle_t& handle,
   }
   RAFT_CUDA_TRY(cudaMemsetAsync(w1, 0, coefCols * coefRows * sizeof(T), stream));
   if (probScale != nullptr) {
-    thrust::device_ptr<thrust::tuple<T, T>> p((thrust::tuple<T, T>*)ps1);
-    thrust::fill(thrust::cuda::par.on(stream), p, p + coefCols, thrust::make_tuple(T(1), T(0)));
+    thrust::device_ptr<cuda::std::tuple<T, T>> p((cuda::std::tuple<T, T>*)ps1);
+    thrust::fill(thrust::cuda::par.on(stream), p, p + coefCols, cuda::std::make_tuple(T(1), T(0)));
   }
 
   // one-vs-rest logic goes over each class

cpp/tests/sg/svc_test.cu

@@ -23,14 +23,14 @@
 #include <rmm/resource_ref.hpp>
 
 #include <cub/cub.cuh>
+#include <cuda/std/tuple>
 #include <thrust/device_ptr.h>
 #include <thrust/execution_policy.h>
 #include <thrust/fill.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/reduce.h>
 #include <thrust/sequence.h>
 #include <thrust/transform.h>
-#include <thrust/tuple.h>
 
 #include <cuvs/distance/distance.hpp>
 #include <cuvs/distance/grammian.hpp>
@@ -1346,7 +1346,7 @@ struct is_same_functor {
   template <typename Tuple>
   __host__ __device__ int operator()(Tuple t)
   {
-    return thrust::get<0>(t) == thrust::get<1>(t);
+    return cuda::std::get<0>(t) == cuda::std::get<1>(t);
   }
 };
 
@@ -1393,8 +1393,8 @@ TYPED_TEST(SmoSolverTest, BlobPredict)
     thrust::device_ptr<TypeParam> ptr1(y_pred.data());
     thrust::device_ptr<TypeParam> ptr2(y_pred2.data());
     thrust::device_ptr<int> ptr3(is_correct.data());
-    auto first = thrust::make_zip_iterator(thrust::make_tuple(ptr1, ptr2));
-    auto last  = thrust::make_zip_iterator(thrust::make_tuple(ptr1 + n_pred, ptr2 + n_pred));
+    auto first = thrust::make_zip_iterator(cuda::std::make_tuple(ptr1, ptr2));
+    auto last  = thrust::make_zip_iterator(cuda::std::make_tuple(ptr1 + n_pred, ptr2 + n_pred));
     thrust::transform(thrust::cuda::par.on(stream), first, last, ptr3, is_same_functor());
     int n_correct = thrust::reduce(thrust::cuda::par.on(stream), ptr3, ptr3 + n_pred);