Skip to content

Commit fc57ba6

Browse files
authored
Add make_initialized_edge_(src|dst_)property utility functions. (#5441)
``` fill_edge_property(handle, graph_view, input) ``` fills the un-maksed edge values to `input`. If we want to fill the entire edge values (both masked and un-masked) with `graph_view` with an attached edge mask, we need to copy a `graph_view` object; clear the edge mask; and call the `fill_edge_property` function with the copy. This is cumbersome and possibly bug inducing. `make_initialized_edge_property` function is added as a utility function. Here, this factory creates an `edge_property_t` object and fills the entire edge property buffer to the user provided `input` value. `make_initialized_edge_src|dst_property` functions are added as well for consistency. Authors: - Seunghwa Kang (https://github.com/seunghwak) Approvers: - Chuck Hastings (https://github.com/ChuckHastings) - Joseph Nke (https://github.com/jnke2016) URL: #5441
1 parent 6bc8593 commit fc57ba6

10 files changed

+173
-75
lines changed

cpp/src/centrality/betweenness_centrality_impl.cuh

Lines changed: 3 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010
#include "prims/extract_transform_if_e.cuh"
1111
#include "prims/extract_transform_if_v_frontier_outgoing_e.cuh"
1212
#include "prims/fill_edge_property.cuh"
13+
#include "prims/make_initialized_edge_property.cuh"
1314
#include "prims/per_v_transform_reduce_incoming_outgoing_e.cuh"
1415
#include "prims/transform_e.cuh"
1516
#include "prims/transform_reduce_if_v_frontier_outgoing_e_by_dst.cuh"
@@ -1502,17 +1503,8 @@ edge_property_t<edge_t, weight_t> edge_betweenness_centrality(
15021503
"Invalid input argument: sources have invalid vertex IDs.");
15031504
}
15041505

1505-
edge_property_t<edge_t, weight_t> centralities(handle, graph_view);
1506-
1507-
if (graph_view.has_edge_mask()) {
1508-
auto unmasked_graph_view = graph_view;
1509-
unmasked_graph_view.clear_edge_mask();
1510-
fill_edge_property(
1511-
handle, unmasked_graph_view, centralities.mutable_view(), weight_t{0}, do_expensive_check);
1512-
} else {
1513-
fill_edge_property(
1514-
handle, graph_view, centralities.mutable_view(), weight_t{0}, do_expensive_check);
1515-
}
1506+
auto centralities =
1507+
make_initialized_edge_property(handle, graph_view, weight_t{0.0}, do_expensive_check);
15161508

15171509
size_t num_sources = cuda::std::distance(vertices_begin, vertices_end);
15181510
std::vector<size_t> source_offsets{{0, num_sources}};

cpp/src/community/approx_weighted_matching_impl.cuh

Lines changed: 3 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@
66

77
#include "detail/shuffle_wrappers.hpp"
88
#include "prims/fill_edge_property.cuh"
9+
#include "prims/make_initialized_edge_property.cuh"
910
#include "prims/reduce_op.cuh"
1011
#include "prims/transform_e.cuh"
1112
#include "prims/transform_reduce_e_by_src_dst_key.cuh"
@@ -38,18 +39,10 @@ std::tuple<rmm::device_uvector<vertex_t>, weight_t> approximate_weighted_matchin
3839
"need to be symmetric");
3940

4041
auto current_graph_view = graph_view;
41-
if (current_graph_view.has_edge_mask()) { current_graph_view.clear_edge_mask(); }
4242

43-
cugraph::edge_property_t<edge_t, bool> edge_masks_even(handle, current_graph_view);
44-
cugraph::fill_edge_property(
45-
handle, current_graph_view, edge_masks_even.mutable_view(), bool{false});
46-
cugraph::edge_property_t<edge_t, bool> edge_masks_odd(handle, current_graph_view);
47-
cugraph::fill_edge_property(
48-
handle, current_graph_view, edge_masks_odd.mutable_view(), bool{false});
43+
auto edge_masks_even = make_initialized_edge_property(handle, current_graph_view, false);
44+
auto edge_masks_odd = make_initialized_edge_property(handle, current_graph_view, false);
4945

