Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Replace graph_view.hpp::number_of_edges with compute_number_of_edges #4026

Merged
Merged
Changes from 16 commits
Commits
Show all changes
28 commits
Select commit Hold shift + click to select a range
345aee7
update transform_e to work with edge masking
seunghwak Nov 15, 2023
0887af9
remove unnecessary checks
seunghwak Nov 16, 2023
69463f4
update fill_edge_property to work with graphs with edge masks
seunghwak Nov 16, 2023
bf733aa
udpate count_if_e & transform_reduce_e to work with edge masking
seunghwak Nov 16, 2023
8d070d8
add missing includes
seunghwak Nov 16, 2023
2d74bfe
Merge branch 'branch-23.12' of https://github.com/rapidsai/cugraph in…
seunghwak Nov 16, 2023
0e6d7c1
update transform_e test with edge mask
seunghwak Nov 16, 2023
6807b0d
add edge_property utility function for prims testing
seunghwak Nov 17, 2023
138b1f4
update count_if_e & transform_reduce_e tests to include edge masking
seunghwak Nov 17, 2023
06c15a2
clang-format
seunghwak Nov 17, 2023
57a5593
Merge branch 'branch-23.12' into fea_transform_e_mask
naimnv Nov 18, 2023
5fcf9e6
Merge branch 'branch-23.12' of https://github.com/rapidsai/cugraph in…
seunghwak Nov 20, 2023
ff2c849
Merge branch 'fea_transform_e_mask' of github.com:seunghwak/cugraph i…
seunghwak Nov 20, 2023
b1f7bc3
Merge branch 'branch-23.12' of https://github.com/rapidsai/cugraph in…
seunghwak Nov 20, 2023
0c10b84
throw an exception when number_of_edges() is called with edge mask
seunghwak Nov 20, 2023
57ea6f6
Merge branch 'upstream_pr4001' into fea_number_of_edges_with_mask
seunghwak Nov 21, 2023
e60d7ff
remove handle and replace raw pointers with raft::device_span in grap…
seunghwak Nov 22, 2023
3a24f30
remove a redundant internal variable in graph_t
seunghwak Nov 22, 2023
2d04112
deprecate number_of_edges() in graph_view.hpp and replace this with c…
seunghwak Nov 29, 2023
26c02ae
Merge branch 'branch-24.02' of https://github.com/rapidsai/cugraph in…
seunghwak Nov 29, 2023
1e3de80
Merge branch 'branch-24.02' of https://github.com/rapidsai/cugraph in…
seunghwak Nov 29, 2023
3e1ce0f
bug fix in tests
seunghwak Nov 29, 2023
f6d39e4
Merge branch 'upstream_pr4001' into fea_number_of_edges_with_mask
seunghwak Nov 29, 2023
93c1d86
resolve merge conflicts
seunghwak Nov 30, 2023
260e86a
bug fix
seunghwak Dec 1, 2023
29a3a5b
Merge branch 'branch-24.02' of https://github.com/rapidsai/cugraph in…
seunghwak Dec 6, 2023
3308c8f
move edge_mask_view_ back to graph_view_t (as graph_base_t is used fo…
seunghwak Dec 6, 2023
f361b81
Merge branch 'branch-24.02' of https://github.com/rapidsai/cugraph in…
seunghwak Dec 9, 2023
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
54 changes: 21 additions & 33 deletions cpp/include/cugraph/graph_view.hpp
Original file line number Diff line number Diff line change
@@ -268,7 +268,11 @@ class graph_base_t {
properties_(properties){};

vertex_t number_of_vertices() const { return number_of_vertices_; }
edge_t number_of_edges() const { return number_of_edges_; }
edge_t number_of_edges() const
{
CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented.");
return number_of_edges_;
}

template <typename vertex_type = vertex_t>
std::enable_if_t<std::is_signed<vertex_type>::value, bool> is_valid_vertex(vertex_type v) const
@@ -285,6 +289,20 @@ class graph_base_t {
bool is_symmetric() const { return properties_.is_symmetric; }
bool is_multigraph() const { return properties_.is_multigraph; }

void attach_edge_mask(edge_property_view_t<edge_t, uint32_t const*, bool> edge_mask_view)
{
edge_mask_view_ = edge_mask_view;
}

void clear_edge_mask() { edge_mask_view_ = std::nullopt; }

bool has_edge_mask() const { return edge_mask_view_.has_value(); }

std::optional<edge_property_view_t<edge_t, uint32_t const*, bool>> edge_mask_view() const
{
return edge_mask_view_;
}

protected:
raft::handle_t const* handle_ptr() const { return handle_ptr_; };
graph_properties_t graph_properties() const { return properties_; }
@@ -296,6 +314,8 @@ class graph_base_t {
edge_t number_of_edges_{0};

graph_properties_t properties_{};

std::optional<edge_property_view_t<edge_t, uint32_t const*, bool>> edge_mask_view_{std::nullopt};
};

} // namespace detail
@@ -731,20 +751,6 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
return local_sorted_unique_edge_dst_vertex_partition_offsets_;
}

void attach_edge_mask(edge_property_view_t<edge_t, uint32_t const*, bool> edge_mask_view)
{
edge_mask_view_ = edge_mask_view;
}

void clear_edge_mask() { edge_mask_view_ = std::nullopt; }

bool has_edge_mask() const { return edge_mask_view_.has_value(); }

std::optional<edge_property_view_t<edge_t, uint32_t const*, bool>> edge_mask_view() const
{
return edge_mask_view_;
}

private:
std::vector<edge_t const*> edge_partition_offsets_{};
std::vector<vertex_t const*> edge_partition_indices_{};
@@ -790,8 +796,6 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
std::optional<raft::host_span<vertex_t const>>,
std::optional<std::byte> /* dummy */>
local_sorted_unique_edge_dst_vertex_partition_offsets_{std::nullopt};

std::optional<edge_property_view_t<edge_t, uint32_t const*, bool>> edge_mask_view_{std::nullopt};
};

