Migrate Thrust zip iterators to CUDA equivalents

cpp/benchmarks/common/generate_input.cu

@@ -43,7 +43,6 @@
 #include <thrust/for_each.h>
 #include <thrust/gather.h>
 #include <thrust/iterator/transform_iterator.h>
-#include <thrust/iterator/zip_iterator.h>
 #include <thrust/random/uniform_int_distribution.h>
 #include <thrust/random/uniform_real_distribution.h>
 #include <thrust/scan.h>
@@ -513,7 +512,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; }),
     cuda::std::logical_not<bool>{});
   auto valid_lengths = thrust::make_transform_iterator(
-    thrust::make_zip_iterator(cuda::std::make_tuple(lengths.begin(), null_mask.begin())),
+    cuda::make_zip_iterator(cuda::std::make_tuple(lengths.begin(), null_mask.begin())),
     valid_or_zero{});
 
   // offsets are created as INT32 or INT64 as appropriate
@@ -523,7 +522,7 @@ std::unique_ptr<cudf::column> create_random_utf8_string_column(data_profile cons
   auto offsets_itr = cudf::detail::offsetalator_factory::make_input_iterator(offsets->view());
   rmm::device_uvector<char> chars(chars_length, cudf::get_default_stream());
   thrust::for_each_n(thrust::device,
-                     thrust::make_zip_iterator(cuda::std::make_tuple(offsets_itr, offsets_itr + 1)),
+                     cuda::make_zip_iterator(cuda::std::make_tuple(offsets_itr, offsets_itr + 1)),
                      num_rows,
                      string_generator<Encoding>{chars.data(), engine});
 

cpp/benchmarks/common/ndsh_data_generator/random_column_generator.cu

@@ -95,7 +95,7 @@ std::unique_ptr<cudf::column> generate_random_string_column(cudf::size_type lowe
   // We generate the strings in parallel into the `chars` vector using the
   // offsets vector generated above.
   thrust::for_each_n(rmm::exec_policy_nosync(stream),
-                     thrust::make_zip_iterator(cuda::std::make_tuple(offset_itr, offset_itr + 1)),
+                     cuda::make_zip_iterator(cuda::std::make_tuple(offset_itr, offset_itr + 1)),
                      num_rows,
                      random_string_generator(chars.data()));
 

cpp/include/cudf/detail/null_mask.cuh

@@ -27,7 +27,6 @@
 #include <cuda/std/tuple>
 #include <thrust/for_each.h>
 #include <thrust/iterator/transform_iterator.h>
-#include <thrust/iterator/zip_iterator.h>
 #include <thrust/transform.h>
 
 #include <algorithm>
@@ -577,8 +576,7 @@ rmm::device_uvector<size_type> segmented_count_bits(bitmask_type const* bitmask,
   if (count_bits == count_bits_policy::UNSET_BITS) {
     // Convert from set bits counts to unset bits by subtracting the number of
     // set bits from the length of the segment.
-    auto segments_begin =
-      thrust::make_zip_iterator(first_bit_indices_begin, last_bit_indices_begin);
+    auto segments_begin = cuda::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 = cuda::std::get<0>(segment);
@@ -795,7 +793,7 @@ std::pair<rmm::device_buffer, size_type> segmented_null_mask_reduction(
   rmm::device_async_resource_ref mr)
 {
   auto const segments_begin =
-    thrust::make_zip_iterator(first_bit_indices_begin, last_bit_indices_begin);
+    cuda::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 = cuda::std::get<0>(segment);
@@ -828,7 +826,7 @@ std::pair<rmm::device_buffer, size_type> segmented_null_mask_reduction(
                                        stream,
                                        cudf::get_current_device_resource_ref());
   auto const length_and_valid_count =
-    thrust::make_zip_iterator(segment_length_iterator, segment_valid_counts.begin());
+    cuda::make_zip_iterator(segment_length_iterator, segment_valid_counts.begin());
   return cudf::detail::valid_if(
     length_and_valid_count,
     length_and_valid_count + num_segments,

cpp/src/groupby/common/m2_var_std.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -15,6 +15,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/iterator>
 #include <cuda/std/cmath>
 #include <cuda/std/functional>
 #include <thrust/tabulate.h>
@@ -133,7 +134,7 @@ std::unique_ptr<column> compute_variance_std(TransformFunc&& transform_fn,
   rmm::device_uvector<bool> validity(size, stream);
 
   auto const out_it =
-    thrust::make_zip_iterator(output->mutable_view().begin<TargetType>(), validity.begin());
+    cuda::make_zip_iterator(output->mutable_view().begin<TargetType>(), validity.begin());
   thrust::tabulate(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
                    out_it,
                    out_it + size,

cpp/src/groupby/sort/group_correlation.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -21,7 +21,6 @@
 
 #include <cuda/iterator>
 #include <cuda/std/tuple>
-#include <thrust/iterator/zip_iterator.h>
 #include <thrust/transform.h>
 
 #include <type_traits>
@@ -177,7 +176,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(cuda::std::make_tuple(stddev0_ptr, stddev1_ptr));
+  auto stddev_iter = cuda::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),

cpp/src/groupby/sort/group_merge_m2.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -13,7 +13,6 @@
 
 #include <cuda/iterator>
 #include <cuda/std/tuple>
-#include <thrust/iterator/zip_iterator.h>
 #include <thrust/transform.h>
 
 namespace cudf {
@@ -87,9 +86,9 @@ std::unique_ptr<column> merge_m2(column_view const& values,
     data_type(type_to_id<result_type>()), num_groups, mask_state::UNALLOCATED, stream, mr);
 
   auto const out_iter =
-    thrust::make_zip_iterator(result_counts->mutable_view().template data<count_type>(),
-                              result_means->mutable_view().template data<result_type>(),
-                              result_M2s->mutable_view().template data<result_type>());
+    cuda::make_zip_iterator(result_counts->mutable_view().template data<count_type>(),
+                            result_means->mutable_view().template data<result_type>(),
+                            result_M2s->mutable_view().template data<result_type>());
 
   auto const count_valid = values.child(0);
   auto const mean_values = values.child(1);

cpp/src/io/comp/compression.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -10,6 +10,7 @@
 #include <rmm/exec_policy.hpp>
 
 #include <cuda/functional>
+#include <cuda/iterator>
 #include <cuda/std/tuple>
 #include <thrust/transform_reduce.h>
 
@@ -33,7 +34,7 @@ writer_compression_statistics collect_compression_statistics(
 
   auto input_size_with_status = [inputs, results, stream](codec_status status) {
     auto const zipped_begin =
-      thrust::make_zip_iterator(cuda::std::make_tuple(inputs.begin(), results.begin()));
+      cuda::make_zip_iterator(cuda::std::make_tuple(inputs.begin(), results.begin()));
     auto const zipped_end = zipped_begin + inputs.size();
 
     return thrust::transform_reduce(

cpp/src/io/comp/gpuinflate.cu

@@ -44,6 +44,7 @@ Mark Adler    madler@alumni.caltech.edu
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/iterator>
 #include <cuda/std/algorithm>
 #include <cuda/std/cmath>
 #include <cuda/std/tuple>
@@ -1211,8 +1212,8 @@ sorted_codec_parameters sort_tasks(device_span<device_span<uint8_t const> const>
   // Precompute costs to avoid repeated computation during sorting
   rmm::device_uvector<double> costs(inputs.size(), stream, mr);
   thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
-                    thrust::make_zip_iterator(inputs.begin(), outputs.begin()),
-                    thrust::make_zip_iterator(inputs.end(), outputs.end()),
+                    cuda::make_zip_iterator(inputs.begin(), outputs.begin()),
+                    cuda::make_zip_iterator(inputs.end(), outputs.end()),
                     costs.begin(),
                     [task_type] __device__(auto const& input_output_pair) {
                       auto const& input  = cuda::std::get<0>(input_output_pair);

cpp/src/io/comp/nvcomp_adapter.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #include "nvcomp_adapter.cuh"
@@ -10,8 +10,8 @@
 #include <rmm/exec_policy.hpp>
 
 #include <cuda/functional>
+#include <cuda/iterator>
 #include <cuda/std/tuple>
-#include <thrust/iterator/zip_iterator.h>
 #include <thrust/transform.h>
 
 namespace cudf::io::detail::nvcomp {
@@ -27,7 +27,7 @@ batched_args create_batched_nvcomp_args(device_span<device_span<uint8_t const> c
   rmm::device_uvector<size_t> output_data_sizes(num_comp_chunks, stream);
 
   // Prepare the input vectors
-  auto ins_it = thrust::make_zip_iterator(input_data_ptrs.begin(), input_data_sizes.begin());
+  auto ins_it = cuda::make_zip_iterator(input_data_ptrs.begin(), input_data_sizes.begin());
   thrust::transform(
     rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
     inputs.begin(),
@@ -36,7 +36,7 @@ batched_args create_batched_nvcomp_args(device_span<device_span<uint8_t const> c
     [] __device__(auto const& in) { return cuda::std::make_tuple(in.data(), in.size()); });
 
   // Prepare the output vectors
-  auto outs_it = thrust::make_zip_iterator(output_data_ptrs.begin(), output_data_sizes.begin());
+  auto outs_it = cuda::make_zip_iterator(output_data_ptrs.begin(), output_data_sizes.begin());
   thrust::transform(
     rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
     outputs.begin(),
@@ -56,7 +56,7 @@ std::pair<rmm::device_uvector<void const*>, rmm::device_uvector<size_t>> create_
   rmm::device_uvector<void const*> input_data_ptrs(inputs.size(), stream);
   rmm::device_uvector<size_t> input_data_sizes(inputs.size(), stream);
 
-  auto ins_it = thrust::make_zip_iterator(input_data_ptrs.begin(), input_data_sizes.begin());
+  auto ins_it = cuda::make_zip_iterator(input_data_ptrs.begin(), input_data_sizes.begin());
   thrust::transform(
     rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
     inputs.begin(),
@@ -108,7 +108,7 @@ void skip_unsupported_inputs(device_span<size_t> input_sizes,
                              rmm::cuda_stream_view stream)
 {
   if (max_valid_input_size.has_value()) {
-    auto status_size_it = thrust::make_zip_iterator(input_sizes.begin(), results.begin());
+    auto status_size_it = cuda::make_zip_iterator(input_sizes.begin(), results.begin());
     thrust::transform_if(
       rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
       results.begin(),

cpp/src/io/json/column_tree_construction.cu

@@ -21,7 +21,6 @@
 #include <cuda/iterator>
 #include <cuda/std/tuple>
 #include <thrust/for_each.h>
-#include <thrust/iterator/zip_iterator.h>
 #include <thrust/scan.h>
 #include <thrust/sort.h>
 #include <thrust/transform.h>
@@ -164,15 +163,14 @@ std::tuple<compressed_sparse_row, column_tree_properties> reduce_to_column_tree(
   rmm::device_uvector<NodeT> column_categories(num_columns, stream);
   thrust::copy_n(
     rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
-    thrust::make_zip_iterator(
+    cuda::make_zip_iterator(
       cuda::make_permutation_iterator(unpermuted_tree.parent_node_ids.begin(),
                                       reordering_index.begin()),
       cuda::make_permutation_iterator(unpermuted_max_row_offsets.begin(), reordering_index.begin()),
       cuda::make_permutation_iterator(unpermuted_tree.node_categories.begin(),
                                       reordering_index.begin())),
     num_columns,
-    thrust::make_zip_iterator(
-      parent_col_ids_it, max_row_offsets.begin(), column_categories.begin()));
+    cuda::make_zip_iterator(parent_col_ids_it, max_row_offsets.begin(), column_categories.begin()));
 
 #ifdef CSR_DEBUG_PRINT
   print<NodeIndexT>(reordering_index, "h_reordering_index", stream);
@@ -212,9 +210,9 @@ std::tuple<compressed_sparse_row, column_tree_properties> reduce_to_column_tree(
     if (num_columns > 1) {
       thrust::transform_inclusive_scan(
         rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
-        thrust::make_zip_iterator(cuda::counting_iterator<NodeIndexT>{1}, row_idx.begin() + 1),
-        thrust::make_zip_iterator(cuda::counting_iterator<NodeIndexT>{1} + num_columns,
-                                  row_idx.end()),
+        cuda::make_zip_iterator(cuda::counting_iterator<NodeIndexT>{1}, row_idx.begin() + 1),
+        cuda::make_zip_iterator(cuda::counting_iterator<NodeIndexT>{1} + num_columns,
+                                row_idx.end()),
         row_idx.begin() + 1,
         cuda::proclaim_return_type<NodeIndexT>([] __device__(auto a) {
           auto n   = cuda::std::get<0>(a);