Migrate transform iterators to CUDA

.pre-commit-config.yaml

@@ -112,6 +112,12 @@ repos:
         entry: 'rmm::exec_policy\('
         language: pygrep
         types_or: [c, c++, cuda]
+      - id: use-cuda-iterators
+        name: use-cuda-iterators
+        description: 'Enforce that CUDA iterators are used instead of equivalent Thrust iterators'
+        entry: 'thrust::(?:make_)?(?:constant|counting|discard|permutation|reverse|strided|tabulate_output|transform|transform_input_output|transform_output|zip)_iterator'
+        language: pygrep
+        types_or: [c, c++, cuda]
       - id: exec-policy-nosync-memory-resource
         name: exec-policy-nosync-memory-resource
         description: 'Enforce that rmm::exec_policy_nosync is always passed a memory resource as its second argument'

cpp/benchmarks/common/generate_input.cu

@@ -42,7 +42,6 @@
 #include <thrust/fill.h>
 #include <thrust/for_each.h>
 #include <thrust/gather.h>
-#include <thrust/iterator/transform_iterator.h>
 #include <thrust/random/uniform_int_distribution.h>
 #include <thrust/random/uniform_real_distribution.h>
 #include <thrust/scan.h>
@@ -419,7 +418,7 @@ rmm::device_uvector<cudf::size_type> sample_indices_with_run_length(cudf::size_t
       thrust::device, run_lens.begin(), run_lens.end(), run_lens.begin(), cuda::std::plus<int>{});
     auto const samples_indices = sample_dist(engine, approx_run_len + 1);
     // This is gather.
