diff --git a/cpp/include/cugraph/edge_src_dst_property.hpp b/cpp/include/cugraph/edge_src_dst_property.hpp index 0724dd329e8..3d61560a300 100644 --- a/cpp/include/cugraph/edge_src_dst_property.hpp +++ b/cpp/include/cugraph/edge_src_dst_property.hpp @@ -294,9 +294,9 @@ class edge_major_property_t { void clear() { - edge_partition_keys_ = std::nullopt; - edge_partition_key_chunk_start_offsets_ = std::nullopt; - key_chunk_size_ = std::nullopt; + edge_partition_keys_.reset(); + edge_partition_key_chunk_start_offsets_.reset(); + key_chunk_size_.reset(); buffers_.clear(); buffers_.shrink_to_fit(); @@ -424,9 +424,9 @@ class edge_minor_property_t { void clear() { - keys_ = std::nullopt; - key_chunk_start_offsets_ = std::nullopt; - key_chunk_size_ = std::nullopt; + keys_.reset(); + key_chunk_start_offsets_.reset(); + key_chunk_size_.reset(); rmm::cuda_stream_view stream{}; if constexpr (std::is_arithmetic_v) { @@ -546,6 +546,12 @@ class edge_src_property_t { public: static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); + using value_type = T; + using value_iterator = + dataframe_buffer_iterator_type_t, uint32_t, T>>; + using const_value_iterator = dataframe_buffer_const_iterator_type_t< + std::conditional_t, uint32_t, T>>; + edge_src_property_t(raft::handle_t const& handle) {} template @@ -599,6 +605,10 @@ class edge_dst_property_t { static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); using value_type = T; + using value_iterator = + dataframe_buffer_iterator_type_t, uint32_t, T>>; + using const_value_iterator = dataframe_buffer_const_iterator_type_t< + std::conditional_t, uint32_t, T>>; edge_dst_property_t(raft::handle_t const& handle) {} @@ -661,6 +671,160 @@ class edge_dst_dummy_property_t { auto view() const { return detail::edge_endpoint_dummy_property_view_t{}; } }; +// SG-only (use a vertex property buffer instead of creating a new edge_src_property_t object to +// save memory) +template +detail::edge_endpoint_property_view_t::const_value_iterator, + T> +make_edge_src_property_view( + GraphViewType const& graph_view, + typename edge_src_property_t::const_value_iterator value_first, + size_t num_values) +{ + using const_value_iterator = typename edge_src_property_t::const_value_iterator; + + CUGRAPH_EXPECTS(!GraphViewType::is_multi_gpu, + "Invalid input argument: this function is only for single-GPU."); + + vertex_t range_size{}; + if constexpr (GraphViewType::is_storage_transposed) { // minor + range_size = graph_view.local_edge_partition_src_range_size(); + } else { // major + range_size = graph_view.local_edge_partition_src_range_size(size_t{0}); + } + auto expected_num_values = std::is_same_v + ? cugraph::packed_bool_size(static_cast(range_size)) + : static_cast(range_size); + CUGRAPH_EXPECTS( + num_values == expected_num_values, + "Invalid input argument: num_values does not match the expected number of values."); + + if constexpr (GraphViewType::is_storage_transposed) { // minor + return detail::edge_endpoint_property_view_t(value_first, + vertex_t{0}); + } else { // major + return detail::edge_endpoint_property_view_t( + std::vector{value_first}, std::vector{vertex_t{0}}); + } +} + +// SG-only (use a vertex property buffer instead of creating a new edge_src_property_t object to +// save memory) +template +detail::edge_endpoint_property_view_t::value_iterator, + T> +make_edge_src_property_mutable_view( + GraphViewType const& graph_view, + typename edge_src_property_t::value_iterator value_first, + size_t num_values) +{ + using value_iterator = typename edge_src_property_t::value_iterator; + + CUGRAPH_EXPECTS(!GraphViewType::is_multi_gpu, + "Invalid input argument: this function is only for single-GPU."); + + vertex_t range_size{}; + if constexpr (GraphViewType::is_storage_transposed) { // minor + range_size = graph_view.local_edge_partition_src_range_size(); + } else { // major + range_size = graph_view.local_edge_partition_src_range_size(size_t{0}); + } + auto expected_num_values = std::is_same_v + ? cugraph::packed_bool_size(static_cast(range_size)) + : static_cast(range_size); + CUGRAPH_EXPECTS( + num_values == expected_num_values, + "Invalid input argument: num_values does not match the expected number of values."); + + if constexpr (GraphViewType::is_storage_transposed) { // minor + return detail::edge_endpoint_property_view_t(value_first, + vertex_t{0}); + } else { // major + return detail::edge_endpoint_property_view_t( + std::vector{value_first}, std::vector{vertex_t{0}}); + } +} + +// SG-only (use a vertex property buffer instead of creating a new edge_dst_property_t object to +// save memory) +template +detail::edge_endpoint_property_view_t::const_value_iterator, + T> +make_edge_dst_property_view( + GraphViewType const& graph_view, + typename edge_dst_property_t::const_value_iterator value_first, + size_t num_values) +{ + using const_value_iterator = typename edge_dst_property_t::const_value_iterator; + + CUGRAPH_EXPECTS(!GraphViewType::is_multi_gpu, + "Invalid input argument: this function is only for single-GPU."); + + vertex_t range_size{}; + if constexpr (GraphViewType::is_storage_transposed) { // major + range_size = graph_view.local_edge_partition_dst_range_size(size_t{0}); + } else { // minor + range_size = graph_view.local_edge_partition_dst_range_size(); + } + auto expected_num_values = std::is_same_v + ? cugraph::packed_bool_size(static_cast(range_size)) + : static_cast(range_size); + CUGRAPH_EXPECTS( + num_values == expected_num_values, + "Invalid input argument: num_values does not match the expected number of values."); + + if constexpr (GraphViewType::is_storage_transposed) { // major + return detail::edge_endpoint_property_view_t( + std::vector{value_first}, std::vector{vertex_t{0}}); + } else { // minor + return detail::edge_endpoint_property_view_t(value_first, + vertex_t{0}); + } +} + +// SG-only (use a vertex property buffer instead of creating a new edge_dst_property_t object to +// save memory) +template +detail::edge_endpoint_property_view_t::value_iterator, + T> +make_edge_dst_property_mutable_view( + GraphViewType const& graph_view, + typename edge_dst_property_t::value_iterator value_first, + size_t num_values) +{ + using value_iterator = typename edge_dst_property_t::value_iterator; + + CUGRAPH_EXPECTS(!GraphViewType::is_multi_gpu, + "Invalid input argument: this function is only for single-GPU."); + + vertex_t range_size{}; + if constexpr (GraphViewType::is_storage_transposed) { // major + range_size = graph_view.local_edge_partition_dst_range_size(size_t{0}); + } else { // minor + range_size = graph_view.local_edge_partition_dst_range_size(); + } + auto expected_num_values = std::is_same_v + ? cugraph::packed_bool_size(static_cast(range_size)) + : static_cast(range_size); + CUGRAPH_EXPECTS( + num_values == expected_num_values, + "Invalid input argument: num_values does not match the expected number of values."); + + if constexpr (GraphViewType::is_storage_transposed) { // major + return detail::edge_endpoint_property_view_t( + std::vector{value_first}, std::vector{vertex_t{0}}); + } else { // minor + return detail::edge_endpoint_property_view_t(value_first, + vertex_t{0}); + } +} + template auto view_concat(detail::edge_endpoint_property_view_t const&... views) { diff --git a/cpp/src/community/approx_weighted_matching_impl.cuh b/cpp/src/community/approx_weighted_matching_impl.cuh index 3bb0961fb9d..3f847e43383 100644 --- a/cpp/src/community/approx_weighted_matching_impl.cuh +++ b/cpp/src/community/approx_weighted_matching_impl.cuh @@ -45,9 +45,7 @@ std::tuple, weight_t> approximate_weighted_matchin "Invalid input arguments: input graph for approximate_weighted_matching must " "need to be symmetric"); - using graph_view_t = cugraph::graph_view_t; - - graph_view_t current_graph_view(graph_view); + auto current_graph_view = graph_view; if (current_graph_view.has_edge_mask()) { current_graph_view.clear_edge_mask(); } cugraph::edge_property_t edge_masks_even(handle, current_graph_view); @@ -100,7 +98,7 @@ std::tuple, weight_t> approximate_weighted_matchin cugraph::edge_src_property_t src_match_flags(handle); cugraph::edge_dst_property_t dst_match_flags(handle); - if constexpr (graph_view_t::is_multi_gpu) { + if constexpr (multi_gpu) { src_key_cache = cugraph::edge_src_property_t(handle, current_graph_view); update_edge_src_property( @@ -129,11 +127,9 @@ std::tuple, weight_t> approximate_weighted_matchin cugraph::edge_src_dummy_property_t{}.view(), cugraph::edge_dst_dummy_property_t{}.view(), edge_weight_view, - graph_view_t::is_multi_gpu - ? src_key_cache.view() - : detail::edge_endpoint_property_view_t( - std::vector{local_vertices.begin()}, - std::vector{vertex_t{0}}), + multi_gpu ? src_key_cache.view() + : make_edge_src_property_view( + current_graph_view, local_vertices.begin(), local_vertices.size()), [] __device__(auto, auto dst, cuda::std::nullopt_t, cuda::std::nullopt_t, auto wt) { return thrust::make_tuple(wt, dst); }, @@ -145,7 +141,7 @@ std::tuple, weight_t> approximate_weighted_matchin // For each target, find the best offer // - if constexpr (graph_view_t::is_multi_gpu) { + if constexpr (multi_gpu) { auto vertex_partition_range_lasts = current_graph_view.vertex_partition_range_lasts(); rmm::device_uvector d_vertex_partition_range_lasts( @@ -225,7 +221,7 @@ std::tuple, weight_t> approximate_weighted_matchin rmm::device_uvector candidates_of_candidates(0, handle.get_stream()); - if (graph_view_t::is_multi_gpu) { + if constexpr (multi_gpu) { auto& comm = handle.get_comms(); auto& major_comm = handle.get_subcomm(cugraph::partition_manager::major_comm_name()); auto const major_comm_size = major_comm.get_size(); @@ -302,7 +298,7 @@ std::tuple, weight_t> approximate_weighted_matchin if (current_graph_view.compute_number_of_edges(handle) == 0) { break; } - if constexpr (graph_view_t::is_multi_gpu) { + if constexpr (multi_gpu) { cugraph::update_edge_src_property( handle, current_graph_view, is_vertex_matched.begin(), src_match_flags.mutable_view()); cugraph::update_edge_dst_property( @@ -310,7 +306,7 @@ std::tuple, weight_t> approximate_weighted_matchin } if (loop_counter % 2 == 0) { - if constexpr (graph_view_t::is_multi_gpu) { + if constexpr (multi_gpu) { cugraph::transform_e( handle, current_graph_view, @@ -326,14 +322,17 @@ std::tuple, weight_t> approximate_weighted_matchin cugraph::transform_e( handle, current_graph_view, - detail::edge_endpoint_property_view_t( - std::vector{is_vertex_matched.begin()}, - std::vector{vertex_t{0}}), - detail::edge_endpoint_property_view_t(is_vertex_matched.begin(), - vertex_t{0}), + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), cugraph::edge_dummy_property_t{}.view(), - [] __device__( - auto src, auto dst, auto is_src_matched, auto is_dst_matched, cuda::std::nullopt_t) { + [is_vertex_matched = raft::device_span( + is_vertex_matched.data(), is_vertex_matched.size())] __device__(auto src, + auto dst, + cuda::std::nullopt_t, + cuda::std::nullopt_t, + cuda::std::nullopt_t) { + auto is_src_matched = is_vertex_matched[src]; + auto is_dst_matched = is_vertex_matched[dst]; return !((is_src_matched == true) || (is_dst_matched == true)); }, edge_masks_odd.mutable_view()); @@ -344,7 +343,7 @@ std::tuple, weight_t> approximate_weighted_matchin handle, current_graph_view, edge_masks_even.mutable_view(), bool{false}); current_graph_view.attach_edge_mask(edge_masks_odd.view()); } else { - if constexpr (graph_view_t::is_multi_gpu) { + if constexpr (multi_gpu) { cugraph::transform_e( handle, current_graph_view, @@ -360,14 +359,17 @@ std::tuple, weight_t> approximate_weighted_matchin cugraph::transform_e( handle, current_graph_view, - detail::edge_endpoint_property_view_t( - std::vector{is_vertex_matched.begin()}, - std::vector{vertex_t{0}}), - detail::edge_endpoint_property_view_t(is_vertex_matched.begin(), - vertex_t{0}), + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), cugraph::edge_dummy_property_t{}.view(), - [] __device__( - auto src, auto dst, auto is_src_matched, auto is_dst_matched, cuda::std::nullopt_t) { + [is_vertex_matched = raft::device_span( + is_vertex_matched.data(), is_vertex_matched.size())] __device__(auto src, + auto dst, + cuda::std::nullopt_t, + cuda::std::nullopt_t, + cuda::std::nullopt_t) { + auto is_src_matched = is_vertex_matched[src]; + auto is_dst_matched = is_vertex_matched[dst]; return !((is_src_matched == true) || (is_dst_matched == true)); }, edge_masks_even.mutable_view()); @@ -385,13 +387,14 @@ std::tuple, weight_t> approximate_weighted_matchin weight_t sum_matched_edge_weights = thrust::reduce( handle.get_thrust_policy(), offers_from_partners.begin(), offers_from_partners.end()); - if constexpr (graph_view_t::is_multi_gpu) { + if constexpr (multi_gpu) { sum_matched_edge_weights = host_scalar_allreduce( handle.get_comms(), sum_matched_edge_weights, raft::comms::op_t::SUM, handle.get_stream()); } return std::make_tuple(std::move(partners), sum_matched_edge_weights / 2.0); } + } // namespace detail template diff --git a/cpp/src/community/detail/common_methods.cuh b/cpp/src/community/detail/common_methods.cuh index 732ba03f4c1..3f189e28a9f 100644 --- a/cpp/src/community/detail/common_methods.cuh +++ b/cpp/src/community/detail/common_methods.cuh @@ -188,13 +188,12 @@ weight_t compute_modularity( weight_t sum_internal = transform_reduce_e( handle, graph_view, - multi_gpu - ? src_clusters_cache.view() - : detail::edge_endpoint_property_view_t( - std::vector{next_clusters.begin()}, std::vector{vertex_t{0}}), + multi_gpu ? src_clusters_cache.view() + : make_edge_src_property_view( + graph_view, next_clusters.begin(), next_clusters.size()), multi_gpu ? dst_clusters_cache.view() - : detail::edge_endpoint_property_view_t( - next_clusters.begin(), vertex_t{0}), + : make_edge_dst_property_view( + graph_view, next_clusters.begin(), next_clusters.size()), *edge_weight_view, cuda::proclaim_return_type( [] __device__(auto, auto, auto src_cluster, auto nbr_cluster, weight_t wt) { @@ -262,7 +261,7 @@ rmm::device_uvector update_clustering_by_delta_modularity( CUGRAPH_EXPECTS(edge_weight_view.has_value(), "Graph must be weighted."); rmm::device_uvector vertex_cluster_weights_v(0, handle.get_stream()); - edge_src_property_t src_cluster_weights(handle); + std::optional> src_cluster_weights{std::nullopt}; if constexpr (multi_gpu) { auto& comm = handle.get_comms(); @@ -291,7 +290,7 @@ rmm::device_uvector update_clustering_by_delta_modularity( src_cluster_weights = edge_src_property_t(handle, graph_view); update_edge_src_property( - handle, graph_view, vertex_cluster_weights_v.begin(), src_cluster_weights.mutable_view()); + handle, graph_view, vertex_cluster_weights_v.begin(), src_cluster_weights->mutable_view()); vertex_cluster_weights_v.resize(0, handle.get_stream()); vertex_cluster_weights_v.shrink_to_fit(handle.get_stream()); } else { @@ -324,13 +323,12 @@ rmm::device_uvector update_clustering_by_delta_modularity( per_v_transform_reduce_outgoing_e( handle, graph_view, - multi_gpu - ? src_clusters_cache.view() - : detail::edge_endpoint_property_view_t( - std::vector{next_clusters_v.data()}, std::vector{vertex_t{0}}), + multi_gpu ? src_clusters_cache.view() + : make_edge_src_property_view( + graph_view, next_clusters_v.begin(), next_clusters_v.size()), multi_gpu ? dst_clusters_cache.view() - : detail::edge_endpoint_property_view_t( - next_clusters_v.data(), vertex_t{0}), + : make_edge_dst_property_view( + graph_view, next_clusters_v.begin(), next_clusters_v.size()), *edge_weight_view, [] __device__(auto src, auto dst, auto src_cluster, auto nbr_cluster, weight_t wt) { weight_t sum{0}; @@ -348,8 +346,8 @@ rmm::device_uvector update_clustering_by_delta_modularity( thrust::make_zip_iterator( thrust::make_tuple(old_cluster_sum_v.begin(), cluster_subtract_v.begin()))); - edge_src_property_t> - src_old_cluster_sum_subtract_pairs(handle); + std::optional>> + src_old_cluster_sum_subtract_pairs{std::nullopt}; if constexpr (multi_gpu) { src_old_cluster_sum_subtract_pairs = @@ -358,7 +356,7 @@ rmm::device_uvector update_clustering_by_delta_modularity( graph_view, thrust::make_zip_iterator(thrust::make_tuple( old_cluster_sum_v.begin(), cluster_subtract_v.begin())), - src_old_cluster_sum_subtract_pairs.mutable_view()); + src_old_cluster_sum_subtract_pairs->mutable_view()); old_cluster_sum_v.resize(0, handle.get_stream()); old_cluster_sum_v.shrink_to_fit(handle.get_stream()); cluster_subtract_v.resize(0, handle.get_stream()); @@ -371,26 +369,19 @@ rmm::device_uvector update_clustering_by_delta_modularity( auto cluster_old_sum_subtract_pair_first = thrust::make_zip_iterator( thrust::make_tuple(old_cluster_sum_v.cbegin(), cluster_subtract_v.cbegin())); auto zipped_src_device_view = - multi_gpu - ? view_concat(src_vertex_weights_cache.view(), - src_clusters_cache.view(), - src_cluster_weights.view(), - src_old_cluster_sum_subtract_pairs.view()) - : view_concat( - detail::edge_endpoint_property_view_t( - std::vector{vertex_weights_v.data()}, - std::vector{vertex_t{0}}), - detail::edge_endpoint_property_view_t( - std::vector{next_clusters_v.data()}, - std::vector{vertex_t{0}}), - detail::edge_endpoint_property_view_t( - std::vector{vertex_cluster_weights_v.data()}, - std::vector{vertex_t{0}}), - detail::edge_endpoint_property_view_t( - std::vector{ - cluster_old_sum_subtract_pair_first}, - std::vector{vertex_t{0}})); + multi_gpu ? view_concat(src_vertex_weights_cache.view(), + src_clusters_cache.view(), + src_cluster_weights->view(), + src_old_cluster_sum_subtract_pairs->view()) + : view_concat( + make_edge_src_property_view( + graph_view, vertex_weights_v.begin(), vertex_weights_v.size()), + make_edge_src_property_view( + graph_view, next_clusters_v.begin(), next_clusters_v.size()), + make_edge_src_property_view( + graph_view, vertex_cluster_weights_v.begin(), vertex_cluster_weights_v.size()), + make_edge_src_property_view>( + graph_view, cluster_old_sum_subtract_pair_first, old_cluster_sum_v.size())); kv_store_t cluster_key_weight_map( cluster_keys_v.begin(), @@ -405,8 +396,8 @@ rmm::device_uvector update_clustering_by_delta_modularity( zipped_src_device_view, *edge_weight_view, multi_gpu ? dst_clusters_cache.view() - : detail::edge_endpoint_property_view_t( - next_clusters_v.data(), vertex_t{0}), + : make_edge_dst_property_view( + graph_view, next_clusters_v.begin(), next_clusters_v.size()), cluster_key_weight_map.view(), detail::key_aggregated_edge_op_t{total_edge_weight, resolution}, thrust::make_tuple(vertex_t{-1}, weight_t{0}), @@ -455,10 +446,9 @@ compute_cluster_keys_and_values( edge_src_dummy_property_t{}.view(), edge_dst_dummy_property_t{}.view(), *edge_weight_view, - multi_gpu - ? src_clusters_cache.view() - : detail::edge_endpoint_property_view_t( - std::vector{next_clusters_v.data()}, std::vector{vertex_t{0}}), + multi_gpu ? src_clusters_cache.view() + : make_edge_src_property_view( + graph_view, next_clusters_v.begin(), next_clusters_v.size()), detail::return_edge_weight_t{}, weight_t{0}, reduce_op::plus{}); diff --git a/cpp/src/community/detail/maximal_independent_moves.cuh b/cpp/src/community/detail/maximal_independent_moves.cuh index cbb71e9c2b9..ff5904c5ac4 100644 --- a/cpp/src/community/detail/maximal_independent_moves.cuh +++ b/cpp/src/community/detail/maximal_independent_moves.cuh @@ -179,12 +179,11 @@ rmm::device_uvector maximal_independent_moves( handle, graph_view, multi_gpu ? src_rank_cache.view() - : detail::edge_endpoint_property_view_t( - std::vector{temporary_ranks.data()}, - std::vector{vertex_t{0}}), + : make_edge_src_property_view( + graph_view, temporary_ranks.begin(), temporary_ranks.size()), multi_gpu ? dst_rank_cache.view() - : detail::edge_endpoint_property_view_t( - temporary_ranks.data(), vertex_t{0}), + : make_edge_dst_property_view( + graph_view, temporary_ranks.begin(), temporary_ranks.size()), edge_dummy_property_t{}.view(), [] __device__(auto src, auto dst, auto src_rank, auto dst_rank, auto wt) { return dst_rank; }, std::numeric_limits::lowest(), @@ -201,12 +200,11 @@ rmm::device_uvector maximal_independent_moves( handle, graph_view, multi_gpu ? src_rank_cache.view() - : detail::edge_endpoint_property_view_t( - std::vector{temporary_ranks.data()}, - std::vector{vertex_t{0}}), + : make_edge_src_property_view( + graph_view, temporary_ranks.begin(), temporary_ranks.size()), multi_gpu ? dst_rank_cache.view() - : detail::edge_endpoint_property_view_t( - temporary_ranks.data(), vertex_t{0}), + : make_edge_dst_property_view( + graph_view, temporary_ranks.begin(), temporary_ranks.size()), edge_dummy_property_t{}.view(), [] __device__(auto src, auto dst, auto src_rank, auto dst_rank, auto wt) { return src_rank; }, std::numeric_limits::lowest(), diff --git a/cpp/src/community/detail/refine_impl.cuh b/cpp/src/community/detail/refine_impl.cuh index 51ca3265739..ad30d6af145 100644 --- a/cpp/src/community/detail/refine_impl.cuh +++ b/cpp/src/community/detail/refine_impl.cuh @@ -28,6 +28,7 @@ #include #include +#include #include #include @@ -213,12 +214,14 @@ refine_clustering( graph_view, GraphViewType::is_multi_gpu ? src_louvain_assignment_cache.view() - : detail::edge_endpoint_property_view_t( - std::vector{louvain_assignment_of_vertices.data()}, - std::vector{vertex_t{0}}), - GraphViewType::is_multi_gpu ? dst_louvain_assignment_cache.view() - : detail::edge_endpoint_property_view_t( - louvain_assignment_of_vertices.data(), vertex_t{0}), + : make_edge_src_property_view(graph_view, + louvain_assignment_of_vertices.begin(), + louvain_assignment_of_vertices.size()), + GraphViewType::is_multi_gpu + ? dst_louvain_assignment_cache.view() + : make_edge_dst_property_view(graph_view, + louvain_assignment_of_vertices.begin(), + louvain_assignment_of_vertices.size()), *edge_weight_view, cuda::proclaim_return_type( [] __device__(auto src, auto dst, auto src_cluster, auto dst_cluster, auto wt) { @@ -355,20 +358,22 @@ refine_clustering( auto src_input_property_values = GraphViewType::is_multi_gpu ? view_concat(src_louvain_assignment_cache.view(), src_leiden_assignment_cache.view()) - : view_concat(detail::edge_endpoint_property_view_t( - std::vector{louvain_assignment_of_vertices.data()}, - std::vector{vertex_t{0}}), - detail::edge_endpoint_property_view_t( - std::vector{leiden_assignment.data()}, - std::vector{vertex_t{0}})); + : view_concat( + make_edge_src_property_view(graph_view, + louvain_assignment_of_vertices.begin(), + louvain_assignment_of_vertices.size()), + make_edge_src_property_view( + graph_view, leiden_assignment.begin(), leiden_assignment.size())); auto dst_input_property_values = GraphViewType::is_multi_gpu ? view_concat(dst_louvain_assignment_cache.view(), dst_leiden_assignment_cache.view()) - : view_concat(detail::edge_endpoint_property_view_t( - louvain_assignment_of_vertices.data(), vertex_t{0}), - detail::edge_endpoint_property_view_t( - leiden_assignment.data(), vertex_t{0})); + : view_concat( + make_edge_dst_property_view(graph_view, + louvain_assignment_of_vertices.begin(), + louvain_assignment_of_vertices.size()), + make_edge_dst_property_view( + graph_view, leiden_assignment.begin(), leiden_assignment.size())); rmm::device_uvector leiden_keys_used_in_edge_reduction(0, handle.get_stream()); rmm::device_uvector refined_community_volumes(0, handle.get_stream()); @@ -391,9 +396,8 @@ refine_clustering( *edge_weight_view, GraphViewType::is_multi_gpu ? dst_leiden_assignment_cache.view() - : detail::edge_endpoint_property_view_t( - leiden_assignment.data(), vertex_t{0}), - + : make_edge_dst_property_view( + graph_view, leiden_assignment.begin(), leiden_assignment.size()), [] __device__(auto src, auto dst, thrust::tuple src_louvain_leidn, @@ -436,24 +440,24 @@ refine_clustering( src_singleton_and_connected_flag_cache.view(), src_leiden_assignment_cache.view(), src_louvain_assignment_cache.view()) - : view_concat(detail::edge_endpoint_property_view_t( - std::vector{weighted_degree_of_vertices.data()}, - std::vector{vertex_t{0}}), - detail::edge_endpoint_property_view_t( - std::vector{weighted_cut_of_vertices_to_louvain.data()}, - std::vector{vertex_t{0}}), - detail::edge_endpoint_property_view_t( - std::vector{vertex_louvain_cluster_weights.data()}, - std::vector{vertex_t{0}}), - detail::edge_endpoint_property_view_t( - std::vector{singleton_and_connected_flags.data()}, - std::vector{vertex_t{0}}), - detail::edge_endpoint_property_view_t( - std::vector{leiden_assignment.data()}, - std::vector{vertex_t{0}}), - detail::edge_endpoint_property_view_t( - std::vector{louvain_assignment_of_vertices.data()}, - std::vector{vertex_t{0}})); + : view_concat( + make_edge_src_property_view( + graph_view, weighted_degree_of_vertices.begin(), weighted_degree_of_vertices.size()), + make_edge_src_property_view( + graph_view, + weighted_cut_of_vertices_to_louvain.begin(), + weighted_cut_of_vertices_to_louvain.size()), + make_edge_src_property_view(graph_view, + vertex_louvain_cluster_weights.begin(), + vertex_louvain_cluster_weights.size()), + make_edge_src_property_view(graph_view, + singleton_and_connected_flags.begin(), + singleton_and_connected_flags.size()), + make_edge_src_property_view( + graph_view, leiden_assignment.begin(), leiden_assignment.size()), + make_edge_src_property_view(graph_view, + louvain_assignment_of_vertices.begin(), + louvain_assignment_of_vertices.size())); rmm::device_uvector louvain_of_leiden_keys_used_in_edge_reduction( 0, handle.get_stream()); @@ -538,8 +542,8 @@ refine_clustering( *edge_weight_view, GraphViewType::is_multi_gpu ? dst_leiden_assignment_cache.view() - : detail::edge_endpoint_property_view_t(leiden_assignment.data(), - vertex_t{0}), + : make_edge_dst_property_view( + graph_view, leiden_assignment.begin(), leiden_assignment.size()), leiden_cluster_key_values_map.view(), detail::leiden_key_aggregated_edge_op_t{ total_edge_weight, resolution, theta, device_state}, diff --git a/cpp/src/components/mis_impl.cuh b/cpp/src/components/mis_impl.cuh index 43563fb6357..d4fab033cb6 100644 --- a/cpp/src/components/mis_impl.cuh +++ b/cpp/src/components/mis_impl.cuh @@ -189,12 +189,11 @@ rmm::device_uvector maximal_independent_set( handle, graph_view, multi_gpu ? src_rank_cache.view() - : detail::edge_endpoint_property_view_t( - std::vector{temporary_ranks.data()}, - std::vector{vertex_t{0}}), + : make_edge_src_property_view( + graph_view, temporary_ranks.begin(), temporary_ranks.size()), multi_gpu ? dst_rank_cache.view() - : detail::edge_endpoint_property_view_t( - temporary_ranks.data(), vertex_t{0}), + : make_edge_dst_property_view( + graph_view, temporary_ranks.begin(), temporary_ranks.size()), edge_dummy_property_t{}.view(), [] __device__(auto src, auto dst, auto src_rank, auto dst_rank, auto wt) { return dst_rank; }, std::numeric_limits::lowest(), @@ -211,12 +210,11 @@ rmm::device_uvector maximal_independent_set( handle, graph_view, multi_gpu ? src_rank_cache.view() - : detail::edge_endpoint_property_view_t( - std::vector{temporary_ranks.data()}, - std::vector{vertex_t{0}}), + : make_edge_src_property_view( + graph_view, temporary_ranks.begin(), temporary_ranks.size()), multi_gpu ? dst_rank_cache.view() - : detail::edge_endpoint_property_view_t( - temporary_ranks.data(), vertex_t{0}), + : make_edge_dst_property_view( + graph_view, temporary_ranks.begin(), temporary_ranks.size()), edge_dummy_property_t{}.view(), [] __device__(auto src, auto dst, auto src_rank, auto dst_rank, auto wt) { return src_rank; }, std::numeric_limits::lowest(), diff --git a/cpp/src/components/vertex_coloring_impl.cuh b/cpp/src/components/vertex_coloring_impl.cuh index 4670c7229d1..31940261878 100644 --- a/cpp/src/components/vertex_coloring_impl.cuh +++ b/cpp/src/components/vertex_coloring_impl.cuh @@ -112,13 +112,14 @@ rmm::device_uvector vertex_coloring( cugraph::transform_e( handle, current_graph_view, - graph_view_t::is_multi_gpu ? src_mis_flags.view() - : detail::edge_endpoint_property_view_t( - std::vector{is_vertex_in_mis.begin()}, - std::vector{vertex_t{0}}), - graph_view_t::is_multi_gpu ? dst_mis_flags.view() - : detail::edge_endpoint_property_view_t( - is_vertex_in_mis.begin(), vertex_t{0}), + graph_view_t::is_multi_gpu + ? src_mis_flags.view() + : make_edge_src_property_view( + current_graph_view, is_vertex_in_mis.begin(), is_vertex_in_mis.size()), + graph_view_t::is_multi_gpu + ? dst_mis_flags.view() + : make_edge_dst_property_view( + current_graph_view, is_vertex_in_mis.begin(), is_vertex_in_mis.size()), cugraph::edge_dummy_property_t{}.view(), [color_id] __device__( auto src, auto dst, auto is_src_in_mis, auto is_dst_in_mis, cuda::std::nullopt_t) { @@ -134,13 +135,14 @@ rmm::device_uvector vertex_coloring( cugraph::transform_e( handle, current_graph_view, - graph_view_t::is_multi_gpu ? src_mis_flags.view() - : detail::edge_endpoint_property_view_t( - std::vector{is_vertex_in_mis.begin()}, - std::vector{vertex_t{0}}), - graph_view_t::is_multi_gpu ? dst_mis_flags.view() - : detail::edge_endpoint_property_view_t( - is_vertex_in_mis.begin(), vertex_t{0}), + graph_view_t::is_multi_gpu + ? src_mis_flags.view() + : make_edge_src_property_view( + current_graph_view, is_vertex_in_mis.begin(), is_vertex_in_mis.size()), + graph_view_t::is_multi_gpu + ? dst_mis_flags.view() + : make_edge_dst_property_view( + current_graph_view, is_vertex_in_mis.begin(), is_vertex_in_mis.size()), cugraph::edge_dummy_property_t{}.view(), [color_id] __device__( auto src, auto dst, auto is_src_in_mis, auto is_dst_in_mis, cuda::std::nullopt_t) { diff --git a/cpp/src/components/weakly_connected_components_impl.cuh b/cpp/src/components/weakly_connected_components_impl.cuh index 505d2ca4564..28b8605d330 100644 --- a/cpp/src/components/weakly_connected_components_impl.cuh +++ b/cpp/src/components/weakly_connected_components_impl.cuh @@ -738,8 +738,10 @@ void weakly_connected_components_impl(raft::handle_t const& handle, ? detail::edge_partition_endpoint_property_device_view_t( edge_dst_components.mutable_view()) : detail::edge_partition_endpoint_property_device_view_t( - detail::edge_endpoint_property_view_t(level_components, - vertex_t{0})), + make_edge_dst_property_mutable_view( + level_graph_view, + level_components, + static_cast(level_graph_view.local_vertex_partition_range_size()))), level_graph_view.local_edge_partition_dst_range_first(), get_dataframe_buffer_begin(edge_buffer), num_edge_inserts.data()}); diff --git a/cpp/src/traversal/sssp_impl.cuh b/cpp/src/traversal/sssp_impl.cuh index b11218665a5..adaec204ffa 100644 --- a/cpp/src/traversal/sssp_impl.cuh +++ b/cpp/src/traversal/sssp_impl.cuh @@ -210,8 +210,8 @@ void sssp(raft::handle_t const& handle, vertex_frontier.bucket(bucket_idx_cur_near), GraphViewType::is_multi_gpu ? edge_src_distances.view() - : detail::edge_endpoint_property_view_t( - std::vector{distances}, std::vector{vertex_t{0}}), + : make_edge_src_property_view( + push_graph_view, distances, push_graph_view.local_vertex_partition_range_size()), edge_dst_dummy_property_t{}.view(), edge_weight_view, e_op_t{},