From 7fe7beab548d66053916fde070d6f6cb58ed2339 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang <45857425+seunghwak@users.noreply.github.com> Date: Wed, 13 Dec 2023 03:41:31 +0900 Subject: [PATCH] Replace graph_view.hpp::number_of_edges with compute_number_of_edges (#4026) Replace graph_view.hpp::number_of_edges (deprecated, throws an exception if an edge mask is attached to the graph view object) with compute_number_of_edges (this function works with or without edge mask) Authors: - Seunghwa Kang (https://github.com/seunghwak) - Naim (https://github.com/naimnv) Approvers: - Joseph Nke (https://github.com/jnke2016) - Chuck Hastings (https://github.com/ChuckHastings) - Naim (https://github.com/naimnv) URL: https://github.com/rapidsai/cugraph/pull/4026 --- cpp/include/cugraph/graph.hpp | 47 ++-- cpp/include/cugraph/graph_view.hpp | 136 +++++----- ...v_transform_reduce_incoming_outgoing_e.cuh | 9 +- .../prims/update_edge_src_dst_property.cuh | 3 +- cpp/src/structure/graph_impl.cuh | 24 +- cpp/src/structure/graph_view_impl.cuh | 237 +++++++++--------- .../traversal/od_shortest_distances_impl.cuh | 4 +- cpp/src/traversal/sssp_impl.cuh | 2 +- cpp/src/utilities/cugraph_ops_utils.hpp | 2 +- cpp/tests/link_analysis/hits_test.cpp | 2 +- ...er_v_random_select_transform_outgoing_e.cu | 4 +- 11 files changed, 237 insertions(+), 233 deletions(-) diff --git a/cpp/include/cugraph/graph.hpp b/cpp/include/cugraph/graph.hpp index 60b9f1a4054..a723fde24df 100644 --- a/cpp/include/cugraph/graph.hpp +++ b/cpp/include/cugraph/graph.hpp @@ -90,24 +90,25 @@ class graph_t meta, bool do_expensive_check = false); + edge_t number_of_edges() const { return this->number_of_edges_; } + graph_view_t view() const { - std::vector offsets(edge_partition_offsets_.size(), nullptr); - std::vector indices(edge_partition_indices_.size(), nullptr); - auto dcs_nzd_vertices = edge_partition_dcs_nzd_vertices_ - ? std::make_optional>( - (*edge_partition_dcs_nzd_vertices_).size(), nullptr) - : std::nullopt; - auto dcs_nzd_vertex_counts = edge_partition_dcs_nzd_vertex_counts_ - ? std::make_optional>( - (*edge_partition_dcs_nzd_vertex_counts_).size(), vertex_t{0}) - : std::nullopt; + std::vector> offsets(edge_partition_offsets_.size()); + std::vector> indices(edge_partition_indices_.size()); + auto dcs_nzd_vertices = edge_partition_dcs_nzd_vertices_ + ? std::make_optional>>( + (*edge_partition_dcs_nzd_vertices_).size()) + : std::nullopt; for (size_t i = 0; i < offsets.size(); ++i) { - offsets[i] = edge_partition_offsets_[i].data(); - indices[i] = edge_partition_indices_[i].data(); + offsets[i] = raft::device_span(edge_partition_offsets_[i].data(), + edge_partition_offsets_[i].size()); + indices[i] = raft::device_span(edge_partition_indices_[i].data(), + edge_partition_indices_[i].size()); if (dcs_nzd_vertices) { - (*dcs_nzd_vertices)[i] = (*edge_partition_dcs_nzd_vertices_)[i].data(); - (*dcs_nzd_vertex_counts)[i] = (*edge_partition_dcs_nzd_vertex_counts_)[i]; + (*dcs_nzd_vertices)[i] = + raft::device_span((*edge_partition_dcs_nzd_vertices_)[i].data(), + (*edge_partition_dcs_nzd_vertices_)[i].size()); } } @@ -196,15 +197,13 @@ class graph_t( - *(this->handle_ptr()), offsets, indices, dcs_nzd_vertices, - dcs_nzd_vertex_counts, graph_view_meta_t{ this->number_of_vertices(), this->number_of_edges(), - this->graph_properties(), + this->properties_, partition_, edge_partition_segment_offsets_, local_sorted_unique_edge_srcs, @@ -224,7 +223,6 @@ class graph_t>> edge_partition_dcs_nzd_vertices_{ std::nullopt}; - std::optional> edge_partition_dcs_nzd_vertex_counts_{std::nullopt}; partition_t partition_{}; // segment offsets within the vertex partition based on vertex degree @@ -283,16 +281,15 @@ class graph_t meta, bool do_expensive_check = false); + edge_t number_of_edges() const { return this->number_of_edges_; } + graph_view_t view() const { return graph_view_t( - *(this->handle_ptr()), - offsets_.data(), - indices_.data(), - graph_view_meta_t{this->number_of_vertices(), - this->number_of_edges(), - this->graph_properties(), - segment_offsets_}); + raft::device_span(offsets_.data(), offsets_.size()), + raft::device_span(indices_.data(), indices_.size()), + graph_view_meta_t{ + this->number_of_vertices(), this->number_of_edges(), this->properties_, segment_offsets_}); } private: diff --git a/cpp/include/cugraph/graph_view.hpp b/cpp/include/cugraph/graph_view.hpp index d79d4635c54..53c66c6483e 100644 --- a/cpp/include/cugraph/graph_view.hpp +++ b/cpp/include/cugraph/graph_view.hpp @@ -258,21 +258,12 @@ class graph_base_t { public: graph_base_t() = default; - graph_base_t(raft::handle_t const& handle, - vertex_t number_of_vertices, - edge_t number_of_edges, - graph_properties_t properties) - : handle_ptr_(&handle), - number_of_vertices_(number_of_vertices), + graph_base_t(vertex_t number_of_vertices, edge_t number_of_edges, graph_properties_t properties) + : number_of_vertices_(number_of_vertices), number_of_edges_(number_of_edges), properties_(properties){}; vertex_t number_of_vertices() const { return number_of_vertices_; } - edge_t number_of_edges() const - { - CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); - return number_of_edges_; - } template std::enable_if_t::value, bool> is_valid_vertex(vertex_type v) const @@ -289,33 +280,12 @@ 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_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_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_; } - - private: - raft::handle_t const* handle_ptr_{nullptr}; - - vertex_t number_of_vertices_{0}; edge_t number_of_edges_{0}; - graph_properties_t properties_{}; - std::optional> edge_mask_view_{std::nullopt}; + private: + vertex_t number_of_vertices_{0}; }; } // namespace detail @@ -405,11 +375,10 @@ class graph_view_t const& edge_partition_offsets, - std::vector const& edge_partition_indices, - std::optional> const& edge_partition_dcs_nzd_vertices, - std::optional> const& edge_partition_dcs_nzd_vertex_counts, + graph_view_t(std::vector> const& edge_partition_offsets, + std::vector> const& edge_partition_indices, + std::optional>> const& + edge_partition_dcs_nzd_vertices, graph_view_meta_t meta); std::vector vertex_partition_range_offsets() const @@ -624,25 +593,16 @@ class graph_view_tlocal_edge_partition_src_value_start_offset(partition_idx); } std::optional major_hypersparse_first{std::nullopt}; - vertex_t offset_size = (major_range_last - major_range_first) + 1; 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]; - offset_size = ((*major_hypersparse_first) - major_range_first) + - (*edge_partition_dcs_nzd_vertex_counts_)[partition_idx] + 1; } return edge_partition_view_t( - raft::device_span(edge_partition_offsets_[partition_idx], - edge_partition_offsets_[partition_idx] + offset_size), - raft::device_span( - edge_partition_indices_[partition_idx], - edge_partition_indices_[partition_idx] + edge_partition_number_of_edges_[partition_idx]), + edge_partition_offsets_[partition_idx], + edge_partition_indices_[partition_idx], edge_partition_dcs_nzd_vertices_ - ? std::make_optional>( - (*edge_partition_dcs_nzd_vertices_)[partition_idx], - (*edge_partition_dcs_nzd_vertices_)[partition_idx] + - (*edge_partition_dcs_nzd_vertex_counts_)[partition_idx]) + ? std::make_optional((*edge_partition_dcs_nzd_vertices_)[partition_idx]) : std::nullopt, major_hypersparse_first, major_range_first, @@ -652,6 +612,16 @@ class graph_view_thas_edge_mask()), "unimplemented."); + return this->number_of_edges_; + } + + edge_t compute_number_of_edges(raft::handle_t const& handle) const; + rmm::device_uvector compute_in_degrees(raft::handle_t const& handle) const; rmm::device_uvector compute_out_degrees(raft::handle_t const& handle) const; @@ -751,15 +721,26 @@ class graph_view_t 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_mask_view() const + { + return edge_mask_view_; + } + private: - std::vector edge_partition_offsets_{}; - std::vector edge_partition_indices_{}; + std::vector> edge_partition_offsets_{}; + std::vector> edge_partition_indices_{}; // relevant only if we use the CSR + DCSR (or CSC + DCSC) hybrid format - std::optional> edge_partition_dcs_nzd_vertices_{}; - std::optional> edge_partition_dcs_nzd_vertex_counts_{}; - - std::vector edge_partition_number_of_edges_{}; + std::optional>> edge_partition_dcs_nzd_vertices_{}; partition_t partition_{}; @@ -796,6 +777,8 @@ class graph_view_t>, std::optional /* dummy */> local_sorted_unique_edge_dst_vertex_partition_offsets_{std::nullopt}; + + std::optional> edge_mask_view_{std::nullopt}; }; // single-GPU version @@ -808,9 +791,8 @@ class graph_view_t offsets, + raft::device_span indices, graph_view_meta_t meta); std::vector vertex_partition_range_offsets() const @@ -924,11 +906,19 @@ class graph_view_t( - raft::device_span(offsets_, offsets_ + (this->number_of_vertices() + 1)), - raft::device_span(indices_, indices_ + this->number_of_edges()), - this->number_of_vertices()); + offsets_, indices_, this->number_of_vertices()); + } + + // FIXME: deprecated, replaced with copmute_number_of_edges (which works with or without edge + // masking) + edge_t number_of_edges() const + { + CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); + return this->number_of_edges_; } + edge_t compute_number_of_edges(raft::handle_t const& handle) const; + rmm::device_uvector compute_in_degrees(raft::handle_t const& handle) const; rmm::device_uvector compute_out_degrees(raft::handle_t const& handle) const; @@ -1016,12 +1006,28 @@ class graph_view_t 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_mask_view() const + { + return edge_mask_view_; + } + private: - edge_t const* offsets_{nullptr}; - vertex_t const* indices_{nullptr}; + raft::device_span offsets_{}; + raft::device_span indices_{}; // segment offsets based on vertex degree, relevant only if vertex IDs are renumbered std::optional> segment_offsets_{std::nullopt}; + + std::optional> edge_mask_view_{std::nullopt}; }; } // namespace cugraph diff --git a/cpp/src/prims/per_v_transform_reduce_incoming_outgoing_e.cuh b/cpp/src/prims/per_v_transform_reduce_incoming_outgoing_e.cuh index 5de2dbc18e2..1a7fc0130c4 100644 --- a/cpp/src/prims/per_v_transform_reduce_incoming_outgoing_e.cuh +++ b/cpp/src/prims/per_v_transform_reduce_incoming_outgoing_e.cuh @@ -593,10 +593,11 @@ void per_v_transform_reduce_e(raft::handle_t const& handle, value_size = sizeof(T); } - auto avg_vertex_degree = graph_view.number_of_vertices() > 0 - ? (static_cast(graph_view.number_of_edges()) / - static_cast(graph_view.number_of_vertices())) - : double{0.0}; + auto avg_vertex_degree = + graph_view.number_of_vertices() > 0 + ? (static_cast(graph_view.compute_number_of_edges(handle)) / + static_cast(graph_view.number_of_vertices())) + : double{0.0}; num_streams = std::min(static_cast(avg_vertex_degree * (static_cast(sizeof(vertex_t)) / diff --git a/cpp/src/prims/update_edge_src_dst_property.cuh b/cpp/src/prims/update_edge_src_dst_property.cuh index 18bdf5bcf2d..0c7058cccb4 100644 --- a/cpp/src/prims/update_edge_src_dst_property.cuh +++ b/cpp/src/prims/update_edge_src_dst_property.cuh @@ -479,7 +479,8 @@ void update_edge_minor_property(raft::handle_t const& handle, bcast_size *= sizeof(typename EdgeMinorPropertyOutputWrapper::value_type); } auto num_concurrent_bcasts = - (static_cast(graph_view.number_of_edges() / comm_size) * sizeof(vertex_t)) / + (static_cast(graph_view.compute_number_of_edges(handle) / comm_size) * + sizeof(vertex_t)) / std::max(bcast_size, size_t{1}); num_concurrent_bcasts = std::max(num_concurrent_bcasts, size_t{1}); num_concurrent_bcasts = std::min(num_concurrent_bcasts, static_cast(major_comm_size)); diff --git a/cpp/src/structure/graph_impl.cuh b/cpp/src/structure/graph_impl.cuh index 75862266789..6568b5e3b9e 100644 --- a/cpp/src/structure/graph_impl.cuh +++ b/cpp/src/structure/graph_impl.cuh @@ -133,8 +133,7 @@ update_local_sorted_unique_edge_majors_minors( graph_meta_t const& meta, std::vector> const& edge_partition_offsets, std::vector> const& edge_partition_indices, - std::optional>> const& edge_partition_dcs_nzd_vertices, - std::optional> const& edge_partition_dcs_nzd_vertex_counts) + std::optional>> const& edge_partition_dcs_nzd_vertices) { auto& comm = handle.get_comms(); auto& major_comm = handle.get_subcomm(cugraph::partition_manager::major_comm_name()); @@ -341,8 +340,7 @@ update_local_sorted_unique_edge_majors_minors( if (use_dcs) { thrust::copy(handle.get_thrust_policy(), (*edge_partition_dcs_nzd_vertices)[i].begin(), - (*edge_partition_dcs_nzd_vertices)[i].begin() + - (*edge_partition_dcs_nzd_vertex_counts)[i], + (*edge_partition_dcs_nzd_vertices)[i].end(), unique_edge_majors.begin() + cur_size); } @@ -390,7 +388,7 @@ graph_t meta, bool do_expensive_check) : detail::graph_base_t( - handle, meta.number_of_vertices, meta.number_of_edges, meta.properties), + meta.number_of_vertices, meta.number_of_edges, meta.properties), partition_(meta.partition) { CUGRAPH_EXPECTS( @@ -408,14 +406,6 @@ graph_t((*edge_partition_dcs_nzd_vertices_).size()); - for (size_t i = 0; i < (*edge_partition_dcs_nzd_vertex_counts_).size(); ++i) { - (*edge_partition_dcs_nzd_vertex_counts_)[i] = - static_cast((*edge_partition_dcs_nzd_vertices_)[i].size()); - } - } // update local sorted unique edge sources/destinations (only if key, value pair will be used) @@ -432,8 +422,7 @@ graph_t meta, bool do_expensive_check) : detail::graph_base_t( - handle, meta.number_of_vertices, static_cast(indices.size()), meta.properties), + meta.number_of_vertices, static_cast(indices.size()), meta.properties), offsets_(std::move(offsets)), indices_(std::move(indices)), segment_offsets_(meta.segment_offsets) diff --git a/cpp/src/structure/graph_view_impl.cuh b/cpp/src/structure/graph_view_impl.cuh index a40747f13f7..da0ecc991df 100644 --- a/cpp/src/structure/graph_view_impl.cuh +++ b/cpp/src/structure/graph_view_impl.cuh @@ -72,44 +72,15 @@ struct out_of_range_t { __device__ bool operator()(vertex_t v) const { return (v < min) || (v >= max); } }; -template -std::vector update_edge_partition_edge_counts( - std::vector const& edge_partition_offsets, - std::optional> const& edge_partition_dcs_nzd_vertex_counts, - partition_t const& partition, - std::vector const& edge_partition_segment_offsets, - cudaStream_t stream) -{ - std::vector edge_partition_edge_counts(partition.number_of_local_edge_partitions(), 0); - auto use_dcs = edge_partition_dcs_nzd_vertex_counts.has_value(); - for (size_t i = 0; i < edge_partition_offsets.size(); ++i) { - auto [major_range_first, major_range_last] = partition.local_edge_partition_major_range(i); - auto segment_offset_size_per_partition = - edge_partition_segment_offsets.size() / edge_partition_offsets.size(); - raft::update_host( - &(edge_partition_edge_counts[i]), - edge_partition_offsets[i] + - (use_dcs - ? (edge_partition_segment_offsets[segment_offset_size_per_partition * i + - detail::num_sparse_segments_per_vertex_partition] + - (*edge_partition_dcs_nzd_vertex_counts)[i]) - : (major_range_last - major_range_first)), - 1, - stream); - } - RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); - return edge_partition_edge_counts; -} - // compute out-degrees (if we are internally storing edges in the sparse 2D matrix using sources as // major indices) or in-degrees (otherwise) template rmm::device_uvector compute_major_degrees( raft::handle_t const& handle, - std::vector const& edge_partition_offsets, - std::optional> const& edge_partition_dcs_nzd_vertices, - std::optional> const& edge_partition_dcs_nzd_vertex_counts, - std::optional> const& edge_partition_masks, + std::vector> const& edge_partition_offsets, + std::optional>> const& + edge_partition_dcs_nzd_vertices, + std::optional>> const& edge_partition_masks, partition_t const& partition, std::vector const& edge_partition_segment_offsets) { @@ -165,35 +136,35 @@ rmm::device_uvector compute_major_degrees( auto local_degree = offsets[i + 1] - offsets[i]; if (masks) { local_degree = static_cast( - detail::count_set_bits(*masks, offsets[i], local_degree)); + detail::count_set_bits((*masks).begin(), offsets[i], local_degree)); } return local_degree; })); if (use_dcs) { - auto dcs_nzd_vertices = (*edge_partition_dcs_nzd_vertices)[i]; - auto dcs_nzd_vertex_count = (*edge_partition_dcs_nzd_vertex_counts)[i]; + auto dcs_nzd_vertices = (*edge_partition_dcs_nzd_vertices)[i]; thrust::fill(execution_policy, local_degrees.begin() + (major_hypersparse_first - major_range_first), local_degrees.begin() + (major_range_last - major_range_first), edge_t{0}); - thrust::for_each(execution_policy, - thrust::make_counting_iterator(vertex_t{0}), - thrust::make_counting_iterator(dcs_nzd_vertex_count), - [offsets, - dcs_nzd_vertices, - masks, - major_range_first, - major_hypersparse_first, - local_degrees = local_degrees.data()] __device__(auto i) { - auto major_idx = (major_hypersparse_first - major_range_first) + i; - auto local_degree = offsets[major_idx + 1] - offsets[major_idx]; - if (masks) { - local_degree = static_cast( - detail::count_set_bits(*masks, offsets[major_idx], local_degree)); - } - auto v = dcs_nzd_vertices[i]; - local_degrees[v - major_range_first] = local_degree; - }); + thrust::for_each( + execution_policy, + thrust::make_counting_iterator(vertex_t{0}), + thrust::make_counting_iterator(static_cast(dcs_nzd_vertices.size())), + [offsets, + dcs_nzd_vertices, + masks, + major_range_first, + major_hypersparse_first, + local_degrees = local_degrees.data()] __device__(auto i) { + auto major_idx = (major_hypersparse_first - major_range_first) + i; + auto local_degree = offsets[major_idx + 1] - offsets[major_idx]; + if (masks) { + local_degree = static_cast( + detail::count_set_bits((*masks).begin(), offsets[major_idx], local_degree)); + } + auto v = dcs_nzd_vertices[i]; + local_degrees[v - major_range_first] = local_degree; + }); } minor_comm.reduce(local_degrees.data(), i == minor_comm_rank ? degrees.data() : static_cast(nullptr), @@ -209,10 +180,11 @@ rmm::device_uvector compute_major_degrees( // compute out-degrees (if we are internally storing edges in the sparse 2D matrix using sources as // major indices) or in-degrees (otherwise) template -rmm::device_uvector compute_major_degrees(raft::handle_t const& handle, - edge_t const* offsets, - std::optional masks, - vertex_t number_of_vertices) +rmm::device_uvector compute_major_degrees( + raft::handle_t const& handle, + raft::device_span offsets, + std::optional> masks, + vertex_t number_of_vertices) { rmm::device_uvector degrees(number_of_vertices, handle.get_stream()); thrust::tabulate( @@ -223,7 +195,7 @@ rmm::device_uvector compute_major_degrees(raft::handle_t const& handle, auto local_degree = offsets[i + 1] - offsets[i]; if (masks) { local_degree = - static_cast(detail::count_set_bits(*masks, offsets[i], local_degree)); + static_cast(detail::count_set_bits((*masks).begin(), offsets[i], local_degree)); } return local_degree; }); @@ -446,24 +418,16 @@ edge_t count_edge_partition_multi_edges( template graph_view_t>:: - graph_view_t(raft::handle_t const& handle, - std::vector const& edge_partition_offsets, - std::vector const& edge_partition_indices, - std::optional> const& edge_partition_dcs_nzd_vertices, - std::optional> const& edge_partition_dcs_nzd_vertex_counts, + graph_view_t(std::vector> const& edge_partition_offsets, + std::vector> const& edge_partition_indices, + std::optional>> const& + edge_partition_dcs_nzd_vertices, graph_view_meta_t meta) : detail::graph_base_t( - handle, meta.number_of_vertices, meta.number_of_edges, meta.properties), + meta.number_of_vertices, meta.number_of_edges, meta.properties), edge_partition_offsets_(edge_partition_offsets), edge_partition_indices_(edge_partition_indices), edge_partition_dcs_nzd_vertices_(edge_partition_dcs_nzd_vertices), - edge_partition_dcs_nzd_vertex_counts_(edge_partition_dcs_nzd_vertex_counts), - edge_partition_number_of_edges_( - update_edge_partition_edge_counts(edge_partition_offsets, - edge_partition_dcs_nzd_vertex_counts, - meta.partition, - meta.edge_partition_segment_offsets, - handle.get_stream())), partition_(meta.partition), edge_partition_segment_offsets_(meta.edge_partition_segment_offsets), local_sorted_unique_edge_srcs_(meta.local_sorted_unique_edge_srcs), @@ -481,51 +445,42 @@ graph_view_thandle_ptr()->get_subcomm(cugraph::partition_manager::minor_comm_name()).get_size(); - auto use_dcs = edge_partition_dcs_nzd_vertices.has_value(); CUGRAPH_EXPECTS(edge_partition_offsets.size() == edge_partition_indices.size(), "Internal Error: edge_partition_offsets.size() and " "edge_partition_indices.size() should coincide."); - CUGRAPH_EXPECTS(edge_partition_dcs_nzd_vertex_counts.has_value() == use_dcs, - "edge_partition_dcs_nzd_vertices.has_value() and " - "edge_partition_dcs_nzd_vertex_counts.has_value() should coincide"); - CUGRAPH_EXPECTS(!use_dcs || ((*edge_partition_dcs_nzd_vertices).size() == - (*edge_partition_dcs_nzd_vertex_counts).size()), - "Internal Error: edge_partition_dcs_nzd_vertices.size() and " - "edge_partition_dcs_nzd_vertex_counts.size() should coincide (if used)."); CUGRAPH_EXPECTS( !use_dcs || ((*edge_partition_dcs_nzd_vertices).size() == edge_partition_offsets.size()), "Internal Error: edge_partition_dcs_nzd_vertices.size() should coincide " "with edge_partition_offsets.size() (if used)."); - CUGRAPH_EXPECTS(edge_partition_offsets.size() == static_cast(minor_comm_size), - "Internal Error: erroneous edge_partition_offsets.size()."); - - CUGRAPH_EXPECTS( - meta.edge_partition_segment_offsets.size() == - minor_comm_size * (detail::num_sparse_segments_per_vertex_partition + (use_dcs ? 3 : 2)), - "Internal Error: invalid edge_partition_segment_offsets.size()."); + CUGRAPH_EXPECTS(meta.edge_partition_segment_offsets.size() == + edge_partition_offsets.size() * + (detail::num_sparse_segments_per_vertex_partition + (use_dcs ? 3 : 2)), + "Internal Error: invalid edge_partition_segment_offsets.size()."); // skip expensive error checks as this function is only called by graph_t } template graph_view_t>:: - graph_view_t(raft::handle_t const& handle, - edge_t const* offsets, - vertex_t const* indices, + graph_view_t(raft::device_span offsets, + raft::device_span indices, graph_view_meta_t meta) : detail::graph_base_t( - handle, meta.number_of_vertices, meta.number_of_edges, meta.properties), + meta.number_of_vertices, meta.number_of_edges, meta.properties), offsets_(offsets), indices_(indices), segment_offsets_(meta.segment_offsets) { // cheap error checks + CUGRAPH_EXPECTS(offsets.size() == static_cast(meta.number_of_vertices + 1), + "Internal Error: offsets.size() returns an invalid value."); + CUGRAPH_EXPECTS(indices.size() == static_cast(meta.number_of_edges), + "Internal Error: indices.size() returns an invalid value."); + CUGRAPH_EXPECTS( !(meta.segment_offsets).has_value() || ((*(meta.segment_offsets)).size() == (detail::num_sparse_segments_per_vertex_partition + 2)), @@ -534,19 +489,62 @@ graph_view_t +edge_t graph_view_t>:: + compute_number_of_edges(raft::handle_t const& handle) const +{ + if (this->has_edge_mask()) { + edge_t ret{}; + auto value_firsts = (*(this->edge_mask_view())).value_firsts(); + auto edge_counts = (*(this->edge_mask_view())).edge_counts(); + for (size_t i = 0; i < value_firsts.size(); ++i) { + ret += static_cast(detail::count_set_bits(handle, value_firsts[i], edge_counts[i])); + } + ret = + host_scalar_allreduce(handle.get_comms(), ret, raft::comms::op_t::SUM, handle.get_stream()); + return ret; + } else { + return this->number_of_edges_; + } +} + +template +edge_t graph_view_t>:: + compute_number_of_edges(raft::handle_t const& handle) const +{ + if (this->has_edge_mask()) { + auto value_firsts = (*(this->edge_mask_view())).value_firsts(); + auto edge_counts = (*(this->edge_mask_view())).edge_counts(); + assert(value_firsts.size() == 0); + assert(edge_counts.size() == 0); + return static_cast(detail::count_set_bits(handle, value_firsts[0], edge_counts[0])); + } else { + return this->number_of_edges_; + } +} + template rmm::device_uvector graph_view_t>:: compute_in_degrees(raft::handle_t const& handle) const { if (store_transposed) { + std::optional>> edge_partition_masks{ + std::nullopt}; + if (this->has_edge_mask()) { + edge_partition_masks = + std::vector>(this->edge_partition_offsets_.size()); + auto value_firsts = (*(this->edge_mask_view())).value_firsts(); + auto edge_counts = (*(this->edge_mask_view())).edge_counts(); + for (size_t i = 0; i < (*edge_partition_masks).size(); ++i) { + (*edge_partition_masks)[i] = + raft::device_span(value_firsts[i], edge_counts[i]); + } + } return compute_major_degrees(handle, this->edge_partition_offsets_, this->edge_partition_dcs_nzd_vertices_, - this->edge_partition_dcs_nzd_vertex_counts_, - this->has_edge_mask() - ? std::make_optional((*(this->edge_mask_view())).value_firsts()) - : std::nullopt, + edge_partition_masks, this->partition_, this->edge_partition_segment_offsets_); } else { @@ -561,12 +559,14 @@ graph_view_toffsets_, - this->has_edge_mask() ? std::make_optional((*(this->edge_mask_view())).value_firsts()[0]) - : std::nullopt, - this->local_vertex_partition_range_size()); + return compute_major_degrees(handle, + this->offsets_, + this->has_edge_mask() + ? std::make_optional(raft::device_span( + (*(this->edge_mask_view())).value_firsts()[0], + (*(this->edge_mask_view())).edge_counts()[0])) + : std::nullopt, + this->local_vertex_partition_range_size()); } else { CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); return compute_minor_degrees(handle, *this); @@ -582,13 +582,22 @@ graph_view_thas_edge_mask()), "unimplemented."); return compute_minor_degrees(handle, *this); } else { + std::optional>> edge_partition_masks{ + std::nullopt}; + if (this->has_edge_mask()) { + edge_partition_masks = + std::vector>(this->edge_partition_offsets_.size()); + auto value_firsts = (*(this->edge_mask_view())).value_firsts(); + auto edge_counts = (*(this->edge_mask_view())).edge_counts(); + for (size_t i = 0; i < (*edge_partition_masks).size(); ++i) { + (*edge_partition_masks)[i] = + raft::device_span(value_firsts[i], edge_counts[i]); + } + } return compute_major_degrees(handle, this->edge_partition_offsets_, this->edge_partition_dcs_nzd_vertices_, - this->edge_partition_dcs_nzd_vertex_counts_, - this->has_edge_mask() - ? std::make_optional((*(this->edge_mask_view())).value_firsts()) - : std::nullopt, + edge_partition_masks, this->partition_, this->edge_partition_segment_offsets_); } @@ -603,12 +612,14 @@ graph_view_thas_edge_mask()), "unimplemented."); return compute_minor_degrees(handle, *this); } else { - return compute_major_degrees( - handle, - this->offsets_, - this->has_edge_mask() ? std::make_optional((*(this->edge_mask_view())).value_firsts()[0]) - : std::nullopt, - this->local_vertex_partition_range_size()); + return compute_major_degrees(handle, + this->offsets_, + this->has_edge_mask() + ? std::make_optional(raft::device_span( + (*(this->edge_mask_view())).value_firsts()[0], + (*(this->edge_mask_view())).edge_counts()[0])) + : std::nullopt, + this->local_vertex_partition_range_size()); } } diff --git a/cpp/src/traversal/od_shortest_distances_impl.cuh b/cpp/src/traversal/od_shortest_distances_impl.cuh index 6a0c5a4a675..cc69cb5f67f 100644 --- a/cpp/src/traversal/od_shortest_distances_impl.cuh +++ b/cpp/src/traversal/od_shortest_distances_impl.cuh @@ -432,7 +432,7 @@ rmm::device_uvector od_shortest_distances( // 1. check input arguments auto const num_vertices = graph_view.number_of_vertices(); - auto const num_edges = graph_view.number_of_edges(); + auto const num_edges = graph_view.compute_number_of_edges(handle); CUGRAPH_EXPECTS(num_vertices != 0 || (origins.size() == 0 && destinations.size() == 0), "Invalid input argument: the input graph is empty but origins.size() > 0 or " @@ -1049,7 +1049,7 @@ rmm::device_uvector od_shortest_distances( CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); auto const num_vertices = graph_view.number_of_vertices(); - auto const num_edges = graph_view.number_of_edges(); + auto const num_edges = graph_view.compute_number_of_edges(handle); weight_t average_vertex_degree = static_cast(num_edges) / static_cast(num_vertices); diff --git a/cpp/src/traversal/sssp_impl.cuh b/cpp/src/traversal/sssp_impl.cuh index c78fa3839e2..5a6d536c6f5 100644 --- a/cpp/src/traversal/sssp_impl.cuh +++ b/cpp/src/traversal/sssp_impl.cuh @@ -93,7 +93,7 @@ void sssp(raft::handle_t const& handle, "GraphViewType should support the push model."); auto const num_vertices = push_graph_view.number_of_vertices(); - auto const num_edges = push_graph_view.number_of_edges(); + auto const num_edges = push_graph_view.compute_number_of_edges(handle); if (num_vertices == 0) { return; } // implements the Near-Far Pile method in diff --git a/cpp/src/utilities/cugraph_ops_utils.hpp b/cpp/src/utilities/cugraph_ops_utils.hpp index 9aea4183866..880a2c8d104 100644 --- a/cpp/src/utilities/cugraph_ops_utils.hpp +++ b/cpp/src/utilities/cugraph_ops_utils.hpp @@ -30,7 +30,7 @@ ops::graph::csc get_graph( ops::graph::csc graph; graph.n_src_nodes = gview.number_of_vertices(); graph.n_dst_nodes = gview.number_of_vertices(); - graph.n_indices = gview.number_of_edges(); + graph.n_indices = gview.local_edge_partition_view().number_of_edges(); // FIXME this is sufficient for now, but if there is a fast (cached) way // of getting max degree, use that instead graph.dst_max_in_degree = std::numeric_limits::max(); diff --git a/cpp/tests/link_analysis/hits_test.cpp b/cpp/tests/link_analysis/hits_test.cpp index 6796761e212..cf35356bb76 100644 --- a/cpp/tests/link_analysis/hits_test.cpp +++ b/cpp/tests/link_analysis/hits_test.cpp @@ -229,7 +229,7 @@ class Tests_Hits : public ::testing::TestWithParam sg_indices(sg_graph_view.number_of_edges(), - handle_->get_stream()); + rmm::device_uvector sg_indices( + sg_graph_view.local_edge_partition_view().indices().size(), handle_->get_stream()); thrust::copy(handle_->get_thrust_policy(), sg_graph_view.local_edge_partition_view().indices().begin(), sg_graph_view.local_edge_partition_view().indices().end(),