Moves more MG graph ETL to libcugraph and re-enables MG tests in CI

cpp/CMakeLists.txt

@@ -203,7 +203,7 @@ set(CUGRAPH_SOURCES
     src/detail/utility_wrappers.cu
     src/structure/graph_view_mg.cu
     src/structure/remove_self_loops.cu
-    src/structure/sort_and_remove_multi_edges.cu
+    src/structure/remove_multi_edges.cu
     src/utilities/path_retrieval.cu
     src/structure/legacy/graph.cu
     src/linear_assignment/legacy/hungarian.cu

cpp/include/cugraph/graph_functions.hpp

@@ -1005,7 +1005,9 @@ remove_self_loops(raft::handle_t const& handle,
                   std::optional<rmm::device_uvector<edge_type_t>>&& edgelist_edge_types);
 
 /**
- * @brief Sort the edges and remove all but one edge when a multi-edge exists
+ * @brief Remove all but one edge when a multi-edge exists.  Note that this function does not use
+ *    stable methods.  When a multi-edge exists, one of the edges will remain, there is no
+ *    guarantee on which one will remain.
  *
  * In an MG context it is assumed that edges have been shuffled to the proper GPU,
  * in which case any multi-edges will be on the same GPU.
@@ -1031,11 +1033,11 @@ std::tuple<rmm::device_uvector<vertex_t>,
            std::optional<rmm::device_uvector<weight_t>>,
            std::optional<rmm::device_uvector<edge_t>>,
            std::optional<rmm::device_uvector<edge_type_t>>>
-sort_and_remove_multi_edges(raft::handle_t const& handle,
-                            rmm::device_uvector<vertex_t>&& edgelist_srcs,
-                            rmm::device_uvector<vertex_t>&& edgelist_dsts,
-                            std::optional<rmm::device_uvector<weight_t>>&& edgelist_weights,
-                            std::optional<rmm::device_uvector<edge_t>>&& edgelist_edge_ids,
-                            std::optional<rmm::device_uvector<edge_type_t>>&& edgelist_edge_types);
+remove_multi_edges(raft::handle_t const& handle,
+                   rmm::device_uvector<vertex_t>&& edgelist_srcs,
+                   rmm::device_uvector<vertex_t>&& edgelist_dsts,
+                   std::optional<rmm::device_uvector<weight_t>>&& edgelist_weights,
+                   std::optional<rmm::device_uvector<edge_t>>&& edgelist_edge_ids,
+                   std::optional<rmm::device_uvector<edge_type_t>>&& edgelist_edge_types);
 
 }  // namespace cugraph

cpp/include/cugraph_c/graph.h

@@ -103,7 +103,8 @@ cugraph_error_code_t cugraph_sg_graph_create(
  *    Note that setting this flag will arbitrarily select one instance of a multi edge to be the
  *    edge that survives.  If the edges have properties that should be honored (e.g. sum the
  weights,
- *    or take the maximum weight), the caller should do that on not rely on this flag.
+ *    or take the maximum weight), the caller should remove specific edges themselves and not rely
+ *    on this flag.
  * @param [in]  do_expensive_check    If true, do expensive checks to validate the input data
  *    is consistent with software assumptions.  If false bypass these checks.
  * @param [out] graph          A pointer to the graph object

cpp/src/c_api/graph_mg.cpp

@@ -212,12 +212,12 @@ struct create_graph_functor : public cugraph::c_api::abstract_functor {
       if (drop_multi_edges_) {
         std::tie(
           edgelist_srcs, edgelist_dsts, edgelist_weights, edgelist_edge_ids, edgelist_edge_types) =
-          cugraph::sort_and_remove_multi_edges(handle_,
-                                               std::move(edgelist_srcs),
-                                               std::move(edgelist_dsts),
-                                               std::move(edgelist_weights),
-                                               std::move(edgelist_edge_ids),
-                                               std::move(edgelist_edge_types));
+          cugraph::remove_multi_edges(handle_,
+                                      std::move(edgelist_srcs),
+                                      std::move(edgelist_dsts),
+                                      std::move(edgelist_weights),
+                                      std::move(edgelist_edge_ids),
+                                      std::move(edgelist_edge_types));
       }
 
       std::tie(*graph, new_edge_weights, new_edge_ids, new_edge_types, new_number_map) =
@@ -324,6 +324,7 @@ extern "C" cugraph_error_code_t cugraph_graph_create_mg(
   //
   cugraph_data_type_id_t vertex_type{cugraph_data_type_id_t::NTYPES};
   cugraph_data_type_id_t weight_type{cugraph_data_type_id_t::NTYPES};
+
   for (size_t i = 0; i < num_arrays; ++i) {
     CAPI_EXPECTS(p_src[i]->size_ == p_dst[i]->size_,
                  CUGRAPH_INVALID_INPUT,
@@ -334,6 +335,7 @@ extern "C" cugraph_error_code_t cugraph_graph_create_mg(
                  CUGRAPH_INVALID_INPUT,
                  "Invalid input arguments: src type != dst type.",
                  *error);
+
     CAPI_EXPECTS((p_vertices == nullptr) || (p_src[i]->type_ == p_vertices[i]->type_),
                  CUGRAPH_INVALID_INPUT,
                  "Invalid input arguments: src type != vertices type.",

cpp/src/c_api/graph_sg.cpp

@@ -195,12 +195,12 @@ struct create_graph_functor : public cugraph::c_api::abstract_functor {
       if (drop_multi_edges_) {
         std::tie(
           edgelist_srcs, edgelist_dsts, edgelist_weights, edgelist_edge_ids, edgelist_edge_types) =
-          cugraph::sort_and_remove_multi_edges(handle_,
-                                               std::move(edgelist_srcs),
-                                               std::move(edgelist_dsts),
-                                               std::move(edgelist_weights),
-                                               std::move(edgelist_edge_ids),
-                                               std::move(edgelist_edge_types));
+          cugraph::remove_multi_edges(handle_,
+                                      std::move(edgelist_srcs),
+                                      std::move(edgelist_dsts),
+                                      std::move(edgelist_weights),
+                                      std::move(edgelist_edge_ids),
+                                      std::move(edgelist_edge_types));
       }
 
       std::tie(*graph, new_edge_weights, new_edge_ids, new_edge_types, new_number_map) =

cpp/src/structure/detail/structure_utils.cuh

@@ -498,63 +498,13 @@ void sort_adjacency_list(raft::handle_t const& handle,
   }
 }
 
-template <typename T>
-struct indirect_array_reference {
-  T const* array_;
-
-  T operator() __host__ __device__(size_t index) { return array_[index]; }
-};
-
-template <typename vertex_t, typename comparison_t>
-std::tuple<size_t, rmm::device_uvector<uint32_t>> mark_edges_for_removal(
-  raft::handle_t const& handle,
-  raft::device_span<vertex_t const> src,
-  raft::device_span<vertex_t const> dst,
-  comparison_t comparison)
-{
-  rmm::device_uvector<uint32_t> remove_flags(packed_bool_size(src.size()), handle.get_stream());
-  thrust::fill(handle.get_thrust_policy(),
-               remove_flags.begin(),
-               remove_flags.end(),
-               cugraph::packed_bool_empty_mask());
-
-  size_t remove_count = thrust::count_if(
-    handle.get_thrust_policy(),
-    thrust::make_counting_iterator(size_t{0}),
-    thrust::make_counting_iterator(src.size()),
-    [comparison, d_remove_flags = remove_flags.data()] __device__(size_t i) {
-      if (comparison(i)) {
-        atomicOr(d_remove_flags + cugraph::packed_bool_offset(i), cugraph::packed_bool_mask(i));
-        return true;
-      }
-      return false;
-    });
-
-  return std::make_tuple(remove_count, std::move(remove_flags));
-}
-
-template <typename T>
-rmm::device_uvector<T> remove_flagged_elements(raft::handle_t const& handle,
-                                               rmm::device_uvector<T>&& vector,
-                                               rmm::device_uvector<uint32_t> const& remove_flags,
-                                               size_t remove_count)
+template <typename comparison_t>
+std::tuple<size_t, rmm::device_uvector<uint32_t>> mark_entries(raft::handle_t const& handle,
+                                                               size_t num_entries,
+                                                               comparison_t comparison)
 {
-  rmm::device_uvector<T> result(vector.size() - remove_count, handle.get_stream());
-
-  thrust::copy_if(
-    handle.get_thrust_policy(),
-    thrust::make_counting_iterator(size_t{0}),
-    thrust::make_counting_iterator(vector.size()),
-    thrust::make_transform_output_iterator(result.begin(),
-                                           indirect_array_reference<T>{vector.data()}),
-    [d_remove_flags = remove_flags.data()] __device__(size_t i) {
-      return !(d_remove_flags[cugraph::packed_bool_offset(i)] & cugraph::packed_bool_mask(i));
-    });
-
-  return result;
-}
-
-}  // namespace detail
+  rmm::device_uvector<uint32_t> marked_entries(cugraph::packed_bool_size(num_entries),
+                                               handle.get_stream());
 
   thrust::tabulate(handle.get_thrust_policy(),
                    marked_entries.begin(),

cpp/src/structure/remove_self_loops_impl.cuh

@@ -44,31 +44,44 @@ remove_self_loops(raft::handle_t const& handle,
                   std::optional<rmm::device_uvector<edge_t>>&& edgelist_edge_ids,
                   std::optional<rmm::device_uvector<edge_type_t>>&& edgelist_edge_types)
 {
-  auto [remove_count, remove_flags] = detail::mark_edges_for_removal(
-    handle,
-    raft::device_span<vertex_t const>{edgelist_srcs.data(), edgelist_srcs.size()},
-    raft::device_span<vertex_t const>{edgelist_dsts.data(), edgelist_dsts.size()},
-    [d_srcs = edgelist_srcs.data(), d_dsts = edgelist_dsts.data()] __device__(size_t i) {
-      return d_srcs[i] == d_dsts[i];
-    });
+  auto [self_loop_count, self_loops_to_delete] =
+    detail::mark_entries(handle,
+                         edgelist_srcs.size(),
+                         [d_srcs = edgelist_srcs.data(), d_dsts = edgelist_dsts.data()] __device__(
+                           size_t i) { return d_srcs[i] == d_dsts[i]; });
 
-  if (remove_count > 0) {
-    edgelist_srcs =
-      detail::remove_flagged_elements(handle, std::move(edgelist_srcs), remove_flags, remove_count);
-    edgelist_dsts =
-      detail::remove_flagged_elements(handle, std::move(edgelist_dsts), remove_flags, remove_count);
+  if (self_loop_count > 0) {
+    edgelist_srcs = detail::remove_flagged_elements(
+      handle,
+      std::move(edgelist_srcs),
+      raft::device_span<uint32_t const>{self_loops_to_delete.data(), self_loops_to_delete.size()},
+      self_loop_count);
+    edgelist_dsts = detail::remove_flagged_elements(
+      handle,
+      std::move(edgelist_dsts),
+      raft::device_span<uint32_t const>{self_loops_to_delete.data(), self_loops_to_delete.size()},
+      self_loop_count);
 
     if (edgelist_weights)
       edgelist_weights = detail::remove_flagged_elements(
-        handle, std::move(*edgelist_weights), remove_flags, remove_count);
+        handle,
+        std::move(*edgelist_weights),
+        raft::device_span<uint32_t const>{self_loops_to_delete.data(), self_loops_to_delete.size()},
+        self_loop_count);
 
     if (edgelist_edge_ids)
       edgelist_edge_ids = detail::remove_flagged_elements(
-        handle, std::move(*edgelist_edge_ids), remove_flags, remove_count);
+        handle,
+        std::move(*edgelist_edge_ids),
+        raft::device_span<uint32_t const>{self_loops_to_delete.data(), self_loops_to_delete.size()},
+        self_loop_count);
 
     if (edgelist_edge_types)
       edgelist_edge_types = detail::remove_flagged_elements(
-        handle, std::move(*edgelist_edge_types), remove_flags, remove_count);
+        handle,
+        std::move(*edgelist_edge_types),
+        raft::device_span<uint32_t const>{self_loops_to_delete.data(), self_loops_to_delete.size()},
+        self_loop_count);
   }
 
   return std::make_tuple(std::move(edgelist_srcs),