3131
3232#include < thrust/adjacent_difference.h>
3333#include < thrust/binary_search.h>
34+ #include < thrust/equal.h>
3435#include < thrust/fill.h>
3536#include < thrust/for_each.h>
3637#include < thrust/gather.h>
@@ -85,6 +86,161 @@ struct rebase_offset_t {
8586 __device__ edge_t operator ()(edge_t offset) const { return offset - base_offset; }
8687};
8788
89+ template <typename vertex_t ,
90+ typename edge_t ,
91+ typename weight_t ,
92+ bool store_transposed,
93+ bool multi_gpu>
94+ bool check_symmetric (raft::handle_t const & handle,
95+ std::vector<edgelist_t <vertex_t , edge_t , weight_t >> const & edgelists)
96+ {
97+ size_t number_of_local_edges{0 };
98+ for (size_t i = 0 ; i < edgelists.size (); ++i) {
99+ number_of_local_edges += edgelists[i].number_of_edges ;
100+ }
101+
102+ auto is_weighted = edgelists[0 ].p_edge_weights .has_value ();
103+
104+ rmm::device_uvector<vertex_t > org_rows (number_of_local_edges, handle.get_stream ());
105+ rmm::device_uvector<vertex_t > org_cols (number_of_local_edges, handle.get_stream ());
106+ auto org_weights = is_weighted ? std::make_optional<rmm::device_uvector<weight_t >>(
107+ number_of_local_edges, handle.get_stream ())
108+ : std::nullopt ;
109+ size_t offset{0 };
110+ for (size_t i = 0 ; i < edgelists.size (); ++i) {
111+ thrust::copy (handle.get_thrust_policy (),
112+ edgelists[i].p_src_vertices ,
113+ edgelists[i].p_src_vertices + edgelists[i].number_of_edges ,
114+ org_rows.begin () + offset);
115+ thrust::copy (handle.get_thrust_policy (),
116+ edgelists[i].p_dst_vertices ,
117+ edgelists[i].p_dst_vertices + edgelists[i].number_of_edges ,
118+ org_cols.begin () + offset);
119+ if (is_weighted) {
120+ thrust::copy (handle.get_thrust_policy (),
121+ *(edgelists[i].p_edge_weights ),
122+ *(edgelists[i].p_edge_weights ) + edgelists[i].number_of_edges ,
123+ (*org_weights).begin () + offset);
124+ }
125+ offset += edgelists[i].number_of_edges ;
126+ }
127+ if constexpr (multi_gpu) {
128+ std::tie (
129+ store_transposed ? org_cols : org_rows, store_transposed ? org_rows : org_cols, org_weights) =
130+ detail::shuffle_edgelist_by_gpu_id (handle,
131+ std::move (store_transposed ? org_cols : org_rows),
132+ std::move (store_transposed ? org_rows : org_cols),
133+ std::move (org_weights));
134+ }
135+
136+ rmm::device_uvector<vertex_t > symmetrized_rows (org_rows.size (), handle.get_stream ());
137+ rmm::device_uvector<vertex_t > symmetrized_cols (org_cols.size (), handle.get_stream ());
138+ auto symmetrized_weights = org_weights ? std::make_optional<rmm::device_uvector<weight_t >>(
139+ (*org_weights).size (), handle.get_stream ())
140+ : std::nullopt ;
141+ thrust::copy (
142+ handle.get_thrust_policy (), org_rows.begin (), org_rows.end (), symmetrized_rows.begin ());
143+ thrust::copy (
144+ handle.get_thrust_policy (), org_cols.begin (), org_cols.end (), symmetrized_cols.begin ());
145+ if (org_weights) {
146+ thrust::copy (handle.get_thrust_policy (),
147+ (*org_weights).begin (),
148+ (*org_weights).end (),
149+ (*symmetrized_weights).begin ());
150+ }
151+ std::tie (symmetrized_rows, symmetrized_cols, symmetrized_weights) =
152+ symmetrize_edgelist<vertex_t , weight_t , store_transposed, multi_gpu>(
153+ handle,
154+ std::move (symmetrized_rows),
155+ std::move (symmetrized_cols),
156+ std::move (symmetrized_weights),
157+ true );
158+
159+ if (org_rows.size () != symmetrized_rows.size ()) { return false ; }
160+
161+ if (org_weights) {
162+ auto org_edge_first = thrust::make_zip_iterator (
163+ thrust::make_tuple (org_rows.begin (), org_cols.begin (), (*org_weights).begin ()));
164+ thrust::sort (handle.get_thrust_policy (), org_edge_first, org_edge_first + org_rows.size ());
165+ auto symmetrized_edge_first = thrust::make_zip_iterator (thrust::make_tuple (
166+ symmetrized_rows.begin (), symmetrized_cols.begin (), (*symmetrized_weights).begin ()));
167+ thrust::sort (handle.get_thrust_policy (),
168+ symmetrized_edge_first,
169+ symmetrized_edge_first + symmetrized_rows.size ());
170+
171+ return thrust::equal (handle.get_thrust_policy (),
172+ org_edge_first,
173+ org_edge_first + org_rows.size (),
174+ symmetrized_edge_first);
175+ } else {
176+ auto org_edge_first =
177+ thrust::make_zip_iterator (thrust::make_tuple (org_rows.begin (), org_cols.begin ()));
178+ thrust::sort (handle.get_thrust_policy (), org_edge_first, org_edge_first + org_rows.size ());
179+ auto symmetrized_edge_first = thrust::make_zip_iterator (
180+ thrust::make_tuple (symmetrized_rows.begin (), symmetrized_cols.begin ()));
181+ thrust::sort (handle.get_thrust_policy (),
182+ symmetrized_edge_first,
183+ symmetrized_edge_first + symmetrized_rows.size ());
184+
185+ return thrust::equal (handle.get_thrust_policy (),
186+ org_edge_first,
187+ org_edge_first + org_rows.size (),
188+ symmetrized_edge_first);
189+ }
190+ }
191+
192+ template <typename vertex_t , typename edge_t , typename weight_t >
193+ bool check_no_parallel_edge (raft::handle_t const & handle,
194+ std::vector<edgelist_t <vertex_t , edge_t , weight_t >> const & edgelists)
195+ {
196+ size_t number_of_local_edges{0 };
197+ for (size_t i = 0 ; i < edgelists.size (); ++i) {
198+ number_of_local_edges += edgelists[i].number_of_edges ;
199+ }
200+
201+ auto is_weighted = edgelists[0 ].p_edge_weights .has_value ();
202+
203+ rmm::device_uvector<vertex_t > edgelist_rows (number_of_local_edges, handle.get_stream ());
204+ rmm::device_uvector<vertex_t > edgelist_cols (number_of_local_edges, handle.get_stream ());
205+ auto edgelist_weights = is_weighted ? std::make_optional<rmm::device_uvector<weight_t >>(
206+ number_of_local_edges, handle.get_stream ())
207+ : std::nullopt ;
208+ size_t offset{0 };
209+ for (size_t i = 0 ; i < edgelists.size (); ++i) {
210+ thrust::copy (handle.get_thrust_policy (),
211+ edgelists[i].p_src_vertices ,
212+ edgelists[i].p_src_vertices + edgelists[i].number_of_edges ,
213+ edgelist_rows.begin () + offset);
214+ thrust::copy (handle.get_thrust_policy (),
215+ edgelists[i].p_dst_vertices ,
216+ edgelists[i].p_dst_vertices + edgelists[i].number_of_edges ,
217+ edgelist_cols.begin () + offset);
218+ if (is_weighted) {
219+ thrust::copy (handle.get_thrust_policy (),
220+ *(edgelists[i].p_edge_weights ),
221+ *(edgelists[i].p_edge_weights ) + edgelists[i].number_of_edges ,
222+ (*edgelist_weights).begin () + offset);
223+ }
224+ offset += edgelists[i].number_of_edges ;
225+ }
226+
227+ if (edgelist_weights) {
228+ auto edge_first = thrust::make_zip_iterator (thrust::make_tuple (
229+ edgelist_rows.begin (), edgelist_cols.begin (), (*edgelist_weights).begin ()));
230+ thrust::sort (handle.get_thrust_policy (), edge_first, edge_first + edgelist_rows.size ());
231+ return thrust::unique (handle.get_thrust_policy (),
232+ edge_first,
233+ edge_first + edgelist_rows.size ()) == (edge_first + edgelist_rows.size ());
234+ } else {
235+ auto edge_first =
236+ thrust::make_zip_iterator (thrust::make_tuple (edgelist_rows.begin (), edgelist_cols.begin ()));
237+ thrust::sort (handle.get_thrust_policy (), edge_first, edge_first + edgelist_rows.size ());
238+ return thrust::unique (handle.get_thrust_policy (),
239+ edge_first,
240+ edge_first + edgelist_rows.size ()) == (edge_first + edgelist_rows.size ());
241+ }
242+ }
243+
88244// compress edge list (COO) to CSR (or CSC) or CSR + DCSR (CSC + DCSC) hybrid
89245template <bool store_transposed, typename vertex_t , typename edge_t , typename weight_t >
90246std::tuple<rmm::device_uvector<edge_t >,
@@ -410,7 +566,7 @@ graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_
410566 " be nullptr if edgelists[].number_of_edges > 0 and edgelists[].p_edge_weights should be "
411567 " neither std::nullopt nor nullptr if weighted and edgelists[].number_of_edges > 0." );
412568
413- // optional expensive checks (part 1/2)
569+ // optional expensive checks
414570
415571 if (do_expensive_check) {
416572 edge_t number_of_local_edges{0 };
@@ -441,6 +597,20 @@ graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_
441597 partition_.get_vertex_partition_last (comm_size - 1 ) == meta.number_of_vertices ,
442598 " Invalid input argument: vertex partition should cover [0, meta.number_of_vertices)." );
443599
600+ if (this ->is_symmetric ()) {
601+ CUGRAPH_EXPECTS (
602+ (check_symmetric<vertex_t , edge_t , weight_t , store_transposed, multi_gpu>(handle,
603+ edgelists)),
604+ " Invalid input argument: meta.property.is_symmetric is true but the input edge list is not "
605+ " symmetric." );
606+ }
607+ if (!this ->is_multigraph ()) {
608+ CUGRAPH_EXPECTS (
609+ check_no_parallel_edge (handle, edgelists),
610+ " Invalid input argument: meta.property.is_multigraph is false but the input edge list has "
611+ " parallel edges." );
612+ }
613+
444614 rmm::device_uvector<vertex_t > majors (number_of_local_edges, handle.get_stream ());
445615 rmm::device_uvector<vertex_t > minors (number_of_local_edges, handle.get_stream ());
446616 size_t cur_size{0 };
@@ -726,16 +896,6 @@ graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_
726896 local_sorted_unique_edge_col_offsets_ = std::move (h_key_offsets);
727897 }
728898 }
729-
730- // optional expensive checks (part 2/2)
731-
732- if (do_expensive_check) {
733- // FIXME: check for symmetricity may better be implemetned with transpose().
734- if (this ->is_symmetric ()) {}
735- // FIXME: check for duplicate edges may better be implemented after deciding whether to sort
736- // neighbor list or not.
737- if (!this ->is_multigraph ()) {}
738- }
739899}
740900
741901template <typename vertex_t ,
@@ -774,7 +934,7 @@ graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_
774934 ((*segment_offsets_).size () == (detail::num_sparse_segments_per_vertex_partition + 1 )),
775935 " Invalid input argument: (*(meta.segment_offsets)).size() returns an invalid value." );
776936
777- // optional expensive checks (part 1/2)
937+ // optional expensive checks
778938
779939 if (do_expensive_check) {
780940 auto edge_first = thrust::make_zip_iterator (
@@ -789,11 +949,20 @@ graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_
789949 0 , this ->get_number_of_vertices (), 0 , this ->get_number_of_vertices ()}) == 0 ,
790950 " Invalid input argument: edgelist have out-of-range values." );
791951
792- // FIXME: check for symmetricity may better be implemetned with transpose().
793- if (this ->is_symmetric ()) {}
794- // FIXME: check for duplicate edges may better be implemented after deciding whether to sort
795- // neighbor list or not.
796- if (!this ->is_multigraph ()) {}
952+ if (this ->is_symmetric ()) {
953+ CUGRAPH_EXPECTS (
954+ (check_symmetric<vertex_t , edge_t , weight_t , store_transposed, multi_gpu>(
955+ handle, std::vector<edgelist_t <vertex_t , edge_t , weight_t >>{edgelist})),
956+ " Invalid input argument: meta.property.is_symmetric is true but the input edge list is not "
957+ " symmetric." );
958+ }
959+ if (!this ->is_multigraph ()) {
960+ CUGRAPH_EXPECTS (
961+ check_no_parallel_edge (handle,
962+ std::vector<edgelist_t <vertex_t , edge_t , weight_t >>{edgelist}),
963+ " Invalid input argument: meta.property.is_multigraph is false but the input edge list has "
964+ " parallel edges." );
965+ }
797966 }
798967
799968 // convert edge list (COO) to compressed sparse format (CSR or CSC)
@@ -815,16 +984,6 @@ graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_
815984 weights_ ? std::optional<weight_t *>{(*weights_).data ()} : std::nullopt ,
816985 static_cast <vertex_t >(offsets_.size () - 1 ),
817986 static_cast <edge_t >(indices_.size ()));
818-
819- // optional expensive checks (part 3/3)
820-
821- if (do_expensive_check) {
822- // FIXME: check for symmetricity may better be implemetned with transpose().
823- if (this ->is_symmetric ()) {}
824- // FIXME: check for duplicate edges may better be implemented after deciding whether to sort
825- // neighbor list or not.
826- if (!this ->is_multigraph ()) {}
827- }
828987}
829988
830989template <typename vertex_t ,
0 commit comments