Skip to content

Commit

Permalink
Skip reduction for zero (in|out-)degree vertices. (rapidsai#2365)
Browse files Browse the repository at this point in the history 
* Added a zero major degree segment when segmenting a vertex partition based on major degrees.
* Added use_dcs() utility function to graph_view_t.
* Added major_hypersparse_first() utility function to edge_partition_device_view_t.
* Skip reduction for zero degree partitions in per_v_transform_reduce_incoming_e (if major == dst) or per_v_transform_reduce_outgoing_e (if major == src).

Authors:
  - Seunghwa Kang (https://github.com/seunghwak)

Approvers:
  - Chuck Hastings (https://github.com/ChuckHastings)
  - Kumar Aatish (https://github.com/kaatish)
  - Joseph Nke (https://github.com/jnke2016)

URL: rapidsai#2365
  • Loading branch information
@seunghwak
seunghwak authored Jul 14, 2022
1 parent 2aad5f2 commit 8ddc7d4
Show file tree
Hide file tree
Showing 21 changed files with 476 additions and 278 deletions.
29 changes: 22 additions & 7 deletions cpp/include/cugraph/edge_partition_device_view.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -112,6 +112,9 @@ class edge_partition_device_view_t<vertex_t,
dcs_nzd_vertex_count_(view.dcs_nzd_vertex_count()
? thrust::optional<vertex_t>{*(view.dcs_nzd_vertex_count())}
: thrust::nullopt),
major_hypersparse_first_(view.major_hypersparse_first()
? thrust::optional<vertex_t>{*(view.major_hypersparse_first())}
: thrust::nullopt),
major_range_first_(view.major_range_first()),
major_range_last_(view.major_range_last()),
minor_range_first_(view.minor_range_first()),
Expand All @@ -120,6 +123,16 @@ class edge_partition_device_view_t<vertex_t,
{
}

__host__ __device__ vertex_t major_value_start_offset() const
{
return major_value_start_offset_;
}

__host__ __device__ thrust::optional<vertex_t> major_hypersparse_first() const noexcept
{
return major_hypersparse_first_;
}

__host__ __device__ vertex_t major_range_first() const noexcept { return major_range_first_; }

__host__ __device__ vertex_t major_range_last() const noexcept { return major_range_last_; }
Expand Down Expand Up @@ -186,11 +199,6 @@ class edge_partition_device_view_t<vertex_t,
return minor_range_first_ + minor_offset;
}

__host__ __device__ vertex_t major_value_start_offset() const
{
return major_value_start_offset_;
}

__host__ __device__ thrust::optional<vertex_t const*> dcs_nzd_vertices() const
{
return dcs_nzd_vertices_;
Expand All @@ -203,8 +211,9 @@ class edge_partition_device_view_t<vertex_t,
private:
// should be trivially copyable to device

thrust::optional<vertex_t const*> dcs_nzd_vertices_{nullptr};
thrust::optional<vertex_t> dcs_nzd_vertex_count_{0};
thrust::optional<vertex_t const*> dcs_nzd_vertices_{thrust::nullopt};
thrust::optional<vertex_t> dcs_nzd_vertex_count_{thrust::nullopt};
thrust::optional<vertex_t> major_hypersparse_first_{thrust::nullopt};

vertex_t major_range_first_{0};
vertex_t major_range_last_{0};
Expand Down Expand Up @@ -232,6 +241,12 @@ class edge_partition_device_view_t<vertex_t,

__host__ __device__ vertex_t major_value_start_offset() const { return vertex_t{0}; }

__host__ __device__ thrust::optional<vertex_t> major_hypersparse_first() const noexcept
{
assert(false);
return thrust::nullopt;
}

__host__ __device__ constexpr vertex_t major_range_first() const noexcept { return vertex_t{0}; }

__host__ __device__ vertex_t major_range_last() const noexcept { return number_of_vertices_; }
Expand Down
9 changes: 7 additions & 2 deletions cpp/include/cugraph/edge_partition_view.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -65,6 +65,7 @@ class edge_partition_view_t<vertex_t, edge_t, weight_t, multi_gpu, std::enable_i
std::optional<weight_t const*> weights,
std::optional<vertex_t const*> dcs_nzd_vertices,
std::optional<vertex_t> dcs_nzd_vertex_count,
std::optional<vertex_t> major_hypersparse_first,
edge_t number_of_edge_partition_edges,
vertex_t major_range_first,
vertex_t major_range_last,
Expand All @@ -75,6 +76,7 @@ class edge_partition_view_t<vertex_t, edge_t, weight_t, multi_gpu, std::enable_i
offsets, indices, weights, number_of_edge_partition_edges),
dcs_nzd_vertices_(dcs_nzd_vertices),
dcs_nzd_vertex_count_(dcs_nzd_vertex_count),
major_hypersparse_first_(major_hypersparse_first),
major_range_first_(major_range_first),
major_range_last_(major_range_last),
minor_range_first_(minor_range_first),
Expand All @@ -85,6 +87,7 @@ class edge_partition_view_t<vertex_t, edge_t, weight_t, multi_gpu, std::enable_i

std::optional<vertex_t const*> dcs_nzd_vertices() const { return dcs_nzd_vertices_; }
std::optional<vertex_t> dcs_nzd_vertex_count() const { return dcs_nzd_vertex_count_; }
std::optional<vertex_t> major_hypersparse_first() const { return major_hypersparse_first_; }

vertex_t major_range_first() const { return major_range_first_; }
vertex_t major_range_last() const { return major_range_last_; }
Expand All @@ -95,8 +98,9 @@ class edge_partition_view_t<vertex_t, edge_t, weight_t, multi_gpu, std::enable_i

private:
// relevant only if we use the CSR + DCSR (or CSC + DCSC) hybrid format
std::optional<vertex_t const*> dcs_nzd_vertices_{};
std::optional<vertex_t> dcs_nzd_vertex_count_{};
std::optional<vertex_t const*> dcs_nzd_vertices_{std::nullopt};
std::optional<vertex_t> dcs_nzd_vertex_count_{std::nullopt};
std::optional<vertex_t> major_hypersparse_first_{std::nullopt};

vertex_t major_range_first_{0};
vertex_t major_range_last_{0};
Expand Down Expand Up @@ -124,6 +128,7 @@ class edge_partition_view_t<vertex_t, edge_t, weight_t, multi_gpu, std::enable_i

std::optional<vertex_t const*> dcs_nzd_vertices() const { return std::nullopt; }
std::optional<vertex_t> dcs_nzd_vertex_count() const { return std::nullopt; }
std::optional<vertex_t> major_hypersparse_first() const { return std::nullopt; }

vertex_t major_range_first() const { return vertex_t{0}; }
vertex_t major_range_last() const { return number_of_vertices_; }
Expand Down
2 changes: 1 addition & 1 deletion cpp/include/cugraph/graph.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -248,7 +248,7 @@ class graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enab
local_sorted_unique_edge_dst_chunk_start_offsets =
std::vector<raft::device_span<vertex_t const>>(
(*local_sorted_unique_edge_dst_chunk_start_offsets_).size());
for (size_t i = 0; (*local_sorted_unique_edge_dsts).size(); ++i) {
for (size_t i = 0; i < (*local_sorted_unique_edge_dsts).size(); ++i) {
(*local_sorted_unique_edge_dsts)[i] =
raft::device_span<vertex_t const>((*local_sorted_unique_edge_dsts_)[i].begin(),
(*local_sorted_unique_edge_dsts_)[i].end());
Expand Down
20 changes: 20 additions & 0 deletions cpp/include/cugraph/graph_view.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -582,6 +582,17 @@ class graph_view_t<vertex_t,
: vertex_t{0};
}

bool use_dcs() const
{
if (edge_partition_segment_offsets_) {
auto size_per_partition =
(*edge_partition_segment_offsets_).size() / edge_partition_offsets_.size();
return size_per_partition > (detail::num_sparse_segments_per_vertex_partition + size_t{2});
} else {
return false;
}
}

std::optional<std::vector<vertex_t>> local_edge_partition_segment_offsets(
size_t partition_idx) const
{
Expand Down Expand Up @@ -626,6 +637,12 @@ class graph_view_t<vertex_t,
major_value_range_start_offset =
this->local_edge_partition_src_value_start_offset(partition_idx);
}
std::optional<vertex_t> major_hypersparse_first{std::nullopt};
if (this->use_dcs()) {
major_hypersparse_first =
major_range_first + (*(this->local_edge_partition_segment_offsets(
partition_idx)))[detail::num_sparse_segments_per_vertex_partition];
}
return edge_partition_view_t<vertex_t, edge_t, weight_t, true>(
edge_partition_offsets_[partition_idx],
edge_partition_indices_[partition_idx],
Expand All @@ -638,6 +655,7 @@ class graph_view_t<vertex_t,
edge_partition_dcs_nzd_vertex_counts_
? std::optional<vertex_t>{(*edge_partition_dcs_nzd_vertex_counts_)[partition_idx]}
: std::nullopt,
major_hypersparse_first,
edge_partition_number_of_edges_[partition_idx],
major_range_first,
major_range_last,
Expand Down Expand Up @@ -926,6 +944,8 @@ class graph_view_t<vertex_t,
return vertex_t{0};
}

bool use_dcs() const { return false; }

std::optional<std::vector<vertex_t>> local_edge_partition_segment_offsets(
size_t partition_idx = 0) const
{
Expand Down
14 changes: 14 additions & 0 deletions cpp/include/cugraph/utilities/device_functors.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -67,6 +67,20 @@ struct check_bit_set_t {
}
};

template <typename T>
struct shift_left_t {
T offset{};

__device__ T operator()(T input) const { return input - offset; }
};

template <typename T>
struct shift_right_t {
T offset{};

__device__ T operator()(T input) const { return input + offset; }
};

template <typename T>
struct multiplier_t {
T multiplier{};
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/detail/utility_wrappers.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -115,7 +115,7 @@ std::tuple<rmm::device_uvector<vertex_t>, rmm::device_uvector<edge_t>> filter_de
auto zip_iter =
thrust::make_zip_iterator(thrust::make_tuple(d_vertices.begin(), d_out_degs.begin()));

CUGRAPH_EXPECTS(d_vertices.size() < std::numeric_limits<int32_t>::max(),
CUGRAPH_EXPECTS(d_vertices.size() < static_cast<size_t>(std::numeric_limits<int32_t>::max()),
"remove_if will fail, d_vertices.size() is too large");

// FIXME: remove_if has a 32-bit overflow issue (https://github.com/NVIDIA/thrust/issues/1302)
Expand Down
8 changes: 4 additions & 4 deletions cpp/src/link_analysis/pagerank_impl.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -105,14 +105,14 @@ void pagerank(
}

if (pull_graph_view.is_weighted()) {
auto num_nonpositive_edge_weights =
auto num_negative_edge_weights =
count_if_e(handle,
pull_graph_view,
dummy_property_t<vertex_t>{}.device_view(),
dummy_property_t<vertex_t>{}.device_view(),
[] __device__(vertex_t, vertex_t, weight_t w, auto, auto) { return w <= 0.0; });
CUGRAPH_EXPECTS(num_nonpositive_edge_weights == 0,
"Invalid input argument: input graph should have postive edge weights.");
[] __device__(vertex_t, vertex_t, weight_t w, auto, auto) { return w < 0.0; });
CUGRAPH_EXPECTS(num_negative_edge_weights == 0,
"Invalid input argument: input graph should have non-negative edge weights.");
}

if (has_initial_guess) {
Expand Down
Loading

0 comments on commit 8ddc7d4

Please sign in to comment.