Transpose

cpp/include/cugraph/algorithms.hpp

@@ -1387,6 +1387,7 @@ random_walks(raft::handle_t const& handle,
  * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false)
  * or multi-GPU (true).
  * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
+ * handles to various CUDA libraries) to run graph algorithms.
  * @param graph_view Graph view object.
  * @param components Pointer to the output component ID array.
  * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).

cpp/include/cugraph/graph.hpp

@@ -97,11 +97,56 @@ class graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enab
           graph_meta_t<vertex_t, edge_t, multi_gpu> meta,
           bool do_expensive_check = false);
 
-  // return a new renumber_map
+  /**
+   * @brief Symmetrize this graph.
+   *
+   * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
+   * handles to various CUDA libraries) to run graph algorithms.
+   * @param renumber_map Renumber map to recover the original vertex IDs from the renumbered vertex
+   * IDs.
+   * @param reciprocal If true, an edge is kept only when the reversed edge also exists. If false,
+   * keep (and symmetrize) all the edges that appear only in one direction.
+   * @return rmm::device_uvector<vertex_t> Return a new renumber map (to recover the original vertex
+   * IDs).
+   */
   rmm::device_uvector<vertex_t> symmetrize(raft::handle_t const& handle,
                                            rmm::device_uvector<vertex_t>&& renumber_map,
                                            bool reciprocal = false);
 
+  /**
+   * @brief Transpose this graph.
+   *
+   * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
+   * handles to various CUDA libraries) to run graph algorithms.
+   * @param renumber_map Renumber map to recover the original vertex IDs from the renumbered vertex
+   * IDs.
+   * @return rmm::device_uvector<vertex_t> Return a new renumber map (to recover the original vertex
+   * IDs).
+   */
+  rmm::device_uvector<vertex_t> transpose(raft::handle_t const& handle,
+                                          rmm::device_uvector<vertex_t>&& renumber_map);
+
+  /**
+   * @brief Transpose the storage format (no change in actual graph).
+   *
+   * In SG, convert between CSR and CSC. In multi-GPU, currently convert between CSR + DCSR hybrid
+   * and CSC + DCSC hybrid (but the internal representation in multi-GPU is subject to change).
+   *
+   * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
+   * handles to various CUDA libraries) to run graph algorithms.
+   * @param renumber_map Renumber map to recover the original vertex IDs from the renumbered vertex
+   * IDs.
+   * @param destroy If true, destroy this graph to free-up memory.
+   * @return std::tuple<graph_t<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>,
+   * rmm::device_uvector<vertex_t>> Return a storage transposed graph and a new renumber map (to
+   * recover the original vertex IDs for the returned graph).
+   */
+  std::tuple<graph_t<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>,
+             rmm::device_uvector<vertex_t>>
+  transpose_storage(raft::handle_t const& handle,
+                    rmm::device_uvector<vertex_t>&& renumber_map,
+                    bool destroy = false);
+
   bool is_weighted() const { return adj_matrix_partition_weights_.has_value(); }
 
   graph_view_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu> view() const
@@ -169,7 +214,7 @@ class graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enab
              std::optional<rmm::device_uvector<weight_t>>>
   decompress_to_edgelist(raft::handle_t const& handle,
                          std::optional<rmm::device_uvector<vertex_t>> const& renumber_map,
-                         bool destroy);
+                         bool destroy = false);
 
  private:
   std::vector<rmm::device_uvector<edge_t>> adj_matrix_partition_offsets_{};
@@ -220,12 +265,61 @@ class graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enab
           graph_meta_t<vertex_t, edge_t, multi_gpu> meta,
           bool do_expensive_check = false);
 
-  // return a new renumber_map if @p renumber_map is valid
+  /**
+   * @brief Symmetrize this graph.
+   *
+   * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
+   * handles to various CUDA libraries) to run graph algorithms.
+   * @param renumber_map Optional renumber map to recover the original vertex IDs from the
+   * renumbered vertex IDs. If @p renuber_map.has_value() is false, this function assumes that
+   * vertex IDs are not renumbered.
+   * @param reciprocal If true, an edge is kept only when the reversed edge also exists. If false,
+   * keep (and symmetrize) all the edges that appear only in one direction.
+   * @return rmm::device_uvector<vertex_t> Return a new renumber map (to recover the original vertex
+   * IDs) if @p renumber_map.has_value() is true.
+   */
   std::optional<rmm::device_uvector<vertex_t>> symmetrize(
     raft::handle_t const& handle,
     std::optional<rmm::device_uvector<vertex_t>>&& renumber_map,
     bool reciprocal = false);
 
