Refactor Graph 500 BFS code and replace thrust::tuple with cuda::std::tuple (#5235)

Graph 500 BFS updates

and

replace thrust::tuple, thrust::make_tuple, thrust::get with cuda::std::tuple, cuda::std::make_tuple, cuda::std::get (#5226 did this migration but there were multiple remaining ones).

Authors:
  - Seunghwa Kang (https://github.com/seunghwak)

Approvers:
  - Chuck Hastings (https://github.com/ChuckHastings)

URL: #5235

Author: seunghwak

cpp/include/cugraph/edge_partition_device_view.cuh

@@ -173,7 +173,7 @@ class edge_partition_device_view_base_t {
     auto edge_offset  = offsets_[major_idx];
     auto local_degree = offsets_[major_idx + 1] - edge_offset;
     auto indices      = indices_.data() + edge_offset;
-    return thrust::make_tuple(indices, edge_offset, local_degree);
+    return cuda::std::make_tuple(indices, edge_offset, local_degree);
   }
 
   // major_idx == major offset if CSR/CSC, major_offset != major_idx if DCSR/DCSC

cpp/include/cugraph/utilities/dataframe_buffer.hpp

@@ -22,8 +22,8 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 
+#include <cuda/std/tuple>
 #include <thrust/iterator/zip_iterator.h>
-#include <thrust/tuple.h>
 
 #include <type_traits>
 

cpp/src/centrality/eigenvector_centrality_impl.cuh

@@ -149,7 +149,9 @@ rmm::device_uvector<weight_t> eigenvector_centrality(
       handle,
       pull_graph_view,
       thrust::make_zip_iterator(centralities.begin(), old_centralities.data()),
-      [] __device__(auto, auto val) { return std::abs(thrust::get<0>(val) - thrust::get<1>(val)); },
+      [] __device__(auto, auto val) {
+        return cuda::std::abs(cuda::std::get<0>(val) - cuda::std::get<1>(val));
+      },
       weight_t{0.0});
 
     iter++;

cpp/src/centrality/katz_centrality_impl.cuh

@@ -143,8 +143,8 @@ void katz_centrality(
                         val_first + pull_graph_view.local_vertex_partition_range_size(),
                         new_katz_centralities,
                         [] __device__(auto val) {
-                          auto const katz_centrality = thrust::get<0>(val);
-                          auto const beta            = thrust::get<1>(val);
+                          auto const katz_centrality = cuda::std::get<0>(val);
+                          auto const beta            = cuda::std::get<1>(val);
                           return katz_centrality + beta;
                         });
     }
@@ -153,7 +153,9 @@ void katz_centrality(
       handle,
       pull_graph_view,
       thrust::make_zip_iterator(new_katz_centralities, old_katz_centralities),
-      [] __device__(auto, auto val) { return std::abs(thrust::get<0>(val) - thrust::get<1>(val)); },
+      [] __device__(auto, auto val) {
+        return std::abs(cuda::std::get<0>(val) - cuda::std::get<1>(val));
+      },
       result_t{0.0});
 
     iter++;

cpp/src/components/mis_impl.cuh

@@ -29,6 +29,7 @@
 
 #include <cuda/functional>
 #include <cuda/std/iterator>
+#include <cuda/std/tuple>
 #include <thrust/count.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/merge.h>
@@ -73,8 +74,8 @@ rmm::device_uvector<vertex_t> maximal_independent_set(
                       thrust::make_zip_iterator(out_degrees.begin(), in_degrees.begin()),
                       remaining_vertices.begin(),
                       [] __device__(auto out_deg_and_in_deg) {
-                        return !((thrust::get<0>(out_deg_and_in_deg) == 0) &&
-                                 (thrust::get<1>(out_deg_and_in_deg) == 0));
+                        return !((cuda::std::get<0>(out_deg_and_in_deg) == 0) &&
+                                 (cuda::std::get<1>(out_deg_and_in_deg) == 0));
                       })),
     handle.get_stream());
 
@@ -90,8 +91,8 @@ rmm::device_uvector<vertex_t> maximal_independent_set(
                        cuda::proclaim_return_type<vertex_t>(
                          [] __device__(auto) { return std::numeric_limits<vertex_t>::max(); }),
                        [] __device__(auto in_out_degree) {
-                         return (thrust::get<0>(in_out_degree) == 0) &&
-                                (thrust::get<1>(in_out_degree) == 0);
+                         return (cuda::std::get<0>(in_out_degree) == 0) &&
+                                (cuda::std::get<1>(in_out_degree) == 0);
                        });
 
   out_degrees.resize(0, handle.get_stream());

cpp/src/cores/core_number_impl.cuh

@@ -136,11 +136,12 @@ void core_number(raft::handle_t const& handle,
       auto in_degrees        = cur_graph_view.compute_in_degrees(handle);
       auto out_degrees       = cur_graph_view.compute_out_degrees(handle);
       auto degree_pair_first = thrust::make_zip_iterator(in_degrees.begin(), out_degrees.begin());
-      thrust::transform(handle.get_thrust_policy(),
-                        degree_pair_first,
-                        degree_pair_first + in_degrees.size(),
-                        core_numbers,
-                        [] __device__(auto p) { return thrust::get<0>(p) + thrust::get<1>(p); });
+      thrust::transform(
+        handle.get_thrust_policy(),
+        degree_pair_first,
+        degree_pair_first + in_degrees.size(),
+        core_numbers,
+        [] __device__(auto p) { return cuda::std::get<0>(p) + cuda::std::get<1>(p); });
     }
   }
 