-    auto avg_repeated_sample_indices_iterator = thrust::make_transform_iterator(
+    auto avg_repeated_sample_indices_iterator = cuda::transform_iterator(
       cuda::counting_iterator<cudf::size_type>{0},
       cuda::proclaim_return_type<cudf::size_type>(
         [rb              = run_lens.begin(),
@@ -441,14 +440,6 @@ rmm::device_uvector<cudf::size_type> sample_indices_with_run_length(cudf::size_t
   }
 }
 
-struct valid_or_zero {
-  template <typename T>
-  __device__ T operator()(cuda::std::tuple<T, bool> len_valid) const
-  {
-    return cuda::std::get<1>(len_valid) ? cuda::std::get<0>(len_valid) : T{0};
-  }
-};
-
 enum class string_encoding {
   ASCII,
   UTF8,
@@ -511,13 +502,9 @@ std::unique_ptr<cudf::column> create_random_utf8_string_column(data_profile cons
     lengths.begin(),
     cuda::proclaim_return_type<cudf::size_type>([] __device__(auto) { return 0; }),
     cuda::std::logical_not<bool>{});
-  auto valid_lengths = thrust::make_transform_iterator(
-    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
   auto [offsets, chars_length] = cudf::strings::detail::make_offsets_child_column(
-    valid_lengths, valid_lengths + num_rows, stream, mr);
+    lengths.begin(), lengths.begin() + num_rows, stream, mr);
   // use the offsetalator to normalize the offset values for use by the string_generator
   auto offsets_itr = cudf::detail::offsetalator_factory::make_input_iterator(offsets->view());
   rmm::device_uvector<char> chars(chars_length, cudf::get_default_stream());

cpp/benchmarks/io/orc/orc_writer_chunks.cpp

@@ -13,7 +13,7 @@
 #include <cudf/table/table.hpp>
 #include <cudf/utilities/default_stream.hpp>
 
-#include <thrust/iterator/transform_iterator.h>
+#include <cuda/iterator>
 
 #include <nvbench/nvbench.cuh>
 
@@ -92,10 +92,10 @@ void nvbench_orc_chunked_write(nvbench::state& state)
 
   auto mem_stats_logger = cudf::memory_stats_logger();
 
-  auto size_iter = thrust::make_transform_iterator(
+  auto size_iter = cuda::transform_iterator(
     tables.begin(), [](auto const& i) { return i->num_columns() * i->num_rows(); });
   auto row_count_iter =
-    thrust::make_transform_iterator(tables.begin(), [](auto const& i) { return i->num_rows(); });
+    cuda::transform_iterator(tables.begin(), [](auto const& i) { return i->num_rows(); });
   auto total_elements = std::accumulate(size_iter, size_iter + num_tables, 0);
   auto total_rows     = std::accumulate(row_count_iter, row_count_iter + num_tables, 0);
 

cpp/benchmarks/iterator/iterator.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -16,7 +16,6 @@
 #include <cub/device/device_reduce.cuh>
 #include <cuda/iterator>
 #include <thrust/execution_policy.h>
-#include <thrust/iterator/transform_iterator.h>
 #include <thrust/reduce.h>
 
 #include <nvbench/nvbench.cuh>

cpp/benchmarks/transpose/transpose.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -12,7 +12,6 @@
 #include <cudf/utilities/default_stream.hpp>
 
 #include <cuda/iterator>
-#include <thrust/iterator/transform_iterator.h>
 
 #include <nvbench/nvbench.cuh>
 
@@ -22,7 +21,7 @@ void bench_transpose(nvbench::state& state)
   constexpr auto column_type_id = cudf::type_id::INT32;
 
   auto int_column_generator =
-    thrust::make_transform_iterator(cuda::counting_iterator<std::size_t>{0}, [count](int i) {
+    cuda::transform_iterator(cuda::counting_iterator<std::size_t>{0}, [count](int i) {
       return cudf::make_numeric_column(
         cudf::data_type{column_type_id}, count, cudf::mask_state::ALL_VALID);
     });

cpp/include/cudf/column/column_device_view.cuh

@@ -20,7 +20,6 @@
 #include <cuda/iterator>
 #include <cuda/std/utility>
 #include <cuda/stream>
-#include <thrust/iterator/transform_iterator.h>
 
 #include <functional>
 
@@ -228,7 +227,7 @@ class alignas(16) column_device_view : public column_device_view_core {
    * @brief Iterator for navigating this column
    */
   template <typename T>
-  using const_iterator = thrust::transform_iterator<detail::value_accessor<T>, count_it>;
+  using const_iterator = cuda::transform_iterator<detail::value_accessor<T>, count_it>;
 
   /**
    * @brief Return an iterator to the first element of the column.
@@ -276,14 +275,14 @@ class alignas(16) column_device_view : public column_device_view_core {
    */
   template <typename T, typename Nullate>
   using const_optional_iterator =
-    thrust::transform_iterator<detail::optional_accessor<T, Nullate>, count_it>;
+    cuda::transform_iterator<detail::optional_accessor<T, Nullate>, count_it>;
 
   /**
    * @brief Pair iterator for navigating this column
    */
   template <typename T, bool has_nulls>
   using const_pair_iterator =
-    thrust::transform_iterator<detail::pair_accessor<T, has_nulls>, count_it>;
+    cuda::transform_iterator<detail::pair_accessor<T, has_nulls>, count_it>;
 
   /**
    * @brief Pair rep iterator for navigating this column
@@ -292,7 +291,7 @@ class alignas(16) column_device_view : public column_device_view_core {
    */
   template <typename T, bool has_nulls>
   using const_pair_rep_iterator =
-    thrust::transform_iterator<detail::pair_rep_accessor<T, has_nulls>, count_it>;
+    cuda::transform_iterator<detail::pair_rep_accessor<T, has_nulls>, count_it>;
 
   /**
    * @brief Return an optional iterator to the first element of the column.
@@ -701,7 +700,7 @@ class alignas(16) mutable_column_device_view : public mutable_column_device_view
    * @brief Iterator for navigating this column
    */
   template <typename T>
-  using iterator = thrust::transform_iterator<detail::mutable_value_accessor<T>, count_it>;
+  using iterator = cuda::transform_iterator<detail::mutable_value_accessor<T>, count_it>;
 
   /**
    * @brief Return first element (accounting for offset) after underlying data

cpp/include/cudf/detail/indexalator.cuh

@@ -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
  */
 
@@ -16,7 +16,6 @@
 #include <cuda/iterator>
 #include <cuda/std/optional>
 #include <cuda/std/utility>
-#include <thrust/iterator/transform_iterator.h>
 
 namespace cudf {
 namespace detail {
@@ -344,8 +343,8 @@ struct indexalator_factory {
    */
   static auto make_input_optional_iterator(scalar const& input, rmm::cuda_stream_view stream)
   {
-    return thrust::make_transform_iterator(cuda::make_constant_iterator<size_type>(0),
-                                           scalar_optional_index_accessor{input, stream});
+    return cuda::transform_iterator(cuda::make_constant_iterator<size_type>(0),
+                                    scalar_optional_index_accessor{input, stream});
   }
 };
 

cpp/include/cudf/detail/iterator.cuh

@@ -30,12 +30,11 @@
 #include <cuda/std/optional>
 #include <cuda/std/type_traits>
 #include <cuda/std/utility>
-#include <thrust/iterator/transform_iterator.h>
 
 namespace cudf {
 namespace detail {
 /**
- * @brief Convenience wrapper for creating a `thrust::transform_iterator` over a
+ * @brief Convenience wrapper for creating a `cuda::transform_iterator` over a
  * `cuda::counting_iterator` within the range [0, INT_MAX].
  *
  *
@@ -65,7 +64,7 @@ CUDF_HOST_DEVICE inline auto make_counting_transform_iterator(CountingIterType s
         cuda::std::numeric_limits<cudf::size_type>::digits,
     "The `start` for the counting_transform_iterator must be size_type or smaller type");
 
-  return thrust::make_transform_iterator(cuda::counting_iterator{start}, f);
+  return cuda::transform_iterator(cuda::counting_iterator{start}, f);
 }
 
 /**
@@ -381,8 +380,8 @@ auto inline make_scalar_iterator(scalar const& scalar_value)
 {
   CUDF_EXPECTS(data_type(type_to_id<Element>()) == scalar_value.type(), "the data type mismatch");
   CUDF_EXPECTS(scalar_value.is_valid(), "the scalar value must be valid");
-  return thrust::make_transform_iterator(cuda::make_constant_iterator<size_type>(0),
-                                         scalar_value_accessor<Element>{scalar_value});
+  return cuda::transform_iterator(cuda::make_constant_iterator<size_type>(0),
+                                  scalar_value_accessor<Element>{scalar_value});
 }
 
 /**
@@ -577,7 +576,7 @@ auto inline make_optional_iterator(scalar const& scalar_value, Nullate has_nulls
 {
   CUDF_EXPECTS(type_id_matches_device_storage_type<Element>(scalar_value.type().id()),
                "the data type mismatch");
-  return thrust::make_transform_iterator(
+  return cuda::transform_iterator(
     cuda::make_constant_iterator<size_type>(0),
     scalar_optional_accessor<Element, Nullate>{scalar_value, has_nulls});
 }
@@ -608,8 +607,8 @@ auto inline make_pair_iterator(scalar const& scalar_value)
 {
   CUDF_EXPECTS(type_id_matches_device_storage_type<Element>(scalar_value.type().id()),
                "the data type mismatch");
-  return thrust::make_transform_iterator(cuda::make_constant_iterator<size_type>(0),
-                                         scalar_pair_accessor<Element>{scalar_value});
+  return cuda::transform_iterator(cuda::make_constant_iterator<size_type>(0),
+                                  scalar_pair_accessor<Element>{scalar_value});
 }
 
 /**

cpp/include/cudf/detail/join/filtered_join.cuh

@@ -89,7 +89,7 @@ class filtered_join {
   struct key_pair_fn {
     CUDF_HOST_DEVICE constexpr key_pair_fn(Hasher const& hasher) : _hasher{hasher} {}
 
-    __device__ __forceinline__ auto operator()(size_type i) const noexcept
+    __device__ __forceinline__ cuco::pair<hash_value_type, T> operator()(size_type i) const noexcept
     {
       return cuco::pair{_hasher(i), T{i}};
     }

cpp/include/cudf/detail/null_mask.cuh

@@ -26,7 +26,6 @@
 #include <cuda/std/tuple>
 #include <cuda/stream>
 #include <thrust/for_each.h>
-#include <thrust/iterator/transform_iterator.h>
 #include <thrust/transform.h>
 
 #include <algorithm>
@@ -536,11 +535,11 @@ rmm::device_uvector<size_type> segmented_count_bits(bitmask_type const* bitmask,
     static_cast<size_type>(std::distance(first_bit_indices_begin, first_bit_indices_end));
   rmm::device_uvector<size_type> d_bit_counts(num_ranges, stream);
 
-  auto num_set_bits_in_word = thrust::make_transform_iterator(bitmask, popcount{});
+  auto num_set_bits_in_word = cuda::transform_iterator(bitmask, popcount{});
   auto first_word_indices =
-    thrust::make_transform_iterator(first_bit_indices_begin, bit_to_word_index{true});
+    cuda::transform_iterator(first_bit_indices_begin, bit_to_word_index{true});
   auto last_word_indices =
-    thrust::make_transform_iterator(last_bit_indices_begin, bit_to_word_index{false});
+    cuda::transform_iterator(last_bit_indices_begin, bit_to_word_index{false});
 
   // Allocate temporary memory.
   size_t temp_storage_bytes{0};
@@ -577,7 +576,7 @@ rmm::device_uvector<size_type> segmented_count_bits(bitmask_type const* bitmask,
     // Convert from set bits counts to unset bits by subtracting the number of
     // set bits from the length of the segment.
     auto segments_begin = cuda::make_zip_iterator(first_bit_indices_begin, last_bit_indices_begin);
-    auto segment_length_iterator = thrust::transform_iterator(
+    auto segment_length_iterator = cuda::transform_iterator(
       segments_begin, cuda::proclaim_return_type<size_type>([] __device__(auto const& segment) {
         auto const begin = cuda::std::get<0>(segment);
         auto const end   = cuda::std::get<1>(segment);
@@ -683,10 +682,10 @@ std::vector<size_type> segmented_count_bits(bitmask_type const* bitmask,
     make_device_uvector_async(h_indices, stream, cudf::get_current_device_resource_ref());
 
   // Compute the bit counts over each segment.
-  auto first_bit_indices_begin = thrust::make_transform_iterator(
+  auto first_bit_indices_begin = cuda::transform_iterator(
     cuda::counting_iterator<cudf::size_type>{0}, index_alternator{false, d_indices.data()});
   auto const first_bit_indices_end = first_bit_indices_begin + num_segments;
-  auto last_bit_indices_begin      = thrust::make_transform_iterator(
+  auto last_bit_indices_begin      = cuda::transform_iterator(
     cuda::counting_iterator<cudf::size_type>{0}, index_alternator{true, d_indices.data()});
   rmm::device_uvector<size_type> d_bit_counts =
     cudf::detail::segmented_count_bits(bitmask,
@@ -794,7 +793,7 @@ std::pair<rmm::device_buffer, size_type> segmented_null_mask_reduction(
 {
   auto const segments_begin =
     cuda::make_zip_iterator(first_bit_indices_begin, last_bit_indices_begin);
-  auto const segment_length_iterator = thrust::make_transform_iterator(
+  auto const segment_length_iterator = cuda::transform_iterator(
     segments_begin, cuda::proclaim_return_type<size_type>([] __device__(auto const& segment) {
       auto const begin = cuda::std::get<0>(segment);
       auto const end   = cuda::std::get<1>(segment);

cpp/include/cudf/detail/row_operator/hashing.cuh

@@ -22,9 +22,9 @@
 #include <cudf/utilities/traits.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
 
+#include <cuda/iterator>
 #include <cuda/std/limits>
 #include <cuda/std/type_traits>
-#include <thrust/iterator/transform_iterator.h>
 
 #include <memory>
 
@@ -126,7 +126,7 @@ class device_row_hasher {
     auto const init        = has_columns ? hasher(_table.column(0)) : _seed;
     auto const start_col   = static_cast<size_type>(has_columns);
 
-    auto it = thrust::make_transform_iterator(_table.begin() + start_col, hasher);
+    auto it = cuda::transform_iterator(_table.begin() + start_col, hasher);
     return detail::accumulate(
       it, it + (_table.num_columns() - start_col), init, [](auto hash, auto h) {
         return cudf::hashing::detail::hash_combine(hash, h);

cpp/include/cudf/detail/scatter.cuh

@@ -28,7 +28,6 @@
 #include <cuda/iterator>
 #include <cuda/stream>
 #include <thrust/count.h>
-#include <thrust/iterator/transform_iterator.h>
 #include <thrust/scatter.h>
 #include <thrust/sequence.h>
 #include <thrust/uninitialized_fill.h>
@@ -380,9 +379,9 @@ std::unique_ptr<table> scatter(table_view const& source,
 
   // Transform negative indices to index + target size.
   auto updated_scatter_map_begin =
-    thrust::make_transform_iterator(scatter_map_begin, index_converter<MapType>{target.num_rows()});
+    cuda::transform_iterator(scatter_map_begin, index_converter<MapType>{target.num_rows()});
   auto updated_scatter_map_end =
-    thrust::make_transform_iterator(scatter_map_end, index_converter<MapType>{target.num_rows()});
+    cuda::transform_iterator(scatter_map_end, index_converter<MapType>{target.num_rows()});
   auto result = std::vector<std::unique_ptr<column>>(target.num_columns());
 
   std::transform(source.begin(),

cpp/include/cudf/detail/utilities/batched_memset.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -16,7 +16,6 @@
 #include <cub/device/device_copy.cuh>
 #include <cuda/functional>
 #include <cuda/iterator>
-#include <thrust/iterator/transform_iterator.h>
 #include <thrust/transform.h>
 
 namespace CUDF_EXPORT cudf {
@@ -43,19 +42,18 @@ void batched_memset(cudf::host_span<cudf::device_span<T> const> host_buffers,
     host_buffers, stream, cudf::get_current_device_resource_ref());
 
   // Vector of sizes of all buffer spans
-  auto sizes = thrust::make_transform_iterator(
-    cuda::counting_iterator<std::size_t>{0},
-    cuda::proclaim_return_type<std::size_t>(
-      [buffers = buffers.data()] __device__(std::size_t i) { return buffers[i].size(); }));
+  auto sizes = cuda::transform_iterator(
+    buffers.begin(), cuda::proclaim_return_type<std::size_t>([] __device__(auto const& buffer) {
+      return buffer.size();
+    }));
 
   // Constant iterator to the value to memset
   auto iter_in = cuda::make_constant_iterator(cuda::make_constant_iterator(value));
 
   // Iterator to each device span pointer
-  auto iter_out = thrust::make_transform_iterator(
-    cuda::counting_iterator<std::size_t>{0},
-    cuda::proclaim_return_type<T*>(
-      [buffers = buffers.data()] __device__(std::size_t i) { return buffers[i].data(); }));
+  auto iter_out = cuda::transform_iterator(
+    buffers.begin(),
+    cuda::proclaim_return_type<T*>([] __device__(auto const& buffer) { return buffer.data(); }));
 
   std::size_t temp_storage_bytes = 0;
   auto const num_buffers         = host_buffers.size();