+  /**
+   * @brief Transpose this graph.
+   *
+   * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
+   * handles to various CUDA libraries) to run graph algorithms.
+   * @param renumber_map Optional renumber map to recover the original vertex IDs from the
+   * renumbered vertex IDs. If @p renuber_map.has_value() is false, this function assumes that
+   * vertex IDs are not renumbered.
+   * @return rmm::device_uvector<vertex_t> Return a new renumber map (to recover the original vertex
+   * IDs) if @p renumber_map.has_value() is true.
+   */
+  std::optional<rmm::device_uvector<vertex_t>> transpose(
+    raft::handle_t const& handle, std::optional<rmm::device_uvector<vertex_t>>&& renumber_map);
+
+  /**
+   * @brief Transpose the storage format (no change in actual graph).
+   *
+   * In SG, convert between CSR and CSC. In multi-GPU, currently convert between CSR + DCSR hybrid
+   * and CSC + DCSC hybrid (but the internal representation in multi-GPU is subject to change).
+   *
+   * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
+   * handles to various CUDA libraries) to run graph algorithms.
+   * @param renumber_map Optional renumber map to recover the original vertex IDs from the
+   * renumbered vertex IDs. If @p renuber_map.has_value() is false, this function assumes that
+   * vertex IDs are not renumbered.
+   * @param destroy If true, destroy this graph to free-up memory.
+   * @return std::tuple<graph_t<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>,
+   * rmm::device_uvector<vertex_t>> Return a storage transposed graph and a optional new renumber
+   * map (to recover the original vertex IDs for the returned graph) The returned optional new
+   * renumber map is valid only if @p renumber_map.has_value() is true.
+   */
+  std::tuple<graph_t<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>,
+             std::optional<rmm::device_uvector<vertex_t>>>
+  transpose_storage(raft::handle_t const& handle,
+                    std::optional<rmm::device_uvector<vertex_t>>&& renumber_map,
+                    bool destroy = false);
+
   bool is_weighted() const { return weights_.has_value(); }
 
   graph_view_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu> view() const
@@ -270,7 +364,7 @@ class graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enab
              std::optional<rmm::device_uvector<weight_t>>>
   decompress_to_edgelist(raft::handle_t const& handle,
                          std::optional<rmm::device_uvector<vertex_t>> const& renumber_map,
-                         bool destroy);
+                         bool destroy = false);
 
  private:
   friend class cugraph::serializer::serializer_t;

cpp/src/structure/graph_impl.cuh

@@ -199,8 +199,6 @@ compress_edgelist(edgelist_t<vertex_t, edge_t, weight_t> const& edgelist,
     }
   }
 
-  // FIXME: need to add an option to sort neighbor lists
-
   return std::make_tuple(
     std::move(offsets), std::move(indices), std::move(weights), std::move(dcs_nzd_vertices));
 }
@@ -748,6 +746,162 @@ graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_
   return std::move(new_renumber_map);
 }
 
