Skip to content

Commit 9f16be4

Browse files
seunghwakseunghwak
authored
Merge becac46 into 1c47f70 
2 parents 1c47f70 + becac46 commit 9f16be4

File tree

9 files changed

+1434
-8
lines changed

9 files changed

+1434
-8
lines changed

‎cpp/include/cugraph/algorithms.hpp‎

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

‎cpp/include/cugraph/graph.hpp‎

Lines changed: 98 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -97,11 +97,56 @@ class graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enab
9797
graph_meta_t<vertex_t, edge_t, multi_gpu> meta,
9898
bool do_expensive_check = false);
9999

100-
// return a new renumber_map
100+
/**
101+
* @brief Symmetrize this graph.
102+
*
103+
* @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
104+
* handles to various CUDA libraries) to run graph algorithms.
105+
* @param renumber_map Renumber map to recover the original vertex IDs from the renumbered vertex
106+
* IDs.
107+
* @param reciprocal If true, an edge is kept only when the reversed edge also exists. If false,
108+
* keep (and symmetrize) all the edges that appear only in one direction.
109+
* @return rmm::device_uvector<vertex_t> Return a new renumber map (to recover the original vertex
110+
* IDs).
111+
*/
101112
rmm::device_uvector<vertex_t> symmetrize(raft::handle_t const& handle,
102113
rmm::device_uvector<vertex_t>&& renumber_map,
103114
bool reciprocal = false);
104115

116+
/**
117+
* @brief Transpose this graph.
118+
*
119+
* @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
120+
* handles to various CUDA libraries) to run graph algorithms.
121+
* @param renumber_map Renumber map to recover the original vertex IDs from the renumbered vertex
122+
* IDs.
123+
* @return rmm::device_uvector<vertex_t> Return a new renumber map (to recover the original vertex
124+
* IDs).
125+
*/
126+
rmm::device_uvector<vertex_t> transpose(raft::handle_t const& handle,
127+
rmm::device_uvector<vertex_t>&& renumber_map);
128+
129+
/**
130+
* @brief Transpose the storage format (no change in actual graph).
131+
*
132+
* In SG, convert between CSR and CSC. In multi-GPU, currently convert between CSR + DCSR hybrid
133+
* and CSC + DCSC hybrid (but the internal representation in multi-GPU is subject to change).
134+
*
135+
* @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
136+
* handles to various CUDA libraries) to run graph algorithms.
137+
* @param renumber_map Renumber map to recover the original vertex IDs from the renumbered vertex
138+
* IDs.
139+
* @param destroy If true, destroy this graph to free-up memory.
140+
* @return std::tuple<graph_t<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>,
141+
* rmm::device_uvector<vertex_t>> Return a storage transposed graph and a new renumber map (to
142+
* recover the original vertex IDs for the returned graph).
143+
*/
144+
std::tuple<graph_t<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>,
145+
rmm::device_uvector<vertex_t>>
146+
transpose_storage(raft::handle_t const& handle,
147+
rmm::device_uvector<vertex_t>&& renumber_map,
148+
bool destroy = false);
149+
105150
bool is_weighted() const { return adj_matrix_partition_weights_.has_value(); }
106151

107152
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
169214
std::optional<rmm::device_uvector<weight_t>>>
170215
decompress_to_edgelist(raft::handle_t const& handle,
171216
std::optional<rmm::device_uvector<vertex_t>> const& renumber_map,
172-
bool destroy);
217+
bool destroy = false);
173218

174219
private:
175220
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
220265
graph_meta_t<vertex_t, edge_t, multi_gpu> meta,
221266
bool do_expensive_check = false);
222267

223-
// return a new renumber_map if @p renumber_map is valid
268+
/**
269+
* @brief Symmetrize this graph.
270+
*
271+
* @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
272+
* handles to various CUDA libraries) to run graph algorithms.
273+
* @param renumber_map Optional renumber map to recover the original vertex IDs from the
274+
* renumbered vertex IDs. If @p renuber_map.has_value() is false, this function assumes that
275+
* vertex IDs are not renumbered.
276+
* @param reciprocal If true, an edge is kept only when the reversed edge also exists. If false,
277+
* keep (and symmetrize) all the edges that appear only in one direction.
278+
* @return rmm::device_uvector<vertex_t> Return a new renumber map (to recover the original vertex
279+
* IDs) if @p renumber_map.has_value() is true.
280+
*/
224281
std::optional<rmm::device_uvector<vertex_t>> symmetrize(
225282
raft::handle_t const& handle,
226283
std::optional<rmm::device_uvector<vertex_t>>&& renumber_map,
227284
bool reciprocal = false);
228285

