diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 57e0aa2d078..2527599fece 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -180,6 +180,7 @@ set(CUGRAPH_SOURCES src/community/detail/refine_sg.cu src/community/detail/refine_mg.cu src/community/edge_triangle_count_sg.cu + src/community/edge_triangle_count_mg.cu src/community/detail/maximal_independent_moves_sg.cu src/community/detail/maximal_independent_moves_mg.cu src/detail/utility_wrappers.cu diff --git a/cpp/include/cugraph/algorithms.hpp b/cpp/include/cugraph/algorithms.hpp index 7c4a978c4b4..cc42399f091 100644 --- a/cpp/include/cugraph/algorithms.hpp +++ b/cpp/include/cugraph/algorithms.hpp @@ -2007,6 +2007,24 @@ void triangle_count(raft::handle_t const& handle, raft::device_span counts, bool do_expensive_check = false); +/* + * @brief Compute edge triangle counts. + * + * Compute edge triangle counts for the entire set of edges. + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam edge_t Type of edge identifiers. Needs to be an integral type. + * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param graph_view Graph view object. + * + * @return edge_property_t containing the edge triangle count + */ +template +edge_property_t, edge_t> edge_triangle_count( + raft::handle_t const& handle, graph_view_t const& graph_view); + /* * @brief Compute K-Truss. * diff --git a/cpp/src/community/edge_triangle_count_impl.cuh b/cpp/src/community/edge_triangle_count_impl.cuh index 1370c1a17f2..c4277e240be 100644 --- a/cpp/src/community/edge_triangle_count_impl.cuh +++ b/cpp/src/community/edge_triangle_count_impl.cuh @@ -17,12 +17,17 @@ #pragma once #include "detail/graph_partition_utils.cuh" +#include "prims/edge_bucket.cuh" +#include "prims/transform_e.cuh" #include "prims/transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v.cuh" +#include #include #include #include +#include + #include #include #include @@ -34,8 +39,9 @@ namespace detail { template struct update_edges_p_r_q_r_num_triangles { - size_t num_edges{}; // rename to num_edges + size_t num_edges{}; const edge_t edge_first_or_second{}; + size_t chunk_start{}; raft::device_span intersection_offsets{}; raft::device_span intersection_indices{}; raft::device_span num_triangles{}; @@ -48,28 +54,22 @@ struct update_edges_p_r_q_r_num_triangles { thrust::seq, intersection_offsets.begin() + 1, intersection_offsets.end(), i); auto idx = thrust::distance(intersection_offsets.begin() + 1, itr); if (edge_first_or_second == 0) { - auto p_r_pair = - thrust::make_tuple(thrust::get<0>(*(edge_first + idx)), intersection_indices[i]); + auto p_r_pair = thrust::make_tuple(thrust::get<0>(*(edge_first + chunk_start + idx)), + intersection_indices[i]); // Find its position in 'edges' auto itr_p_r_p_q = - thrust::lower_bound(thrust::seq, - edge_first, - edge_first + num_edges, // pass the number of vertex pairs - p_r_pair); + thrust::lower_bound(thrust::seq, edge_first, edge_first + num_edges, p_r_pair); assert(*itr_p_r_p_q == p_r_pair); idx = thrust::distance(edge_first, itr_p_r_p_q); } else { - auto p_r_pair = - thrust::make_tuple(thrust::get<1>(*(edge_first + idx)), intersection_indices[i]); + auto p_r_pair = thrust::make_tuple(thrust::get<1>(*(edge_first + chunk_start + idx)), + intersection_indices[i]); // Find its position in 'edges' auto itr_p_r_p_q = - thrust::lower_bound(thrust::seq, - edge_first, - edge_first + num_edges, // pass the number of vertex pairs - p_r_pair); + thrust::lower_bound(thrust::seq, edge_first, edge_first + num_edges, p_r_pair); assert(*itr_p_r_p_q == p_r_pair); idx = thrust::distance(edge_first, itr_p_r_p_q); } @@ -78,77 +78,296 @@ struct update_edges_p_r_q_r_num_triangles { } }; +template +struct extract_p_r_q_r { + size_t chunk_start{}; + size_t p_r_or_q_r{}; + raft::device_span intersection_offsets{}; + raft::device_span intersection_indices{}; + EdgeIterator edge_first; + + __device__ thrust::tuple operator()(edge_t i) const + { + auto itr = thrust::upper_bound( + thrust::seq, intersection_offsets.begin() + 1, intersection_offsets.end(), i); + auto idx = thrust::distance(intersection_offsets.begin() + 1, itr); + + if (p_r_or_q_r == 0) { + return thrust::make_tuple(thrust::get<0>(*(edge_first + chunk_start + idx)), + intersection_indices[i]); + } else { + return thrust::make_tuple(thrust::get<1>(*(edge_first + chunk_start + idx)), + intersection_indices[i]); + } + } +}; + +template +struct extract_q_r { + size_t chunk_start{}; + raft::device_span intersection_offsets{}; + raft::device_span intersection_indices{}; + EdgeIterator edge_first; + + __device__ thrust::tuple operator()(edge_t i) const + { + auto itr = thrust::upper_bound( + thrust::seq, intersection_offsets.begin() + 1, intersection_offsets.end(), i); + auto idx = thrust::distance(intersection_offsets.begin() + 1, itr); + auto pair = thrust::make_tuple(thrust::get<1>(*(edge_first + chunk_start + idx)), + intersection_indices[i]); + + return pair; + } +}; + template -std::enable_if_t> edge_triangle_count_impl( +edge_property_t, edge_t> edge_triangle_count_impl( raft::handle_t const& handle, - graph_view_t const& graph_view, - raft::device_span edgelist_srcs, - raft::device_span edgelist_dsts) + graph_view_t const& graph_view) { - auto edge_first = thrust::make_zip_iterator(edgelist_srcs.begin(), edgelist_dsts.begin()); + using weight_t = float; + rmm::device_uvector edgelist_srcs(0, handle.get_stream()); + rmm::device_uvector edgelist_dsts(0, handle.get_stream()); + std::tie(edgelist_srcs, edgelist_dsts, std::ignore, std::ignore, std::ignore) = + decompress_to_edgelist( + handle, graph_view, std::nullopt, std::nullopt, std::nullopt, std::nullopt); - thrust::sort(handle.get_thrust_policy(), edge_first, edge_first + edgelist_srcs.size()); + auto edge_first = thrust::make_zip_iterator(edgelist_srcs.begin(), edgelist_dsts.begin()); - // FIXME: Perform 'nbr_intersection' in chunks to reduce peak memory. - auto [intersection_offsets, intersection_indices] = - detail::nbr_intersection(handle, - graph_view, - cugraph::edge_dummy_property_t{}.view(), - edge_first, - edge_first + edgelist_srcs.size(), - std::array{true, true}, - false /*FIXME: pass 'do_expensive_check' as argument*/); + size_t edges_to_intersect_per_iteration = + static_cast(handle.get_device_properties().multiProcessorCount) * (1 << 17); + auto num_chunks = + raft::div_rounding_up_safe(edgelist_srcs.size(), edges_to_intersect_per_iteration); + size_t prev_chunk_size = 0; + auto num_remaining_edges = edgelist_srcs.size(); rmm::device_uvector num_triangles(edgelist_srcs.size(), handle.get_stream()); - // Update the number of triangles of each (p, q) edges by looking at their intersection - // size - thrust::adjacent_difference(handle.get_thrust_policy(), - intersection_offsets.begin() + 1, - intersection_offsets.end(), - num_triangles.begin()); - - // Given intersection offsets and indices that are used to update the number of - // triangles of (p, q) edges where `r`s are the intersection indices, update - // the number of triangles of the pairs (p, r) and (q, r). - - thrust::for_each( - handle.get_thrust_policy(), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(intersection_indices.size()), - update_edges_p_r_q_r_num_triangles{ - edgelist_srcs.size(), - 0, - raft::device_span(intersection_offsets.data(), intersection_offsets.size()), - raft::device_span(intersection_indices.data(), intersection_indices.size()), - raft::device_span(num_triangles.data(), num_triangles.size()), - edge_first}); - - thrust::for_each( - handle.get_thrust_policy(), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(intersection_indices.size()), - update_edges_p_r_q_r_num_triangles{ - edgelist_srcs.size(), - 1, - raft::device_span(intersection_offsets.data(), intersection_offsets.size()), - raft::device_span(intersection_indices.data(), intersection_indices.size()), - raft::device_span(num_triangles.data(), num_triangles.size()), - edge_first}); - - return num_triangles; + // auto my_rank = handle.get_comms().get_rank(); + if constexpr (multi_gpu) { + num_chunks = host_scalar_allreduce( + handle.get_comms(), num_chunks, raft::comms::op_t::MAX, handle.get_stream()); + } + + // Need to ensure that the vector has its values initialized to 0 before incrementing + thrust::fill(handle.get_thrust_policy(), num_triangles.begin(), num_triangles.end(), 0); + + for (size_t i = 0; i < num_chunks; ++i) { + auto chunk_size = std::min(edges_to_intersect_per_iteration, num_remaining_edges); + num_remaining_edges -= chunk_size; + // Perform 'nbr_intersection' in chunks to reduce peak memory. + auto [intersection_offsets, intersection_indices] = + detail::nbr_intersection(handle, + graph_view, + cugraph::edge_dummy_property_t{}.view(), + edge_first + prev_chunk_size, + edge_first + prev_chunk_size + chunk_size, + std::array{true, true}, + false /*FIXME: pass 'do_expensive_check' as argument*/); + + // Update the number of triangles of each (p, q) edges by looking at their intersection + // size + thrust::for_each( + handle.get_thrust_policy(), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(chunk_size), + [chunk_start = prev_chunk_size, + num_triangles = raft::device_span(num_triangles.data(), num_triangles.size()), + intersection_offsets = raft::device_span( + intersection_offsets.data(), intersection_offsets.size())] __device__(auto i) { + num_triangles[chunk_start + i] += (intersection_offsets[i + 1] - intersection_offsets[i]); + }); + + if constexpr (multi_gpu) { + // stores all the pairs (p, r) and (q, r) + auto vertex_pair_buffer_tmp = allocate_dataframe_buffer>( + intersection_indices.size() * 2, handle.get_stream()); + + // tabulate with the size of intersection_indices, and call binary search on + // intersection_offsets to get (p, r). + thrust::tabulate( + handle.get_thrust_policy(), + get_dataframe_buffer_begin(vertex_pair_buffer_tmp), + get_dataframe_buffer_begin(vertex_pair_buffer_tmp) + intersection_indices.size(), + extract_p_r_q_r{ + prev_chunk_size, + 0, + raft::device_span(intersection_offsets.data(), intersection_offsets.size()), + raft::device_span(intersection_indices.data(), + intersection_indices.size()), + edge_first}); + // FIXME: Consolidate both functions + thrust::tabulate( + handle.get_thrust_policy(), + get_dataframe_buffer_begin(vertex_pair_buffer_tmp) + intersection_indices.size(), + get_dataframe_buffer_begin(vertex_pair_buffer_tmp) + (2 * intersection_indices.size()), + extract_p_r_q_r{ + prev_chunk_size, + 1, + raft::device_span(intersection_offsets.data(), intersection_offsets.size()), + raft::device_span(intersection_indices.data(), + intersection_indices.size()), + edge_first}); + + thrust::sort(handle.get_thrust_policy(), + get_dataframe_buffer_begin(vertex_pair_buffer_tmp), + get_dataframe_buffer_end(vertex_pair_buffer_tmp)); + + rmm::device_uvector increase_count_tmp(2 * intersection_indices.size(), + handle.get_stream()); + thrust::fill(handle.get_thrust_policy(), + increase_count_tmp.begin(), + increase_count_tmp.end(), + size_t{1}); + + auto count_p_r_q_r = thrust::unique_count(handle.get_thrust_policy(), + get_dataframe_buffer_begin(vertex_pair_buffer_tmp), + get_dataframe_buffer_end(vertex_pair_buffer_tmp)); + + rmm::device_uvector increase_count(count_p_r_q_r, handle.get_stream()); + + auto vertex_pair_buffer = allocate_dataframe_buffer>( + count_p_r_q_r, handle.get_stream()); + thrust::reduce_by_key(handle.get_thrust_policy(), + get_dataframe_buffer_begin(vertex_pair_buffer_tmp), + get_dataframe_buffer_end(vertex_pair_buffer_tmp), + increase_count_tmp.begin(), + get_dataframe_buffer_begin(vertex_pair_buffer), + increase_count.begin(), + thrust::equal_to>{}); + + rmm::device_uvector pair_srcs(0, handle.get_stream()); + rmm::device_uvector pair_dsts(0, handle.get_stream()); + std::optional> pair_count{std::nullopt}; + + std::optional> opt_increase_count = + std::make_optional(rmm::device_uvector(increase_count.size(), handle.get_stream())); + + raft::copy((*opt_increase_count).begin(), + increase_count.begin(), + increase_count.size(), + handle.get_stream()); + + // There are still multiple copies here but is it worth sorting and reducing again? + std::tie(pair_srcs, pair_dsts, std::ignore, pair_count, std::ignore) = + shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( + handle, + std::move(std::get<0>(vertex_pair_buffer)), + std::move(std::get<1>(vertex_pair_buffer)), + std::nullopt, + // FIXME: Add general purpose function for shuffling vertex pairs and arbitrary attributes + std::move(opt_increase_count), + std::nullopt, + graph_view.vertex_partition_range_lasts()); + + thrust::for_each( + handle.get_thrust_policy(), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(pair_srcs.size()), + [num_edges = edgelist_srcs.size(), + num_triangles = num_triangles.data(), + pair_srcs = pair_srcs.data(), + pair_dsts = pair_dsts.data(), + pair_count = (*pair_count).data(), + edge_first] __device__(auto idx) { + auto src = pair_srcs[idx]; + auto dst = pair_dsts[idx]; + auto p_r_q_r_pair = thrust::make_tuple(src, dst); + + // Find its position in 'edges' + auto itr_p_r_q_r = + thrust::lower_bound(thrust::seq, edge_first, edge_first + num_edges, p_r_q_r_pair); + + assert(*itr_p_r_q_r == p_r_q_r_pair); + auto idx_p_r_q_r = thrust::distance(edge_first, itr_p_r_q_r); + + cuda::atomic_ref atomic_counter( + num_triangles[idx_p_r_q_r]); + auto r = atomic_counter.fetch_add(pair_count[idx], cuda::std::memory_order_relaxed); + }); + + } else { + // Given intersection offsets and indices that are used to update the number of + // triangles of (p, q) edges where `r`s are the intersection indices, update + // the number of triangles of the pairs (p, r) and (q, r). + thrust::for_each( + handle.get_thrust_policy(), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(intersection_indices.size()), + update_edges_p_r_q_r_num_triangles{ + edgelist_srcs.size(), + 0, + prev_chunk_size, + raft::device_span(intersection_offsets.data(), intersection_offsets.size()), + raft::device_span(intersection_indices.data(), + intersection_indices.size()), + raft::device_span(num_triangles.data(), num_triangles.size()), + edge_first}); + + thrust::for_each( + handle.get_thrust_policy(), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(intersection_indices.size()), + update_edges_p_r_q_r_num_triangles{ + edgelist_srcs.size(), + 1, + prev_chunk_size, + raft::device_span(intersection_offsets.data(), intersection_offsets.size()), + raft::device_span(intersection_indices.data(), + intersection_indices.size()), + raft::device_span(num_triangles.data(), num_triangles.size()), + edge_first}); + } + prev_chunk_size += chunk_size; + } + + cugraph::edge_property_t, edge_t> counts( + handle, graph_view); + + cugraph::edge_bucket_t valid_edges(handle); + valid_edges.insert(edgelist_srcs.begin(), edgelist_srcs.end(), edgelist_dsts.begin()); + + auto cur_graph_view = graph_view; + + cugraph::transform_e( + handle, + graph_view, + valid_edges, + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + cugraph::edge_dummy_property_t{}.view(), + [edge_first, + edge_last = edge_first + edgelist_srcs.size(), + num_edges = edgelist_srcs.size(), + num_triangles = num_triangles.data()] __device__(auto src, + auto dst, + thrust::nullopt_t, + thrust::nullopt_t, + thrust::nullopt_t) { + auto pair = thrust::make_tuple(src, dst); + + // Find its position in 'edges' + auto itr_pair = thrust::lower_bound(thrust::seq, edge_first, edge_last, pair); + auto idx_pair = thrust::distance(edge_first, itr_pair); + return num_triangles[idx_pair]; + }, + counts.mutable_view(), + false); + + return counts; } } // namespace detail -template -rmm::device_uvector edge_triangle_count( - raft::handle_t const& handle, - graph_view_t const& graph_view, - raft::device_span edgelist_srcs, - raft::device_span edgelist_dsts) +template +edge_property_t, edge_t> edge_triangle_count( + raft::handle_t const& handle, graph_view_t const& graph_view) { - return detail::edge_triangle_count_impl(handle, graph_view, edgelist_srcs, edgelist_dsts); + return detail::edge_triangle_count_impl(handle, graph_view); } } // namespace cugraph diff --git a/cpp/src/community/edge_triangle_count_mg.cu b/cpp/src/community/edge_triangle_count_mg.cu new file mode 100644 index 00000000000..254a0807e56 --- /dev/null +++ b/cpp/src/community/edge_triangle_count_mg.cu @@ -0,0 +1,33 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include "community/edge_triangle_count_impl.cuh" + +namespace cugraph { + +// SG instantiation +template edge_property_t, int32_t> edge_triangle_count( + raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view); + +template edge_property_t, int64_t> edge_triangle_count( + raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view); + +template edge_property_t, int64_t> edge_triangle_count( + raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view); + +} // namespace cugraph diff --git a/cpp/src/community/edge_triangle_count_sg.cu b/cpp/src/community/edge_triangle_count_sg.cu index c4b7e71b967..4ccb968458d 100644 --- a/cpp/src/community/edge_triangle_count_sg.cu +++ b/cpp/src/community/edge_triangle_count_sg.cu @@ -18,22 +18,16 @@ namespace cugraph { // SG instantiation -template rmm::device_uvector edge_triangle_count( +template edge_property_t, int32_t> edge_triangle_count( raft::handle_t const& handle, - cugraph::graph_view_t const& graph_view, - raft::device_span edgelist_srcs, - raft::device_span edgelist_dsts); + cugraph::graph_view_t const& graph_view); -template rmm::device_uvector edge_triangle_count( +template edge_property_t, int64_t> edge_triangle_count( raft::handle_t const& handle, - cugraph::graph_view_t const& graph_view, - raft::device_span edgelist_srcs, - raft::device_span edgelist_dsts); + cugraph::graph_view_t const& graph_view); -template rmm::device_uvector edge_triangle_count( +template edge_property_t, int64_t> edge_triangle_count( raft::handle_t const& handle, - cugraph::graph_view_t const& graph_view, - raft::device_span edgelist_srcs, - raft::device_span edgelist_dsts); + cugraph::graph_view_t const& graph_view); } // namespace cugraph diff --git a/cpp/src/community/k_truss_impl.cuh b/cpp/src/community/k_truss_impl.cuh index 7f96312703d..f830e6a7700 100644 --- a/cpp/src/community/k_truss_impl.cuh +++ b/cpp/src/community/k_truss_impl.cuh @@ -27,6 +27,8 @@ #include #include +#include + #include #include #include @@ -39,14 +41,6 @@ namespace cugraph { -// FIXME : This will be deleted once edge_triangle_count becomes public -template -rmm::device_uvector edge_triangle_count( - raft::handle_t const& handle, - graph_view_t const& graph_view, - raft::device_span edgelist_srcs, - raft::device_span edgelist_dsts); - template struct unroll_edge { size_t num_valid_edges{}; @@ -442,6 +436,7 @@ struct extract_low_to_high_degree_edges_t { template struct generate_p_r_or_q_r_from_p_q { + size_t chunk_start{}; raft::device_span intersection_offsets{}; raft::device_span intersection_indices{}; raft::device_span invalid_srcs{}; @@ -454,10 +449,10 @@ struct generate_p_r_or_q_r_from_p_q { auto idx = thrust::distance(intersection_offsets.begin() + 1, itr); if constexpr (generate_p_r) { - return thrust::make_tuple(invalid_srcs[idx], intersection_indices[i]); + return thrust::make_tuple(invalid_srcs[chunk_start + idx], intersection_indices[i]); } else { - return thrust::make_tuple(invalid_dsts[idx], intersection_indices[i]); + return thrust::make_tuple(invalid_dsts[chunk_start + idx], intersection_indices[i]); } } }; @@ -491,6 +486,7 @@ k_truss(raft::handle_t const& handle, std::optional> renumber_map{std::nullopt}; std::optional, weight_t>> edge_weight{std::nullopt}; + std::optional> wgts{std::nullopt}; if (graph_view.count_self_loops(handle) > edge_t{0}) { auto [srcs, dsts] = extract_transform_e(handle, @@ -524,31 +520,30 @@ k_truss(raft::handle_t const& handle, modified_graph_view = (*modified_graph).view(); } - // FIXME: Investigate k-1 core failure to yield correct results. // 3. Find (k-1)-core and exclude edges that do not belong to (k-1)-core - /* { auto cur_graph_view = modified_graph_view ? *modified_graph_view : graph_view; + auto vertex_partition_range_lasts = renumber_map ? std::make_optional>(cur_graph_view.vertex_partition_range_lasts()) : std::nullopt; - rmm::device_uvector d_core_numbers(cur_graph_view.local_vertex_partition_range_size(), - handle.get_stream()); - raft::device_span core_number_span{d_core_numbers.data(), d_core_numbers.size()}; + rmm::device_uvector core_numbers(cur_graph_view.number_of_vertices(), + handle.get_stream()); + core_number( + handle, cur_graph_view, core_numbers.data(), k_core_degree_type_t::OUT, size_t{2}, size_t{2}); + + raft::device_span core_number_span{core_numbers.data(), core_numbers.size()}; rmm::device_uvector srcs{0, handle.get_stream()}; rmm::device_uvector dsts{0, handle.get_stream()}; - std::tie(srcs, dsts, std::ignore) = - k_core(handle, - cur_graph_view, - std::optional>{std::nullopt}, - size_t{k - 1}, - std::make_optional(k_core_degree_type_t::OUT), - // Seems like the below argument is required. passing a std::nullopt - // create a compiler error - std::make_optional(core_number_span)); + std::tie(srcs, dsts, wgts) = k_core(handle, + cur_graph_view, + edge_weight_view, + k - 1, + std::make_optional(k_core_degree_type_t::OUT), + std::make_optional(core_number_span)); if constexpr (multi_gpu) { std::tie(srcs, dsts, std::ignore, std::ignore, std::ignore) = @@ -561,17 +556,17 @@ k_truss(raft::handle_t const& handle, std::optional> tmp_renumber_map{std::nullopt}; - std::tie(*modified_graph, std::ignore, std::ignore, std::ignore, tmp_renumber_map) = + std::tie(*modified_graph, edge_weight, std::ignore, std::ignore, tmp_renumber_map) = create_graph_from_edgelist( handle, std::nullopt, std::move(srcs), std::move(dsts), - std::nullopt, + std::move(wgts), std::nullopt, std::nullopt, cugraph::graph_properties_t{true, graph_view.is_multigraph()}, - true); + false); modified_graph_view = (*modified_graph).view(); @@ -584,7 +579,6 @@ k_truss(raft::handle_t const& handle, } renumber_map = std::move(tmp_renumber_map); } - */ // 4. Keep only the edges from a low-degree vertex to a high-degree vertex. @@ -606,7 +600,10 @@ k_truss(raft::handle_t const& handle, rmm::device_uvector srcs(0, handle.get_stream()); rmm::device_uvector dsts(0, handle.get_stream()); - std::optional> wgts{std::nullopt}; + + edge_weight_view = + edge_weight ? std::make_optional((*edge_weight).view()) + : std::optional>{std::nullopt}; if (edge_weight_view) { std::tie(srcs, dsts, wgts) = extract_transform_e( handle, @@ -666,38 +663,36 @@ k_truss(raft::handle_t const& handle, auto cur_graph_view = modified_graph_view ? *modified_graph_view : graph_view; rmm::device_uvector edgelist_srcs(0, handle.get_stream()); rmm::device_uvector edgelist_dsts(0, handle.get_stream()); + std::optional> num_triangles{std::nullopt}; std::optional> edgelist_wgts{std::nullopt}; edge_weight_view = edge_weight ? std::make_optional((*edge_weight).view()) : std::optional>{std::nullopt}; - std::tie(edgelist_srcs, edgelist_dsts, edgelist_wgts, std::ignore, std::ignore) = + + auto prop_num_triangles = edge_triangle_count(handle, cur_graph_view); + + std::tie(edgelist_srcs, edgelist_dsts, edgelist_wgts, num_triangles, std::ignore) = decompress_to_edgelist( handle, cur_graph_view, edge_weight_view, - std::optional>{std::nullopt}, + // FIXME: Update 'decompress_edgelist' to support int32_t and int64_t values + std::make_optional(prop_num_triangles.view()), std::optional>{std::nullopt}, std::optional>(std::nullopt)); - - auto num_triangles = edge_triangle_count( - handle, - cur_graph_view, - raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), - raft::device_span(edgelist_dsts.data(), edgelist_dsts.size())); - auto transposed_edge_first = thrust::make_zip_iterator(edgelist_dsts.begin(), edgelist_srcs.begin()); auto edge_first = thrust::make_zip_iterator(edgelist_srcs.begin(), edgelist_dsts.begin()); auto transposed_edge_triangle_count_pair_first = - thrust::make_zip_iterator(transposed_edge_first, num_triangles.begin()); + thrust::make_zip_iterator(transposed_edge_first, (*num_triangles).begin()); thrust::sort_by_key(handle.get_thrust_policy(), transposed_edge_first, transposed_edge_first + edgelist_srcs.size(), - num_triangles.begin()); + (*num_triangles).begin()); cugraph::edge_property_t edge_mask(handle, cur_graph_view); cugraph::fill_edge_property(handle, cur_graph_view, true, edge_mask); @@ -728,92 +723,115 @@ k_truss(raft::handle_t const& handle, // nbr_intersection requires the edges to be sort by 'src' // sort the invalid edges by src for nbr intersection - thrust::sort_by_key(handle.get_thrust_policy(), - edge_first + num_valid_edges, - edge_first + edgelist_srcs.size(), - num_triangles.begin() + num_valid_edges); - - auto [intersection_offsets, intersection_indices] = - detail::nbr_intersection(handle, - cur_graph_view, - cugraph::edge_dummy_property_t{}.view(), - edge_first + num_valid_edges, - edge_first + edgelist_srcs.size(), - std::array{true, true}, - do_expensive_check); - - // Update the number of triangles of each (p, q) edges by looking at their intersection - // size. - thrust::for_each( - handle.get_thrust_policy(), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(num_invalid_edges), - [num_triangles = - raft::device_span(num_triangles.data() + num_valid_edges, num_invalid_edges), - intersection_offsets = raft::device_span( - intersection_offsets.data(), intersection_offsets.size())] __device__(auto i) { - num_triangles[i] -= intersection_offsets[i + 1] - intersection_offsets[i]; - }); - - // FIXME: Find a way to not have to maintain a dataframe_buffer - auto vertex_pair_buffer_p_r_edge_p_q = - allocate_dataframe_buffer>(intersection_indices.size(), - handle.get_stream()); - - thrust::tabulate( - handle.get_thrust_policy(), - get_dataframe_buffer_begin(vertex_pair_buffer_p_r_edge_p_q), - get_dataframe_buffer_end(vertex_pair_buffer_p_r_edge_p_q), - generate_p_r_or_q_r_from_p_q{ - raft::device_span(intersection_offsets.data(), intersection_offsets.size()), - raft::device_span(intersection_indices.data(), - intersection_indices.size()), - raft::device_span(edgelist_srcs.data() + num_valid_edges, num_invalid_edges), - raft::device_span(edgelist_dsts.data() + num_valid_edges, num_invalid_edges)}); - - auto vertex_pair_buffer_q_r_edge_p_q = - allocate_dataframe_buffer>(intersection_indices.size(), - handle.get_stream()); - thrust::tabulate( - handle.get_thrust_policy(), - get_dataframe_buffer_begin(vertex_pair_buffer_q_r_edge_p_q), - get_dataframe_buffer_end(vertex_pair_buffer_q_r_edge_p_q), - generate_p_r_or_q_r_from_p_q{ - raft::device_span(intersection_offsets.data(), intersection_offsets.size()), - raft::device_span(intersection_indices.data(), - intersection_indices.size()), - raft::device_span(edgelist_srcs.data() + num_valid_edges, num_invalid_edges), - raft::device_span(edgelist_dsts.data() + num_valid_edges, num_invalid_edges)}); - - // Unrolling the edges require the edges to be sorted by destination - // re-sort the invalid edges by 'dst' - thrust::sort_by_key(handle.get_thrust_policy(), - transposed_edge_first + num_valid_edges, - transposed_edge_first + edgelist_srcs.size(), - num_triangles.begin() + num_valid_edges); - - thrust::for_each(handle.get_thrust_policy(), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(intersection_indices.size()), - unroll_edge{ - num_valid_edges, - raft::device_span(num_triangles.data(), num_triangles.size()), - get_dataframe_buffer_begin(vertex_pair_buffer_p_r_edge_p_q), - transposed_edge_first, - transposed_edge_first + num_valid_edges, - transposed_edge_first + edgelist_srcs.size()}); - - thrust::for_each(handle.get_thrust_policy(), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(intersection_indices.size()), - unroll_edge{ - num_valid_edges, - raft::device_span(num_triangles.data(), num_triangles.size()), - get_dataframe_buffer_begin(vertex_pair_buffer_q_r_edge_p_q), - transposed_edge_first, - transposed_edge_first + num_valid_edges, - transposed_edge_first + edgelist_srcs.size()}); - + size_t edges_to_intersect_per_iteration = + static_cast(handle.get_device_properties().multiProcessorCount) * (1 << 17); + + size_t prev_chunk_size = 0; + size_t chunk_num_invalid_edges = num_invalid_edges; + + auto num_chunks = + raft::div_rounding_up_safe(edgelist_srcs.size(), edges_to_intersect_per_iteration); + + for (size_t i = 0; i < num_chunks; ++i) { + auto chunk_size = std::min(edges_to_intersect_per_iteration, chunk_num_invalid_edges); + thrust::sort_by_key(handle.get_thrust_policy(), + edge_first + num_valid_edges, + edge_first + edgelist_srcs.size(), + (*num_triangles).begin() + num_valid_edges); + + auto [intersection_offsets, intersection_indices] = + detail::nbr_intersection(handle, + cur_graph_view, + cugraph::edge_dummy_property_t{}.view(), + edge_first + num_valid_edges + prev_chunk_size, + edge_first + num_valid_edges + prev_chunk_size + chunk_size, + std::array{true, true}, + do_expensive_check); + + // Update the number of triangles of each (p, q) edges by looking at their intersection + // size. + thrust::for_each( + handle.get_thrust_policy(), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(chunk_size), + [chunk_start = prev_chunk_size, + num_triangles = raft::device_span((*num_triangles).data() + num_valid_edges, + num_invalid_edges), + intersection_offsets = raft::device_span( + intersection_offsets.data(), intersection_offsets.size())] __device__(auto i) { + num_triangles[chunk_start + i] -= + (intersection_offsets[i + 1] - intersection_offsets[i]); + }); + + // FIXME: Find a way to not have to maintain a dataframe_buffer + auto vertex_pair_buffer_p_r_edge_p_q = + allocate_dataframe_buffer>(intersection_indices.size(), + handle.get_stream()); + thrust::tabulate( + handle.get_thrust_policy(), + get_dataframe_buffer_begin(vertex_pair_buffer_p_r_edge_p_q), + get_dataframe_buffer_end(vertex_pair_buffer_p_r_edge_p_q), + generate_p_r_or_q_r_from_p_q{ + prev_chunk_size, + raft::device_span(intersection_offsets.data(), + intersection_offsets.size()), + raft::device_span(intersection_indices.data(), + intersection_indices.size()), + raft::device_span(edgelist_srcs.data() + num_valid_edges, num_invalid_edges), + raft::device_span(edgelist_dsts.data() + num_valid_edges, + num_invalid_edges)}); + + auto vertex_pair_buffer_q_r_edge_p_q = + allocate_dataframe_buffer>(intersection_indices.size(), + handle.get_stream()); + thrust::tabulate( + handle.get_thrust_policy(), + get_dataframe_buffer_begin(vertex_pair_buffer_q_r_edge_p_q), + get_dataframe_buffer_end(vertex_pair_buffer_q_r_edge_p_q), + generate_p_r_or_q_r_from_p_q{ + prev_chunk_size, + raft::device_span(intersection_offsets.data(), + intersection_offsets.size()), + raft::device_span(intersection_indices.data(), + intersection_indices.size()), + raft::device_span(edgelist_srcs.data() + num_valid_edges, num_invalid_edges), + raft::device_span(edgelist_dsts.data() + num_valid_edges, + num_invalid_edges)}); + + // Unrolling the edges require the edges to be sorted by destination + // re-sort the invalid edges by 'dst' + thrust::sort_by_key(handle.get_thrust_policy(), + transposed_edge_first + num_valid_edges, + transposed_edge_first + edgelist_srcs.size(), + (*num_triangles).begin() + num_valid_edges); + + thrust::for_each( + handle.get_thrust_policy(), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(intersection_indices.size()), + unroll_edge{ + num_valid_edges, + raft::device_span((*num_triangles).data(), (*num_triangles).size()), + get_dataframe_buffer_begin(vertex_pair_buffer_p_r_edge_p_q), + transposed_edge_first, + transposed_edge_first + num_valid_edges, + transposed_edge_first + edgelist_srcs.size()}); + + thrust::for_each( + handle.get_thrust_policy(), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(intersection_indices.size()), + unroll_edge{ + num_valid_edges, + raft::device_span((*num_triangles).data(), (*num_triangles).size()), + get_dataframe_buffer_begin(vertex_pair_buffer_q_r_edge_p_q), + transposed_edge_first, + transposed_edge_first + num_valid_edges, + transposed_edge_first + edgelist_srcs.size()}); + + prev_chunk_size += chunk_size; + chunk_num_invalid_edges -= chunk_size; + } // case 2: unroll (q, r) // For each (q, r) edges to unroll, find the incoming edges to 'r' let's say from 'p' and // create the pair (p, q) @@ -824,7 +842,7 @@ k_truss(raft::handle_t const& handle, num_valid_edges, raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), raft::device_span(edgelist_dsts.data(), edgelist_dsts.size()), - raft::device_span(num_triangles.data(), num_triangles.size())); + raft::device_span((*num_triangles).data(), (*num_triangles).size())); // case 3: unroll (p, r) cugraph::unroll_p_r_or_q_r_edges( @@ -834,18 +852,18 @@ k_truss(raft::handle_t const& handle, num_valid_edges, raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), raft::device_span(edgelist_dsts.data(), edgelist_dsts.size()), - raft::device_span(num_triangles.data(), num_triangles.size())); + raft::device_span((*num_triangles).data(), (*num_triangles).size())); // Remove edges that have a triangle count of zero. Those should not be accounted // for during the unroling phase. - auto edges_with_triangle_last = - thrust::stable_partition(handle.get_thrust_policy(), - transposed_edge_triangle_count_pair_first, - transposed_edge_triangle_count_pair_first + num_triangles.size(), - [] __device__(auto e) { - auto num_triangles = thrust::get<1>(e); - return num_triangles > 0; - }); + auto edges_with_triangle_last = thrust::stable_partition( + handle.get_thrust_policy(), + transposed_edge_triangle_count_pair_first, + transposed_edge_triangle_count_pair_first + (*num_triangles).size(), + [] __device__(auto e) { + auto num_triangles = thrust::get<1>(e); + return num_triangles > 0; + }); auto num_edges_with_triangles = static_cast( thrust::distance(transposed_edge_triangle_count_pair_first, edges_with_triangle_last)); @@ -893,7 +911,7 @@ k_truss(raft::handle_t const& handle, edgelist_srcs.resize(num_edges_with_triangles, handle.get_stream()); edgelist_dsts.resize(num_edges_with_triangles, handle.get_stream()); - num_triangles.resize(num_edges_with_triangles, handle.get_stream()); + (*num_triangles).resize(num_edges_with_triangles, handle.get_stream()); } std::tie(edgelist_srcs, edgelist_dsts, edgelist_wgts, std::ignore, std::ignore) = diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index ced3b7bedb1..d1dd2dec069 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -490,6 +490,11 @@ ConfigureTest(K_TRUSS_TEST community/k_truss_test.cpp) # - Triangle Count tests -------------------------------------------------------------------------- ConfigureTest(TRIANGLE_COUNT_TEST community/triangle_count_test.cpp) +################################################################################################### +# - Edge Triangle Count tests -------------------------------------------------------------------------- +ConfigureTest(EDGE_TRIANGLE_COUNT_TEST community/edge_triangle_count_test.cpp) + + ################################################################################################### # - K-hop Neighbors tests ------------------------------------------------------------------------- ConfigureTest(K_HOP_NBRS_TEST traversal/k_hop_nbrs_test.cpp) @@ -590,6 +595,10 @@ if(BUILD_CUGRAPH_MG_TESTS) # - MG LOUVAIN tests -------------------------------------------------------------------------- ConfigureTestMG(MG_EGONET_TEST community/mg_egonet_test.cu) + ############################################################################################### + # - MG EDGE TRIANGLE COUNT tests -------------------------------------------------------------------------- + ConfigureTestMG(MG_EDGE_TRIANGLE_COUNT_TEST community/mg_edge_triangle_count_test.cpp) + ############################################################################################### # - MG WEAKLY CONNECTED COMPONENTS tests ------------------------------------------------------ ConfigureTestMG(MG_WEAKLY_CONNECTED_COMPONENTS_TEST diff --git a/cpp/tests/centrality/mg_betweenness_centrality_test.cpp b/cpp/tests/centrality/mg_betweenness_centrality_test.cpp index 7924d449897..798e767085e 100644 --- a/cpp/tests/centrality/mg_betweenness_centrality_test.cpp +++ b/cpp/tests/centrality/mg_betweenness_centrality_test.cpp @@ -152,13 +152,15 @@ class Tests_MGBetweennessCentrality std::optional< cugraph::edge_property_t, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = - cugraph::test::mg_graph_to_sg_graph(*handle_, - mg_graph_view, - mg_edge_weight_view, - std::make_optional>( - (*mg_renumber_map).data(), (*mg_renumber_map).size()), - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/centrality/mg_edge_betweenness_centrality_test.cpp b/cpp/tests/centrality/mg_edge_betweenness_centrality_test.cpp index c3417e96c03..1703f198a4c 100644 --- a/cpp/tests/centrality/mg_edge_betweenness_centrality_test.cpp +++ b/cpp/tests/centrality/mg_edge_betweenness_centrality_test.cpp @@ -142,12 +142,14 @@ class Tests_MGEdgeBetweennessCentrality std::optional< cugraph::edge_property_t, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - mg_edge_weight_view, - std::optional>{std::nullopt}, - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + false); if (handle_->get_comms().get_rank() == 0) { auto sg_edge_weights_view = diff --git a/cpp/tests/centrality/mg_eigenvector_centrality_test.cpp b/cpp/tests/centrality/mg_eigenvector_centrality_test.cpp index ed24bee0923..76c52d52bfd 100644 --- a/cpp/tests/centrality/mg_eigenvector_centrality_test.cpp +++ b/cpp/tests/centrality/mg_eigenvector_centrality_test.cpp @@ -144,13 +144,15 @@ class Tests_MGEigenvectorCentrality std::optional< cugraph::edge_property_t, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = - cugraph::test::mg_graph_to_sg_graph(*handle_, - mg_graph_view, - mg_edge_weight_view, - std::make_optional>( - (*mg_renumber_map).data(), (*mg_renumber_map).size()), - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == int{0}) { // 3-2. run SG Eigenvector Centrality diff --git a/cpp/tests/centrality/mg_katz_centrality_test.cpp b/cpp/tests/centrality/mg_katz_centrality_test.cpp index abe02b2287b..e38f87749b8 100644 --- a/cpp/tests/centrality/mg_katz_centrality_test.cpp +++ b/cpp/tests/centrality/mg_katz_centrality_test.cpp @@ -151,13 +151,15 @@ class Tests_MGKatzCentrality std::optional< cugraph::edge_property_t, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = - cugraph::test::mg_graph_to_sg_graph(*handle_, - mg_graph_view, - mg_edge_weight_view, - std::make_optional>( - (*mg_renumber_map).data(), (*mg_renumber_map).size()), - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == int{0}) { // 4-2. run SG Katz Centrality diff --git a/cpp/tests/community/edge_triangle_count_test.cpp b/cpp/tests/community/edge_triangle_count_test.cpp new file mode 100644 index 00000000000..8cefc2c31f4 --- /dev/null +++ b/cpp/tests/community/edge_triangle_count_test.cpp @@ -0,0 +1,260 @@ +/* + * Copyright (c) 2022-2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governin_from_mtxg permissions and + * limitations under the License. + */ + +#include "utilities/base_fixture.hpp" +#include "utilities/check_utilities.hpp" +#include "utilities/conversion_utilities.hpp" +#include "utilities/property_generator_utilities.hpp" +#include "utilities/test_graphs.hpp" +#include "utilities/thrust_wrapper.hpp" + +#include +#include +#include +#include +#include + +#include +#include + +#include + +#include +#include +#include +#include +#include +#include +#include + +struct EdgeTriangleCount_Usecase { + bool edge_masking_{false}; + bool check_correctness_{true}; +}; + +template +class Tests_EdgeTriangleCount + : public ::testing::TestWithParam> { + public: + Tests_EdgeTriangleCount() {} + + static void SetUpTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + // FIXME: There is an utility equivalent functor not + // supporting host vectors. + template + struct host_nearly_equal { + const type_t threshold_ratio; + const type_t threshold_magnitude; + + bool operator()(type_t lhs, type_t rhs) const + { + return std::abs(lhs - rhs) < + std::max(std::max(lhs, rhs) * threshold_ratio, threshold_magnitude); + } + }; + + template + std::vector edge_triangle_count_reference(std::vector h_srcs, + std::vector h_dsts) + { + std::vector edge_triangle_counts(h_srcs.size()); + std::uninitialized_fill(edge_triangle_counts.begin(), edge_triangle_counts.end(), 0); + + for (int i = 0; i < h_srcs.size(); ++i) { // edge centric implementation + // for each edge, find the intersection + auto src = h_srcs[i]; + auto dst = h_dsts[i]; + auto it_src_start = std::lower_bound(h_srcs.begin(), h_srcs.end(), src); + auto src_start = std::distance(h_srcs.begin(), it_src_start); + + auto src_end = + src_start + std::distance(it_src_start, std::upper_bound(it_src_start, h_srcs.end(), src)); + + auto it_dst_start = std::lower_bound(h_srcs.begin(), h_srcs.end(), dst); + auto dst_start = std::distance(h_srcs.begin(), it_dst_start); + auto dst_end = + dst_start + std::distance(it_dst_start, std::upper_bound(it_dst_start, h_srcs.end(), dst)); + + std::set nbr_intersection; + std::set_intersection(h_dsts.begin() + src_start, + h_dsts.begin() + src_end, + h_dsts.begin() + dst_start, + h_dsts.begin() + dst_end, + std::inserter(nbr_intersection, nbr_intersection.end())); + // Find the supporting edges + for (auto v : nbr_intersection) { + auto it_edge = std::lower_bound(h_dsts.begin() + src_start, h_dsts.begin() + src_end, v); + auto idx_edge = std::distance(h_dsts.begin(), it_edge); + edge_triangle_counts[idx_edge] += 1; + + it_edge = std::lower_bound(h_dsts.begin() + dst_start, h_dsts.begin() + dst_end, v); + idx_edge = std::distance(h_dsts.begin(), it_edge); + } + } + + std::transform(edge_triangle_counts.begin(), + edge_triangle_counts.end(), + edge_triangle_counts.begin(), + [](auto count) { return count * 3; }); + return std::move(edge_triangle_counts); + } + + template + void run_current_test( + std::tuple const& param) + { + constexpr bool renumber = false; + auto [edge_triangle_count_usecase, input_usecase] = param; + raft::handle_t handle{}; + HighResTimer hr_timer{}; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_timer.start("SG Construct graph"); + } + + auto [graph, edge_weight, d_renumber_map_labels] = + cugraph::test::construct_graph( + handle, input_usecase, false, renumber, true, true); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + auto graph_view = graph.view(); + + std::optional> edge_mask{std::nullopt}; + if (edge_triangle_count_usecase.edge_masking_) { + edge_mask = + cugraph::test::generate::edge_property(handle, graph_view, 2); + graph_view.attach_edge_mask((*edge_mask).view()); + } + + rmm::device_uvector edgelist_srcs(0, handle.get_stream()); + rmm::device_uvector edgelist_dsts(0, handle.get_stream()); + std::optional> d_edge_triangle_counts{std::nullopt}; + + auto d_cugraph_results = + cugraph::edge_triangle_count(handle, graph_view); + + std::tie(edgelist_srcs, edgelist_dsts, std::ignore, d_edge_triangle_counts, std::ignore) = + cugraph::decompress_to_edgelist( + handle, + graph_view, + std::optional>{std::nullopt}, + std::make_optional(d_cugraph_results.view()), + std::optional>{std::nullopt}, + std::optional>{std::nullopt}); // FIXME: No longer needed + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_timer.start("EdgeTriangleCount"); + } + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + if (edge_triangle_count_usecase.check_correctness_) { + std::optional> modified_graph{std::nullopt}; + std::vector h_srcs(edgelist_srcs.size()); + std::vector h_dsts(edgelist_dsts.size()); + std::tie(h_srcs, h_dsts, std::ignore) = cugraph::test::graph_to_host_coo( + handle, + graph_view, + edge_weight ? std::make_optional((*edge_weight).view()) : std::nullopt, + std::optional>(std::nullopt)); + + auto h_cugraph_edge_triangle_counts = cugraph::test::to_host(handle, *d_edge_triangle_counts); + + auto h_reference_edge_triangle_counts = + edge_triangle_count_reference(h_srcs, h_dsts); + + for (size_t i = 0; i < h_srcs.size(); ++i) { + ASSERT_EQ(h_cugraph_edge_triangle_counts[i], h_reference_edge_triangle_counts[i]) + << "Edge triangle count values do not match with the reference values."; + } + } + } +}; + +using Tests_EdgeTriangleCount_File = Tests_EdgeTriangleCount; +using Tests_EdgeTriangleCount_Rmat = Tests_EdgeTriangleCount; + +TEST_P(Tests_EdgeTriangleCount_File, CheckInt32Int32Float) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} +TEST_P(Tests_EdgeTriangleCount_Rmat, CheckInt32Int32Float) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} +TEST_P(Tests_EdgeTriangleCount_File, CheckInt64Int64Float) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} +TEST_P(Tests_EdgeTriangleCount_Rmat, CheckInt64Int64Float) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +INSTANTIATE_TEST_SUITE_P( + simple_test, + Tests_EdgeTriangleCount_File, + ::testing::Combine( + // enable correctness checks + ::testing::Values(EdgeTriangleCount_Usecase{false, true}, + EdgeTriangleCount_Usecase{true, true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), + cugraph::test::File_Usecase("test/datasets/dolphins.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + rmat_small_test, + Tests_EdgeTriangleCount_Rmat, + // enable correctness checks + ::testing::Combine( + ::testing::Values(EdgeTriangleCount_Usecase{false, true}, + EdgeTriangleCount_Usecase{true, true}), + ::testing::Values(cugraph::test::Rmat_Usecase(10, 16, 0.57, 0.19, 0.19, 0, true, false)))); + +INSTANTIATE_TEST_SUITE_P( + rmat_benchmark_test, /* note that scale & edge factor can be overridden in benchmarking (with + --gtest_filter to select only the rmat_benchmark_test with a specific + vertex & edge type combination) by command line arguments and do not + include more than one Rmat_Usecase that differ only in scale or edge + factor (to avoid running same benchmarks more than once) */ + Tests_EdgeTriangleCount_Rmat, + // disable correctness checks for large graphs + // FIXME: High memory footprint. Perform nbr_intersection in chunks. + ::testing::Combine( + ::testing::Values(EdgeTriangleCount_Usecase{false, false}, + EdgeTriangleCount_Usecase{true, false}), + ::testing::Values(cugraph::test::Rmat_Usecase(16, 16, 0.57, 0.19, 0.19, 0, true, false)))); + +CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/community/mg_ecg_test.cpp b/cpp/tests/community/mg_ecg_test.cpp index a5e02c4f532..c99f83fa2e8 100644 --- a/cpp/tests/community/mg_ecg_test.cpp +++ b/cpp/tests/community/mg_ecg_test.cpp @@ -127,12 +127,14 @@ class Tests_MGEcg : public ::testing::TestWithParam, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - mg_edge_weight_view, - std::optional>{std::nullopt}, - false); // crate a SG graph with MG graph vertex IDs + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + false); // crate a SG graph with MG graph vertex IDs auto const comm_rank = handle_->get_comms().get_rank(); if (comm_rank == 0) { diff --git a/cpp/tests/community/mg_edge_triangle_count_test.cpp b/cpp/tests/community/mg_edge_triangle_count_test.cpp new file mode 100644 index 00000000000..89bdf870ccd --- /dev/null +++ b/cpp/tests/community/mg_edge_triangle_count_test.cpp @@ -0,0 +1,253 @@ +/* + * Copyright (c) 2022-2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "utilities/base_fixture.hpp" +#include "utilities/conversion_utilities.hpp" +#include "utilities/device_comm_wrapper.hpp" +#include "utilities/mg_utilities.hpp" +#include "utilities/property_generator_utilities.hpp" +#include "utilities/test_graphs.hpp" +#include "utilities/thrust_wrapper.hpp" + +#include +#include +#include +#include +#include + +#include +#include +#include + +#include +#include + +#include + +#include + +struct EdgeTriangleCount_Usecase { + bool edge_masking_{false}; + bool check_correctness_{true}; +}; + +template +class Tests_MGEdgeTriangleCount + : public ::testing::TestWithParam> { + public: + Tests_MGEdgeTriangleCount() {} + + static void SetUpTestCase() { handle_ = cugraph::test::initialize_mg_handle(); } + + static void TearDownTestCase() { handle_.reset(); } + + virtual void SetUp() {} + virtual void TearDown() {} + + // Compare the results of running EdgeTriangleCount on multiple GPUs to that of a single-GPU run + template + void run_current_test(EdgeTriangleCount_Usecase const& edge_triangle_count_usecase, + input_usecase_t const& input_usecase) + { + using weight_t = float; + + HighResTimer hr_timer{}; + + // 1. create MG graph + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle_->get_comms().barrier(); + hr_timer.start("MG Construct graph"); + } + + cugraph::graph_t mg_graph(*handle_); + std::optional> mg_renumber_map{std::nullopt}; + std::tie(mg_graph, std::ignore, mg_renumber_map) = + cugraph::test::construct_graph( + *handle_, input_usecase, false, true, false, true); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle_->get_comms().barrier(); + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + auto mg_graph_view = mg_graph.view(); + + std::optional> edge_mask{std::nullopt}; + if (edge_triangle_count_usecase.edge_masking_) { + edge_mask = cugraph::test::generate::edge_property( + *handle_, mg_graph_view, 2); + mg_graph_view.attach_edge_mask((*edge_mask).view()); + } + + // 2. run MG EdgeTriangleCount + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle_->get_comms().barrier(); + hr_timer.start("MG EdgeTriangleCount"); + } + + auto d_mg_cugraph_results = + cugraph::edge_triangle_count(*handle_, mg_graph_view); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle_->get_comms().barrier(); + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + // 3. Compare SG & MG results + + if (edge_triangle_count_usecase.check_correctness_) { + // 3-1. Convert to SG graph + + cugraph::graph_t sg_graph(*handle_); + std::optional< + cugraph::edge_property_t, edge_t>> + d_sg_cugraph_results{std::nullopt}; + std::tie(sg_graph, std::ignore, d_sg_cugraph_results, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + // FIXME: Update 'create_graph_from_edgelist' to support int32_t and int64_t values + std::make_optional(d_mg_cugraph_results.view()), + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); + + if (handle_->get_comms().get_rank() == int{0}) { + // 3-2. Convert the MG triangle counts stored as 'edge_property_t' to device vector + + auto [edgelist_srcs, + edgelist_dsts, + d_edgelist_weights, + d_edge_triangle_counts, + d_edgelist_type] = + cugraph::decompress_to_edgelist( + *handle_, + sg_graph.view(), + std::optional>{std::nullopt}, + // FIXME: Update 'decompress_edgelist' to support int32_t and int64_t values + std::make_optional((*d_sg_cugraph_results).view()), + std::optional>{std::nullopt}, + std::optional>{ + std::nullopt}); // FIXME: No longer needed + + // 3-3. Run SG EdgeTriangleCount + + auto ref_d_sg_cugraph_results = + cugraph::edge_triangle_count(*handle_, sg_graph.view()); + auto [ref_edgelist_srcs, + ref_edgelist_dsts, + ref_d_edgelist_weights, + ref_d_edge_triangle_counts, + ref_d_edgelist_type] = + cugraph::decompress_to_edgelist( + *handle_, + sg_graph.view(), + std::optional>{std::nullopt}, + std::make_optional(ref_d_sg_cugraph_results.view()), + std::optional>{std::nullopt}, + std::optional>{ + std::nullopt}); // FIXME: No longer needed + + // 3-4. Compare + + auto h_mg_edge_triangle_counts = cugraph::test::to_host(*handle_, *d_edge_triangle_counts); + auto h_sg_edge_triangle_counts = + cugraph::test::to_host(*handle_, *ref_d_edge_triangle_counts); + + ASSERT_TRUE(std::equal(h_mg_edge_triangle_counts.begin(), + h_mg_edge_triangle_counts.end(), + h_sg_edge_triangle_counts.begin())); + } + } + } + + private: + static std::unique_ptr handle_; +}; + +template +std::unique_ptr Tests_MGEdgeTriangleCount::handle_ = nullptr; + +using Tests_MGEdgeTriangleCount_File = Tests_MGEdgeTriangleCount; +using Tests_MGEdgeTriangleCount_Rmat = Tests_MGEdgeTriangleCount; + +TEST_P(Tests_MGEdgeTriangleCount_File, CheckInt32Int32) +{ + auto param = GetParam(); + run_current_test(std::get<0>(param), std::get<1>(param)); +} + +TEST_P(Tests_MGEdgeTriangleCount_Rmat, CheckInt32Int32) +{ + auto param = GetParam(); + run_current_test( + std::get<0>(param), override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); +} + +TEST_P(Tests_MGEdgeTriangleCount_Rmat, CheckInt32Int64) +{ + auto param = GetParam(); + run_current_test( + std::get<0>(param), override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); +} + +TEST_P(Tests_MGEdgeTriangleCount_Rmat, CheckInt64Int64) +{ + auto param = GetParam(); + run_current_test( + std::get<0>(param), override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); +} + +INSTANTIATE_TEST_SUITE_P( + file_tests, + Tests_MGEdgeTriangleCount_File, + ::testing::Combine( + // enable correctness checks + ::testing::Values(EdgeTriangleCount_Usecase{false, true}, + EdgeTriangleCount_Usecase{true, true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), + cugraph::test::File_Usecase("test/datasets/dolphins.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + rmat_small_tests, + Tests_MGEdgeTriangleCount_Rmat, + ::testing::Combine( + ::testing::Values(EdgeTriangleCount_Usecase{false, true}, + EdgeTriangleCount_Usecase{true, true}), + ::testing::Values(cugraph::test::Rmat_Usecase(10, 16, 0.57, 0.19, 0.19, 0, true, false)))); + +INSTANTIATE_TEST_SUITE_P( + rmat_benchmark_test, /* note that scale & edge factor can be overridden in benchmarking (with + --gtest_filter to select only the rmat_benchmark_test with a specific + vertex & edge type combination) by command line arguments and do not + include more than one Rmat_Usecase that differ only in scale or edge + factor (to avoid running same benchmarks more than once) */ + Tests_MGEdgeTriangleCount_Rmat, + ::testing::Combine( + ::testing::Values(EdgeTriangleCount_Usecase{false, false}, + EdgeTriangleCount_Usecase{true, false}), + ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, true, false)))); + +CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/community/mg_egonet_test.cu b/cpp/tests/community/mg_egonet_test.cu index 66ab1f47312..ac363df3ec5 100644 --- a/cpp/tests/community/mg_egonet_test.cu +++ b/cpp/tests/community/mg_egonet_test.cu @@ -199,13 +199,15 @@ class Tests_MGEgonet triplet_first + d_mg_aggregate_edgelist_src.size()); } - auto [sg_graph, sg_edge_weights, sg_number_map] = - cugraph::test::mg_graph_to_sg_graph(*handle_, - mg_graph_view, - mg_edge_weight_view, - std::make_optional>( - (*mg_renumber_map).data(), (*mg_renumber_map).size()), - false); + auto [sg_graph, sg_edge_weights, sg_edge_ids, sg_number_map] = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto d_mg_aggregate_edgelist_offsets = diff --git a/cpp/tests/community/mg_leiden_test.cpp b/cpp/tests/community/mg_leiden_test.cpp index f1a2fc83192..65f4827ba06 100644 --- a/cpp/tests/community/mg_leiden_test.cpp +++ b/cpp/tests/community/mg_leiden_test.cpp @@ -87,12 +87,14 @@ class Tests_MGLeiden std::optional< cugraph::edge_property_t, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - mg_edge_weight_view, - std::optional>{std::nullopt}, - false); // crate an SG graph with MG graph vertex IDs + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + false); // crate an SG graph with MG graph vertex IDs // FIXME: We need to figure out how to test each iteration of // SG vs MG Leiden, possibly by passing results of refinement phase diff --git a/cpp/tests/community/mg_louvain_test.cpp b/cpp/tests/community/mg_louvain_test.cpp index 733ee9368ac..106ad2562f7 100644 --- a/cpp/tests/community/mg_louvain_test.cpp +++ b/cpp/tests/community/mg_louvain_test.cpp @@ -85,12 +85,14 @@ class Tests_MGLouvain std::optional< cugraph::edge_property_t, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - mg_edge_weight_view, - std::optional>{std::nullopt}, - false); // crate an SG graph with MG graph vertex IDs + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + false); // crate an SG graph with MG graph vertex IDs weight_t sg_modularity{-1.0}; diff --git a/cpp/tests/community/mg_triangle_count_test.cpp b/cpp/tests/community/mg_triangle_count_test.cpp index ca3e0b2ac8f..932ff5050f1 100644 --- a/cpp/tests/community/mg_triangle_count_test.cpp +++ b/cpp/tests/community/mg_triangle_count_test.cpp @@ -178,13 +178,15 @@ class Tests_MGTriangleCount d_mg_triangle_counts.size())); cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == int{0}) { // 4-2. run SG TriangleCount diff --git a/cpp/tests/community/mg_weighted_matching_test.cpp b/cpp/tests/community/mg_weighted_matching_test.cpp index 21963922ab1..4f36ee36902 100644 --- a/cpp/tests/community/mg_weighted_matching_test.cpp +++ b/cpp/tests/community/mg_weighted_matching_test.cpp @@ -130,12 +130,14 @@ class Tests_MGWeightedMatching std::optional< cugraph::edge_property_t, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - mg_edge_weight_view, - std::optional>(std::nullopt), - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::optional>(std::nullopt), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/components/mg_weakly_connected_components_test.cpp b/cpp/tests/components/mg_weakly_connected_components_test.cpp index c510e3139fb..368fea68877 100644 --- a/cpp/tests/components/mg_weakly_connected_components_test.cpp +++ b/cpp/tests/components/mg_weakly_connected_components_test.cpp @@ -125,13 +125,15 @@ class Tests_MGWeaklyConnectedComponents raft::device_span(d_mg_components.data(), d_mg_components.size())); cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == int{0}) { // 3-2. run SG weakly connected components diff --git a/cpp/tests/cores/mg_core_number_test.cpp b/cpp/tests/cores/mg_core_number_test.cpp index ac99d7d4a93..f8294d81fdf 100644 --- a/cpp/tests/cores/mg_core_number_test.cpp +++ b/cpp/tests/cores/mg_core_number_test.cpp @@ -143,13 +143,15 @@ class Tests_MGCoreNumber raft::device_span(d_mg_core_numbers.data(), d_mg_core_numbers.size())); cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == int{0}) { // 3-2. run SG CoreNumber diff --git a/cpp/tests/cores/mg_k_core_test.cpp b/cpp/tests/cores/mg_k_core_test.cpp index 100c7fa3bcf..28bc445bda8 100644 --- a/cpp/tests/cores/mg_k_core_test.cpp +++ b/cpp/tests/cores/mg_k_core_test.cpp @@ -160,13 +160,15 @@ class Tests_MGKCore : public ::testing::TestWithParam>{std::nullopt}, raft::device_span(d_mg_core_numbers.data(), d_mg_core_numbers.size())); - auto [sg_graph, sg_edge_weights, sg_number_map] = - cugraph::test::mg_graph_to_sg_graph(*handle_, - mg_graph_view, - mg_edge_weight_view, - std::make_optional>( - (*mg_renumber_map).data(), (*mg_renumber_map).size()), - false); + auto [sg_graph, sg_edge_weights, sg_edge_ids, sg_number_map] = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/link_analysis/mg_hits_test.cpp b/cpp/tests/link_analysis/mg_hits_test.cpp index 101a4fe1557..40a439ffc4c 100644 --- a/cpp/tests/link_analysis/mg_hits_test.cpp +++ b/cpp/tests/link_analysis/mg_hits_test.cpp @@ -186,13 +186,15 @@ class Tests_MGHits : public ::testing::TestWithParam, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == int{0}) { // 3-3. run SG Hits diff --git a/cpp/tests/link_analysis/mg_pagerank_test.cpp b/cpp/tests/link_analysis/mg_pagerank_test.cpp index 6be451ac5fd..26136c8c9d2 100644 --- a/cpp/tests/link_analysis/mg_pagerank_test.cpp +++ b/cpp/tests/link_analysis/mg_pagerank_test.cpp @@ -202,13 +202,15 @@ class Tests_MGPageRank std::optional< cugraph::edge_property_t, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - mg_edge_weight_view, - std::make_optional>((*d_mg_renumber_map).data(), - (*d_mg_renumber_map).size()), - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::make_optional>((*d_mg_renumber_map).data(), + (*d_mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == int{0}) { // 4-2. run SG PageRank diff --git a/cpp/tests/mtmg/threaded_test_louvain.cu b/cpp/tests/mtmg/threaded_test_louvain.cu index ab51d701b57..b9c8f621ab8 100644 --- a/cpp/tests/mtmg/threaded_test_louvain.cu +++ b/cpp/tests/mtmg/threaded_test_louvain.cu @@ -384,12 +384,13 @@ class Tests_Multithreaded auto thread_handle = instance_manager->get_handle(); if (thread_handle.get_rank() == 0) { - std::tie(sg_graph, sg_edge_weights, std::ignore) = + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( thread_handle.raft_handle(), graph_view.get(thread_handle), edge_weights ? std::make_optional(edge_weights->get(thread_handle).view()) : std::nullopt, + std::optional>{std::nullopt}, std::optional>{std::nullopt}, false); // create an SG graph with MG graph vertex IDs } else { @@ -398,6 +399,7 @@ class Tests_Multithreaded graph_view.get(thread_handle), edge_weights ? std::make_optional(edge_weights->get(thread_handle).view()) : std::nullopt, + std::optional>{std::nullopt}, std::optional>{std::nullopt}, false); // create an SG graph with MG graph vertex IDs } diff --git a/cpp/tests/prims/mg_count_if_e.cu b/cpp/tests/prims/mg_count_if_e.cu index 8ad1a20e585..137f7db8625 100644 --- a/cpp/tests/prims/mg_count_if_e.cu +++ b/cpp/tests/prims/mg_count_if_e.cu @@ -149,13 +149,15 @@ class Tests_MGCountIfE if (prims_usecase.check_correctness) { cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/prims/mg_count_if_v.cu b/cpp/tests/prims/mg_count_if_v.cu index eb0e8cf9835..e3f30e37729 100644 --- a/cpp/tests/prims/mg_count_if_v.cu +++ b/cpp/tests/prims/mg_count_if_v.cu @@ -123,13 +123,15 @@ class Tests_MGCountIfV if (prims_usecase.check_correctness) { cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/prims/mg_extract_transform_e.cu b/cpp/tests/prims/mg_extract_transform_e.cu index 48b893f6fea..20e87070fa5 100644 --- a/cpp/tests/prims/mg_extract_transform_e.cu +++ b/cpp/tests/prims/mg_extract_transform_e.cu @@ -253,13 +253,15 @@ class Tests_MGExtractTransformE } cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*d_mg_renumber_map_labels).data(), - (*d_mg_renumber_map_labels).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*d_mg_renumber_map_labels).data(), + (*d_mg_renumber_map_labels).size()), + false); rmm::device_uvector sg_vertex_prop(0, handle_->get_stream()); std::tie(std::ignore, sg_vertex_prop) = cugraph::test::mg_vertex_property_values_to_sg_vertex_property_values( diff --git a/cpp/tests/prims/mg_extract_transform_v_frontier_outgoing_e.cu b/cpp/tests/prims/mg_extract_transform_v_frontier_outgoing_e.cu index 3611a250afd..9e7611190ae 100644 --- a/cpp/tests/prims/mg_extract_transform_v_frontier_outgoing_e.cu +++ b/cpp/tests/prims/mg_extract_transform_v_frontier_outgoing_e.cu @@ -283,13 +283,15 @@ class Tests_MGExtractTransformVFrontierOutgoingE } cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*d_mg_renumber_map_labels).data(), - (*d_mg_renumber_map_labels).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*d_mg_renumber_map_labels).data(), + (*d_mg_renumber_map_labels).size()), + false); rmm::device_uvector sg_vertex_prop(0, handle_->get_stream()); std::tie(std::ignore, sg_vertex_prop) = cugraph::test::mg_vertex_property_values_to_sg_vertex_property_values( diff --git a/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu b/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu index 762da62eeb8..75b711fbd9c 100644 --- a/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu +++ b/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu @@ -226,13 +226,15 @@ class Tests_MGPerVPairTransformDstNbrIntersection *handle_, std::get<1>(mg_result_buffer).data(), std::get<1>(mg_result_buffer).size()); cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_weighted_intersection.cu b/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_weighted_intersection.cu index de78b42603d..48bbc6176d8 100644 --- a/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_weighted_intersection.cu +++ b/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_weighted_intersection.cu @@ -258,15 +258,17 @@ class Tests_MGPerVPairTransformDstNbrIntersection weight_t>> sg_edge_weight{std::nullopt}; - std::tie(sg_graph, sg_edge_weight, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - mg_edge_weight - ? std::make_optional(mg_edge_weight_view) - : std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, sg_edge_weight, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight + ? std::make_optional(mg_edge_weight_view) + : std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/prims/mg_per_v_random_select_transform_outgoing_e.cu b/cpp/tests/prims/mg_per_v_random_select_transform_outgoing_e.cu index 97c7333cd2e..b99dbf16107 100644 --- a/cpp/tests/prims/mg_per_v_random_select_transform_outgoing_e.cu +++ b/cpp/tests/prims/mg_per_v_random_select_transform_outgoing_e.cu @@ -320,13 +320,15 @@ class Tests_MGPerVRandomSelectTransformOutgoingE } cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { std::optional> mg_aggregate_sample_offsets{std::nullopt}; diff --git a/cpp/tests/prims/mg_per_v_transform_reduce_dst_key_aggregated_outgoing_e.cu b/cpp/tests/prims/mg_per_v_transform_reduce_dst_key_aggregated_outgoing_e.cu index efcfee9fc66..fd9192dcce5 100644 --- a/cpp/tests/prims/mg_per_v_transform_reduce_dst_key_aggregated_outgoing_e.cu +++ b/cpp/tests/prims/mg_per_v_transform_reduce_dst_key_aggregated_outgoing_e.cu @@ -297,13 +297,15 @@ class Tests_MGPerVTransformReduceDstKeyAggregatedOutgoingE std::optional< cugraph::edge_property_t, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); for (size_t i = 0; i < reduction_types.size(); ++i) { auto mg_aggregate_results = diff --git a/cpp/tests/prims/mg_per_v_transform_reduce_incoming_outgoing_e.cu b/cpp/tests/prims/mg_per_v_transform_reduce_incoming_outgoing_e.cu index e3eb56d5a6e..be29c793ad5 100644 --- a/cpp/tests/prims/mg_per_v_transform_reduce_incoming_outgoing_e.cu +++ b/cpp/tests/prims/mg_per_v_transform_reduce_incoming_outgoing_e.cu @@ -271,13 +271,15 @@ class Tests_MGPerVTransformReduceIncomingOutgoingE if (prims_usecase.check_correctness) { cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); for (size_t i = 0; i < reduction_types.size(); ++i) { auto mg_aggregate_in_results = diff --git a/cpp/tests/prims/mg_reduce_v.cu b/cpp/tests/prims/mg_reduce_v.cu index 1449e8f9910..e91db5fa6ad 100644 --- a/cpp/tests/prims/mg_reduce_v.cu +++ b/cpp/tests/prims/mg_reduce_v.cu @@ -163,13 +163,15 @@ class Tests_MGReduceV if (prims_usecase.check_correctness) { cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/prims/mg_transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v.cu b/cpp/tests/prims/mg_transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v.cu index 71cdf27fda1..4fac6ef3be7 100644 --- a/cpp/tests/prims/mg_transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v.cu +++ b/cpp/tests/prims/mg_transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v.cu @@ -174,13 +174,15 @@ class Tests_MGTransformReduceDstNbrIntersectionOfEEndpointsByV raft::device_span(mg_result_buffer.data(), mg_result_buffer.size())); cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/prims/mg_transform_reduce_e.cu b/cpp/tests/prims/mg_transform_reduce_e.cu index a086571d6e0..4785a8bb01b 100644 --- a/cpp/tests/prims/mg_transform_reduce_e.cu +++ b/cpp/tests/prims/mg_transform_reduce_e.cu @@ -159,13 +159,15 @@ class Tests_MGTransformReduceE if (prims_usecase.check_correctness) { cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/prims/mg_transform_reduce_e_by_src_dst_key.cu b/cpp/tests/prims/mg_transform_reduce_e_by_src_dst_key.cu index a66c70ff586..9950b5bdbf4 100644 --- a/cpp/tests/prims/mg_transform_reduce_e_by_src_dst_key.cu +++ b/cpp/tests/prims/mg_transform_reduce_e_by_src_dst_key.cu @@ -237,13 +237,15 @@ class Tests_MGTransformReduceEBySrcDstKey cugraph::get_dataframe_buffer_begin(mg_aggregate_by_dst_values)); cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/prims/mg_transform_reduce_v.cu b/cpp/tests/prims/mg_transform_reduce_v.cu index c26085a55c4..f6f07bc03ab 100644 --- a/cpp/tests/prims/mg_transform_reduce_v.cu +++ b/cpp/tests/prims/mg_transform_reduce_v.cu @@ -169,13 +169,15 @@ class Tests_MGTransformReduceV if (prims_usecase.check_correctness) { cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/prims/mg_transform_reduce_v_frontier_outgoing_e_by_dst.cu b/cpp/tests/prims/mg_transform_reduce_v_frontier_outgoing_e_by_dst.cu index 07a0f7e7aab..335a7ec879c 100644 --- a/cpp/tests/prims/mg_transform_reduce_v_frontier_outgoing_e_by_dst.cu +++ b/cpp/tests/prims/mg_transform_reduce_v_frontier_outgoing_e_by_dst.cu @@ -292,13 +292,15 @@ class Tests_MGTransformReduceVFrontierOutgoingEByDst } cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == int{0}) { if constexpr (std::is_same_v) { diff --git a/cpp/tests/structure/mg_coarsen_graph_test.cpp b/cpp/tests/structure/mg_coarsen_graph_test.cpp index 1da30869545..471773d71bd 100644 --- a/cpp/tests/structure/mg_coarsen_graph_test.cpp +++ b/cpp/tests/structure/mg_coarsen_graph_test.cpp @@ -330,23 +330,26 @@ class Tests_MGCoarsenGraph cugraph::edge_property_t, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - mg_edge_weight_view, - std::optional>{std::nullopt}, - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + false); cugraph::graph_t sg_coarse_graph(*handle_); std::optional< cugraph::edge_property_t, weight_t>> sg_coarse_edge_weights{std::nullopt}; - std::tie(sg_coarse_graph, sg_coarse_edge_weights, std::ignore) = + std::tie(sg_coarse_graph, sg_coarse_edge_weights, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( *handle_, mg_coarse_graph_view, mg_coarse_edge_weight_view, + std::optional>{std::nullopt}, std::optional>{std::nullopt}, false); diff --git a/cpp/tests/structure/mg_count_self_loops_and_multi_edges_test.cpp b/cpp/tests/structure/mg_count_self_loops_and_multi_edges_test.cpp index 45fac884f49..61f40049e31 100644 --- a/cpp/tests/structure/mg_count_self_loops_and_multi_edges_test.cpp +++ b/cpp/tests/structure/mg_count_self_loops_and_multi_edges_test.cpp @@ -126,13 +126,15 @@ class Tests_MGCountSelfLoopsAndMultiEdges // 3-1. aggregate MG results cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/structure/mg_has_edge_and_compute_multiplicity_test.cpp b/cpp/tests/structure/mg_has_edge_and_compute_multiplicity_test.cpp index 0ee72726294..3d3d881fb23 100644 --- a/cpp/tests/structure/mg_has_edge_and_compute_multiplicity_test.cpp +++ b/cpp/tests/structure/mg_has_edge_and_compute_multiplicity_test.cpp @@ -204,13 +204,15 @@ class Tests_MGHasEdgeAndComputeMultiplicity d_mg_edge_multiplicities.size())); cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); diff --git a/cpp/tests/structure/mg_induced_subgraph_test.cu b/cpp/tests/structure/mg_induced_subgraph_test.cu index 3b32c15bf9f..2ed909b9955 100644 --- a/cpp/tests/structure/mg_induced_subgraph_test.cu +++ b/cpp/tests/structure/mg_induced_subgraph_test.cu @@ -214,12 +214,14 @@ class Tests_MGInducedSubgraph true, handle_->get_stream()); - auto [sg_graph, sg_edge_weights, sg_number_map] = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - mg_edge_weight_view, - std::optional>{std::nullopt}, - false); + auto [sg_graph, sg_edge_weights, sg_edge_ids, sg_number_map] = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + false); if (my_rank == 0) { auto d_sg_subgraph_offsets = cugraph::test::to_device(*handle_, h_sg_subgraph_offsets); diff --git a/cpp/tests/structure/mg_symmetrize_test.cpp b/cpp/tests/structure/mg_symmetrize_test.cpp index e607370f62a..7f1e4f04dc7 100644 --- a/cpp/tests/structure/mg_symmetrize_test.cpp +++ b/cpp/tests/structure/mg_symmetrize_test.cpp @@ -88,13 +88,15 @@ class Tests_MGSymmetrize weight_t>> sg_edge_weights{std::nullopt}; if (symmetrize_usecase.check_correctness) { - std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph.view(), - mg_edge_weights ? std::make_optional((*mg_edge_weights).view()) : std::nullopt, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph.view(), + mg_edge_weights ? std::make_optional((*mg_edge_weights).view()) : std::nullopt, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); } // 3. run MG symmetrize diff --git a/cpp/tests/structure/mg_transpose_storage_test.cpp b/cpp/tests/structure/mg_transpose_storage_test.cpp index c8b4f70f1e2..e870f648039 100644 --- a/cpp/tests/structure/mg_transpose_storage_test.cpp +++ b/cpp/tests/structure/mg_transpose_storage_test.cpp @@ -87,13 +87,15 @@ class Tests_MGTransposeStorage weight_t>> sg_edge_weights{std::nullopt}; if (transpose_storage_usecase.check_correctness) { - std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph.view(), - mg_edge_weights ? std::make_optional((*mg_edge_weights).view()) : std::nullopt, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph.view(), + mg_edge_weights ? std::make_optional((*mg_edge_weights).view()) : std::nullopt, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); } // 2. run MG transpose storage diff --git a/cpp/tests/structure/mg_transpose_test.cpp b/cpp/tests/structure/mg_transpose_test.cpp index 4428f8430d5..921cef42595 100644 --- a/cpp/tests/structure/mg_transpose_test.cpp +++ b/cpp/tests/structure/mg_transpose_test.cpp @@ -87,13 +87,15 @@ class Tests_MGTranspose weight_t>> sg_edge_weights{std::nullopt}; if (transpose_usecase.check_correctness) { - std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph.view(), - mg_edge_weights ? std::make_optional((*mg_edge_weights).view()) : std::nullopt, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph.view(), + mg_edge_weights ? std::make_optional((*mg_edge_weights).view()) : std::nullopt, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); } // 3. run MG transpose diff --git a/cpp/tests/traversal/mg_bfs_test.cpp b/cpp/tests/traversal/mg_bfs_test.cpp index 431ed75c82d..1b63ad3b085 100644 --- a/cpp/tests/traversal/mg_bfs_test.cpp +++ b/cpp/tests/traversal/mg_bfs_test.cpp @@ -183,13 +183,15 @@ class Tests_MGBFS : public ::testing::TestWithParam sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == int{0}) { // 3-3. run SG BFS diff --git a/cpp/tests/traversal/mg_extract_bfs_paths_test.cu b/cpp/tests/traversal/mg_extract_bfs_paths_test.cu index 8484066c6a0..476a6ffab8f 100644 --- a/cpp/tests/traversal/mg_extract_bfs_paths_test.cu +++ b/cpp/tests/traversal/mg_extract_bfs_paths_test.cu @@ -237,13 +237,15 @@ class Tests_MGExtractBFSPaths cugraph::test::device_gatherv(*handle_, d_mg_paths.data(), d_mg_paths.size()); cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == int{0}) { // run SG extract_bfs_paths diff --git a/cpp/tests/traversal/mg_k_hop_nbrs_test.cpp b/cpp/tests/traversal/mg_k_hop_nbrs_test.cpp index 07ea107a2ed..64674fb3799 100644 --- a/cpp/tests/traversal/mg_k_hop_nbrs_test.cpp +++ b/cpp/tests/traversal/mg_k_hop_nbrs_test.cpp @@ -178,13 +178,15 @@ class Tests_MGKHopNbrs *handle_, raft::device_span(d_mg_nbrs.data(), d_mg_nbrs.size())); cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( - *handle_, - mg_graph_view, - std::optional>{std::nullopt}, - std::make_optional>((*mg_renumber_map).data(), - (*mg_renumber_map).size()), - false); + std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == int{0}) { // 3-3. run SG K-hop neighbors diff --git a/cpp/tests/traversal/mg_sssp_test.cpp b/cpp/tests/traversal/mg_sssp_test.cpp index 188d0eca115..9ad16d1c947 100644 --- a/cpp/tests/traversal/mg_sssp_test.cpp +++ b/cpp/tests/traversal/mg_sssp_test.cpp @@ -176,13 +176,15 @@ class Tests_MGSSSP : public ::testing::TestWithParam, weight_t>> sg_edge_weights{std::nullopt}; - std::tie(sg_graph, sg_edge_weights, std::ignore) = - cugraph::test::mg_graph_to_sg_graph(*handle_, - mg_graph_view, - mg_edge_weight_view, - std::make_optional>( - (*mg_renumber_map).data(), (*mg_renumber_map).size()), - false); + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = + cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>{std::nullopt}, + std::make_optional>((*mg_renumber_map).data(), + (*mg_renumber_map).size()), + false); if (handle_->get_comms().get_rank() == int{0}) { // 3-3. run SG SSSP diff --git a/cpp/tests/utilities/conversion_utilities.hpp b/cpp/tests/utilities/conversion_utilities.hpp index 9b55f45d5bd..24a8ecbe4fd 100644 --- a/cpp/tests/utilities/conversion_utilities.hpp +++ b/cpp/tests/utilities/conversion_utilities.hpp @@ -216,15 +216,20 @@ graph_to_host_csc( // Only the rank 0 GPU holds the valid data template -std::tuple, - std::optional, - weight_t>>, - std::optional>> +std::tuple< + cugraph::graph_t, + std::optional< + cugraph::edge_property_t, + weight_t>>, + std::optional< + cugraph::edge_property_t, + edge_t>>, + std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber); diff --git a/cpp/tests/utilities/conversion_utilities_impl.cuh b/cpp/tests/utilities/conversion_utilities_impl.cuh index 6eb7357eedd..748a5731b89 100644 --- a/cpp/tests/utilities/conversion_utilities_impl.cuh +++ b/cpp/tests/utilities/conversion_utilities_impl.cuh @@ -283,23 +283,26 @@ template , std::optional, weight_t>>, + std::optional, edge_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber) { rmm::device_uvector d_src(0, handle.get_stream()); rmm::device_uvector d_dst(0, handle.get_stream()); std::optional> d_wgt{std::nullopt}; + std::optional> d_edge_id{std::nullopt}; - std::tie(d_src, d_dst, d_wgt, std::ignore, std::ignore) = cugraph::decompress_to_edgelist( + std::tie(d_src, d_dst, d_wgt, d_edge_id, std::ignore) = cugraph::decompress_to_edgelist( handle, graph_view, edge_weight_view, - std::optional>{std::nullopt}, + edge_id_view, std::optional>{std::nullopt}, renumber_map); @@ -310,6 +313,9 @@ mg_graph_to_sg_graph( if (d_wgt) *d_wgt = cugraph::test::device_gatherv( handle, raft::device_span{d_wgt->data(), d_wgt->size()}); + if (d_edge_id) + *d_edge_id = cugraph::test::device_gatherv( + handle, raft::device_span{d_edge_id->data(), d_edge_id->size()}); rmm::device_uvector vertices(0, handle.get_stream()); if (renumber_map) { vertices = cugraph::test::device_gatherv(handle, *renumber_map); } @@ -317,6 +323,8 @@ mg_graph_to_sg_graph( graph_t sg_graph(handle); std::optional, weight_t>> sg_edge_weights{std::nullopt}; + std::optional, edge_t>> + sg_edge_ids{std::nullopt}; std::optional> sg_number_map; if (handle.get_comms().get_rank() == 0) { if (!renumber_map) { @@ -325,7 +333,7 @@ mg_graph_to_sg_graph( handle.get_stream(), vertices.data(), vertices.size(), vertex_t{0}); } - std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore, sg_number_map) = + std::tie(sg_graph, sg_edge_weights, sg_edge_ids, std::ignore, sg_number_map) = cugraph::create_graph_from_edgelist diff --git a/cpp/tests/utilities/conversion_utilities_mg.cu b/cpp/tests/utilities/conversion_utilities_mg.cu index d657f868497..cb4703ec89b 100644 --- a/cpp/tests/utilities/conversion_utilities_mg.cu +++ b/cpp/tests/utilities/conversion_utilities_mg.cu @@ -381,132 +381,156 @@ graph_to_host_csc( template std::tuple< cugraph::graph_t, std::optional, float>>, + std::optional, int32_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber); template std::tuple< cugraph::graph_t, std::optional, float>>, + std::optional, int64_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber); template std::tuple< cugraph::graph_t, std::optional, float>>, + std::optional, int64_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber); template std::tuple< cugraph::graph_t, std::optional, double>>, + std::optional, int32_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber); template std::tuple< cugraph::graph_t, std::optional, double>>, + std::optional, int64_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber); template std::tuple< cugraph::graph_t, std::optional, double>>, + std::optional, int64_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber); template std::tuple< cugraph::graph_t, std::optional, float>>, + std::optional, int32_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber); template std::tuple< cugraph::graph_t, std::optional, float>>, + std::optional, int64_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber); template std::tuple< cugraph::graph_t, std::optional, float>>, + std::optional, int64_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber); template std::tuple< cugraph::graph_t, std::optional, double>>, + std::optional, int32_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber); template std::tuple< cugraph::graph_t, std::optional, double>>, + std::optional, int64_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber); template std::tuple< cugraph::graph_t, std::optional, double>>, + std::optional, int64_t>>, std::optional>> mg_graph_to_sg_graph( raft::handle_t const& handle, cugraph::graph_view_t const& graph_view, std::optional> edge_weight_view, + std::optional> edge_id_view, std::optional> renumber_map, bool renumber);