Replace Thrust sequential algorithms with libcu++ ones

Author: miscco

cpp/examples/developers/vertex_and_edge_partition/vertex_and_edge_partition.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024, NVIDIA CORPORATION.
+ * Copyright (c) 2024-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -24,7 +24,7 @@
 #include <raft/core/handle.hpp>
 #include <raft/random/rng_state.hpp>
 
-#include <thrust/for_each.h>
+#include <cuda/std/__algorithm_>
 
 #include <iostream>
 #include <string>
@@ -237,21 +237,20 @@ void look_into_vertex_and_edge_partitions(
   //
 
   if (renumber_map) {
-    thrust::for_each(thrust::host,
-                     thrust::make_zip_iterator(thrust::make_tuple(
-                       h_vertices_in_this_proces.begin(),
-                       thrust::make_counting_iterator(renumbered_vertex_id_of_local_first))),
-                     thrust::make_zip_iterator(thrust::make_tuple(
-                       h_vertices_in_this_proces.end(),
-                       thrust::make_counting_iterator(renumbered_vertex_id_of_local_last))),
-                     [comm_rank](auto old_and_new_id_pair) {
-                       auto old_id = thrust::get<0>(old_and_new_id_pair);
-                       auto new_id = thrust::get<1>(old_and_new_id_pair);
-                       printf("owner rank = %d, original vertex id %d is renumbered to  %d\n",
-                              comm_rank,
-                              static_cast<int>(old_id),
-                              static_cast<int>(new_id));
-                     });
+    cuda::std::for_each(thrust::make_zip_iterator(thrust::make_tuple(
+                          h_vertices_in_this_proces.begin(),
+                          thrust::make_counting_iterator(renumbered_vertex_id_of_local_first))),
+                        thrust::make_zip_iterator(thrust::make_tuple(
+                          h_vertices_in_this_proces.end(),
+                          thrust::make_counting_iterator(renumbered_vertex_id_of_local_last))),
+                        [comm_rank](auto old_and_new_id_pair) {
+                          auto old_id = thrust::get<0>(old_and_new_id_pair);
+                          auto new_id = thrust::get<1>(old_and_new_id_pair);
+                          printf("owner rank = %d, original vertex id %d is renumbered to  %d\n",
+                                 comm_rank,
+                                 static_cast<int>(old_id),
+                                 static_cast<int>(new_id));
+                        });
   }
 
   //
@@ -307,8 +306,7 @@ void look_into_vertex_and_edge_partitions(
         auto v                               = major_range_first + i;
         auto deg_of_v_in_this_edge_partition = offsets[i + 1] - offsets[i];
 
-        thrust::for_each(
-          thrust::seq,
+        cuda::std::for_each(
           thrust::make_counting_iterator(edge_t{offsets[i]}),
           thrust::make_counting_iterator(edge_t{offsets[i + 1]}),
           [comm_rank, ep_idx, v, indices, is_weighted, weights] __device__(auto pos) {
@@ -353,8 +351,7 @@ void look_into_vertex_and_edge_partitions(
           auto major_idx                       = (major_hypersparse_first - major_range_first) + i;
           auto deg_of_v_in_this_edge_partition = offsets[major_idx + 1] - offsets[major_idx];
 
-          thrust::for_each(
-            thrust::seq,
+          cuda::std::for_each(
             thrust::make_counting_iterator(edge_t{offsets[major_idx]}),
             thrust::make_counting_iterator(edge_t{offsets[major_idx + 1]}),
             [comm_rank, ep_idx, v, indices, is_weighted, weights] __device__(auto pos) {

cpp/include/cugraph/detail/decompress_edge_partition.cuh

@@ -24,9 +24,9 @@
 
 #include <rmm/device_uvector.hpp>
 
+#include <cuda/std/__algorithm_>
 #include <thrust/copy.h>
 #include <thrust/execution_policy.h>
-#include <thrust/fill.h>
 #include <thrust/for_each.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/sequence.h>
@@ -159,10 +159,9 @@ void decompress_edge_partition_to_fill_edgelist_majors(
           auto major_offset = edge_partition.major_offset_from_major_nocheck(major);
           auto local_degree = edge_partition.local_degree(major_offset);
           auto local_offset = edge_partition.local_offset(major_offset);
-          thrust::fill(thrust::seq,
-                       output_buffer.begin() + local_offset,
-                       output_buffer.begin() + local_offset + local_degree,
-                       major);
+          cuda::std::fill(output_buffer.begin() + local_offset,
+                          output_buffer.begin() + local_offset + local_degree,
+                          major);
         });
     }
     if (edge_partition.dcs_nzd_vertex_count() && (*(edge_partition.dcs_nzd_vertex_count()) > 0)) {
@@ -177,10 +176,9 @@ void decompress_edge_partition_to_fill_edgelist_majors(
             major_start_offset + idx;  // major_offset != major_idx in the hypersparse region
           auto local_degree = edge_partition.local_degree(major_idx);
           auto local_offset = edge_partition.local_offset(major_idx);
-          thrust::fill(thrust::seq,
-                       output_buffer.begin() + local_offset,
-                       output_buffer.begin() + local_offset + local_degree,
-                       major);
+          cuda::std::fill(output_buffer.begin() + local_offset,
+                          output_buffer.begin() + local_offset + local_degree,
+                          major);
         });
     }
   } else {
@@ -192,10 +190,9 @@ void decompress_edge_partition_to_fill_edgelist_majors(
                        auto major_offset = edge_partition.major_offset_from_major_nocheck(major);
                        auto local_degree = edge_partition.local_degree(major_offset);
                        auto local_offset = edge_partition.local_offset(major_offset);
-                       thrust::fill(thrust::seq,
-                                    output_buffer.begin() + local_offset,
-                                    output_buffer.begin() + local_offset + local_degree,
-                                    major);
+                       cuda::std::fill(output_buffer.begin() + local_offset,
+                                       output_buffer.begin() + local_offset + local_degree,
+                                       major);
                      });
   }
 

cpp/include/cugraph/edge_partition_device_view.cuh

@@ -26,9 +26,9 @@
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/std/__algorithm_>
 #include <cuda/std/iterator>
 #include <cuda/std/optional>
-#include <thrust/binary_search.h>
 #include <thrust/execution_policy.h>
 #include <thrust/transform.h>
 #include <thrust/transform_reduce.h>
@@ -48,8 +48,7 @@ __device__ cuda::std::optional<vertex_t> major_hypersparse_idx_from_major_nochec
 {
   // we can avoid binary search (and potentially improve performance) if we add an auxiliary array
   // or cuco::static_map (at the expense of additional memory)
-  auto it =
-    thrust::lower_bound(thrust::seq, dcs_nzd_vertices.begin(), dcs_nzd_vertices.end(), major);
+  auto it = cuda::std::lower_bound(dcs_nzd_vertices.begin(), dcs_nzd_vertices.end(), major);
   return it != dcs_nzd_vertices.end()
            ? (*it == major ? cuda::std::optional<vertex_t>{static_cast<vertex_t>(
                                cuda::std::distance(dcs_nzd_vertices.begin(), it))}
@@ -163,7 +162,7 @@ class edge_partition_device_view_base_t {
   {
     return static_cast<vertex_t>(cuda::std::distance(
       offsets_.begin() + 1,
-      thrust::upper_bound(thrust::seq, offsets_.begin() + 1, offsets_.end(), local_edge_idx)));
+      cuda::std::upper_bound(offsets_.begin() + 1, offsets_.end(), local_edge_idx)));
   }
 
   // major_idx == major offset if CSR/CSC, major_offset != major_idx if DCSR/DCSC

cpp/include/cugraph/edge_partition_endpoint_property_device_view.cuh

@@ -23,9 +23,9 @@
 
 #include <raft/core/device_span.hpp>
 
+#include <cuda/std/__algorithm_>
 #include <cuda/std/iterator>
 #include <cuda/std/optional>
-#include <thrust/binary_search.h>
 #include <thrust/execution_policy.h>
 #include <thrust/fill.h>
 #include <thrust/iterator/iterator_traits.h>
@@ -197,10 +197,10 @@ class edge_partition_endpoint_property_device_view_t {
     auto val_offset = offset;
     if (keys_) {
       auto chunk_idx = static_cast<size_t>(offset) / (*key_chunk_size_);
-      auto it        = thrust::lower_bound(thrust::seq,
-                                    (*keys_).begin() + (*key_chunk_start_offsets_)[chunk_idx],
-                                    (*keys_).begin() + (*key_chunk_start_offsets_)[chunk_idx + 1],
-                                    range_first_ + offset);
+      auto it =
+        cuda::std::lower_bound((*keys_).begin() + (*key_chunk_start_offsets_)[chunk_idx],
+                               (*keys_).begin() + (*key_chunk_start_offsets_)[chunk_idx + 1],
+                               range_first_ + offset);
       assert((it != (*keys_).begin() + (*key_chunk_start_offsets_)[chunk_idx + 1]) &&
              (*it == (range_first_ + offset)));
       val_offset = (*key_chunk_start_offsets_)[chunk_idx] +

cpp/include/cugraph/utilities/mask_utils.cuh

@@ -21,7 +21,9 @@
 #include <raft/core/handle.hpp>
 
 #include <cuda/functional>
+#include <cuda/std/__algorithm_>
 #include <cuda/std/iterator>
+#include <cuda/std/numeric>
 #include <thrust/copy.h>
 #include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
@@ -53,19 +55,18 @@ __device__ size_t count_set_bits(MaskIterator mask_first, size_t start_offset, s
     ++mask_first;
   }
 
-  return thrust::transform_reduce(
-    thrust::seq,
+  return cuda::std::transform_reduce(
     thrust::make_counting_iterator(size_t{0}),
     thrust::make_counting_iterator(packed_bool_size(num_bits)),
+    ret,
+    thrust::plus<size_t>{},
     [mask_first, num_bits] __device__(size_t i) {
       auto word = *(mask_first + i);
       if ((i + 1) * packed_bools_per_word() > num_bits) {
         word &= packed_bool_partial_mask(num_bits % packed_bools_per_word());
       }
       return static_cast<size_t>(__popc(word));
-    },
-    ret,
-    thrust::plus<size_t>{});
+    });
 }
 
 // @p n starts from 1

cpp/include/cugraph/utilities/shuffle_comm.cuh

@@ -26,6 +26,7 @@
 
 #include <cuda/atomic>
 #include <cuda/functional>
+#include <cuda/std/__algorithm_>
 #include <cuda/std/iterator>
 #include <thrust/binary_search.h>
 #include <thrust/copy.h>
@@ -63,9 +64,8 @@ struct compute_group_id_count_pair_t {
   {
     static_assert(
       std::is_same_v<typename thrust::iterator_traits<GroupIdIterator>::value_type, int>);
-    auto lower_it =
-      thrust::lower_bound(thrust::seq, group_id_first, group_id_last, static_cast<int>(i));
-    auto upper_it = thrust::upper_bound(thrust::seq, lower_it, group_id_last, static_cast<int>(i));
+    auto lower_it = cuda::std::lower_bound(group_id_first, group_id_last, static_cast<int>(i));
+    auto upper_it = cuda::std::upper_bound(lower_it, group_id_last, static_cast<int>(i));
     return thrust::make_tuple(static_cast<int>(i),
                               static_cast<size_t>(cuda::std::distance(lower_it, upper_it)));
   }

cpp/src/c_api/capi_helper.cu

@@ -20,6 +20,7 @@
 #include <cugraph/shuffle_functions.hpp>
 #include <cugraph/utilities/misc_utils.cuh>
 
+#include <cuda/std/__algorithm_>
 #include <cuda/std/iterator>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/sort.h>
@@ -141,8 +142,7 @@ reorder_extracted_egonets(raft::handle_t const& handle,
      source_indices = raft::device_span<size_t const>(source_indices.data(),
                                                       source_indices.size())] __device__(size_t i) {
       auto idx = static_cast<size_t>(cuda::std::distance(
-        offset_lasts.begin(),
-        thrust::upper_bound(thrust::seq, offset_lasts.begin(), offset_lasts.end(), i)));
+        offset_lasts.begin(), cuda::std::upper_bound(offset_lasts.begin(), offset_lasts.end(), i)));
       return source_indices[idx];
     });
   source_indices.resize(0, handle.get_stream());
@@ -166,8 +166,7 @@ reorder_extracted_egonets(raft::handle_t const& handle,
     [sort_indices = raft::device_span<size_t const>(sort_indices.data(),
                                                     sort_indices.size())] __device__(size_t i) {
       return static_cast<size_t>(cuda::std::distance(
-        sort_indices.begin(),
-        thrust::upper_bound(thrust::seq, sort_indices.begin(), sort_indices.end(), i)));
+        sort_indices.begin(), cuda::std::upper_bound(sort_indices.begin(), sort_indices.end(), i)));
     });
 
   return std::make_tuple(

cpp/src/c_api/renumber_arbitrary_edgelist.cu

@@ -23,8 +23,8 @@
 #include <cugraph/graph.hpp>
 #include <cugraph/utilities/error.hpp>
 
+#include <cuda/std/__algorithm_>
 #include <cuda/std/iterator>
-#include <thrust/binary_search.h>
 #include <thrust/iterator/counting_iterator.h>
 
 namespace {
@@ -87,8 +87,8 @@ cugraph_error_code_t renumber_arbitrary_edgelist(
          raft::device_span<vertex_t const>{renumber_chunk.data(), renumber_chunk.size()},
        vertices_span,
        ids_span] __device__(size_t idx) {
-        auto pos = thrust::lower_bound(
-          thrust::seq, vertices_span.begin(), vertices_span.end(), renumber_chunk_span[idx]);
+        auto pos = cuda::std::lower_bound(
+          vertices_span.begin(), vertices_span.end(), renumber_chunk_span[idx]);
         if ((pos != vertices_span.end()) && (*pos == renumber_chunk_span[idx])) {
           ids_span[cuda::std::distance(vertices_span.begin(), pos)] =
             static_cast<vertex_t>(chunk_base_offset + idx);
@@ -102,27 +102,25 @@ cugraph_error_code_t renumber_arbitrary_edgelist(
                                 cugraph::invalid_vertex_id<vertex_t>::value) == 0,
                   "some vertices were not renumbered");
 
-  thrust::transform(
-    handle.get_thrust_policy(),
-    srcs->as_type<vertex_t>(),
-    srcs->as_type<vertex_t>() + srcs->size_,
-    srcs->as_type<vertex_t>(),
-    [vertices_span, ids_span] __device__(vertex_t v) {
-      return ids_span[cuda::std::distance(
-        vertices_span.begin(),
-        thrust::lower_bound(thrust::seq, vertices_span.begin(), vertices_span.end(), v))];
-    });
-
-  thrust::transform(
-    handle.get_thrust_policy(),
-    dsts->as_type<vertex_t>(),
-    dsts->as_type<vertex_t>() + srcs->size_,
-    dsts->as_type<vertex_t>(),
-    [vertices_span, ids_span] __device__(vertex_t v) {
-      return ids_span[cuda::std::distance(
-        vertices_span.begin(),
-        thrust::lower_bound(thrust::seq, vertices_span.begin(), vertices_span.end(), v))];
-    });
+  thrust::transform(handle.get_thrust_policy(),
+                    srcs->as_type<vertex_t>(),
+                    srcs->as_type<vertex_t>() + srcs->size_,
+                    srcs->as_type<vertex_t>(),
+                    [vertices_span, ids_span] __device__(vertex_t v) {
+                      return ids_span[cuda::std::distance(
+                        vertices_span.begin(),
+                        cuda::std::lower_bound(vertices_span.begin(), vertices_span.end(), v))];
+                    });
+
+  thrust::transform(handle.get_thrust_policy(),
+                    dsts->as_type<vertex_t>(),
+                    dsts->as_type<vertex_t>() + srcs->size_,
+                    dsts->as_type<vertex_t>(),
+                    [vertices_span, ids_span] __device__(vertex_t v) {
+                      return ids_span[cuda::std::distance(
+                        vertices_span.begin(),
+                        cuda::std::lower_bound(vertices_span.begin(), vertices_span.end(), v))];
+                    });
 
   return CUGRAPH_SUCCESS;
 }

cpp/src/community/detail/common_methods.cuh

@@ -30,8 +30,8 @@
 #include <cugraph/graph_functions.hpp>
 
 #include <cuda/functional>
+#include <cuda/std/__algorithm_>
 #include <cuda/std/optional>
-#include <thrust/binary_search.h>
 #include <thrust/execution_policy.h>
 #include <thrust/functional.h>
 #include <thrust/iterator/zip_iterator.h>
@@ -319,8 +319,8 @@ rmm::device_uvector<vertex_t> update_clustering_by_delta_modularity(
                       [d_cluster_weights = cluster_weights_v.data(),
                        d_cluster_keys    = cluster_keys_v.data(),
                        num_clusters      = cluster_keys_v.size()] __device__(vertex_t cluster) {
-                        auto pos = thrust::lower_bound(
-                          thrust::seq, d_cluster_keys, d_cluster_keys + num_clusters, cluster);
+                        auto pos = cuda::std::lower_bound(
+                          d_cluster_keys, d_cluster_keys + num_clusters, cluster);
                         return d_cluster_weights[pos - d_cluster_keys];
                       });
   }