286+
/**
287+
* @brief Transpose this graph.
288+
*
289+
* @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
290+
* handles to various CUDA libraries) to run graph algorithms.
291+
* @param renumber_map Optional renumber map to recover the original vertex IDs from the
292+
* renumbered vertex IDs. If @p renuber_map.has_value() is false, this function assumes that
293+
* vertex IDs are not renumbered.
294+
* @return rmm::device_uvector<vertex_t> Return a new renumber map (to recover the original vertex
295+
* IDs) if @p renumber_map.has_value() is true.
296+
*/
297+
std::optional<rmm::device_uvector<vertex_t>> transpose(
298+
raft::handle_t const& handle, std::optional<rmm::device_uvector<vertex_t>>&& renumber_map);
299+
300+
/**
301+
* @brief Transpose the storage format (no change in actual graph).
302+
*
303+
* In SG, convert between CSR and CSC. In multi-GPU, currently convert between CSR + DCSR hybrid
304+
* and CSC + DCSC hybrid (but the internal representation in multi-GPU is subject to change).
305+
*
306+
* @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
307+
* handles to various CUDA libraries) to run graph algorithms.
308+
* @param renumber_map Optional renumber map to recover the original vertex IDs from the
309+
* renumbered vertex IDs. If @p renuber_map.has_value() is false, this function assumes that
310+
* vertex IDs are not renumbered.
311+
* @param destroy If true, destroy this graph to free-up memory.
312+
* @return std::tuple<graph_t<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>,
313+
* rmm::device_uvector<vertex_t>> Return a storage transposed graph and a optional new renumber
314+
* map (to recover the original vertex IDs for the returned graph) The returned optional new
315+
* renumber map is valid only if @p renumber_map.has_value() is true.
316+
*/
317+
std::tuple<graph_t<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>,
318+
std::optional<rmm::device_uvector<vertex_t>>>
319+
transpose_storage(raft::handle_t const& handle,
320+
std::optional<rmm::device_uvector<vertex_t>>&& renumber_map,
321+
bool destroy = false);
322+
229323
bool is_weighted() const { return weights_.has_value(); }
230324

231325
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
270364
std::optional<rmm::device_uvector<weight_t>>>
271365
decompress_to_edgelist(raft::handle_t const& handle,
272366
std::optional<rmm::device_uvector<vertex_t>> const& renumber_map,
273-
bool destroy);
367+
bool destroy = false);
274368

275369
private:
276370
friend class cugraph::serializer::serializer_t;

‎cpp/src/structure/graph_impl.cuh‎

Lines changed: 156 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -199,8 +199,6 @@ compress_edgelist(edgelist_t<vertex_t, edge_t, weight_t> const& edgelist,
199199
}
200200
}
201201

202-
// FIXME: need to add an option to sort neighbor lists
203-
204202
return std::make_tuple(
205203
std::move(offsets), std::move(indices), std::move(weights), std::move(dcs_nzd_vertices));
206204
}
@@ -748,6 +746,162 @@ graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_
748746
return std::move(new_renumber_map);
749747
}
750748

