Drop thrust::constant_iterator in favor of cuda::constant_iterator

cpp/src/c_api/leiden.cpp

@@ -82,7 +82,7 @@ struct leiden_functor : public cugraph::c_api::abstract_functor {
                                              handle_.get_stream());
 
       // FIXME: Revisit the constant edge property idea.  We could consider an alternate
-      // implementation (perhaps involving the thrust::constant_iterator), or we
+      // implementation (perhaps involving the cuda::constant_iterator), or we
       // could add support in Leiden for std::nullopt as the edge weights behaving
       // as desired and only instantiating a real edge_property_view_t for the
       // coarsened graphs.

cpp/src/c_api/louvain.cpp

@@ -79,7 +79,7 @@ struct louvain_functor : public cugraph::c_api::abstract_functor {
                                              handle_.get_stream());
 
       // FIXME: Revisit the constant edge property idea.  We could consider an alternate
-      // implementation (perhaps involving the thrust::constant_iterator), or we
+      // implementation (perhaps involving the cuda::constant_iterator), or we
       // could add support in Louvain for std::nullopt as the edge weights behaving
       // as desired and only instantiating a real edge_property_view_t for the
       // coarsened graphs.

cpp/src/centrality/eigenvector_centrality_impl.cuh

@@ -25,7 +25,6 @@
 #include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/fill.h>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/transform.h>
 

cpp/src/centrality/katz_centrality_impl.cuh

@@ -22,7 +22,6 @@
 #include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/fill.h>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/transform.h>
 

cpp/src/community/k_truss_impl.cuh

@@ -21,6 +21,7 @@
 #include <raft/util/integer_utils.hpp>
 
 #include <cuda/functional>
+#include <cuda/iterator>
 #include <cuda/std/iterator>
 #include <cuda/std/optional>
 #include <cuda/std/tuple>
@@ -490,7 +491,7 @@ k_truss(raft::handle_t const& handle,
       thrust::reduce_by_key(handle.get_thrust_policy(),
                             get_dataframe_buffer_begin(edgelist_to_update_count),
                             get_dataframe_buffer_end(edgelist_to_update_count),
-                            thrust::make_constant_iterator(size_t{1}),
+                            cuda::make_constant_iterator(size_t{1}),
                             get_dataframe_buffer_begin(vertex_pair_buffer_unique),
                             decrease_count.begin(),
                             thrust::equal_to<cuda::std::tuple<vertex_t, vertex_t>>{});

cpp/src/components/weakly_connected_components_impl.cuh

@@ -28,14 +28,14 @@
 #include <rmm/device_uvector.hpp>
 
 #include <cuda/functional>
+#include <cuda/iterator>
 #include <cuda/std/iterator>
 #include <cuda/std/optional>
 #include <cuda/std/tuple>
 #include <thrust/binary_search.h>
 #include <thrust/copy.h>
 #include <thrust/for_each.h>
 #include <thrust/functional.h>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
@@ -753,7 +753,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle,
           vertex_frontier,
           raft::host_span<size_t const>(next_frontier_bucket_indices.data(),
                                         next_frontier_bucket_indices.size()),
-          thrust::make_constant_iterator(0) /* dummy */,
+          cuda::make_constant_iterator(0) /* dummy */,
           thrust::make_discard_iterator() /* dummy */,
           v_op_t<GraphViewType>{
             vertex_partition, level_components, bucket_idx_next, bucket_idx_conflict});