+template <typename vertex_t,
+          typename edge_t,
+          typename weight_t,
+          bool store_transposed,
+          bool multi_gpu>
+rmm::device_uvector<vertex_t>
+graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_t<multi_gpu>>::
+  transpose(raft::handle_t const& handle, rmm::device_uvector<vertex_t>&& renumber_map)
+{
+  if (this->is_symmetric()) { return std::move(renumber_map); }
+
+  auto is_multigraph = this->is_multigraph();
+
+  auto wrapped_renumber_map = std::optional<rmm::device_uvector<vertex_t>>(std::move(renumber_map));
+
+  auto [edgelist_rows, edgelist_cols, edgelist_weights] =
+    this->decompress_to_edgelist(handle, wrapped_renumber_map, true);
+
+  std::tie(store_transposed ? edgelist_rows : edgelist_cols,
+           store_transposed ? edgelist_cols : edgelist_rows,
+           edgelist_weights) =
+    detail::shuffle_edgelist_by_gpu_id(handle,
+                                       std::move(store_transposed ? edgelist_rows : edgelist_cols),
+                                       std::move(store_transposed ? edgelist_cols : edgelist_rows),
+                                       std::move(edgelist_weights));
+
+  auto [transposed_graph, new_renumber_map] =
+    create_graph_from_edgelist<vertex_t, edge_t, weight_t, store_transposed, multi_gpu>(
+      handle,
+      std::move(*wrapped_renumber_map),
+      std::move(edgelist_cols),
+      std::move(edgelist_rows),
+      std::move(edgelist_weights),
+      graph_properties_t{is_multigraph, false},
+      true);
+  *this = std::move(transposed_graph);
+
+  return std::move(*new_renumber_map);
+}
+
+template <typename vertex_t,
+          typename edge_t,
+          typename weight_t,
+          bool store_transposed,
+          bool multi_gpu>
+std::optional<rmm::device_uvector<vertex_t>>
+graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_t<!multi_gpu>>::
+  transpose(raft::handle_t const& handle,
+            std::optional<rmm::device_uvector<vertex_t>>&& renumber_map)
+{
+  if (this->is_symmetric()) { return std::move(renumber_map); }
+
+  auto number_of_vertices = this->get_number_of_vertices();
+  auto is_multigraph      = this->is_multigraph();
+  bool renumber           = renumber_map.has_value();
+
+  auto [edgelist_rows, edgelist_cols, edgelist_weights] =
+    this->decompress_to_edgelist(handle, renumber_map, true);
+  auto vertex_span = renumber ? std::move(renumber_map)
+                              : std::make_optional<rmm::device_uvector<vertex_t>>(
+                                  number_of_vertices, handle.get_stream());
+  if (!renumber) {
+    thrust::sequence(
+      handle.get_thrust_policy(), (*vertex_span).begin(), (*vertex_span).end(), vertex_t{0});
+  }
+
+  auto [transposed_graph, new_renumber_map] =
+    create_graph_from_edgelist<vertex_t, edge_t, weight_t, store_transposed, multi_gpu>(
+      handle,
+      std::move(vertex_span),
+      std::move(edgelist_cols),
+      std::move(edgelist_rows),
+      std::move(edgelist_weights),
+      graph_properties_t{is_multigraph, false},
+      renumber);
+  *this = std::move(transposed_graph);
+
+  return std::move(new_renumber_map);
+}
+
+template <typename vertex_t,
+          typename edge_t,
+          typename weight_t,
+          bool store_transposed,
+          bool multi_gpu>
+std::tuple<graph_t<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>,
+           rmm::device_uvector<vertex_t>>
+graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_t<multi_gpu>>::
+  transpose_storage(raft::handle_t const& handle,
+                    rmm::device_uvector<vertex_t>&& renumber_map,
+                    bool destroy)
+{
+  auto is_multigraph = this->is_multigraph();
+
+  auto wrapped_renumber_map = std::optional<rmm::device_uvector<vertex_t>>(std::move(renumber_map));
+
+  auto [edgelist_rows, edgelist_cols, edgelist_weights] =
+    this->decompress_to_edgelist(handle, wrapped_renumber_map, destroy);
+
+  std::tie(!store_transposed ? edgelist_cols : edgelist_rows,
+           !store_transposed ? edgelist_rows : edgelist_cols,
+           edgelist_weights) =
+    detail::shuffle_edgelist_by_gpu_id(handle,
+                                       std::move(!store_transposed ? edgelist_cols : edgelist_rows),
+                                       std::move(!store_transposed ? edgelist_rows : edgelist_cols),
+                                       std::move(edgelist_weights));
+
+  auto [storage_transposed_graph, new_renumber_map] =
+    create_graph_from_edgelist<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>(
+      handle,
+      std::move(*wrapped_renumber_map),
+      std::move(edgelist_rows),
+      std::move(edgelist_cols),
+      std::move(edgelist_weights),
+      graph_properties_t{is_multigraph, false},
+      true);
+
+  return std::make_tuple(std::move(storage_transposed_graph), std::move(*new_renumber_map));
+}
+
+template <typename vertex_t,
+          typename edge_t,
+          typename weight_t,
+          bool store_transposed,
+          bool multi_gpu>
+std::tuple<graph_t<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>,
+           std::optional<rmm::device_uvector<vertex_t>>>
+graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_t<!multi_gpu>>::
+  transpose_storage(raft::handle_t const& handle,
+                    std::optional<rmm::device_uvector<vertex_t>>&& renumber_map,
+                    bool destroy)
+{
+  auto number_of_vertices = this->get_number_of_vertices();
+  auto is_multigraph      = this->is_multigraph();
+  bool renumber           = renumber_map.has_value();
+
+  auto [edgelist_rows, edgelist_cols, edgelist_weights] =
+    this->decompress_to_edgelist(handle, renumber_map, destroy);
+  auto vertex_span = renumber ? std::move(renumber_map)
+                              : std::make_optional<rmm::device_uvector<vertex_t>>(
+                                  number_of_vertices, handle.get_stream());
+  if (!renumber) {
+    thrust::sequence(
+      handle.get_thrust_policy(), (*vertex_span).begin(), (*vertex_span).end(), vertex_t{0});
+  }
+
+  return create_graph_from_edgelist<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>(
+    handle,
+    std::move(vertex_span),
+    std::move(edgelist_cols),
+    std::move(edgelist_rows),
+    std::move(edgelist_weights),
+    graph_properties_t{is_multigraph, false},
+    renumber);
+}
+
 template <typename vertex_t,
           typename edge_t,
           typename weight_t,