749+
template <typename vertex_t,
750+
typename edge_t,
751+
typename weight_t,
752+
bool store_transposed,
753+
bool multi_gpu>
754+
rmm::device_uvector<vertex_t>
755+
graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_t<multi_gpu>>::
756+
transpose(raft::handle_t const& handle, rmm::device_uvector<vertex_t>&& renumber_map)
757+
{
758+
if (this->is_symmetric()) { return std::move(renumber_map); }
759+
760+
auto is_multigraph = this->is_multigraph();
761+
762+
auto wrapped_renumber_map = std::optional<rmm::device_uvector<vertex_t>>(std::move(renumber_map));
763+
764+
auto [edgelist_rows, edgelist_cols, edgelist_weights] =
765+
this->decompress_to_edgelist(handle, wrapped_renumber_map, true);
766+
767+
std::tie(store_transposed ? edgelist_rows : edgelist_cols,
768+
store_transposed ? edgelist_cols : edgelist_rows,
769+
edgelist_weights) =
770+
detail::shuffle_edgelist_by_gpu_id(handle,
771+
std::move(store_transposed ? edgelist_rows : edgelist_cols),
772+
std::move(store_transposed ? edgelist_cols : edgelist_rows),
773+
std::move(edgelist_weights));
774+
775+
auto [transposed_graph, new_renumber_map] =
776+
create_graph_from_edgelist<vertex_t, edge_t, weight_t, store_transposed, multi_gpu>(
777+
handle,
778+
std::move(*wrapped_renumber_map),
779+
std::move(edgelist_cols),
780+
std::move(edgelist_rows),
781+
std::move(edgelist_weights),
782+
graph_properties_t{is_multigraph, false},
783+
true);
784+
*this = std::move(transposed_graph);
785+
786+
return std::move(*new_renumber_map);
787+
}
788+
789+
template <typename vertex_t,
790+
typename edge_t,
791+
typename weight_t,
792+
bool store_transposed,
793+
bool multi_gpu>
794+
std::optional<rmm::device_uvector<vertex_t>>
795+
graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_t<!multi_gpu>>::
796+
transpose(raft::handle_t const& handle,
797+
std::optional<rmm::device_uvector<vertex_t>>&& renumber_map)
798+
{
799+
if (this->is_symmetric()) { return std::move(renumber_map); }
800+
801+
auto number_of_vertices = this->get_number_of_vertices();
802+
auto is_multigraph = this->is_multigraph();
803+
bool renumber = renumber_map.has_value();
804+
805+
auto [edgelist_rows, edgelist_cols, edgelist_weights] =
806+
this->decompress_to_edgelist(handle, renumber_map, true);
807+
auto vertex_span = renumber ? std::move(renumber_map)
808+
: std::make_optional<rmm::device_uvector<vertex_t>>(
809+
number_of_vertices, handle.get_stream());
810+
if (!renumber) {
811+
thrust::sequence(
812+
handle.get_thrust_policy(), (*vertex_span).begin(), (*vertex_span).end(), vertex_t{0});
813+
}
814+
815+
auto [transposed_graph, new_renumber_map] =
816+
create_graph_from_edgelist<vertex_t, edge_t, weight_t, store_transposed, multi_gpu>(
817+
handle,
818+
std::move(vertex_span),
819+
std::move(edgelist_cols),
820+
std::move(edgelist_rows),
821+
std::move(edgelist_weights),
822+
graph_properties_t{is_multigraph, false},
823+
renumber);
824+
*this = std::move(transposed_graph);
825+
826+
return std::move(new_renumber_map);
827+
}
828+
829+
template <typename vertex_t,
830+
typename edge_t,
831+
typename weight_t,
832+
bool store_transposed,
833+
bool multi_gpu>
834+
std::tuple<graph_t<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>,
835+
rmm::device_uvector<vertex_t>>
836+
graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_t<multi_gpu>>::
837+
transpose_storage(raft::handle_t const& handle,
838+
rmm::device_uvector<vertex_t>&& renumber_map,
839+
bool destroy)
840+
{
841+
auto is_multigraph = this->is_multigraph();
842+
843+
auto wrapped_renumber_map = std::optional<rmm::device_uvector<vertex_t>>(std::move(renumber_map));
844+
845+
auto [edgelist_rows, edgelist_cols, edgelist_weights] =
846+
this->decompress_to_edgelist(handle, wrapped_renumber_map, destroy);
847+
848+
std::tie(!store_transposed ? edgelist_cols : edgelist_rows,
849+
!store_transposed ? edgelist_rows : edgelist_cols,
850+
edgelist_weights) =
851+
detail::shuffle_edgelist_by_gpu_id(handle,
852+
std::move(!store_transposed ? edgelist_cols : edgelist_rows),
853+
std::move(!store_transposed ? edgelist_rows : edgelist_cols),
854+
std::move(edgelist_weights));
855+
856+
auto [storage_transposed_graph, new_renumber_map] =
857+
create_graph_from_edgelist<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>(
858+
handle,
859+
std::move(*wrapped_renumber_map),
860+
std::move(edgelist_rows),
861+
std::move(edgelist_cols),
862+
std::move(edgelist_weights),
863+
graph_properties_t{is_multigraph, false},
864+
true);
865+
866+
return std::make_tuple(std::move(storage_transposed_graph), std::move(*new_renumber_map));
867+
}
868+
869+
template <typename vertex_t,
870+
typename edge_t,
871+
typename weight_t,
872+
bool store_transposed,
873+
bool multi_gpu>
874+
std::tuple<graph_t<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>,
875+
std::optional<rmm::device_uvector<vertex_t>>>
876+
graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_t<!multi_gpu>>::
877+
transpose_storage(raft::handle_t const& handle,
878+
std::optional<rmm::device_uvector<vertex_t>>&& renumber_map,
879+
bool destroy)
880+
{
881+
auto number_of_vertices = this->get_number_of_vertices();
882+
auto is_multigraph = this->is_multigraph();
883+
bool renumber = renumber_map.has_value();
884+
885+
auto [edgelist_rows, edgelist_cols, edgelist_weights] =
886+
this->decompress_to_edgelist(handle, renumber_map, destroy);
887+
auto vertex_span = renumber ? std::move(renumber_map)
888+
: std::make_optional<rmm::device_uvector<vertex_t>>(
889+
number_of_vertices, handle.get_stream());
890+
if (!renumber) {
891+
thrust::sequence(
892+
handle.get_thrust_policy(), (*vertex_span).begin(), (*vertex_span).end(), vertex_t{0});
893+
}
894+
895+
return create_graph_from_edgelist<vertex_t, edge_t, weight_t, !store_transposed, multi_gpu>(
896+
handle,
897+
std::move(vertex_span),
898+
std::move(edgelist_cols),
899+
std::move(edgelist_rows),
900+
std::move(edgelist_weights),
901+
graph_properties_t{is_multigraph, false},
902+
renumber);
903+
}
904+
751905
template <typename vertex_t,
752906
typename edge_t,
753907
typename weight_t,

