Drop thrust::zip_iterator in favor of cuda::zip_iterator

cpp/include/raft/comms/detail/std_comms.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -16,7 +16,6 @@
 #include <rmm/device_uvector.hpp>
 
 #include <cuda_runtime.h>
-#include <thrust/iterator/zip_iterator.h>
 
 #include <nccl.h>
 #include <stdlib.h>

cpp/include/raft/sparse/op/detail/sort.h

@@ -13,10 +13,10 @@
 
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/iterator>
 #include <cuda/std/tuple>
 #include <cuda_runtime.h>
 #include <thrust/device_ptr.h>
-#include <thrust/iterator/zip_iterator.h>
 #include <thrust/scan.h>
 #include <thrust/sort.h>
 
@@ -60,7 +60,7 @@ struct TupleComp {
 template <typename T, typename IdxT = int, typename nnz_t>
 void coo_sort(IdxT m, IdxT n, nnz_t nnz, IdxT* rows, IdxT* cols, T* vals, cudaStream_t stream)
 {
-  auto coo_indices = thrust::make_zip_iterator(cuda::std::make_tuple(rows, cols));
+  auto coo_indices = cuda::make_zip_iterator(cuda::std::make_tuple(rows, cols));
 
   // get all the colors in contiguous locations so we can map them to warps.
   thrust::sort_by_key(rmm::exec_policy(stream), coo_indices, coo_indices + nnz, vals, TupleComp());
@@ -95,7 +95,7 @@ void coo_sort_by_weight(
 {
   thrust::device_ptr<value_t> t_data = thrust::device_pointer_cast(data);
 
-  auto first = thrust::make_zip_iterator(cuda::std::make_tuple(rows, cols));
+  auto first = cuda::make_zip_iterator(cuda::std::make_tuple(rows, cols));
 
   thrust::sort_by_key(rmm::exec_policy(stream), t_data, t_data + nnz, first);
 }

cpp/include/raft/sparse/solver/detail/mst_solver_inl.cuh

@@ -15,13 +15,13 @@
 #include <rmm/device_uvector.hpp>
 
 #include <cuda/functional>
+#include <cuda/iterator>
 #include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/device_ptr.h>
 #include <thrust/execution_policy.h>
 #include <thrust/fill.h>
 #include <thrust/host_vector.h>
-#include <thrust/iterator/zip_iterator.h>
 #include <thrust/reduce.h>
 #include <thrust/sequence.h>
 #include <thrust/sort.h>
@@ -194,8 +194,8 @@ alteration_t MST_solver<vertex_t, edge_t, weight_t, alteration_t>::alteration_ma
   auto new_end = thrust::unique(policy, tmp.begin(), tmp.end());
 
   // min(a[i+1]-a[i])/2
-  auto begin = thrust::make_zip_iterator(cuda::std::make_tuple(tmp.begin(), tmp.begin() + 1));
-  auto end   = thrust::make_zip_iterator(cuda::std::make_tuple(new_end - 1, new_end));
+  auto begin = cuda::make_zip_iterator(cuda::std::make_tuple(tmp.begin(), tmp.begin() + 1));
+  auto end   = cuda::make_zip_iterator(cuda::std::make_tuple(new_end - 1, new_end));
   auto init  = tmp.element(1, stream) - tmp.element(0, stream);
   auto max   = thrust::transform_reduce(
     policy, begin, end, alteration_functor<weight_t>(), init, cuda::minimum<weight_t>());
@@ -383,14 +383,14 @@ void MST_solver<vertex_t, edge_t, weight_t, alteration_t>::append_src_dst_pair(
 
   // iterator to end of mst edges added to final output in previous iteration
   auto src_dst_zip_end =
-    thrust::make_zip_iterator(cuda::std::make_tuple(mst_src + curr_mst_edge_count,
-                                                    mst_dst + curr_mst_edge_count,
-                                                    mst_weights + curr_mst_edge_count));
+    cuda::make_zip_iterator(cuda::std::make_tuple(mst_src + curr_mst_edge_count,
+                                                  mst_dst + curr_mst_edge_count,
+                                                  mst_weights + curr_mst_edge_count));
 
   // iterator to new mst edges found
-  auto temp_src_dst_zip_begin = thrust::make_zip_iterator(
+  auto temp_src_dst_zip_begin = cuda::make_zip_iterator(
     cuda::std::make_tuple(temp_src.begin(), temp_dst.begin(), temp_weights.begin()));
-  auto temp_src_dst_zip_end = thrust::make_zip_iterator(
+  auto temp_src_dst_zip_end = cuda::make_zip_iterator(
     cuda::std::make_tuple(temp_src.end(), temp_dst.end(), temp_weights.end()));
 
   // copy new mst edges to final output

cpp/include/raft/spectral/detail/spectral_util.cuh

@@ -19,7 +19,6 @@
 #include <thrust/device_ptr.h>
 #include <thrust/fill.h>
 #include <thrust/for_each.h>
-#include <thrust/iterator/zip_iterator.h>
 #include <thrust/reduce.h>
 #include <thrust/transform.h>
 
@@ -140,10 +139,10 @@ bool construct_indicator(
 
   thrust::for_each(
     thrust_exec_policy,
-    thrust::make_zip_iterator(cuda::std::make_tuple(thrust::device_pointer_cast(clusters),
-                                                    thrust::device_pointer_cast(part_i.raw()))),
-    thrust::make_zip_iterator(cuda::std::make_tuple(thrust::device_pointer_cast(clusters + n),
-                                                    thrust::device_pointer_cast(part_i.raw() + n))),
+    cuda::make_zip_iterator(cuda::std::make_tuple(thrust::device_pointer_cast(clusters),
+                                                  thrust::device_pointer_cast(part_i.raw()))),
+    cuda::make_zip_iterator(cuda::std::make_tuple(thrust::device_pointer_cast(clusters + n),
+                                                  thrust::device_pointer_cast(part_i.raw() + n))),
     equal_to_i_op<vertex_t, weight_t>(index));
   RAFT_CHECK_CUDA(stream);