@@ -291,8 +292,8 @@ void core_number(raft::handle_t const& handle,
               auto new_core_number = v_val >= pushed_val ? v_val - pushed_val : edge_t{0};
               new_core_number      = new_core_number < (k - delta) ? (k - delta) : new_core_number;
               new_core_number      = new_core_number < k_first ? edge_t{0} : new_core_number;
-              return thrust::make_tuple(cuda::std::optional<size_t>{bucket_idx_next},
-                                        cuda::std::optional<edge_t>{new_core_number});
+              return cuda::std::make_tuple(cuda::std::optional<size_t>{bucket_idx_next},
+                                           cuda::std::optional<edge_t>{new_core_number});
             });
         }
 

cpp/src/cores/k_core_impl.cuh

@@ -68,7 +68,7 @@ k_core(raft::handle_t const& handle,
       thrust::make_counting_iterator(graph_view.local_vertex_partition_range_last()),
       core_numbers->end()),
     thrust::make_zip_iterator(subgraph_vertices.begin(), thrust::make_discard_iterator()),
-    [k] __device__(auto tuple) { return (k <= thrust::get<1>(tuple)); });
+    [k] __device__(auto tuple) { return (k <= cuda::std::get<1>(tuple)); });
 
   subgraph_vertices.resize(
     cuda::std::distance(

cpp/src/detail/groupby_and_count.cuh

@@ -91,7 +91,7 @@ rmm::device_uvector<size_t> groupby_and_count_edgelist_by_local_partition_id(
        cugraph::detail::compute_vertex_partition_id_from_ext_vertex_t<vertex_t>{
          comm_size}] __device__(auto pair) {
       auto local_edge_partition_id = local_edge_partition_id_key_func(pair);
-      auto vertex_partition_id     = vertex_partition_id_key_func(thrust::get<1>(pair));
+      auto vertex_partition_id     = vertex_partition_id_key_func(cuda::std::get<1>(pair));
       return (local_edge_partition_id * major_comm_size) +
              ((vertex_partition_id) % major_comm_size);
     };

cpp/src/detail/utility_wrappers_impl.cuh

@@ -138,11 +138,11 @@ vertex_t compute_maximum_vertex_id(rmm::cuda_stream_view const& stream_view,
                                    vertex_t const* d_edgelist_dsts,
                                    size_t num_edges)
 {
-  auto max_v_first =
-    thrust::make_transform_iterator(thrust::make_zip_iterator(d_edgelist_srcs, d_edgelist_dsts),
-                                    cuda::proclaim_return_type<vertex_t>([] __device__(auto e) {
-                                      return cuda::std::max(thrust::get<0>(e), thrust::get<1>(e));
-                                    }));
+  auto max_v_first = thrust::make_transform_iterator(
+    thrust::make_zip_iterator(d_edgelist_srcs, d_edgelist_dsts),
+    cuda::proclaim_return_type<vertex_t>([] __device__(auto e) {
+      return cuda::std::max(cuda::std::get<0>(e), cuda::std::get<1>(e));
+    }));
   return thrust::reduce(rmm::exec_policy(stream_view),
                         max_v_first,
                         max_v_first + num_edges,

cpp/src/generators/generate_bipartite_rmat_edgelist.cuh

@@ -110,7 +110,7 @@ generate_bipartite_rmat_edgelist(raft::handle_t const& handle,
               dst_bit_set ? static_cast<vertex_t>(vertex_t{1} << (dst_scale - (level + 1))) : 0;
           }
         }
-        return thrust::make_tuple(src, dst);
+        return cuda::std::make_tuple(src, dst);
       });
     num_edges_generated += num_edges_to_generate;
   }