cpp/tests/CMakeLists.txt

@@ -340,6 +340,14 @@ ConfigureTest(GRAPH_TEST structure/graph_test.cpp)
 # - Symmetrize tests ------------------------------------------------------------------------------
 ConfigureTest(SYMMETRIZE_TEST structure/symmetrize_test.cpp)
 
+###################################################################################################
+# - Transpose tests ------------------------------------------------------------------------------
+ConfigureTest(TRANSPOSE_TEST structure/transpose_test.cpp)
+
+###################################################################################################
+# - Transpose Storage tests -----------------------------------------------------------------------
+ConfigureTest(TRANSPOSE_STORAGE_TEST structure/transpose_storage_test.cpp)
+
 ###################################################################################################
 # - Weight-sum tests ------------------------------------------------------------------------------
 ConfigureTest(WEIGHT_SUM_TEST structure/weight_sum_test.cpp)
@@ -440,9 +448,17 @@ if(BUILD_CUGRAPH_MG_TESTS)
 
     if(MPI_CXX_FOUND)
         ###########################################################################################
-        # - MG SYMMETRIZE tests ---------------------------------------------------------------------
+        # - MG SYMMETRIZE tests -------------------------------------------------------------------
         ConfigureTestMG(MG_SYMMETRIZE_TEST structure/mg_symmetrize_test.cpp)
 
+        ###########################################################################################
+        # - MG Transpose tests --------------------------------------------------------------------
+        ConfigureTestMG(MG_TRANSPOSE_TEST structure/mg_transpose_test.cpp)
+
+        ###########################################################################################
+        # - MG Transpose Storage tests ------------------------------------------------------------
+        ConfigureTestMG(MG_TRANSPOSE_STORAGE_TEST structure/mg_transpose_storage_test.cpp)
+
         ###########################################################################################
         # - MG PAGERANK tests ---------------------------------------------------------------------
         ConfigureTestMG(MG_PAGERANK_TEST link_analysis/mg_pagerank_test.cpp)

cpp/tests/structure/mg_symmetrize_test.cpp

@@ -139,7 +139,7 @@ class Tests_MGSymmetrize
           cugraph::test::construct_graph<vertex_t, edge_t, weight_t, store_transposed, false>(
             handle, input_usecase, symmetrize_usecase.test_weighted, false);
 
-        // 4-4. run SG symmetrize Centrality
+        // 4-4. run SG symmetrize
 
         auto d_sg_renumber_map_labels =
           sg_graph.symmetrize(handle, std::nullopt, symmetrize_usecase.reciprocal);