cpp/src/converters/legacy/COOtoCSR.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 /*
@@ -20,14 +20,14 @@
 
 #include <cub/device/device_radix_sort.cuh>
 #include <cub/device/device_run_length_encode.cuh>
+#include <cuda/std/iterator>
 #include <cuda/std/tuple>
 #include <thrust/device_ptr.h>
 #include <thrust/extrema.h>
 #include <thrust/fill.h>
 #include <thrust/for_each.h>
 #include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/reverse_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/scan.h>
 #include <thrust/sort.h>
@@ -109,7 +109,7 @@ void fill_offset(VT* source,
   thrust::device_ptr<ET> off = thrust::device_pointer_cast(offsets);
   off[src[0]]                = ET{0};
 
-  auto iter = thrust::make_reverse_iterator(offsets + number_of_vertices + 1);
+  auto iter = cuda::std::make_reverse_iterator(offsets + number_of_vertices + 1);
   thrust::inclusive_scan(rmm::exec_policy(stream_view),
                          iter,
                          iter + number_of_vertices + 1,

cpp/src/detail/permute_range.cuh

@@ -17,8 +17,8 @@
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/iterator>
 #include <cuda/std/iterator>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/reduce.h>
 #include <thrust/sort.h>
 
@@ -84,7 +84,7 @@ rmm::device_uvector<vertex_t> permute_range(raft::handle_t const& handle,
       auto output_end = thrust::reduce_by_key(handle.get_thrust_policy(),
                                               d_target_ranks.begin(),
                                               d_target_ranks.end(),
-                                              thrust::make_constant_iterator(1),
+                                              cuda::make_constant_iterator(1),
                                               d_reduced_ranks.begin(),
                                               d_reduced_counts.begin(),
                                               thrust::equal_to<int>());

cpp/src/link_analysis/hits_impl.cuh

@@ -16,11 +16,11 @@
 #include <cugraph/edge_src_dst_property.hpp>
 #include <cugraph/graph_view.hpp>
 
+#include <cuda/iterator>
 #include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/fill.h>
 #include <thrust/functional.h>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/transform.h>
 
@@ -38,7 +38,7 @@ void normalize(raft::handle_t const& handle,
   thrust::transform(handle.get_thrust_policy(),
                     hubs,
                     hubs + graph_view.local_vertex_partition_range_size(),
-                    thrust::make_constant_iterator(hubs_norm),
+                    cuda::make_constant_iterator(hubs_norm),
                     hubs,
                     thrust::divides<result_t>());
 }

cpp/src/link_analysis/pagerank_impl.cuh

@@ -28,7 +28,6 @@
 #include <thrust/count.h>
 #include <thrust/fill.h>
 #include <thrust/for_each.h>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/reduce.h>
 #include <thrust/transform.h>

cpp/src/prims/detail/partition_v_frontier.cuh

@@ -24,7 +24,6 @@
 #include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/count.h>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/remove.h>
 #include <thrust/sort.h>

cpp/src/prims/detail/sample_and_compute_local_nbr_indices.cuh

@@ -33,7 +33,6 @@
 #include <thrust/adjacent_difference.h>
 #include <thrust/copy.h>
 #include <thrust/count.h>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/remove.h>
 #include <thrust/sort.h>

cpp/src/prims/edge_bucket.cuh

@@ -17,7 +17,6 @@
 #include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/fill.h>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/merge.h>
 #include <thrust/partition.h>

cpp/src/prims/fill_edge_property.cuh

@@ -13,11 +13,11 @@
 
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/iterator>
 #include <cuda/std/functional>
 #include <cuda/std/optional>
 #include <cuda/std/tuple>
 #include <thrust/fill.h>
-#include <thrust/iterator/constant_iterator.h>
 
 #include <cstddef>
 
@@ -93,8 +93,8 @@ void fill_edge_property(raft::handle_t const& handle,
     } else {
       if (edge_partition_e_mask) {
         thrust::transform_if(handle.get_thrust_policy(),
-                             thrust::make_constant_iterator(input),
-                             thrust::make_constant_iterator(input) + edge_counts[i],
+                             cuda::make_constant_iterator(input),
+                             cuda::make_constant_iterator(input) + edge_counts[i],
                              thrust::make_counting_iterator(edge_t{0}),
                              value_firsts[i],
                              cuda::std::identity{},

cpp/src/prims/fill_edge_src_dst_property.cuh

@@ -19,11 +19,10 @@
 
 #include <rmm/exec_policy.hpp>
 
-#include <cuda/std/iterator>
+#include <cuda/iterator>
 #include <cuda/std/tuple>
 #include <thrust/count.h>
 #include <thrust/fill.h>
-#include <thrust/iterator/constant_iterator.h>
 
 #include <cstddef>
 #include <utility>
@@ -212,7 +211,7 @@ void fill_edge_major_property(raft::handle_t const& handle,
             cuda::proclaim_return_type<vertex_t>([edge_partition] __device__(auto v) {
               return edge_partition.major_offset_from_major_nocheck(v);
             }));
-          auto val_first = thrust::make_constant_iterator(input);
+          auto val_first = cuda::make_constant_iterator(input);
           // FIXME: this scatter is unnecessary if NCCL directly takes a permutation iterator (and
           // directly scatters from the internal buffer)
           thrust::scatter(handle.get_thrust_policy(),
@@ -235,7 +234,7 @@ void fill_edge_major_property(raft::handle_t const& handle,
                        [input, output_value_first = edge_partition_value_firsts[0]] __device__(
                          auto v) { packed_bool_atomic_set(output_value_first, v, input); });
     } else {
-      auto val_first = thrust::make_constant_iterator(input);
+      auto val_first = cuda::make_constant_iterator(input);
       thrust::scatter(
         handle.get_thrust_policy(),
         val_first,
@@ -834,7 +833,7 @@ void fill_edge_minor_property(raft::handle_t const& handle,
                          local_v_list_range_firsts[partition_idx]] __device__(auto v_offset) {
                         return static_cast<vertex_t>(v_offset + (range_first - minor_range_first));
                       }));
-                  auto val_first = thrust::make_constant_iterator(input);
+                  auto val_first = cuda::make_constant_iterator(input);
                   thrust::scatter(rmm::exec_policy_nosync(loop_stream),
                                   val_first,
                                   val_first + local_v_list_sizes[partition_idx],
@@ -845,7 +844,7 @@ void fill_edge_minor_property(raft::handle_t const& handle,
                     std::get<0>(edge_partition_v_buffers[j]).begin(),
                     cuda::proclaim_return_type<vertex_t>(
                       [minor_range_first] __device__(auto v) { return v - minor_range_first; }));
-                  auto val_first = thrust::make_constant_iterator(input);
+                  auto val_first = cuda::make_constant_iterator(input);
                   thrust::scatter(rmm::exec_policy_nosync(loop_stream),
                                   val_first,
                                   val_first + local_v_list_sizes[partition_idx],
@@ -913,7 +912,7 @@ void fill_edge_minor_property(raft::handle_t const& handle,
                 fill_scalar_or_thrust_tuple(output_value_first, minor_offset, input);
               });
           } else {
-            auto val_first = thrust::make_constant_iterator(input);
+            auto val_first = cuda::make_constant_iterator(input);
             if (compressed_v_list) {
               auto map_first = thrust::make_transform_iterator(
                 thrust::make_counting_iterator(vertex_t{0}),
@@ -974,7 +973,7 @@ void fill_edge_minor_property(raft::handle_t const& handle,
                          fill_scalar_or_thrust_tuple(output_value_first, v, input);
                        });
     } else {
-      auto val_first = thrust::make_constant_iterator(input);
+      auto val_first = cuda::make_constant_iterator(input);
       thrust::scatter(
         handle.get_thrust_policy(),
         val_first,

cpp/src/prims/per_v_random_select_transform_outgoing_e.cuh

@@ -28,7 +28,6 @@
 #include <cuda/std/tuple>
 #include <thrust/copy.h>
 #include <thrust/count.h>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/remove.h>
 #include <thrust/sort.h>

cpp/src/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh

@@ -36,7 +36,6 @@
 #include <thrust/count.h>
 #include <thrust/fill.h>
 #include <thrust/functional.h>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/iterator/zip_iterator.h>

cpp/src/prims/transform_reduce_if_v_frontier_outgoing_e_by_dst.cuh

@@ -40,7 +40,6 @@
 #include <thrust/count.h>
 #include <thrust/execution_policy.h>
 #include <thrust/functional.h>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/iterator/zip_iterator.h>

cpp/src/prims/vertex_frontier.cuh

@@ -21,12 +21,12 @@
 
 #include <cuda/atomic>
 #include <cuda/functional>
+#include <cuda/iterator>
 #include <cuda/std/iterator>
 #include <cuda/std/tuple>
 #include <thrust/binary_search.h>
 #include <thrust/copy.h>
 #include <thrust/fill.h>
-#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/merge.h>
 #include <thrust/partition.h>
@@ -777,7 +777,7 @@ class vertex_frontier_t {
       auto it = thrust::reduce_by_key(handle_ptr_->get_thrust_policy(),
                                       bucket_idx_first,
                                       bucket_idx_last,
-                                      thrust::make_constant_iterator(size_t{1}),
+                                      cuda::make_constant_iterator(size_t{1}),
                                       d_indices.begin(),
                                       d_counts.begin());
       d_indices.resize(cuda::std::distance(d_indices.begin(), cuda::std::get<0>(it)),

cpp/src/sampling/negative_sampling_impl.cuh

@@ -24,8 +24,8 @@
 
 #include <rmm/device_scalar.hpp>
 
+#include <cuda/iterator>
 #include <cuda/std/functional>
-#include <cuda/std/iterator>
 #include <thrust/adjacent_difference.h>
 #include <thrust/binary_search.h>
 #include <thrust/copy.h>
@@ -106,7 +106,7 @@ normalize_biases(raft::handle_t const& handle,
 
     // FIXME: conclusion of above.  Using 1.1 since it is > 1.0 and easy to type
     thrust::copy_n(handle.get_thrust_policy(),
-                   thrust::make_constant_iterator<weight_t>(1.1),
+                   cuda::make_constant_iterator<weight_t>(1.1),
                    trailing_zeros + 1,
                    gpu_biases->begin() + gpu_biases->size() - trailing_zeros - 1);
   }