‎cpp/tests/CMakeLists.txt‎

Lines changed: 17 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -340,6 +340,14 @@ ConfigureTest(GRAPH_TEST structure/graph_test.cpp)
340340
# - Symmetrize tests ------------------------------------------------------------------------------
341341
ConfigureTest(SYMMETRIZE_TEST structure/symmetrize_test.cpp)
342342

343+
###################################################################################################
344+
# - Transpose tests ------------------------------------------------------------------------------
345+
ConfigureTest(TRANSPOSE_TEST structure/transpose_test.cpp)
346+
347+
###################################################################################################
348+
# - Transpose Storage tests -----------------------------------------------------------------------
349+
ConfigureTest(TRANSPOSE_STORAGE_TEST structure/transpose_storage_test.cpp)
350+
343351
###################################################################################################
344352
# - Weight-sum tests ------------------------------------------------------------------------------
345353
ConfigureTest(WEIGHT_SUM_TEST structure/weight_sum_test.cpp)
@@ -440,9 +448,17 @@ if(BUILD_CUGRAPH_MG_TESTS)
440448

441449
if(MPI_CXX_FOUND)
442450
###########################################################################################
443-
# - MG SYMMETRIZE tests ---------------------------------------------------------------------
451+
# - MG SYMMETRIZE tests -------------------------------------------------------------------
444452
ConfigureTestMG(MG_SYMMETRIZE_TEST structure/mg_symmetrize_test.cpp)
445453

454+
###########################################################################################
455+
# - MG Transpose tests --------------------------------------------------------------------
456+
ConfigureTestMG(MG_TRANSPOSE_TEST structure/mg_transpose_test.cpp)
457+
458+
###########################################################################################
459+
# - MG Transpose Storage tests ------------------------------------------------------------
460+
ConfigureTestMG(MG_TRANSPOSE_STORAGE_TEST structure/mg_transpose_storage_test.cpp)
461+
446462
###########################################################################################
447463
# - MG PAGERANK tests ---------------------------------------------------------------------
448464
ConfigureTestMG(MG_PAGERANK_TEST link_analysis/mg_pagerank_test.cpp)

‎cpp/tests/structure/mg_symmetrize_test.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -139,7 +139,7 @@ class Tests_MGSymmetrize
139139
cugraph::test::construct_graph<vertex_t, edge_t, weight_t, store_transposed, false>(
140140
handle, input_usecase, symmetrize_usecase.test_weighted, false);
141141

142-
// 4-4. run SG symmetrize Centrality
142+
// 4-4. run SG symmetrize
143143

144144
auto d_sg_renumber_map_labels =
145145
sg_graph.symmetrize(handle, std::nullopt, symmetrize_usecase.reciprocal);

0 commit comments

Comments
 (0)