Improve MG Louvain scalability (#2062)

Pulls updates from PRs #2044, better be reviewed and merged after PR #2044

Mainly experimenting, not ready for review.

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

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

URL: #2062

Author: seunghwak

cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -485,7 +485,7 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle,
   [[maybe_unused]] std::conditional_t<GraphViewType::is_adj_matrix_transposed,
                                       row_properties_t<GraphViewType, T>,
                                       col_properties_t<GraphViewType, T>>
-    minor_tmp_buffer{};  // relevant only when (GraphViewType::is_multi_gpu && !update_major
+    minor_tmp_buffer(handle);  // relevant only when (GraphViewType::is_multi_gpu && !update_major
   if constexpr (GraphViewType::is_multi_gpu && !update_major) {
     if constexpr (GraphViewType::is_adj_matrix_transposed) {
       minor_tmp_buffer = row_properties_t<GraphViewType, T>(handle, graph_view);

cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -173,6 +173,18 @@ class minor_properties_device_view_t {
   {
   }
 
+  std::optional<vertex_t const*> key_data() const
+  {
+    return key_first_ ? std::optional<vertex_t const*>{*key_first_} : std::nullopt;
+  }
+
+  std::optional<vertex_t> number_of_keys() const
+  {
+    return key_first_ ? std::optional<vertex_t>{static_cast<vertex_t>(
+                          thrust::distance(*key_first_, *key_last_))}
+                      : std::nullopt;
+  }
+
   ValueIterator value_data() const { return value_first_; }
 
   __device__ ValueIterator get_iter(vertex_t offset) const
@@ -199,7 +211,10 @@ class minor_properties_device_view_t {
 template <typename vertex_t, typename T>
 class major_properties_t {
  public:
-  major_properties_t() : buffer_(allocate_dataframe_buffer<T>(0, rmm::cuda_stream_view{})) {}
+  major_properties_t(raft::handle_t const& handle)
+    : buffer_(allocate_dataframe_buffer<T>(size_t{0}, handle.get_stream()))
+  {
+  }
 
   major_properties_t(raft::handle_t const& handle, vertex_t buffer_size)
     : buffer_(allocate_dataframe_buffer<T>(buffer_size, handle.get_stream()))
@@ -227,6 +242,19 @@ class major_properties_t {
   {
   }
 
+  void clear(raft::handle_t const& handle)
+  {
+    key_first_ = std::nullopt;
+
+    resize_dataframe_buffer(buffer_, size_t{0}, handle.get_stream());
+    shrink_to_fit_dataframe_buffer(buffer_, handle.get_stream());
+
+    matrix_partition_key_offsets_  = std::nullopt;
+    matrix_partition_major_firsts_ = std::nullopt;
+
+    matrix_partition_major_value_start_offsets_ = std::nullopt;
+  }
+
   void fill(T value, rmm::cuda_stream_view stream)
   {
     thrust::fill(
@@ -279,7 +307,7 @@ class major_properties_t {
  private:
   std::optional<vertex_t const*> key_first_{std::nullopt};
 
-  decltype(allocate_dataframe_buffer<T>(0, rmm::cuda_stream_view{})) buffer_;
+  decltype(allocate_dataframe_buffer<T>(size_t{0}, rmm::cuda_stream_view{})) buffer_;
 
   std::optional<std::vector<vertex_t>> matrix_partition_key_offsets_{std::nullopt};
   std::optional<std::vector<vertex_t>> matrix_partition_major_firsts_{std::nullopt};
@@ -290,7 +318,10 @@ class major_properties_t {
 template <typename vertex_t, typename T>
 class minor_properties_t {
  public:
-  minor_properties_t() : buffer_(allocate_dataframe_buffer<T>(0, rmm::cuda_stream_view{})) {}
+  minor_properties_t(raft::handle_t const& handle)
+    : buffer_(allocate_dataframe_buffer<T>(size_t{0}, handle.get_stream()))
+  {
+  }
 
   minor_properties_t(raft::handle_t const& handle, vertex_t buffer_size)
     : buffer_(allocate_dataframe_buffer<T>(buffer_size, handle.get_stream()))
@@ -309,6 +340,16 @@ class minor_properties_t {
   {
   }
 
+  void clear(raft::handle_t const& handle)
+  {
+    key_first_   = std::nullopt;
+    key_last_    = std::nullopt;
+    minor_first_ = std::nullopt;
+
+    resize_dataframe_buffer(buffer_, size_t{0}, handle.get_stream());
+    shrink_to_fit_dataframe_buffer(buffer_, handle.get_stream());
+  }
+
   void fill(T value, rmm::cuda_stream_view stream)
   {
     thrust::fill(
@@ -346,7 +387,7 @@ class minor_properties_t {
   std::optional<vertex_t const*> key_last_{std::nullopt};
   std::optional<vertex_t> minor_first_{std::nullopt};
 
-  decltype(allocate_dataframe_buffer<T>(0, rmm::cuda_stream_view{})) buffer_;
+  decltype(allocate_dataframe_buffer<T>(size_t{0}, rmm::cuda_stream_view{})) buffer_;
 };
 
 template <typename Iterator,
@@ -380,9 +421,10 @@ class row_properties_t {
 
   static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic<T>::value);
 
-  row_properties_t() = default;
+  row_properties_t(raft::handle_t const& handle) : properties_(handle) {}
 
   row_properties_t(raft::handle_t const& handle, GraphViewType const& graph_view)
+    : properties_(handle)
   {
     using vertex_t = typename GraphViewType::vertex_type;
 
@@ -433,6 +475,8 @@ class row_properties_t {
     }
   }
 
+  void clear(raft::handle_t const& handle) { properties_.clear(handle); }
+
   void fill(T value, rmm::cuda_stream_view stream) { properties_.fill(value, stream); }
 
   auto key_first() { return properties_.key_first(); }
@@ -447,7 +491,7 @@ class row_properties_t {
   std::conditional_t<GraphViewType::is_adj_matrix_transposed,
                      detail::minor_properties_t<typename GraphViewType::vertex_type, T>,
                      detail::major_properties_t<typename GraphViewType::vertex_type, T>>
-    properties_{};
+    properties_;
 };
 
 template <typename GraphViewType, typename T>
@@ -457,9 +501,10 @@ class col_properties_t {
 
   static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic<T>::value);
 
-  col_properties_t() = default;
+  col_properties_t(raft::handle_t const& handle) : properties_(handle) {}
 
   col_properties_t(raft::handle_t const& handle, GraphViewType const& graph_view)
+    : properties_(handle)
   {
     using vertex_t = typename GraphViewType::vertex_type;
 
@@ -510,6 +555,8 @@ class col_properties_t {
     }
   }
 
+  void clear(raft::handle_t const& handle) { properties_.clear(handle); }
+
   void fill(T value, rmm::cuda_stream_view stream) { properties_.fill(value, stream); }
 
   auto key_first() { return properties_.key_first(); }
@@ -524,7 +571,7 @@ class col_properties_t {
   std::conditional_t<GraphViewType::is_adj_matrix_transposed,
                      detail::major_properties_t<typename GraphViewType::vertex_type, T>,
                      detail::minor_properties_t<typename GraphViewType::vertex_type, T>>
-    properties_{};
+    properties_;
 };
 
 template <typename vertex_t>