50-
if (graph_view.has_edge_mask()) {
51-
current_graph_view.attach_edge_mask(*(graph_view.edge_mask_view()));
52-
}
5346
// Mask out self-loop
5447
cugraph::transform_e(
5548
handle,

cpp/src/community/edge_triangle_count_impl.cuh

Lines changed: 2 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,7 @@
88
#include "detail/graph_partition_utils.cuh"
99
#include "prims/edge_bucket.cuh"
1010
#include "prims/fill_edge_property.cuh"
11+
#include "prims/make_initialized_edge_property.cuh"
1112
#include "prims/per_v_pair_dst_nbr_intersection.cuh"
1213
#include "prims/transform_e.cuh"
1314

@@ -338,12 +339,7 @@ edge_property_t<edge_t, edge_t> edge_triangle_count_impl(
338339
prev_chunk_size += chunk_size;
339340
}
340341

341-
cugraph::edge_property_t<edge_t, edge_t> counts(handle, cur_graph_view);
342-
{
343-
auto unmasked_graph_view = cur_graph_view;
344-
if (unmasked_graph_view.has_edge_mask()) { unmasked_graph_view.clear_edge_mask(); }
345-
cugraph::fill_edge_property(handle, unmasked_graph_view, counts.mutable_view(), edge_t{0});
346-
}
342+
auto counts = make_initialized_edge_property(handle, cur_graph_view, edge_t{0});
347343

348344
cugraph::edge_bucket_t<vertex_t, edge_t, true, multi_gpu, true> valid_edges(
349345
handle, false /* multigraph */);

cpp/src/community/k_truss_impl.cuh

Lines changed: 6 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,7 @@
77
#include "prims/edge_bucket.cuh"
88
#include "prims/extract_transform_if_e.cuh"
99
#include "prims/fill_edge_property.cuh"
10+
#include "prims/make_initialized_edge_property.cuh"
1011
#include "prims/per_v_pair_dst_nbr_intersection.cuh"
1112
#include "prims/transform_e.cuh"
1213
#include "prims/transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v.cuh"
@@ -179,10 +180,8 @@ k_truss(raft::handle_t const& handle,
179180

180181
// 2. Exclude self-loops and edges that do not belong to (k-1)-core
181182

182-
auto cur_graph_view = graph_view;
183-
auto unmasked_cur_graph_view = cur_graph_view;
183+
auto cur_graph_view = graph_view;
184184

185-
if (unmasked_cur_graph_view.has_edge_mask()) { unmasked_cur_graph_view.clear_edge_mask(); }
186185
// mask for self-loops and edges not part of k-1 core
187186
cugraph::edge_property_t<edge_t, bool> undirected_mask(handle);
188187
{
@@ -191,9 +190,7 @@ k_truss(raft::handle_t const& handle,
191190
if (cur_graph_view.count_self_loops(handle) > edge_t{0}) {
192191
// 2.1. Exclude self-loops
193192

194-
cugraph::edge_property_t<edge_t, bool> self_loop_edge_mask(handle, cur_graph_view);
195-
cugraph::fill_edge_property(
196-
handle, unmasked_cur_graph_view, self_loop_edge_mask.mutable_view(), false);
193+
auto self_loop_edge_mask = make_initialized_edge_property(handle, cur_graph_view, false);
197194

198195
transform_e(handle,
199196
cur_graph_view,
@@ -238,9 +235,8 @@ k_truss(raft::handle_t const& handle,
238235
in_k_minus_1_core_flags.begin(),
239236
edge_dst_in_k_minus_1_cores.mutable_view());
240237

241-
cugraph::edge_property_t<edge_t, bool> in_k_minus_1_core_edge_mask(handle, cur_graph_view);
242-
cugraph::fill_edge_property(
243-
handle, unmasked_cur_graph_view, in_k_minus_1_core_edge_mask.mutable_view(), false);
238+
auto in_k_minus_1_core_edge_mask =
239+
make_initialized_edge_property(handle, cur_graph_view, false);
244240

245241
transform_e(
246242
handle,
@@ -265,17 +261,14 @@ k_truss(raft::handle_t const& handle,
265261
edge_src_property_t<vertex_t, edge_t> edge_src_out_degrees(handle, cur_graph_view);
266262
edge_dst_property_t<vertex_t, edge_t> edge_dst_out_degrees(handle, cur_graph_view);
267263

268-
cugraph::edge_property_t<edge_t, bool> dodg_mask(handle, cur_graph_view);
264+
auto dodg_mask = make_initialized_edge_property(handle, cur_graph_view, false);
269265
{
270266
auto out_degrees = cur_graph_view.compute_out_degrees(handle);
271267
update_edge_src_property(
272268
handle, cur_graph_view, out_degrees.begin(), edge_src_out_degrees.mutable_view());
273269
update_edge_dst_property(
274270
handle, cur_graph_view, out_degrees.begin(), edge_dst_out_degrees.mutable_view());
275271

276-
cugraph::fill_edge_property(
277-
handle, unmasked_cur_graph_view, dodg_mask.mutable_view(), bool{false});
278-
279272
cugraph::transform_e(
280273
handle,
281274
cur_graph_view,

cpp/src/community/triangle_count_impl.cuh

Lines changed: 4 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,7 @@
77
#include "detail/shuffle_wrappers.hpp"
88
#include "prims/extract_transform_if_e.cuh"
99
#include "prims/fill_edge_property.cuh"
10+
#include "prims/make_initialized_edge_property.cuh"
1011
#include "prims/transform_e.cuh"
1112
#include "prims/transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v.cuh"
1213
#include "prims/update_edge_src_dst_property.cuh"
@@ -182,18 +183,13 @@ void triangle_count(raft::handle_t const& handle,
182183

183184
auto cur_graph_view = graph_view;
184185

185-
auto unmasked_cur_graph_view = cur_graph_view;
186-
if (unmasked_cur_graph_view.has_edge_mask()) { unmasked_cur_graph_view.clear_edge_mask(); }
187-
188186
// 2. Mask out the edges that has source or destination that cannot be reached from vertices
189187
// within two hop (if vertices.has_value() is true).
190188

191189
cugraph::edge_property_t<edge_t, bool> edge_mask(handle);
192190

193191
if (vertices) {
194-
cugraph::edge_property_t<edge_t, bool> within_two_hop_edge_mask(handle, cur_graph_view);
195-
cugraph::fill_edge_property(
196-
handle, unmasked_cur_graph_view, within_two_hop_edge_mask.mutable_view(), false);
192+
auto within_two_hop_edge_mask = make_initialized_edge_property(handle, cur_graph_view, false);
197193

198194
rmm::device_uvector<vertex_t> unique_vertices((*vertices).size(), handle.get_stream());
199195
thrust::copy(
@@ -353,9 +349,7 @@ void triangle_count(raft::handle_t const& handle,
353349
// 3. Exclude self-loops
354350

355351
if (cur_graph_view.count_self_loops(handle) > edge_t{0}) {
356-
cugraph::edge_property_t<edge_t, bool> self_loop_edge_mask(handle, cur_graph_view);
357-
cugraph::fill_edge_property(
358-
handle, unmasked_cur_graph_view, self_loop_edge_mask.mutable_view(), false);
352+
auto self_loop_edge_mask = make_initialized_edge_property(handle, cur_graph_view, false);
359353

360354
transform_e(handle,
361355
cur_graph_view,
@@ -374,9 +368,7 @@ void triangle_count(raft::handle_t const& handle,
374368
// 4. Find 2-core and exclude edges that do not belong to 2-core add masking support).
375369

376370
{
377-
cugraph::edge_property_t<edge_t, bool> in_two_core_edge_mask(handle, cur_graph_view);
378-
cugraph::fill_edge_property(
379-
handle, unmasked_cur_graph_view, in_two_core_edge_mask.mutable_view(), false);
371+
auto in_two_core_edge_mask = make_initialized_edge_property(handle, cur_graph_view, false);
380372

381373
rmm::device_uvector<edge_t> core_numbers(cur_graph_view.number_of_vertices(),
382374
handle.get_stream());

cpp/src/components/vertex_coloring_impl.cuh

Lines changed: 4 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,11 @@
11
/*
2-
* SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION.
2+
* SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
33
* SPDX-License-Identifier: Apache-2.0
44
*/
55
#pragma once
66

77
#include "prims/fill_edge_property.cuh"
8+
#include "prims/make_initialized_edge_property.cuh"
89
#include "prims/transform_e.cuh"
910
#include "prims/update_edge_src_dst_property.cuh"
1011

@@ -29,13 +30,8 @@ rmm::device_uvector<vertex_t> vertex_coloring(
2930
graph_view_t current_graph_view(graph_view);
3031

3132
// edge mask
32-
cugraph::edge_property_t<edge_t, bool> edge_masks_even(handle, current_graph_view);
33-
cugraph::fill_edge_property(
34-
handle, current_graph_view, edge_masks_even.mutable_view(), bool{false});
35-
36-
cugraph::edge_property_t<edge_t, bool> edge_masks_odd(handle, current_graph_view);
37-
cugraph::fill_edge_property(
38-
handle, current_graph_view, edge_masks_odd.mutable_view(), bool{false});
33+
auto edge_masks_even = make_initialized_edge_property(handle, current_graph_view, false);
34+
auto edge_masks_odd = make_initialized_edge_property(handle, current_graph_view, false);
3935

4036
cugraph::transform_e(
4137
handle,

cpp/src/components/weakly_connected_components_impl.cuh

Lines changed: 3 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@
66

77
#include "prims/fill_edge_property.cuh"
88
#include "prims/fill_edge_src_dst_property.cuh"
9+
#include "prims/make_initialized_edge_property.cuh"
910
#include "prims/transform_e.cuh"
1011
#include "prims/transform_reduce_if_v_frontier_outgoing_e_by_dst.cuh"
1112
#include "prims/update_edge_src_dst_property.cuh"
@@ -457,15 +458,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle,
457458
handle, level_graph_view, visited.begin(), edge_src_visited.mutable_view());
458459
{
459460
auto tmp_graph_view = level_graph_view;
460-
edge_property_t<edge_t, bool> edge_mask(handle, level_graph_view);
461-
if (tmp_graph_view.has_edge_mask()) {
462-
tmp_graph_view.clear_edge_mask();
463-
cugraph::fill_edge_property(handle,
464-
tmp_graph_view,
465-
edge_mask.mutable_view(),
466-
false); // if level_graph_view has an attached edge mask,
467-
// masked out values of edge_mask won't be updated.
468-
}
461+
auto edge_mask = make_initialized_edge_property(handle, level_graph_view, false);
469462
transform_e(handle,
470463
level_graph_view,
471464
edge_src_visited.view(),
@@ -477,6 +470,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle,
477470
!src_visited;
478471
}),
479472
edge_mask.mutable_view());
473+
if (tmp_graph_view.has_edge_mask()) { tmp_graph_view.clear_edge_mask(); }
480474
tmp_graph_view.attach_edge_mask(edge_mask.view());
481475
std::tie(std::get<0>(edge_buffer),
482476
std::get<1>(edge_buffer),

cpp/src/cores/core_number_impl.cuh

Lines changed: 4 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@
66

77
#include "prims/fill_edge_property.cuh"
88
#include "prims/fill_edge_src_dst_property.cuh"
9+
#include "prims/make_initialized_edge_property.cuh"
910
#include "prims/reduce_v.cuh"
1011
#include "prims/transform_e.cuh"
1112
#include "prims/transform_reduce_if_v_frontier_outgoing_e_by_dst.cuh"
@@ -81,19 +82,16 @@ void core_number(raft::handle_t const& handle,
8182
std::optional<cugraph::edge_property_t<edge_t, bool>> self_loop_edge_mask{std::nullopt};
8283
auto cur_graph_view = graph_view;
8384
if (cur_graph_view.count_self_loops(handle) > edge_t{0}) {
84-
self_loop_edge_mask = cugraph::edge_property_t<edge_t, bool>(handle, cur_graph_view);
85-
if (cur_graph_view.has_edge_mask()) { cur_graph_view.clear_edge_mask(); }
86-
cugraph::fill_edge_property(handle, cur_graph_view, self_loop_edge_mask->mutable_view(), false);
87-
85+
self_loop_edge_mask = make_initialized_edge_property(handle, cur_graph_view, false);
8886
transform_e(handle,
89-
graph_view,
87+
cur_graph_view,
9088
edge_src_dummy_property_t{}.view(),
9189
edge_dst_dummy_property_t{}.view(),
9290
edge_dummy_property_t{}.view(),
9391
cuda::proclaim_return_type<bool>(
9492
[] __device__(auto src, auto dst, auto, auto, auto) { return src != dst; }),
9593
self_loop_edge_mask->mutable_view());
96-
94+
if (cur_graph_view.has_edge_mask()) { cur_graph_view.clear_edge_mask(); }
9795
cur_graph_view.attach_edge_mask(self_loop_edge_mask->view());
9896
}
9997

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,56 @@
1+
/*
2+
* SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
3+
* SPDX-License-Identifier: Apache-2.0
4+
*/
5+
#pragma once
6+
7+
#include "prims/fill_edge_property.cuh"
8+
9+
#include <cugraph/edge_property.hpp>
10+
#include <cugraph/graph_view.hpp>
11+
12+
#include <raft/core/handle.hpp>
13+
14+
namespace cugraph {
15+
16+
/**
17+
* @brief Create an edge property and fill the entire buffer with the given value.
18+
*
19+
* Create an edge property object for @p graph_view and initialize the edge property values to @p
20+
* input. Note that the entire edge buffer is initialized to the provided @p
21+
* input (ignoring the edge mask if @p graph_view has an attached edge mask); this contrasts
22+
* with the fill_edge_property function which only fills the edge values for un-masked edges if @p
23+
* graph_view has an attached edge mask.
24+
*
25+
* @tparam GraphViewType Type of the passed non-owning graph object.
26+
* @tparam T Type of the edge property values.
27+
* @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
28+
* handles to various CUDA libraries) to run graph algorithms.
29+
* @param graph_view Non-owning graph object.
30+
* @param input Edge property values will be set to @p input.
31+
* @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
32+
* @return edge_property_t<typename GraphViewType::edge_type, T> with all entries set to @p input.
33+
*/
34+
template <typename GraphViewType, typename T>
35+
edge_property_t<typename GraphViewType::edge_type, T> make_initialized_edge_property(
36+
raft::handle_t const& handle,
37+
GraphViewType const& graph_view,
38+
T input,
39+
bool do_expensive_check = false)
40+
{
41+
using edge_t = typename GraphViewType::edge_type;
42+
43+
edge_property_t<edge_t, T> edge_prop(handle, graph_view);
44+
if (graph_view.has_edge_mask()) {
45+
auto graph_view_copy = graph_view;
46+
graph_view_copy.clear_edge_mask();
47+
fill_edge_property(
48+
handle, graph_view_copy, edge_prop.mutable_view(), input, do_expensive_check);
49+
} else {
50+
fill_edge_property(handle, graph_view, edge_prop.mutable_view(), input, do_expensive_check);
51+
}
52+
53+
return edge_prop;
54+
}
55+
56+
} // namespace cugraph

0 commit comments

Comments
 (0)