// single-GPU version
@@ -1012,28 +1016,12 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
return std::nullopt;
}

void attach_edge_mask(edge_property_view_t<edge_t, uint32_t const*, bool> edge_mask_view)
{
edge_mask_view_ = edge_mask_view;
}

void clear_edge_mask() { edge_mask_view_ = std::nullopt; }

bool has_edge_mask() const { return edge_mask_view_.has_value(); }

std::optional<edge_property_view_t<edge_t, uint32_t const*, bool>> edge_mask_view() const
{
return edge_mask_view_;
}

private:
edge_t const* offsets_{nullptr};
vertex_t const* indices_{nullptr};

// segment offsets based on vertex degree, relevant only if vertex IDs are renumbered
std::optional<std::vector<vertex_t>> segment_offsets_{std::nullopt};

std::optional<edge_property_view_t<edge_t, uint32_t const*, bool>> edge_mask_view_{std::nullopt};
};

} // namespace cugraph
1 change: 1 addition & 0 deletions cpp/include/cugraph/utilities/misc_utils.cuh
Original file line number Diff line number Diff line change
@@ -19,6 +19,7 @@
#include <raft/util/cudart_utils.hpp>
#include <rmm/device_uvector.hpp>

#include <cuda/atomic>
#include <thrust/binary_search.h>
#include <thrust/gather.h>
#include <thrust/iterator/counting_iterator.h>
2 changes: 0 additions & 2 deletions cpp/src/prims/count_if_e.cuh
Original file line number Diff line number Diff line change
@@ -74,8 +74,6 @@ typename GraphViewType::edge_type count_if_e(raft::handle_t const& handle,
using vertex_t = typename GraphViewType::vertex_type;
using edge_t = typename GraphViewType::edge_type;

CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented.");

if (do_expensive_check) {
// currently, nothing to do
}
73 changes: 65 additions & 8 deletions cpp/src/prims/fill_edge_property.cuh
Original file line number Diff line number Diff line change
@@ -15,6 +15,7 @@
*/
#pragma once

#include <cugraph/edge_partition_edge_property_device_view.cuh>
#include <cugraph/edge_property.hpp>
#include <cugraph/graph_view.hpp>
#include <cugraph/utilities/error.hpp>
@@ -23,6 +24,7 @@
#include <rmm/exec_policy.hpp>

#include <thrust/fill.h>
#include <thrust/iterator/constant_iterator.h>

#include <cstddef>

@@ -38,21 +40,78 @@ void fill_edge_property(raft::handle_t const& handle,
{
static_assert(std::is_same_v<T, typename EdgePropertyOutputWrapper::value_type>);

using edge_t = typename GraphViewType::edge_type;

auto edge_mask_view = graph_view.edge_mask_view();

auto value_firsts = edge_property_output.value_firsts();
auto edge_counts = edge_property_output.edge_counts();
for (size_t i = 0; i < graph_view.number_of_local_edge_partitions(); ++i) {
auto edge_partition_e_mask =
edge_mask_view
? thrust::make_optional<
detail::edge_partition_edge_property_device_view_t<edge_t, uint32_t const*, bool>>(
*edge_mask_view, i)
: thrust::nullopt;

if constexpr (cugraph::has_packed_bool_element<
std::remove_reference_t<decltype(value_firsts[i])>,
T>()) {
static_assert(std::is_arithmetic_v<T>, "unimplemented for thrust::tuple types.");
auto packed_input = input ? packed_bool_full_mask() : packed_bool_empty_mask();
thrust::fill_n(handle.get_thrust_policy(),
value_firsts[i],
packed_bool_size(static_cast<size_t>(edge_counts[i])),
packed_input);
auto rem = edge_counts[i] % packed_bools_per_word();
if (edge_partition_e_mask) {
auto input_first =
thrust::make_zip_iterator(value_firsts[i], (*edge_partition_e_mask).value_first());
thrust::transform(handle.get_thrust_policy(),
input_first,
input_first + packed_bool_size(static_cast<size_t>(edge_counts[i] - rem)),
value_firsts[i],
[packed_input] __device__(thrust::tuple<T, uint32_t> pair) {
auto old_value = thrust::get<0>(pair);
auto mask = thrust::get<1>(pair);
return (old_value & ~mask) | (packed_input & mask);
});
if (rem > 0) {
thrust::transform(
handle.get_thrust_policy(),
input_first + packed_bool_size(static_cast<size_t>(edge_counts[i] - rem)),
input_first + packed_bool_size(static_cast<size_t>(edge_counts[i])),
value_firsts[i] + packed_bool_size(static_cast<size_t>(edge_counts[i] - rem)),
[packed_input, rem] __device__(thrust::tuple<T, uint32_t> pair) {
auto old_value = thrust::get<0>(pair);
auto mask = thrust::get<1>(pair);
return ((old_value & ~mask) | (packed_input & mask)) & packed_bool_partial_mask(rem);
});
}
} else {
thrust::fill_n(handle.get_thrust_policy(),
value_firsts[i],
packed_bool_size(static_cast<size_t>(edge_counts[i] - rem)),
packed_input);
if (rem > 0) {
thrust::fill_n(
handle.get_thrust_policy(),
value_firsts[i] + packed_bool_size(static_cast<size_t>(edge_counts[i] - rem)),
1,
packed_input & packed_bool_partial_mask(rem));
}
}
} else {
thrust::fill_n(
handle.get_thrust_policy(), value_firsts[i], static_cast<size_t>(edge_counts[i]), input);
if (edge_partition_e_mask) {
thrust::transform_if(handle.get_thrust_policy(),
thrust::make_constant_iterator(input),
thrust::make_constant_iterator(input) + edge_counts[i],
thrust::make_counting_iterator(edge_t{0}),
value_firsts[i],
thrust::identity<T>{},
[edge_partition_e_mask = *edge_partition_e_mask] __device__(edge_t i) {
return edge_partition_e_mask.get(i);
});
} else {
thrust::fill_n(
handle.get_thrust_policy(), value_firsts[i], static_cast<size_t>(edge_counts[i]), input);
}
}
}
}
@@ -79,8 +138,6 @@ void fill_edge_property(raft::handle_t const& handle,
edge_property_t<GraphViewType, T>& edge_property_output,
bool do_expensive_check = false)
{
CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented.");

if (do_expensive_check) {
// currently, nothing to do
}
Loading