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.

Sign up for GitHub

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

Update count_if_e, transform_reduce_e, and transform_e to support edge masking #4001

Merged
merged 17 commits into from
Nov 30, 2023
Merged

Update count_if_e, transform_reduce_e, and transform_e to support edge masking #4001

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
17 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
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
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
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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; }
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I have no preference. Is there a reason to prefer this syntax over edge_mask_view_.reset()?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I assume they are same. I guess mainly personal preference.

According to
https://stackoverflow.com/questions/47441623/how-to-assign-nothing-to-stdoptionalt

It seems like = {} or = std::nullopt were initially recommended and .reset() is added later.

Copy link
Contributor Author

@seunghwak seunghwak Nov 30, 2023 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

But consistency matters, so if we are mainly using .reset() elsewhere in our codebase, I am OK to make a change.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That was my assessment. A quick check in the code, I see a few calls to .reset() and a bunch more assignments to std::nullopt.

I'm going to merge this PR. We should probably pick a way and create a PR to update things to be consistent.


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_; }
Expand All @@ -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
Expand Down Expand Up @@ -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_{};
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>
Expand Down
2 changes: 0 additions & 2 deletions cpp/src/prims/count_if_e.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
}
Expand Down
73 changes: 65 additions & 8 deletions cpp/src/prims/fill_edge_property.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>
Expand All @@ -23,6 +24,7 @@
#include <rmm/exec_policy.hpp>

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

#include <cstddef>

Expand All @@ -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);
}
}
}
}
Expand All @@ -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
}
Expand Down
Loading