Optimize K-Truss

cpp/CMakeLists.txt

@@ -284,6 +284,9 @@ set(CUGRAPH_SOURCES
     src/community/k_truss_sg_v64_e64.cu
     src/community/k_truss_sg_v32_e32.cu
     src/community/k_truss_sg_v32_e64.cu
+    src/community/k_truss_mg_v64_e64.cu
+    src/community/k_truss_mg_v32_e32.cu
+    src/community/k_truss_mg_v32_e64.cu
     src/lookup/lookup_src_dst_mg_v32_e32.cu
     src/lookup/lookup_src_dst_mg_v32_e64.cu
     src/lookup/lookup_src_dst_mg_v64_e64.cu
@@ -628,7 +631,7 @@ add_library(cugraph_c
         src/c_api/eigenvector_centrality.cpp
         src/c_api/betweenness_centrality.cpp
         src/c_api/core_number.cpp
-        src/c_api/k_truss.cpp
+        #src/c_api/k_truss.cpp
         src/c_api/core_result.cpp
         src/c_api/extract_ego.cpp
         src/c_api/ecg.cpp

cpp/include/cugraph/detail/shuffle_wrappers.hpp

@@ -53,7 +53,8 @@ std::tuple<rmm::device_uvector<vertex_t>,
            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_id_t>>>
+           std::optional<rmm::device_uvector<edge_type_id_t>>,
+           std::vector<size_t>>
 shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning(
   raft::handle_t const& handle,
   rmm::device_uvector<vertex_t>&& majors,
@@ -93,7 +94,8 @@ std::tuple<rmm::device_uvector<vertex_t>,
            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_id_t>>>
+           std::optional<rmm::device_uvector<edge_type_id_t>>,
+           std::vector<size_t>>
 shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning(
   raft::handle_t const& handle,
   rmm::device_uvector<vertex_t>&& majors,

cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp

@@ -251,7 +251,8 @@ class per_device_edgelist_t {
              store_transposed ? src_[0] : dst_[0],
              tmp_wgt,
              tmp_edge_id,
-             tmp_edge_type) =
+             tmp_edge_type,
+             std::ignore) =
       cugraph::detail::shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning(
         handle.raft_handle(),
         store_transposed ? std::move(dst_[0]) : std::move(src_[0]),

cpp/src/c_api/graph_functions.cpp

@@ -72,7 +72,7 @@ struct create_vertex_pairs_functor : public cugraph::c_api::abstract_functor {
         second_copy.data(), second_->as_type<vertex_t>(), second_->size_, handle_.get_stream());
 
       if constexpr (multi_gpu) {
-        std::tie(first_copy, second_copy, std::ignore, std::ignore, std::ignore) =
+        std::tie(first_copy, second_copy, std::ignore, std::ignore, std::ignore, std::ignore) =
           cugraph::detail::shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning<
             vertex_t,
             edge_t,

cpp/src/c_api/graph_mg.cpp

@@ -167,7 +167,8 @@ struct create_graph_functor : public cugraph::c_api::abstract_functor {
                store_transposed ? edgelist_srcs : edgelist_dsts,
                edgelist_weights,
                edgelist_edge_ids,
-               edgelist_edge_types) =
+               edgelist_edge_types,
+               std::ignore) =
         cugraph::detail::shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning(
           handle_,
           std::move(store_transposed ? edgelist_dsts : edgelist_srcs),

cpp/src/community/detail/refine_impl.cuh

@@ -627,6 +627,7 @@ refine_clustering(
                store_transposed ? d_srcs : d_dsts,
                d_weights,
                std::ignore,
+               std::ignore,
                std::ignore) =
         cugraph::detail::shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning<
           vertex_t,

cpp/src/community/edge_triangle_count_impl.cuh

@@ -136,7 +136,7 @@ edge_property_t<graph_view_t<vertex_t, edge_t, false, multi_gpu>, edge_t> edge_t
   auto edge_first = thrust::make_zip_iterator(edgelist_srcs.begin(), edgelist_dsts.begin());
 
   size_t edges_to_intersect_per_iteration =
-    static_cast<size_t>(handle.get_device_properties().multiProcessorCount) * (1 << 17);
+    static_cast<size_t>(handle.get_device_properties().multiProcessorCount) * (1 << 13);
 
   auto num_chunks =
     raft::div_rounding_up_safe(edgelist_srcs.size(), edges_to_intersect_per_iteration);
@@ -250,7 +250,7 @@ edge_property_t<graph_view_t<vertex_t, edge_t, false, multi_gpu>, edge_t> edge_t
                          handle.get_stream());
 
       // There are still multiple copies here but is it worth sorting and reducing again?
-      std::tie(pair_srcs, pair_dsts, std::ignore, pair_count, std::ignore) =
+      std::tie(pair_srcs, pair_dsts, std::ignore, pair_count, std::ignore, std::ignore) =
         shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning<vertex_t,
                                                                                edge_t,
                                                                                weight_t,