15
15
*/
16
16
#pragma once
17
17
18
+ #include " prims/fill_edge_src_dst_property.cuh"
18
19
#include " prims/reduce_v.cuh"
19
20
#include " prims/transform_reduce_if_v_frontier_outgoing_e_by_dst.cuh"
20
- #include " prims/update_edge_src_dst_property.cuh"
21
21
#include " prims/update_v_frontier.cuh"
22
22
#include " prims/vertex_frontier.cuh"
23
23
@@ -91,7 +91,7 @@ void core_number(raft::handle_t const& handle,
91
91
" Invalid input argument: graph_view has self-loops." );
92
92
}
93
93
94
- // initialize core_numbers to degrees
94
+ // initialize core_numbers to degrees (upper-bound)
95
95
96
96
if (graph_view.is_symmetric ()) { // in-degree == out-degree
97
97
auto out_degrees = graph_view.compute_out_degrees (handle);
@@ -163,9 +163,32 @@ void core_number(raft::handle_t const& handle,
163
163
164
164
vertex_frontier_t <vertex_t , void , multi_gpu, true > vertex_frontier (handle, num_buckets);
165
165
166
- edge_dst_property_t <graph_view_t <vertex_t , edge_t , false , multi_gpu>, edge_t > dst_core_numbers (
166
+ edge_dst_property_t <graph_view_t <vertex_t , edge_t , false , multi_gpu>, bool > edge_dst_valids (
167
167
handle, graph_view);
168
- update_edge_dst_property (handle, graph_view, core_numbers, dst_core_numbers.mutable_view ());
168
+ fill_edge_dst_property (handle, graph_view, edge_dst_valids.mutable_view (), true );
169
+ if (!graph_view.is_symmetric () &&
170
+ degree_type != k_core_degree_type_t ::INOUT) { // 0 core number vertex may have non-zero
171
+ // in|out-degrees (so still can be accessed)
172
+ rmm::device_uvector<vertex_t > zero_degree_vertices (
173
+ graph_view.local_vertex_partition_range_size () - remaining_vertices.size (),
174
+ handle.get_stream ());
175
+ thrust::copy_if (
176
+ handle.get_thrust_policy (),
177
+ thrust::make_counting_iterator (graph_view.local_vertex_partition_range_first ()),
178
+ thrust::make_counting_iterator (graph_view.local_vertex_partition_range_last ()),
179
+ zero_degree_vertices.begin (),
180
+ [core_numbers, v_first = graph_view.local_vertex_partition_range_first ()] __device__ (auto v) {
181
+ return core_numbers[v - v_first] == edge_t {0 };
182
+ });
183
+ fill_edge_dst_property (handle,
184
+ graph_view,
185
+ zero_degree_vertices.begin (),
186
+ zero_degree_vertices.end (),
187
+ edge_dst_valids.mutable_view (),
188
+ false );
189
+ } else { // skip invaliding 0 core number vertices as they won't be accessed anyways
190
+ /* nothing to do */
191
+ }
169
192
170
193
auto k = std::max (k_first, size_t {2 }); // degree 0|1 vertices belong to 0|1-core
171
194
if (graph_view.is_symmetric () && (degree_type == k_core_degree_type_t ::INOUT) &&
@@ -202,27 +225,35 @@ void core_number(raft::handle_t const& handle,
202
225
: edge_t {1 };
203
226
if (vertex_frontier.bucket (bucket_idx_cur).aggregate_size () > 0 ) {
204
227
do {
228
+ fill_edge_dst_property (handle,
229
+ graph_view,
230
+ vertex_frontier.bucket (bucket_idx_cur).begin (),
231
+ vertex_frontier.bucket (bucket_idx_cur).end (),
232
+ edge_dst_valids.mutable_view (),
233
+ false );
234
+
205
235
// FIXME: If most vertices have core numbers less than k, (dst_val >= k) will be mostly
206
236
// false leading to too many unnecessary edge traversals (this is especially problematic if
207
237
// the number of distinct core numbers in [k_first, std::min(max_degree, k_last)] is large).
208
238
// There are two potential solutions: 1) extract a sub-graph and work on the sub-graph & 2)
209
239
// mask-out/delete edges.
210
- if (graph_view.is_symmetric () || ((degree_type == k_core_degree_type_t ::IN) ||
211
- (degree_type == k_core_degree_type_t ::INOUT))) {
240
+ if (graph_view.is_symmetric () ||
241
+ ((degree_type == k_core_degree_type_t ::IN) ||
242
+ (degree_type == k_core_degree_type_t ::INOUT))) { // decrement in-degrees
212
243
auto [new_frontier_vertex_buffer, delta_buffer] =
213
244
cugraph::transform_reduce_if_v_frontier_outgoing_e_by_dst (
214
245
handle,
215
246
graph_view,
216
247
vertex_frontier.bucket (bucket_idx_cur),
217
248
edge_src_dummy_property_t {}.view (),
218
- dst_core_numbers .view (),
249
+ edge_dst_valids .view (),
219
250
edge_dummy_property_t {}.view (),
220
251
cuda::proclaim_return_type<edge_t >(
221
252
[k, delta] __device__ (auto , auto , auto , auto , auto ) { return delta; }),
222
253
reduce_op::plus<edge_t >(),
223
254
cuda::proclaim_return_type<bool >(
224
- [k, delta] __device__ (auto , auto , auto , edge_t dst_val , auto ) {
225
- return dst_val >= k ? true : false ;
255
+ [k, delta] __device__ (auto , auto , auto , bool dst_valid , auto ) {
256
+ return dst_valid ;
226
257
}));
227
258
228
259
update_v_frontier (
@@ -249,20 +280,14 @@ void core_number(raft::handle_t const& handle,
249
280
});
250
281
}
251
282
252
- if (!graph_view.is_symmetric () && ((degree_type == k_core_degree_type_t ::OUT) ||
253
- (degree_type == k_core_degree_type_t ::INOUT))) {
283
+ if (!graph_view.is_symmetric () &&
284
+ ((degree_type == k_core_degree_type_t ::OUT) ||
285
+ (degree_type == k_core_degree_type_t ::INOUT))) { // decrement out-degrees
254
286
// FIXME: we can create a transposed copy of the input graph (note that currently,
255
287
// transpose works only on graph_t (and does not work on graph_view_t)).
256
288
CUGRAPH_FAIL (" unimplemented." );
257
289
}
258
290
259
- update_edge_dst_property (handle,
260
- graph_view,
261
- vertex_frontier.bucket (bucket_idx_next).begin (),
262
- vertex_frontier.bucket (bucket_idx_next).end (),
263
- core_numbers,
264
- dst_core_numbers.mutable_view ());
265
-
266
291
vertex_frontier.bucket (bucket_idx_next)
267
292
.resize (static_cast <size_t >(thrust::distance (
268
293
vertex_frontier.bucket (bucket_idx_next).begin (),
0 commit comments