Implement a random selection primitive

cpp/include/cugraph/edge_partition_device_view.cuh

@@ -17,6 +17,7 @@
 
 #include <cugraph/edge_partition_view.hpp>
 #include <cugraph/utilities/error.hpp>
+#include <cugraph/utilities/misc_utils.cuh>
 
 #include <raft/core/device_span.hpp>
 #include <rmm/cuda_stream_view.hpp>
@@ -165,13 +166,8 @@ class edge_partition_device_view_t<vertex_t,
   edge_partition_device_view_t(edge_partition_view_t<vertex_t, edge_t, weight_t, multi_gpu> view)
     : detail::edge_partition_device_view_base_t<vertex_t, edge_t, weight_t>(
         view.offsets(), view.indices(), view.weights()),
-      dcs_nzd_vertices_(
-        view.dcs_nzd_vertices()
-          ? thrust::optional<raft::device_span<vertex_t const>>{*(view.dcs_nzd_vertices())}
-          : thrust::nullopt),
-      major_hypersparse_first_(view.major_hypersparse_first()
-                                 ? thrust::optional<vertex_t>{*(view.major_hypersparse_first())}
-                                 : thrust::nullopt),
+      dcs_nzd_vertices_(detail::to_thrust_optional(view.dcs_nzd_vertices())),
+      major_hypersparse_first_(detail::to_thrust_optional(view.major_hypersparse_first())),
       major_range_first_(view.major_range_first()),
       major_range_last_(view.major_range_last()),
       minor_range_first_(view.minor_range_first()),
@@ -211,13 +207,16 @@ class edge_partition_device_view_t<vertex_t,
   {
     rmm::device_uvector<edge_t> local_degrees(this->major_range_size(), stream);
     if (dcs_nzd_vertices_) {
-      thrust::transform(
-        rmm::exec_policy(stream),
-        thrust::make_counting_iterator(this->major_range_first()),
-        thrust::make_counting_iterator(this->major_range_last()),
-        local_degrees.begin(),
-        detail::local_degree_op_t<vertex_t, edge_t, multi_gpu, true>{
-          this->offsets_, major_range_first_, *dcs_nzd_vertices_, *major_hypersparse_first_});
+      assert(major_hypersparse_first_);
+      thrust::transform(rmm::exec_policy(stream),
+                        thrust::make_counting_iterator(this->major_range_first()),
+                        thrust::make_counting_iterator(this->major_range_last()),
+                        local_degrees.begin(),
+                        detail::local_degree_op_t<vertex_t, edge_t, multi_gpu, true>{
+                          this->offsets_,
+                          major_range_first_,
+                          *dcs_nzd_vertices_,
+                          major_hypersparse_first_.value_or(vertex_t{0})});
     } else {
       thrust::transform(
         rmm::exec_policy(stream),
@@ -235,13 +234,16 @@ class edge_partition_device_view_t<vertex_t,
   {
     rmm::device_uvector<edge_t> local_degrees(majors.size(), stream);
     if (dcs_nzd_vertices_) {
-      thrust::transform(
-        rmm::exec_policy(stream),
-        majors.begin(),
-        majors.end(),
-        local_degrees.begin(),
-        detail::local_degree_op_t<vertex_t, edge_t, multi_gpu, true>{
-          this->offsets_, major_range_first_, *dcs_nzd_vertices_, *major_hypersparse_first_});
+      assert(major_hypersparse_first_);
+      thrust::transform(rmm::exec_policy(stream),
+                        majors.begin(),
+                        majors.end(),
+                        local_degrees.begin(),
+                        detail::local_degree_op_t<vertex_t, edge_t, multi_gpu, true>{
+                          this->offsets_,
+                          major_range_first_,
+                          dcs_nzd_vertices_.value(),
+                          major_hypersparse_first_.value_or(vertex_t{0})});
     } else {
       thrust::transform(
         rmm::exec_policy(stream),

cpp/include/cugraph/utilities/device_functors.cuh

@@ -75,6 +75,22 @@ struct check_in_range_t {
   __device__ bool operator()(T val) const { return (val >= min) && (val < max); }
 };
 
+template <typename T>
+struct strided_sum_t {
+  T const* values{nullptr};
+  size_t stride{0};
+  size_t count{0};
+
+  __device__ T operator()(size_t start_offset) const
+  {
+    T sum{0};
+    for (size_t j = 0; j < count; ++j) {
+      sum += values[start_offset + stride * j];
+    }
+    return sum;
+  }
+};
+
 template <typename T>
 struct shift_left_t {
   T offset{};
@@ -104,6 +120,13 @@ struct multiply_and_add_t {
   __device__ T operator()(T input) const { return input * multiplier + adder; }
 };
 
+template <typename T>
+struct divider_t {
+  T divisor{};
+
+  __device__ T operator()(T input) const { return input / divisor; }
+};
+
 }  // namespace detail
 
 }  // namespace cugraph

cpp/include/cugraph/utilities/misc_utils.cuh

@@ -23,7 +23,9 @@
 #include <thrust/gather.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
+#include <thrust/optional.h>
 
+#include <optional>
 #include <tuple>
 #include <vector>
 
@@ -72,6 +74,14 @@ std::tuple<std::vector<vertex_t>, std::vector<edge_t>> compute_offset_aligned_ed
   return std::make_tuple(h_vertex_offsets, h_edge_offsets);
 }
 
+template <typename T>
+thrust::optional<T> to_thrust_optional(std::optional<T> val)
+{
+  thrust::optional<T> ret{thrust::nullopt};
+  if (val) { ret = *val; }
+  return ret;
+}
+
 }  // namespace detail
 
 }  // namespace cugraph

cpp/include/cugraph/utilities/thrust_tuple_utils.hpp

@@ -184,6 +184,18 @@ auto std_tuple_to_thrust_tuple(TupleType tup)
     tup, std::make_index_sequence<std::tuple_size_v<TupleType>>{});
 }
 
+template <typename T>
+auto to_thrust_tuple(T scalar_value)
+{
+  return thrust::make_tuple(scalar_value);
+}
+
+template <typename... Ts>
+auto to_thrust_tuple(thrust::tuple<Ts...> tuple_value)
+{
+  return tuple_value;
+}
+
 // a temporary function to emulate thrust::tuple_cat (not supported) using std::tuple_cat (should
 // retire once thrust::tuple is replaced with cuda::std::tuple)
 template <typename... TupleTypes>