diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 2911be3046e..54a8b2cde93 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -199,7 +199,8 @@ set(CUGRAPH_SOURCES src/community/legacy/ktruss.cu src/community/legacy/ecg.cu src/community/legacy/extract_subgraph_by_vertex.cu - src/community/legacy/egonet.cu + src/community/egonet_sg.cu + src/community/egonet_mg.cu src/sampling/random_walks.cu src/sampling/random_walks_sg.cu src/sampling/detail/sampling_utils_mg.cu @@ -365,6 +366,7 @@ add_library(cugraph_c src/c_api/betweenness_centrality.cpp src/c_api/core_number.cpp src/c_api/core_result.cpp + src/c_api/extract_ego.cpp src/c_api/k_core.cpp src/c_api/induced_subgraph.cpp src/c_api/induced_subgraph_helper.cu diff --git a/cpp/include/cugraph/algorithms.hpp b/cpp/include/cugraph/algorithms.hpp index bf02887dc04..7715ce30882 100644 --- a/cpp/include/cugraph/algorithms.hpp +++ b/cpp/include/cugraph/algorithms.hpp @@ -1349,10 +1349,14 @@ void katz_centrality(raft::handle_t const& handle, bool has_initial_guess = false, bool normalize = false, bool do_expensive_check = false); + /** * @brief returns induced EgoNet subgraph(s) of neighbors centered at nodes in source_vertex within * a given radius. * + * @deprecated This algorithm will be deprecated to replaced by the new version + * that uses the raft::device_span. + * * @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 weight_t Type of edge weights. Needs to be a floating point type. @@ -1380,6 +1384,37 @@ extract_ego(raft::handle_t const& handle, vertex_t n_subgraphs, vertex_t radius); +/** + * @brief returns induced EgoNet subgraph(s) of neighbors centered at nodes in source_vertex within + * a given radius. + * + * @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 weight_t Type of edge weights. Needs to be a floating point type. + * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * or multi-GPU (true). + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. Must have at least one worker stream. + * @param graph_view Graph view object of, we extract induced egonet subgraphs from @p graph_view. + * @param source_vertex Pointer to egonet center vertices (size == @p n_subgraphs). + * @param n_subgraphs Number of induced EgoNet subgraphs to extract (ie. number of elements in @p + * source_vertex). + * @param radius Include all neighbors of distance <= radius from @p source_vertex. + * @return std::tuple, rmm::device_uvector, + * rmm::device_uvector, rmm::device_uvector> Quadraplet of edge source vertices, + * edge destination vertices, edge weights, and edge offsets for each induced EgoNet subgraph. + */ +template +std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertices, + vertex_t radius, + bool do_expensive_check = false); + /** * @brief returns random walks (RW) from starting sources, where each path is of given maximum * length. Uniform distribution is assumed for the random engine. diff --git a/cpp/include/cugraph_c/community_algorithms.h b/cpp/include/cugraph_c/community_algorithms.h index 4921630a82f..9ec17a20486 100644 --- a/cpp/include/cugraph_c/community_algorithms.h +++ b/cpp/include/cugraph_c/community_algorithms.h @@ -18,6 +18,7 @@ #include #include +#include #include /** @defgroup community Community algorithms @@ -135,6 +136,30 @@ double cugraph_heirarchical_clustering_result_get_modularity( */ void cugraph_heirarchical_clustering_result_free(cugraph_heirarchical_clustering_result_t* result); +/** + * @brief Extract ego graphs + * + * @param [in] handle Handle for accessing resources + * @param [in] graph Pointer to graph. NOTE: Graph might be modified if the storage + * needs to be transposed + * @param [in] source_vertices Device array of vertices we want to extract egonets for. + * @param [in] radius The number of hops to go out from each source vertex + * @param [in] do_expensive_check + * A flag to run expensive checks for input arguments (if set to true) + * @param [out] result Opaque object containing the extracted subgraph + * @param [out] error Pointer to an error object storing details of any error. Will + * be populated if error code is not CUGRAPH_SUCCESS + * @return error code + */ +cugraph_error_code_t cugraph_extract_ego( + const cugraph_resource_handle_t* handle, + cugraph_graph_t* graph, + const cugraph_type_erased_device_array_view_t* source_vertices, + size_t radius, + bool_t do_expensive_check, + cugraph_induced_subgraph_result_t** result, + cugraph_error_t** error); + #ifdef __cplusplus } #endif diff --git a/cpp/src/c_api/extract_ego.cpp b/cpp/src/c_api/extract_ego.cpp new file mode 100644 index 00000000000..cf4009c965a --- /dev/null +++ b/cpp/src/c_api/extract_ego.cpp @@ -0,0 +1,154 @@ +/* + * Copyright (c) 2022, 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 + +#include +#include +#include +#include +#include + +#include +#include +#include +#include + +#include + +namespace { + +struct extract_ego_functor : public cugraph::c_api::abstract_functor { + raft::handle_t const& handle_; + cugraph::c_api::cugraph_graph_t* graph_; + cugraph::c_api::cugraph_type_erased_device_array_view_t const* source_vertices_; + size_t radius_; + bool do_expensive_check_; + cugraph::c_api::cugraph_induced_subgraph_result_t* result_{}; + + extract_ego_functor(::cugraph_resource_handle_t const* handle, + ::cugraph_graph_t* graph, + ::cugraph_type_erased_device_array_view_t const* source_vertices, + size_t radius, + bool do_expensive_check) + : abstract_functor(), + handle_(*reinterpret_cast(handle)->handle_), + graph_(reinterpret_cast(graph)), + source_vertices_( + reinterpret_cast( + source_vertices)), + radius_(radius), + do_expensive_check_(do_expensive_check) + { + } + + template + void operator()() + { + if constexpr (!cugraph::is_candidate::value) { + unsupported(); + } else { + // extract ego expects store_transposed == false + if constexpr (store_transposed) { + error_code_ = cugraph::c_api:: + transpose_storage( + handle_, graph_, error_.get()); + if (error_code_ != CUGRAPH_SUCCESS) return; + } + + auto graph = + reinterpret_cast*>( + graph_->graph_); + + auto graph_view = graph->view(); + + auto number_map = reinterpret_cast*>(graph_->number_map_); + + rmm::device_uvector source_vertices(source_vertices_->size_, handle_.get_stream()); + raft::copy(source_vertices.data(), + source_vertices_->as_type(), + source_vertices.size(), + handle_.get_stream()); + + if constexpr (multi_gpu) { + source_vertices = + cugraph::detail::shuffle_ext_vertices_by_gpu_id(handle_, std::move(source_vertices)); + } + + cugraph::renumber_ext_vertices( + handle_, + source_vertices.data(), + source_vertices.size(), + number_map->data(), + graph_view.local_vertex_partition_range_first(), + graph_view.local_vertex_partition_range_last(), + do_expensive_check_); + + auto [src, dst, wgt, edge_offsets] = + cugraph::extract_ego( + handle_, + graph_view, + raft::device_span{source_vertices.data(), source_vertices.size()}, + radius_, + do_expensive_check_); + + cugraph::unrenumber_int_vertices( + handle_, + src.data(), + src.size(), + number_map->data(), + graph_view.vertex_partition_range_lasts(), + do_expensive_check_); + + cugraph::unrenumber_int_vertices( + handle_, + dst.data(), + dst.size(), + number_map->data(), + graph_view.vertex_partition_range_lasts(), + do_expensive_check_); + + result_ = new cugraph::c_api::cugraph_induced_subgraph_result_t{ + new cugraph::c_api::cugraph_type_erased_device_array_t(src, graph_->vertex_type_), + new cugraph::c_api::cugraph_type_erased_device_array_t(dst, graph_->edge_type_), + wgt ? new cugraph::c_api::cugraph_type_erased_device_array_t(*wgt, graph_->edge_type_) + : NULL, + new cugraph::c_api::cugraph_type_erased_device_array_t(edge_offsets, + data_type_id_t::SIZE_T)}; + } + } +}; + +} // namespace + +extern "C" cugraph_error_code_t cugraph_extract_ego( + const cugraph_resource_handle_t* handle, + cugraph_graph_t* graph, + const cugraph_type_erased_device_array_view_t* source_vertices, + size_t radius, + bool_t do_expensive_check, + cugraph_induced_subgraph_result_t** result, + cugraph_error_t** error) +{ + extract_ego_functor functor(handle, graph, source_vertices, radius, do_expensive_check); + + return cugraph::c_api::run_algorithm(graph, functor, result, error); +} diff --git a/cpp/src/community/legacy/egonet.cu b/cpp/src/community/egonet_impl.cuh similarity index 50% rename from cpp/src/community/legacy/egonet.cu rename to cpp/src/community/egonet_impl.cuh index 7d1719bbc14..9b855ac1bcd 100644 --- a/cpp/src/community/legacy/egonet.cu +++ b/cpp/src/community/egonet_impl.cuh @@ -13,8 +13,15 @@ * See the License for the specific language governing permissions and * limitations under the License. */ +#pragma once + +#include + +#include +#include +#include +#include -// Alex Fender afender@nvida.com #include #include #include @@ -33,15 +40,6 @@ #include #include -#include - -#include -#include -#include - -#include -#include - #include namespace { @@ -59,30 +57,30 @@ Rather than doing custom one/two hops features, we propose a generic k-hops solu cutoff and subgraph extraction */ -template +template std::tuple, rmm::device_uvector, std::optional>, rmm::device_uvector> extract(raft::handle_t const& handle, - cugraph::graph_view_t const& csr_view, - vertex_t* source_vertex, - vertex_t n_subgraphs, - vertex_t radius) + cugraph::graph_view_t const& graph_view, + raft::device_span source_vertex, + vertex_t radius, + bool do_expensive_check) { - auto v = csr_view.number_of_vertices(); auto user_stream_view = handle.get_stream(); - rmm::device_vector neighbors_offsets(n_subgraphs + 1); - rmm::device_vector neighbors; + rmm::device_uvector neighbors_offsets(source_vertex.size() + 1, user_stream_view); + rmm::device_uvector neighbors(0, user_stream_view); - std::vector h_neighbors_offsets(n_subgraphs + 1); + std::vector h_neighbors_offsets(source_vertex.size() + 1); // Streams will allocate concurrently later std::vector> reached{}; - reached.reserve(n_subgraphs); - for (vertex_t i = 0; i < n_subgraphs; i++) { + reached.reserve(source_vertex.size()); + for (size_t i = 0; i < source_vertex.size(); i++) { // Allocations and operations are attached to the worker stream - rmm::device_uvector local_reach(v, handle.get_next_usable_stream(i)); + rmm::device_uvector local_reach(graph_view.local_vertex_partition_range_size(), + handle.get_next_usable_stream(i)); reached.push_back(std::move(local_reach)); } @@ -92,67 +90,74 @@ extract(raft::handle_t const& handle, hr_timer.start("ego_neighbors"); #endif - for (vertex_t i = 0; i < n_subgraphs; i++) { + // FIXME: Explore the performance here. Single-seed BFS + // has a slow parallelism ramp up, would we be better off + // using the technique in induced subgraph where we tag + // the vertices and search for matches until the frontiers + // are large enough to use this approach? + + for (size_t i = 0; i < source_vertex.size(); i++) { // get light handle from worker pool raft::handle_t light_handle(handle.get_next_usable_stream(i)); - auto worker_stream_view = light_handle.get_stream(); + auto worker_stream_view = multi_gpu ? handle.get_stream() : light_handle.get_stream(); // BFS with cutoff // consider adding a device API to BFS (ie. accept source on the device) - rmm::device_uvector predecessors(v, worker_stream_view); // not used bool direction_optimizing = false; thrust::fill(rmm::exec_policy(worker_stream_view), reached[i].begin(), reached[i].end(), std::numeric_limits::max()); - thrust::fill( - rmm::exec_policy(worker_stream_view), reached[i].begin(), reached[i].begin() + 100, 1.0); - - cugraph::bfs(light_handle, - csr_view, - reached[i].data(), - predecessors.data(), - source_vertex + i, - 1, - direction_optimizing, - radius); + + cugraph::bfs(multi_gpu ? handle : light_handle, + graph_view, + reached[i].data(), + nullptr, + source_vertex.data() + i, + 1, + direction_optimizing, + radius, + do_expensive_check); // identify reached vertex ids from distance array - thrust::transform(rmm::exec_policy(worker_stream_view), - thrust::make_counting_iterator(vertex_t{0}), - thrust::make_counting_iterator(v), - reached[i].begin(), - reached[i].begin(), - [sentinel = std::numeric_limits::max()] __device__( - auto id, auto val) { return val < sentinel ? id : sentinel; }); + thrust::transform( + rmm::exec_policy(worker_stream_view), + thrust::make_counting_iterator(graph_view.local_vertex_partition_range_first()), + thrust::make_counting_iterator(graph_view.local_vertex_partition_range_last()), + reached[i].begin(), + reached[i].begin(), + [sentinel = std::numeric_limits::max()] __device__(auto id, auto val) { + return val < sentinel ? id : sentinel; + }); // removes unreached data auto reached_end = thrust::remove(rmm::exec_policy(worker_stream_view), reached[i].begin(), reached[i].end(), std::numeric_limits::max()); + // release temp storage reached[i].resize(thrust::distance(reached[i].begin(), reached_end), worker_stream_view); reached[i].shrink_to_fit(worker_stream_view); } - // wait on every one to identify their neighboors before proceeding to concatenation + // wait on every one to identify their neighbors before proceeding to concatenation handle.sync_stream_pool(); - // Construct neighboors offsets (just a scan on neighborhod vector sizes) + // Construct neighbors offsets (just a scan on neighborhod vector sizes) h_neighbors_offsets[0] = 0; - for (vertex_t i = 0; i < n_subgraphs; i++) { + for (size_t i = 0; i < source_vertex.size(); i++) { h_neighbors_offsets[i + 1] = h_neighbors_offsets[i] + reached[i].size(); } - raft::update_device(neighbors_offsets.data().get(), + raft::update_device(neighbors_offsets.data(), &h_neighbors_offsets[0], - n_subgraphs + 1, + source_vertex.size() + 1, user_stream_view.value()); - neighbors.resize(h_neighbors_offsets[n_subgraphs]); + neighbors.resize(h_neighbors_offsets[source_vertex.size()], user_stream_view.value()); user_stream_view.synchronize(); - // Construct the neighboors list concurrently - for (vertex_t i = 0; i < n_subgraphs; i++) { + // Construct the neighbors list concurrently + for (size_t i = 0; i < source_vertex.size(); i++) { auto worker_stream_view = handle.get_next_usable_stream(i); thrust::copy(rmm::exec_policy(worker_stream_view), reached[i].begin(), @@ -175,10 +180,10 @@ extract(raft::handle_t const& handle, // extract return cugraph::extract_induced_subgraphs( handle, - csr_view, - raft::device_span(neighbors_offsets.data().get(), neighbors_offsets.size()), - raft::device_span(neighbors.data().get(), neighbors.size()), - false); + graph_view, + raft::device_span(neighbors_offsets.data(), neighbors_offsets.size()), + raft::device_span(neighbors.data(), neighbors.size()), + do_expensive_check); } } // namespace @@ -196,13 +201,6 @@ extract_ego(raft::handle_t const& handle, vertex_t n_subgraphs, vertex_t radius) { - if (multi_gpu) { - CUGRAPH_FAIL("Unimplemented."); - return std::make_tuple(rmm::device_uvector(0, handle.get_stream()), - rmm::device_uvector(0, handle.get_stream()), - rmm::device_uvector(0, handle.get_stream()), - rmm::device_uvector(0, handle.get_stream())); - } CUGRAPH_EXPECTS(n_subgraphs > 0, "Need at least one source to extract the egonet from"); CUGRAPH_EXPECTS(n_subgraphs < graph_view.number_of_vertices(), "Can't have more sources to extract from than vertices in the graph"); @@ -210,65 +208,31 @@ extract_ego(raft::handle_t const& handle, CUGRAPH_EXPECTS(radius < graph_view.number_of_vertices(), "radius is too large"); // source_vertex range is checked in bfs. - return extract( - handle, graph_view, source_vertex, n_subgraphs, radius); + return extract(handle, + graph_view, + raft::device_span{source_vertex, static_cast(n_subgraphs)}, + radius, + false); +} + +template +std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + vertex_t radius, + bool do_expensive_check) +{ + CUGRAPH_EXPECTS(source_vertex.size() > 0, "Need at least one source to extract the egonet from"); + CUGRAPH_EXPECTS(source_vertex.size() < static_cast(graph_view.number_of_vertices()), + "Can't have more sources to extract from than vertices in the graph"); + CUGRAPH_EXPECTS(radius > 0, "Radius should be at least 1"); + CUGRAPH_EXPECTS(radius < graph_view.number_of_vertices(), "radius is too large"); + + return extract(handle, graph_view, source_vertex, radius, do_expensive_check); } -// SG FP32 -template std::tuple, - rmm::device_uvector, - std::optional>, - rmm::device_uvector> -extract_ego(raft::handle_t const&, - graph_view_t const&, - int32_t*, - int32_t, - int32_t); -template std::tuple, - rmm::device_uvector, - std::optional>, - rmm::device_uvector> -extract_ego(raft::handle_t const&, - graph_view_t const&, - int32_t*, - int32_t, - int32_t); -template std::tuple, - rmm::device_uvector, - std::optional>, - rmm::device_uvector> -extract_ego(raft::handle_t const&, - graph_view_t const&, - int64_t*, - int64_t, - int64_t); - -// SG FP64 -template std::tuple, - rmm::device_uvector, - std::optional>, - rmm::device_uvector> -extract_ego(raft::handle_t const&, - graph_view_t const&, - int32_t*, - int32_t, - int32_t); -template std::tuple, - rmm::device_uvector, - std::optional>, - rmm::device_uvector> -extract_ego(raft::handle_t const&, - graph_view_t const&, - int32_t*, - int32_t, - int32_t); -template std::tuple, - rmm::device_uvector, - std::optional>, - rmm::device_uvector> -extract_ego(raft::handle_t const&, - graph_view_t const&, - int64_t*, - int64_t, - int64_t); } // namespace cugraph diff --git a/cpp/src/community/egonet_mg.cu b/cpp/src/community/egonet_mg.cu new file mode 100644 index 00000000000..4e043e54320 --- /dev/null +++ b/cpp/src/community/egonet_mg.cu @@ -0,0 +1,143 @@ +/* + * Copyright (c) 2021-2022, 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 + +namespace cugraph { + +// MG FP32 +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const&, + graph_view_t const&, + int32_t*, + int32_t, + int32_t); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> + +extract_ego(raft::handle_t const&, + graph_view_t const&, + int32_t*, + int32_t, + int32_t); +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> + +extract_ego(raft::handle_t const&, + graph_view_t const&, + int64_t*, + int64_t, + int64_t); + +// MG FP64 +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const&, + graph_view_t const&, + int32_t*, + int32_t, + int32_t); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const&, + graph_view_t const&, + int32_t*, + int32_t, + int32_t); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const&, + graph_view_t const&, + int64_t*, + int64_t, + int64_t); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + int32_t radius, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + int32_t radius, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + int64_t radius, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + int32_t radius, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + int32_t radius, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + int64_t radius, + bool do_expensive_check); + +} // namespace cugraph diff --git a/cpp/src/community/egonet_sg.cu b/cpp/src/community/egonet_sg.cu new file mode 100644 index 00000000000..02150f8e791 --- /dev/null +++ b/cpp/src/community/egonet_sg.cu @@ -0,0 +1,142 @@ +/* + * Copyright (c) 2021-2022, 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 + +namespace cugraph { + +// SG FP32 +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + int32_t* source_vertex, + int32_t n_subgraphs, + int32_t radius); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + int32_t* source_vertex, + int32_t n_subgraphs, + int32_t radius); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + int64_t* source_vertex, + int64_t n_subgraphs, + int64_t radius); + +// SG FP64 +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const&, + graph_view_t const& graph_view, + int32_t* source_vertex, + int32_t n_subgraphs, + int32_t radius); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + int32_t* source_vertex, + int32_t n_subgraphs, + int32_t radius); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + int64_t* source_vertex, + int64_t n_subgraphs, + int64_t radius); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + int32_t radius, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + int32_t radius, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + int64_t radius, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + int32_t radius, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + int32_t radius, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +extract_ego(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span source_vertex, + int64_t radius, + bool do_expensive_check); + +} // namespace cugraph diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 1f4b56e36a3..95954d59340 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -27,6 +27,7 @@ add_library(cugraphtestutil STATIC utilities/test_utilities_mg.cu link_prediction/similarity_compare.cpp centrality/betweenness_centrality_validate.cu + community/egonet_validate.cu structure/induced_subgraph_validate.cu components/wcc_graphs.cu sampling/random_walks_check_sg.cu @@ -229,7 +230,7 @@ ConfigureTest(BALANCED_TEST community/balanced_edge_test.cpp) ################################################################################################### # - EGO tests -------------------------------------------------------------------------------- -ConfigureTest(EGO_TEST community/egonet_test.cu) +ConfigureTest(EGO_TEST community/egonet_test.cpp) ################################################################################################### # - FORCE ATLAS 2 tests -------------------------------------------------------------------------- @@ -473,6 +474,10 @@ if(BUILD_CUGRAPH_MG_TESTS) community/mg_louvain_helper.cu community/mg_louvain_test.cpp) + ############################################################################################### + # - MG LOUVAIN tests -------------------------------------------------------------------------- + ConfigureTestMG(MG_EGO_TEST community/mg_egonet_test.cpp) + ############################################################################################### # - MG WEAKLY CONNECTED COMPONENTS tests ------------------------------------------------------ ConfigureTestMG(MG_WEAKLY_CONNECTED_COMPONENTS_TEST @@ -591,6 +596,7 @@ if(BUILD_CUGRAPH_MG_TESTS) ConfigureCTestMG(MG_CAPI_SIMILARITY_TEST c_api/mg_similarity_test.c c_api/mg_test_utils.cpp) ConfigureCTestMG(MG_CAPI_K_CORE_TEST c_api/mg_k_core_test.c c_api/mg_test_utils.cpp) ConfigureCTestMG(MG_CAPI_INDUCED_SUBGRAPH_TEST c_api/mg_induced_subgraph_test.c c_api/mg_test_utils.cpp) + ConfigureCTestMG(MG_CAPI_EGONET_TEST c_api/mg_egonet_test.c c_api/mg_test_utils.cpp) ConfigureCTestMG(MG_CAPI_TWO_HOP_NEIGHBORS_TEST c_api/mg_two_hop_neighbors_test.c c_api/mg_test_utils.cpp) endif() @@ -646,6 +652,7 @@ ConfigureCTest(CAPI_CORE_NUMBER_TEST c_api/core_number_test.c) ConfigureCTest(CAPI_SIMILARITY_TEST c_api/similarity_test.c) ConfigureCTest(CAPI_K_CORE_TEST c_api/k_core_test.c) ConfigureCTest(CAPI_INDUCED_SUBGRAPH_TEST c_api/induced_subgraph_test.c) +ConfigureCTest(CAPI_EGONET_TEST c_api/egonet_test.c) ConfigureCTest(CAPI_TWO_HOP_NEIGHBORS_TEST c_api/two_hop_neighbors_test.c) ################################################################################################### diff --git a/cpp/tests/c_api/egonet_test.c b/cpp/tests/c_api/egonet_test.c new file mode 100644 index 00000000000..fac9815c150 --- /dev/null +++ b/cpp/tests/c_api/egonet_test.c @@ -0,0 +1,195 @@ +/* + * Copyright (c) 2022, 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 "c_test_utils.h" /* RUN_TEST */ + +#include +#include + +#include + +typedef int32_t vertex_t; +typedef int32_t edge_t; +typedef float weight_t; + +int generic_egonet_test(vertex_t* h_src, + vertex_t* h_dst, + weight_t* h_wgt, + vertex_t* h_seeds, + vertex_t* h_expected_src, + vertex_t* h_expected_dst, + size_t* h_expected_offsets, + size_t num_vertices, + size_t num_edges, + size_t num_seeds, + size_t radius, + bool_t store_transposed) +{ + int test_ret_value = 0; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error; + + cugraph_resource_handle_t* resource_handle = NULL; + cugraph_graph_t* graph = NULL; + cugraph_type_erased_device_array_t* seeds = NULL; + cugraph_type_erased_device_array_view_t* seeds_view = NULL; + cugraph_induced_subgraph_result_t* result = NULL; + + resource_handle = cugraph_create_resource_handle(NULL); + TEST_ASSERT(test_ret_value, resource_handle != NULL, "resource handle creation failed."); + + ret_code = create_test_graph(resource_handle, + h_src, + h_dst, + h_wgt, + num_edges, + store_transposed, + FALSE, + FALSE, + &graph, + &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "create_test_graph failed."); + TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + ret_code = + cugraph_type_erased_device_array_create(resource_handle, num_seeds, INT32, &seeds, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "seeds create failed."); + + seeds_view = cugraph_type_erased_device_array_view(seeds); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + resource_handle, seeds_view, (byte_t*)h_seeds, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src copy_from_host failed."); + + ret_code = + cugraph_extract_ego(resource_handle, graph, seeds_view, radius, FALSE, &result, &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, "cugraph_egonet failed."); + + if (test_ret_value == 0) { + cugraph_type_erased_device_array_view_t* src; + cugraph_type_erased_device_array_view_t* dst; + cugraph_type_erased_device_array_view_t* wgt; + cugraph_type_erased_device_array_view_t* offsets; + + src = cugraph_induced_subgraph_get_sources(result); + dst = cugraph_induced_subgraph_get_destinations(result); + wgt = cugraph_induced_subgraph_get_edge_weights(result); + offsets = cugraph_induced_subgraph_get_subgraph_offsets(result); + + size_t num_result_edges = cugraph_type_erased_device_array_view_size(src); + size_t num_result_offsets = cugraph_type_erased_device_array_view_size(offsets); + + vertex_t h_result_src[num_result_edges]; + vertex_t h_result_dst[num_result_edges]; + weight_t h_result_wgt[num_result_edges]; + size_t h_result_offsets[num_result_offsets]; + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + resource_handle, (byte_t*)h_result_src, src, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + resource_handle, (byte_t*)h_result_dst, dst, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + resource_handle, (byte_t*)h_result_wgt, wgt, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + resource_handle, (byte_t*)h_result_offsets, offsets, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + TEST_ASSERT( + test_ret_value, (num_seeds + 1) == num_result_offsets, "number of offsets doesn't match"); + + for (int i = 0; (i < num_result_offsets) && (test_ret_value == 0); ++i) { + TEST_ASSERT( + test_ret_value, h_result_offsets[i] == h_expected_offsets[i], "offsets don't match"); + } + + weight_t M[num_vertices][num_vertices]; + + for (int i = 0; (i < num_seeds) && (test_ret_value == 0); ++i) { + for (int r = 0 ; r < num_vertices ; ++r) + for (int c = 0 ; c < num_vertices ; ++c) + M[r][c] = 0; + + for (size_t e = h_expected_offsets[i] ; e < h_expected_offsets[i+1] ; ++e) + M[h_expected_src[e]][h_expected_dst[e]] = 1; + + for (size_t e = h_result_offsets[i] ; (e < h_result_offsets[i+1]) && (test_ret_value == 0) ; ++e) { + TEST_ASSERT(test_ret_value, (M[h_result_src[e]][h_result_dst[e]] > 0), "found different edges"); + } + } + + cugraph_type_erased_device_array_view_free(src); + cugraph_type_erased_device_array_view_free(dst); + cugraph_type_erased_device_array_view_free(wgt); + cugraph_type_erased_device_array_view_free(offsets); + cugraph_induced_subgraph_result_free(result); + } + + cugraph_sg_graph_free(graph); + cugraph_free_resource_handle(resource_handle); + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int test_egonet() +{ + size_t num_edges = 9; + size_t num_vertices = 6; + size_t radius = 2; + size_t num_seeds = 2; + + vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 3, 4}; + vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 4, 5, 5}; + weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 6.1f}; + vertex_t h_seeds[] = {0, 1}; + + vertex_t h_result_src[] = {0, 1, 1, 3, 1, 1, 3, 3, 4}; + vertex_t h_result_dst[] = {1, 3, 4, 4, 3, 4, 4, 5, 5}; + size_t h_result_offsets[] = {0, 4, 9}; + + // Egonet wants store_transposed = FALSE + return generic_egonet_test(h_src, + h_dst, + h_wgt, + h_seeds, + h_result_src, + h_result_dst, + h_result_offsets, + num_vertices, + num_edges, + num_seeds, + radius, + FALSE); +} + +/******************************************************************************/ + +int main(int argc, char** argv) +{ + int result = 0; + result |= RUN_TEST(test_egonet); + return result; +} diff --git a/cpp/tests/c_api/mg_egonet_test.c b/cpp/tests/c_api/mg_egonet_test.c new file mode 100644 index 00000000000..8f0dda42a57 --- /dev/null +++ b/cpp/tests/c_api/mg_egonet_test.c @@ -0,0 +1,242 @@ +/* + * Copyright (c) 2022, 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 "mg_test_utils.h" /* RUN_TEST */ + +#include +#include + +#include + +typedef int32_t vertex_t; +typedef int32_t edge_t; +typedef float weight_t; + +int generic_egonet_test(const cugraph_resource_handle_t* resource_handle, + vertex_t* h_src, + vertex_t* h_dst, + weight_t* h_wgt, + vertex_t* h_seeds, + vertex_t* h_expected_src, + vertex_t* h_expected_dst, + size_t* h_expected_offsets, + size_t num_vertices, + size_t num_edges, + size_t num_seeds, + size_t radius, + bool_t store_transposed) +{ + int test_ret_value = 0; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error; + + cugraph_graph_t* graph = NULL; + cugraph_type_erased_device_array_t* seeds = NULL; + cugraph_type_erased_device_array_view_t* seeds_view = NULL; + cugraph_induced_subgraph_result_t* result = NULL; + + int rank = cugraph_resource_handle_get_rank(resource_handle); + + resource_handle = cugraph_create_resource_handle(NULL); + TEST_ASSERT(test_ret_value, resource_handle != NULL, "resource handle creation failed."); + + ret_code = create_test_graph(resource_handle, + h_src, + h_dst, + h_wgt, + num_edges, + store_transposed, + FALSE, + FALSE, + &graph, + &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "create_test_graph failed."); + TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + if (rank != 0) { + num_seeds = 0; + } + + ret_code = + cugraph_type_erased_device_array_create(resource_handle, num_seeds, INT32, &seeds, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "seeds create failed."); + + seeds_view = cugraph_type_erased_device_array_view(seeds); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + resource_handle, seeds_view, (byte_t*)h_seeds, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src copy_from_host failed."); + + ret_code = + cugraph_extract_ego(resource_handle, graph, seeds_view, radius, FALSE, &result, &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, "cugraph_egonet failed."); + + if (test_ret_value == 0) { + cugraph_type_erased_device_array_view_t* src; + cugraph_type_erased_device_array_view_t* dst; + cugraph_type_erased_device_array_view_t* wgt; + cugraph_type_erased_device_array_view_t* offsets; + + src = cugraph_induced_subgraph_get_sources(result); + dst = cugraph_induced_subgraph_get_destinations(result); + wgt = cugraph_induced_subgraph_get_edge_weights(result); + offsets = cugraph_induced_subgraph_get_subgraph_offsets(result); + + size_t num_result_edges = cugraph_type_erased_device_array_view_size(src); + size_t num_result_offsets = cugraph_type_erased_device_array_view_size(offsets); + + vertex_t h_result_src[num_result_edges]; + vertex_t h_result_dst[num_result_edges]; + weight_t h_result_wgt[num_result_edges]; + size_t h_result_offsets[num_result_offsets]; + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + resource_handle, (byte_t*)h_result_src, src, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + resource_handle, (byte_t*)h_result_dst, dst, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + resource_handle, (byte_t*)h_result_wgt, wgt, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + resource_handle, (byte_t*)h_result_offsets, offsets, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + TEST_ASSERT( + test_ret_value, (num_seeds + 1) == num_result_offsets, "number of offsets doesn't match"); + + for (int i = 0; (i < num_result_offsets) && (test_ret_value == 0); ++i) { + TEST_ASSERT( + test_ret_value, h_result_offsets[i] == h_expected_offsets[i], "offsets don't match"); + } + + weight_t M[num_vertices][num_vertices]; + + for (int i = 0; (i < num_seeds) && (test_ret_value == 0); ++i) { + for (int r = 0 ; r < num_vertices ; ++r) + for (int c = 0 ; c < num_vertices ; ++c) + M[r][c] = 0; + + for (size_t e = h_expected_offsets[i] ; e < h_expected_offsets[i+1] ; ++e) + M[h_expected_src[e]][h_expected_dst[e]] = 1; + + for (size_t e = h_result_offsets[i] ; (e < h_result_offsets[i+1]) && (test_ret_value == 0) ; ++e) { + TEST_ASSERT(test_ret_value, (M[h_result_src[e]][h_result_dst[e]] > 0), "found different edges"); + } + } + + cugraph_type_erased_device_array_view_free(src); + cugraph_type_erased_device_array_view_free(dst); + cugraph_type_erased_device_array_view_free(wgt); + cugraph_type_erased_device_array_view_free(offsets); + cugraph_induced_subgraph_result_free(result); + } + + cugraph_sg_graph_free(graph); + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int test_egonet(const cugraph_resource_handle_t *resource_handle) +{ + size_t num_edges = 9; + size_t num_vertices = 6; + size_t radius = 2; + size_t num_seeds = 2; + + vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 3, 4}; + vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 4, 5, 5}; + weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 6.1f}; + vertex_t h_seeds[] = {0, 1}; + + vertex_t h_result_src[] = {0, 1, 1, 3, 1, 1, 3, 3, 4}; + vertex_t h_result_dst[] = {1, 3, 4, 4, 3, 4, 4, 5, 5}; + size_t h_result_offsets[] = {0, 4, 9}; + + // Egonet wants store_transposed = FALSE + return generic_egonet_test(resource_handle, + h_src, + h_dst, + h_wgt, + h_seeds, + h_result_src, + h_result_dst, + h_result_offsets, + num_vertices, + num_edges, + num_seeds, + radius, + FALSE); +} + +/******************************************************************************/ + +int main(int argc, char** argv) +{ + // Set up MPI: + int comm_rank; + int comm_size; + int num_gpus_per_node; + cudaError_t status; + int mpi_status; + int result = 0; + cugraph_resource_handle_t* handle = NULL; + cugraph_error_t* ret_error; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + int prows = 1; + + C_MPI_TRY(MPI_Init(&argc, &argv)); + C_MPI_TRY(MPI_Comm_rank(MPI_COMM_WORLD, &comm_rank)); + C_MPI_TRY(MPI_Comm_size(MPI_COMM_WORLD, &comm_size)); + C_CUDA_TRY(cudaGetDeviceCount(&num_gpus_per_node)); + C_CUDA_TRY(cudaSetDevice(comm_rank % num_gpus_per_node)); + +#if 0 + // TODO: Need something a bit more sophisticated for bigger systems + prows = (int)sqrt((double)comm_size); + while (comm_size % prows != 0) { + --prows; + } + + ret_code = cugraph_resource_handle_init_comms(handle, prows, &ret_error); + TEST_ASSERT(result, ret_code == CUGRAPH_SUCCESS, "handle create failed."); + TEST_ASSERT(result, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); +#endif + + void* raft_handle = create_raft_handle(prows); + handle = cugraph_create_resource_handle(raft_handle); + + if (result == 0) { + result |= RUN_MG_TEST(test_egonet, handle); + + cugraph_free_resource_handle(handle); + } + + free_raft_handle(raft_handle); + + C_MPI_TRY(MPI_Finalize()); + + return result; +} diff --git a/cpp/tests/community/egonet_test.cpp b/cpp/tests/community/egonet_test.cpp new file mode 100644 index 00000000000..57dc1384f40 --- /dev/null +++ b/cpp/tests/community/egonet_test.cpp @@ -0,0 +1,199 @@ +/* + * Copyright (c) 2021-2022, 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 + +#include +#include +#include +#include + +#include +#include +#include +#include + +#include +#include +#include +#include + +#include + +struct Egonet_Usecase { + std::vector ego_sources_{}; + int32_t radius_{1}; + bool test_weighted_{false}; + bool check_correctness_{false}; +}; + +template +class Tests_Egonet : public ::testing::TestWithParam> { + public: + Tests_Egonet() {} + + static void SetUpTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(std::tuple const& param) + { + auto [egonet_usecase, input_usecase] = param; + + HighResClock hr_clock{}; + + auto n_streams = std::min(egonet_usecase.ego_sources_.size(), size_t{128}); + auto stream_pool = std::make_shared(n_streams); + raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); + + bool renumber = true; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_clock.start(); + } + + auto [graph, d_renumber_map_labels] = + cugraph::test::construct_graph( + handle, input_usecase, egonet_usecase.test_weighted_, renumber); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + double elapsed_time{0.0}; + hr_clock.stop(&elapsed_time); + std::cout << "construct_graph took " << elapsed_time * 1e-6 << " s.\n"; + } + + auto graph_view = graph.view(); + + rmm::device_uvector d_ego_sources(egonet_usecase.ego_sources_.size(), + handle.get_stream()); + + raft::update_device(d_ego_sources.data(), + egonet_usecase.ego_sources_.data(), + egonet_usecase.ego_sources_.size(), + handle.get_stream()); + + cugraph::renumber_ext_vertices(handle, + d_ego_sources.data(), + d_ego_sources.size(), + d_renumber_map_labels->data(), + graph_view.local_vertex_partition_range_first(), + graph_view.local_vertex_partition_range_last()); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_clock.start(); + } + + auto [d_ego_edgelist_src, d_ego_edgelist_dst, d_ego_edgelist_wgt, d_ego_edgelist_offsets] = + cugraph::extract_ego( + handle, + graph_view, + raft::device_span{d_ego_sources.data(), egonet_usecase.ego_sources_.size()}, + egonet_usecase.radius_); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + double elapsed_time{0.0}; + hr_clock.stop(&elapsed_time); + std::cout << "Egonet took " << elapsed_time * 1e-6 << " s.\n"; + } + + if (egonet_usecase.check_correctness_) { + auto [d_reference_src, d_reference_dst, d_reference_wgt, d_reference_offsets] = + cugraph::test::egonet_reference( + handle, + graph_view, + raft::device_span{d_ego_sources.data(), d_ego_sources.size()}, + egonet_usecase.radius_); + + cugraph::test::egonet_validate(handle, + d_ego_edgelist_src, + d_ego_edgelist_dst, + d_ego_edgelist_wgt, + d_ego_edgelist_offsets, + d_reference_src, + d_reference_dst, + d_reference_wgt, + d_reference_offsets); + } + } +}; + +using Tests_Egonet_File = Tests_Egonet; +using Tests_Egonet_Rmat = Tests_Egonet; + +TEST_P(Tests_Egonet_File, CheckInt32Int32Float) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_Egonet_Rmat, CheckInt32Int32Float) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +INSTANTIATE_TEST_SUITE_P( + simple_test, + Tests_Egonet_File, + ::testing::Combine( + ::testing::Values(Egonet_Usecase{std::vector{0}, 1, false, true}, + Egonet_Usecase{std::vector{0}, 2, false, true}, + Egonet_Usecase{std::vector{1}, 3, false, true}, + Egonet_Usecase{std::vector{10, 0, 5}, 2, false, true}, + Egonet_Usecase{std::vector{9, 3, 10}, 2, false, true}, + Egonet_Usecase{std::vector{5, 9, 3, 10, 12, 13}, 2, true, true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), + cugraph::test::File_Usecase("test/datasets/web-Google.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + file_benchmark_test, + Tests_Egonet_File, + ::testing::Combine( + ::testing::Values(Egonet_Usecase{std::vector{5, 9, 3, 10, 12, 13}, 2, true, false}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + rmat_small_test, + Tests_Egonet_Rmat, + // enable correctness checks + ::testing::Combine( + ::testing::Values(Egonet_Usecase{std::vector{0}, 1, false, true}, + Egonet_Usecase{std::vector{0}, 2, false, true}, + Egonet_Usecase{std::vector{1}, 3, false, true}, + Egonet_Usecase{std::vector{10, 0, 5}, 2, false, true}, + Egonet_Usecase{std::vector{9, 3, 10}, 2, false, true}, + Egonet_Usecase{std::vector{5, 9, 3, 10, 12, 13}, 2, 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_Egonet_Rmat, + // disable correctness checks for large graphs + ::testing::Combine( + ::testing::Values(Egonet_Usecase{std::vector{5, 9, 3, 10, 12, 13}, 2, true, false}), + ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); + +CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/community/egonet_test.cu b/cpp/tests/community/egonet_test.cu deleted file mode 100644 index 7a9343d3a84..00000000000 --- a/cpp/tests/community/egonet_test.cu +++ /dev/null @@ -1,279 +0,0 @@ -/* - * Copyright (c) 2021-2022, 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 -#include -#include - -#include -#include -#include - -#include -#include -#include -#include - -#include - -#include -#include -#include -#include -#include - -#include - -typedef struct InducedEgo_Usecase_t { - std::string graph_file_full_path{}; - std::vector ego_sources{}; - int32_t radius; - bool test_weighted{false}; - - InducedEgo_Usecase_t(std::string const& graph_file_path, - std::vector const& ego_sources, - int32_t radius, - bool test_weighted) - : ego_sources(ego_sources), radius(radius), test_weighted(test_weighted) - { - if ((graph_file_path.length() > 0) && (graph_file_path[0] != '/')) { - graph_file_full_path = cugraph::test::get_rapids_dataset_root_dir() + "/" + graph_file_path; - } else { - graph_file_full_path = graph_file_path; - } - }; -} InducedEgo_Usecase; - -class Tests_InducedEgo : public ::testing::TestWithParam { - public: - Tests_InducedEgo() {} - - static void SetUpTestCase() {} - static void TearDownTestCase() {} - - virtual void SetUp() {} - virtual void TearDown() {} - - template - void run_current_test(InducedEgo_Usecase const& configuration) - { - int n_streams = std::min(configuration.ego_sources.size(), static_cast(128)); - auto stream_pool = std::make_shared(n_streams); - raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); - - cugraph::graph_t graph(handle); - std::tie(graph, std::ignore) = cugraph::test:: - read_graph_from_matrix_market_file( - handle, configuration.graph_file_full_path, configuration.test_weighted, false); - auto graph_view = graph.view(); - - rmm::device_uvector d_ego_sources(configuration.ego_sources.size(), - handle.get_stream()); - - raft::update_device(d_ego_sources.data(), - configuration.ego_sources.data(), - configuration.ego_sources.size(), - handle.get_stream()); - - HighResTimer hr_timer; - hr_timer.start("egonet"); - cudaProfilerStart(); - auto [d_ego_edgelist_src, d_ego_edgelist_dst, d_ego_edgelist_weights, d_ego_edge_offsets] = - cugraph::extract_ego(handle, - graph_view, - d_ego_sources.data(), - static_cast(configuration.ego_sources.size()), - configuration.radius); - cudaProfilerStop(); - hr_timer.stop(); - hr_timer.display(std::cout); - auto h_cugraph_ego_edgelist_src = cugraph::test::to_host(handle, d_ego_edgelist_src); - auto h_cugraph_ego_edgelist_dst = cugraph::test::to_host(handle, d_ego_edgelist_dst); - auto h_cugraph_ego_edge_offsets = cugraph::test::to_host(handle, d_ego_edge_offsets); - ASSERT_TRUE(d_ego_edge_offsets.size() == (configuration.ego_sources.size() + 1)); - ASSERT_TRUE(d_ego_edgelist_src.size() == d_ego_edgelist_dst.size()); - if (configuration.test_weighted) - ASSERT_TRUE(d_ego_edgelist_src.size() == (*d_ego_edgelist_weights).size()); - ASSERT_TRUE(h_cugraph_ego_edge_offsets[configuration.ego_sources.size()] == - d_ego_edgelist_src.size()); - for (size_t i = 0; i < configuration.ego_sources.size(); i++) - ASSERT_TRUE(h_cugraph_ego_edge_offsets[i] <= h_cugraph_ego_edge_offsets[i + 1]); - auto n_vertices = graph_view.number_of_vertices(); - for (size_t i = 0; i < d_ego_edgelist_src.size(); i++) { - ASSERT_TRUE(cugraph::is_valid_vertex(n_vertices, h_cugraph_ego_edgelist_src[i])); - ASSERT_TRUE(cugraph::is_valid_vertex(n_vertices, h_cugraph_ego_edgelist_dst[i])); - } - } -}; - -TEST_P(Tests_InducedEgo, CheckInt32Int32FloatTransposeFalse) -{ - run_current_test(GetParam()); -} - -INSTANTIATE_TEST_SUITE_P( - simple_test, - Tests_InducedEgo, - ::testing::Values( - InducedEgo_Usecase("test/datasets/karate.mtx", std::vector{0}, 1, false), - InducedEgo_Usecase("test/datasets/karate.mtx", std::vector{0}, 2, false), - InducedEgo_Usecase("test/datasets/karate.mtx", std::vector{1}, 3, false), - InducedEgo_Usecase("test/datasets/karate.mtx", std::vector{10, 0, 5}, 2, false), - InducedEgo_Usecase("test/datasets/karate.mtx", std::vector{9, 3, 10}, 2, false), - InducedEgo_Usecase( - "test/datasets/karate.mtx", std::vector{5, 9, 3, 10, 12, 13}, 2, true))); - -// For perf analysis -/* -INSTANTIATE_TEST_SUITE_P( - simple_test, - Tests_InducedEgo, - ::testing::Values( - InducedEgo_Usecase("test/datasets/soc-LiveJournal1.mtx", std::vector{0}, 1, false), - InducedEgo_Usecase("test/datasets/soc-LiveJournal1.mtx", std::vector{0}, 2, false), - InducedEgo_Usecase("test/datasets/soc-LiveJournal1.mtx", std::vector{0}, 3, false), - InducedEgo_Usecase("test/datasets/soc-LiveJournal1.mtx", std::vector{0}, 4, false), - InducedEgo_Usecase("test/datasets/soc-LiveJournal1.mtx", std::vector{0}, 5, false), - InducedEgo_Usecase( - "test/datasets/soc-LiveJournal1.mtx", std::vector{363617}, 2, false), - InducedEgo_Usecase( - "test/datasets/soc-LiveJournal1.mtx", - std::vector{ - 363617, 722214, 2337449, 2510183, 2513389, 225853, 2035807, 3836330, 1865496, 28755}, - 2, - false), - InducedEgo_Usecase( - "test/datasets/soc-LiveJournal1.mtx", - std::vector{ - 363617, 722214, 2337449, 2510183, 2513389, 225853, 2035807, 3836330, 1865496, 28755, - 2536834, 3070144, 3888415, 3131712, 2382526, 1040771, 2631543, 4607218, 4465829, 3341686, - 2772973, 2611175, 4526129, 2624421, 1220593, 2593137, 3270705, 1503899, 1213033, 4840102, - 4529036, 3421116, 4264831, 4089751, 4272322, 3486998, 2830318, 320953, 2388331, 520808, - 3023094, 1600294, 3631119, 1716614, 4829213, 1175844, 960680, 847662, 3277365, 3957318, - 3455123, 2454259, 670953, 4465677, 1027332, 2560721, 89061, 1163406, 3109528, 3221856, - 4714426, 2382774, 37828, 4433616, 3283229, 591911, 4200188, 442522, 872207, 2437601, - 741003, 266241, 914618, 3626195, 2021080, 4679624, 777476, 2527796, 1114017, 640142, - 49259, 4069879, 3869098, 1105040, 4707804, 3208582, 3325885, 1450601, 4072548, 2037062, - 2029646, 4575891, 1488598, 79105, 4827273, 3795434, 4647518, 4733397, 3980718, 1184627}, - 2, - false), - InducedEgo_Usecase( - "test/datasets/soc-LiveJournal1.mtx", - std::vector{ - 363617, 722214, 2337449, 2510183, 2513389, 225853, 2035807, 3836330, 1865496, 28755, - 2536834, 3070144, 3888415, 3131712, 2382526, 1040771, 2631543, 4607218, 4465829, 3341686, - 2772973, 2611175, 4526129, 2624421, 1220593, 2593137, 3270705, 1503899, 1213033, 4840102, - 4529036, 3421116, 4264831, 4089751, 4272322, 3486998, 2830318, 320953, 2388331, 520808, - 3023094, 1600294, 3631119, 1716614, 4829213, 1175844, 960680, 847662, 3277365, 3957318, - 3455123, 2454259, 670953, 4465677, 1027332, 2560721, 89061, 1163406, 3109528, 3221856, - 4714426, 2382774, 37828, 4433616, 3283229, 591911, 4200188, 442522, 872207, 2437601, - 741003, 266241, 914618, 3626195, 2021080, 4679624, 777476, 2527796, 1114017, 640142, - 49259, 4069879, 3869098, 1105040, 4707804, 3208582, 3325885, 1450601, 4072548, 2037062, - 2029646, 4575891, 1488598, 79105, 4827273, 3795434, 4647518, 4733397, 3980718, 1184627, - 984983, 3114832, 1967741, 1599818, 144593, 2698770, 2889449, 2495550, 1053813, 1193622, - 686026, 3989015, 2040719, 4693428, 3190376, 2926728, 3399030, 1664419, 662429, 4526841, - 2186957, 3752558, 2440046, 2930226, 3633006, 4058166, 3137060, 3499296, 2126343, 148971, - 2199672, 275811, 2813976, 2274536, 1189239, 1335942, 2465624, 2596042, 829684, 193400, - 2682845, 3691697, 4022437, 4051170, 4195175, 2876420, 3984220, 2174475, 326134, 2606530, - 2493046, 4706121, 1498980, 4576225, 1271339, 44832, 1875673, 4664940, 134931, 736397, - 4333554, 2751031, 2163610, 2879676, 3174153, 3317403, 2052464, 1881883, 4757859, 3596257, - 2358088, 2578758, 447504, 590720, 1717038, 1869795, 1133885, 3027521, 840312, 2818881, - 3654321, 2730947, 353585, 1134903, 2223378, 1508824, 3662521, 1363776, 2712071, 288441, - 1204581, 3502242, 4645567, 2767267, 1514366, 3956099, 1422145, 1216608, 2253360, 189132, - 4238225, 1345783, 451571, 1599442, 3237284, 4711405, 929446, 1857675, 150759, 1277633, - 761210, 138628, 1026833, 2599544, 2464737, 989203, 3399615, 2144292, 216142, 637312, - 2044964, 716256, 1660632, 1762919, 4784357, 2213415, 2764769, 291806, 609772, 3264819, - 1870953, 1516385, 235647, 1045474, 2664957, 819095, 1824119, 4045271, 4448109, 1676788, - 4285177, 1580502, 3546548, 2771971, 3927086, 1339779, 3156204, 1730998, 1172522, 2433024, - 4533449, 479930, 2010695, 672994, 3542039, 3176455, 26352, 2137735, 866910, 4410835, - 2623982, 3603159, 2555625, 2765653, 267865, 2015523, 1009052, 4713994, 1600667, 2176195, - 3179631, 4570390, 2018424, 3356384, 1784287, 894861, 3622099, 1647273, 3044136, 950354, - 1491760, 3416929, 3757300, 2244912, 4129215, 1600848, 3867343, 72329, 919189, 992521, - 3445975, 4712557, 4680974, 188419, 2612093, 1991268, 3566207, 2281468, 3859078, 2492806, - 3398628, 763441, 2679107, 2554420, 2130132, 4664374, 1182901, 3890770, 4714667, 4209303, - 4013060, 3617653, 2040022, 3296519, 4190671, 1693353, 2678411, 3788834, 2781815, 191965, - 1083926, 503974, 3529226, 1650522, 1900976, 542080, 3423929, 3418905, 878165, 4701703, - 3022790, 4316365, 76365, 4053672, 1358185, 3830478, 4445661, 3210024, 1895915, 4541133, - 2938808, 562788, 3920065, 1458776, 4052046, 2967475, 1092809, 3203538, 159626, 3399464, - 214467, 3343982, 1811854, 3189045, 4272117, 4701563, 424807, 4341116, 760545, 4674683, - 1538018, 386762, 194237, 2162719, 1694433, 943728, 2389036, 2196653, 3085571, 1513424, - 3689413, 3278747, 4197291, 3324063, 3651090, 1737936, 2768803, 2768889, 3108096, 4311775, - 3569480, 886705, 733256, 2477493, 1735412, 2960895, 1983781, 1861797, 3566460, 4537673, - 1164093, 3499764, 4553071, 3518985, 847658, 918948, 2922351, 1056144, 652895, 1013195, - 780505, 1702928, 3562838, 1432719, 2405207, 1054920, 641647, 2240939, 3617702, 383165, - 652641, 879593, 1810739, 2096385, 4497865, 4768530, 1743968, 3582014, 1025009, 3002122, - 2422190, 527647, 1251821, 2571153, 4095874, 3705333, 3637407, 1385567, 4043855, 4041930, - 2433139, 1710383, 1127734, 4362316, 711588, 817839, 3214775, 910077, 1313768, 2382229, - 16864, 2081770, 3095420, 3195272, 548711, 2259860, 1167323, 2435974, 425238, 2085179, - 2630042, 2632881, 2867923, 3703565, 1037695, 226617, 4379130, 1541468, 3581937, 605965, - 1137674, 4655221, 4769963, 1394370, 4425315, 2990132, 2364485, 1561137, 2713384, 481509, - 2900382, 934766, 2986774, 1767669, 298593, 2502539, 139296, 3794229, 4002180, 4718138, - 2909238, 423691, 3023810, 2784924, 2760160, 1971980, 316683, 3828090, 3253691, 4839313, - 1203624, 584938, 3901482, 1747543, 1572737, 3533226, 774708, 1691195, 1037110, 1557763, - 225120, 4424243, 3524086, 1717663, 4332507, 3513592, 4274932, 1232118, 873498, 1416042, - 2488925, 111391, 4704545, 4492545, 445317, 1584812, 2187737, 2471948, 3731678, 219255, - 2282627, 2589971, 2372185, 4609096, 3673961, 2524410, 12823, 2437155, 3015974, 4188352, - 3184084, 3690756, 1222341, 1278376, 3652030, 4162647, 326548, 3930062, 3926100, 1551222, - 2722165, 4526695, 3997534, 4815513, 3139056, 2547644, 3028915, 4149092, 3656554, 2691582, - 2676699, 1878842, 260174, 3129900, 4379993, 182347, 2189338, 3783616, 2616666, 2596952, - 243007, 4179282, 2730, 1939894, 2332032, 3335636, 182332, 3112260, 2174584, 587481, - 4527368, 3154106, 3403059, 673206, 2150292, 446521, 1600204, 4819428, 2591357, 48490, - 2917012, 2285923, 1072926, 2824281, 4364250, 956033, 311938, 37251, 3729300, 2726300, - 644966, 1623020, 1419070, 4646747, 2417222, 2680238, 2561083, 1793801, 2349366, 339747, - 611366, 4684147, 4356907, 1277161, 4510381, 3218352, 4161658, 3200733, 1172372, 3997786, - 3169266, 3353418, 2248955, 2875885, 2365369, 498208, 2968066, 2681505, 2059048, 2097106, - 3607540, 1121504, 2016789, 1762605, 3138431, 866081, 3705757, 3833066, 2599788, 760816, - 4046672, 1544367, 2983906, 4842911, 209599, 1250954, 3333704, 561212, 4674336, 2831841, - 3690724, 2929360, 4830834, 1177524, 2487687, 3525137, 875283, 651241, 2110742, 1296646, - 1543739, 4349417, 2384725, 1931751, 1519208, 1520034, 3385008, 3219962, 734912, 170230, - 1741419, 729913, 2860117, 2362381, 1199807, 2424230, 177824, 125948, 2722701, 4687548, - 1140771, 3232742, 4522020, 4376360, 1125603, 590312, 2481884, 138951, 4086775, 615155, - 3395781, 4587272, 283209, 568470, 4296185, 4344150, 2454321, 2672602, 838828, 4051647, - 1709120, 3074610, 693235, 4356087, 3018806, 239410, 2431497, 691186, 766276, 4462126, - 859155, 2370304, 1571808, 1938673, 1694955, 3871296, 4245059, 3987376, 301524, 2512461, - 3410437, 3300380, 684922, 4581995, 3599557, 683515, 1850634, 3704678, 1937490, 2035591, - 3718533, 2065879, 3160765, 1467884, 1912241, 2501509, 3668572, 3390469, 2501150, 612319, - 713633, 1976262, 135946, 3641535, 632083, 13414, 4217765, 4137712, 2550250, 3281035, - 4179598, 961045, 2020694, 4380006, 1345936, 289162, 1359035, 770872, 4509911, 3947317, - 4719693, 248568, 2625660, 1237232, 2153208, 4814282, 1259954, 3677369, 861222, 2883506, - 3339149, 3998335, 491017, 1609022, 2648112, 742132, 649609, 4206953, 3131106, 3504814, - 3344486, 611721, 3215620, 2856233, 4447505, 1949222, 1868345, 712710, 6966, 4730666, - 3181872, 2972889, 3038521, 3525444, 4385208, 1845613, 1124187, 2030476, 4468651, 2478792, - 3473580, 3783357, 1852991, 1648485, 871319, 1670723, 4458328, 3218600, 1811100, 3443356, - 2233873, 3035207, 2548692, 3337891, 3773674, 1552957, 4782811, 3144712, 3523466, 1491315, - 3955852, 1838410, 3164028, 1092543, 776459, 2959379, 2541744, 4064418, 3908320, 2854145, - 3960709, 1348188, 977678, 853619, 1304291, 2848702, 1657913, 1319826, 3322665, 788037, - 2913686, 4471279, 1766285, 348304, 56570, 1892118, 4017244, 401006, 3524539, 4310134, - 1624693, 4081113, 957511, 849400, 129975, 2616130, 378537, 1556787, 3916162, 1039980, - 4407778, 2027690, 4213675, 839863, 683134, 75805, 2493150, 4215796, 81587, 751845, - 1255588, 1947964, 1950470, 859401, 3077088, 3931110, 2316256, 1523761, 4527477, 4237511, - 1123513, 4209796, 3584772, 4250563, 2091754, 1618766, 2139944, 4525352, 382159, 2955887, - 41760, 2313998, 496912, 3791570, 3904792, 3613654, 873959, 127076, 2537797, 2458107, - 4543265, 3661909, 26828, 271816, 17854, 2461269, 1776042, 1573899, 3409957, 4335712, - 4534313, 3392751, 1230124, 2159031, 4444015, 3373087, 3848014, 2026600, 1382747, 3537242, - 4536743, 4714155, 3788371, 3570849, 173741, 211962, 4377778, 119369, 2856973, 2945854, - 1508054, 4503932, 3141566, 1842177, 3448683, 3384614, 2886508, 1573965, 990618, 3053734, - 2918742, 4508753, 1032149, 60943, 4291620, 722607, 2883224, 169359, 4356585, 3725543, - 3678729, 341673, 3592828, 4077251, 3382936, 3885685, 4630994, 1286698, 4449616, 1138430, - 3113385, 4660578, 2539973, 4562286, 4085089, 494737, 3967610, 2130702, 1823755, 1369324, - 3796951, 956299, 141730, 935144, 4381893, 4412545, 1382250, 3024476, 2364546, 3396164, - 3573511, 314081, 577688, 4154135, 1567018, 4047761, 2446220, 1148833, 4842497, 3967186, - 1175290, 3749667, 1209593, 3295627, 3169065, 2460328, 1838486, 1436923, 2843887, 3676426, - 2079145, 2975635, 535071, 4287509, 3281107, 39606, 3115500, 3204573, 722131, 3124073}, - 2, - false))); -*/ -CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/community/egonet_validate.cu b/cpp/tests/community/egonet_validate.cu new file mode 100644 index 00000000000..c0b82f8036b --- /dev/null +++ b/cpp/tests/community/egonet_validate.cu @@ -0,0 +1,451 @@ +/* + * Copyright (c) 2022, 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 + +#include +#include + +#include + +#include +#include +#include +#include +#include +#include + +namespace cugraph { +namespace test { + +template +std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +egonet_reference(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span ego_sources, + int radius) +{ +#if 1 + auto [d_coo_src, d_coo_dst, d_coo_wgt] = cugraph::decompress_to_edgelist( + handle, graph_view, std::optional>{std::nullopt}); +#else + // FIXME: This should be faster (smaller list of edges to operate on), but uniform_nbr_sample + // doesn't preserve multi-edges (which is probably a bug) + std::vector fan_out(radius + 1, -1); + + rmm::device_uvector local_ego_sources(ego_sources.size(), handle.get_stream()); + raft::copy(local_ego_sources.data(), ego_sources.data(), ego_sources.size(), handle.get_stream()); + + auto [d_coo_src, d_coo_dst, d_coo_wgt, d_coo_counts] = cugraph::uniform_nbr_sample( + handle, + graph_view, + raft::device_span{local_ego_sources.data(), local_ego_sources.size()}, + raft::host_span{fan_out.data(), fan_out.size()}, + false); + + d_coo_counts.resize(0, handle.get_stream()); + d_coo_counts.shrink_to_fit(handle.get_stream()); +#endif + + rmm::device_uvector d_reference_src(0, handle.get_stream()); + rmm::device_uvector d_reference_dst(0, handle.get_stream()); + + auto d_reference_wgt = + graph_view.is_weighted() + ? std::make_optional(rmm::device_uvector(0, handle.get_stream())) + : std::nullopt; + + std::vector h_reference_offsets; + + size_t offset{0}; + + for (size_t idx = 0; idx < ego_sources.size(); ++idx) { + h_reference_offsets.push_back(offset); + + rmm::device_uvector frontier(1, handle.get_stream()); + rmm::device_uvector visited(1, handle.get_stream()); + raft::copy(frontier.data(), ego_sources.data() + idx, 1, handle.get_stream()); + raft::copy(visited.data(), ego_sources.data() + idx, 1, handle.get_stream()); + + for (int hop = 0; hop < radius; ++hop) { + size_t new_entries = thrust::count_if( + handle.get_thrust_policy(), + d_coo_src.begin(), + d_coo_src.end(), + [frontier_begin = frontier.begin(), + frontier_end = frontier.end()] __device__(vertex_t src) { + return thrust::binary_search(thrust::seq, frontier_begin, frontier_end, src); + }); + + size_t old_size = d_reference_src.size(); + + d_reference_src.resize(old_size + new_entries, handle.get_stream()); + d_reference_dst.resize(old_size + new_entries, handle.get_stream()); + + if (graph_view.is_weighted()) { + d_reference_wgt->resize(old_size + new_entries, handle.get_stream()); + + thrust::copy_if( + handle.get_thrust_policy(), + thrust::make_zip_iterator(d_coo_src.begin(), d_coo_dst.begin(), d_coo_wgt->begin()), + thrust::make_zip_iterator(d_coo_src.end(), d_coo_dst.end(), d_coo_wgt->end()), + thrust::make_zip_iterator(d_reference_src.begin() + old_size, + d_reference_dst.begin() + old_size, + d_reference_wgt->begin() + old_size), + [frontier_begin = frontier.begin(), + frontier_end = frontier.end()] __device__(auto tuple) { + vertex_t src = thrust::get<0>(tuple); + return thrust::binary_search(thrust::seq, frontier_begin, frontier_end, src); + }); + } else { + thrust::copy_if(handle.get_thrust_policy(), + thrust::make_zip_iterator(d_coo_src.begin(), d_coo_dst.begin()), + thrust::make_zip_iterator(d_coo_src.end(), d_coo_dst.end()), + thrust::make_zip_iterator(d_reference_src.begin() + old_size, + d_reference_dst.begin() + old_size), + [frontier_begin = frontier.begin(), + frontier_end = frontier.end()] __device__(auto tuple) { + vertex_t src = thrust::get<0>(tuple); + return thrust::binary_search( + thrust::seq, frontier_begin, frontier_end, src); + }); + } + + frontier.resize(new_entries, handle.get_stream()); + auto new_end = thrust::copy_if( + handle.get_thrust_policy(), + d_reference_dst.begin() + old_size, + d_reference_dst.end(), + frontier.begin(), + [visited_begin = visited.begin(), visited_end = visited.end()] __device__(auto v) { + return !thrust::binary_search(thrust::seq, visited_begin, visited_end, v); + }); + frontier.resize(thrust::distance(frontier.begin(), new_end), handle.get_stream()); + thrust::sort(handle.get_thrust_policy(), frontier.begin(), frontier.end()); + new_end = thrust::unique(handle.get_thrust_policy(), frontier.begin(), frontier.end()); + frontier.resize(thrust::distance(frontier.begin(), new_end), handle.get_stream()); + + size_t old_visited_size = visited.size(); + visited.resize(old_visited_size + frontier.size(), handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), + frontier.begin(), + frontier.end(), + visited.begin() + old_visited_size); + thrust::sort(handle.get_thrust_policy(), visited.begin(), visited.end()); + + offset += new_entries; + } + + if (frontier.size() > 0) { + size_t new_entries = thrust::count_if( + handle.get_thrust_policy(), + thrust::make_zip_iterator(d_coo_src.begin(), d_coo_dst.begin()), + thrust::make_zip_iterator(d_coo_src.end(), d_coo_dst.end()), + [frontier_begin = frontier.begin(), + frontier_end = frontier.end(), + visited_begin = visited.begin(), + visited_end = visited.end()] __device__(auto tuple) { + vertex_t src = thrust::get<0>(tuple); + vertex_t dst = thrust::get<1>(tuple); + return thrust::binary_search(thrust::seq, frontier_begin, frontier_end, src) && + thrust::binary_search(thrust::seq, visited_begin, visited_end, dst); + }); + + size_t old_size = d_reference_src.size(); + + d_reference_src.resize(old_size + new_entries, handle.get_stream()); + d_reference_dst.resize(old_size + new_entries, handle.get_stream()); + + if (graph_view.is_weighted()) { + d_reference_wgt->resize(old_size + new_entries, handle.get_stream()); + + thrust::copy_if( + handle.get_thrust_policy(), + thrust::make_zip_iterator(d_coo_src.begin(), d_coo_dst.begin(), d_coo_wgt->begin()), + thrust::make_zip_iterator(d_coo_src.end(), d_coo_dst.end(), d_coo_wgt->end()), + thrust::make_zip_iterator(d_reference_src.begin() + old_size, + d_reference_dst.begin() + old_size, + d_reference_wgt->begin() + old_size), + [frontier_begin = frontier.begin(), + frontier_end = frontier.end(), + visited_begin = visited.begin(), + visited_end = visited.end()] __device__(auto tuple) { + vertex_t src = thrust::get<0>(tuple); + vertex_t dst = thrust::get<1>(tuple); + return thrust::binary_search(thrust::seq, frontier_begin, frontier_end, src) && + thrust::binary_search(thrust::seq, visited_begin, visited_end, dst); + }); + } else { + thrust::copy_if( + handle.get_thrust_policy(), + thrust::make_zip_iterator(d_coo_src.begin(), d_coo_dst.begin()), + thrust::make_zip_iterator(d_coo_src.end(), d_coo_dst.end()), + thrust::make_zip_iterator(d_reference_src.begin() + old_size, + d_reference_dst.begin() + old_size), + [frontier_begin = frontier.begin(), + frontier_end = frontier.end(), + visited_begin = visited.begin(), + visited_end = visited.end()] __device__(auto tuple) { + vertex_t src = thrust::get<0>(tuple); + vertex_t dst = thrust::get<1>(tuple); + return thrust::binary_search(thrust::seq, frontier_begin, frontier_end, src) && + thrust::binary_search(thrust::seq, visited_begin, visited_end, dst); + }); + } + + offset += new_entries; + } + } + + h_reference_offsets.push_back(offset); + + auto d_reference_offsets = cugraph::test::to_device(handle, h_reference_offsets); + + return std::make_tuple(std::move(d_reference_src), + std::move(d_reference_dst), + std::move(d_reference_wgt), + std::move(d_reference_offsets)); +} + +template +void egonet_validate(raft::handle_t const& handle, + rmm::device_uvector& d_cugraph_egonet_src, + rmm::device_uvector& d_cugraph_egonet_dst, + std::optional>& d_cugraph_egonet_wgt, + rmm::device_uvector& d_cugraph_egonet_offsets, + rmm::device_uvector& d_reference_egonet_src, + rmm::device_uvector& d_reference_egonet_dst, + std::optional>& d_reference_egonet_wgt, + rmm::device_uvector& d_reference_egonet_offsets) +{ + ASSERT_EQ(d_reference_egonet_offsets.size(), d_cugraph_egonet_offsets.size()) + << "Returned edge offset vector has an invalid size."; + + ASSERT_TRUE(thrust::equal(handle.get_thrust_policy(), + d_reference_egonet_offsets.begin(), + d_reference_egonet_offsets.end(), + d_cugraph_egonet_offsets.begin())) + << "Returned egonet edge offset values do not match with the reference values."; + + ASSERT_EQ(d_cugraph_egonet_wgt.has_value(), d_reference_egonet_wgt.has_value()); + + auto h_offsets = + cugraph::test::to_host(handle, + raft::device_span{d_cugraph_egonet_offsets.data(), + d_cugraph_egonet_offsets.size()}); + + for (size_t i = 0; i < (h_offsets.size() - 1); ++i) { + if (d_reference_egonet_wgt) { + thrust::sort(handle.get_thrust_policy(), + thrust::make_zip_iterator(d_reference_egonet_src.begin(), + d_reference_egonet_dst.begin(), + d_reference_egonet_wgt->begin()) + + h_offsets[i], + thrust::make_zip_iterator(d_reference_egonet_src.begin(), + d_reference_egonet_dst.begin(), + d_reference_egonet_wgt->begin()) + + h_offsets[i + 1]); + thrust::sort(handle.get_thrust_policy(), + thrust::make_zip_iterator(d_cugraph_egonet_src.begin(), + d_cugraph_egonet_dst.begin(), + d_cugraph_egonet_wgt->begin()) + + h_offsets[i], + thrust::make_zip_iterator(d_cugraph_egonet_src.begin(), + d_cugraph_egonet_dst.begin(), + d_cugraph_egonet_wgt->begin()) + + h_offsets[i + 1]); + + ASSERT_TRUE(thrust::equal(handle.get_thrust_policy(), + thrust::make_zip_iterator(d_reference_egonet_src.begin(), + d_reference_egonet_dst.begin(), + d_reference_egonet_wgt->begin()) + + h_offsets[i], + thrust::make_zip_iterator(d_reference_egonet_src.begin(), + d_reference_egonet_dst.begin(), + d_reference_egonet_wgt->begin()) + + h_offsets[i + 1], + thrust::make_zip_iterator(d_cugraph_egonet_src.begin(), + d_cugraph_egonet_dst.begin(), + d_cugraph_egonet_wgt->begin()) + + h_offsets[i], + [] __device__(auto left, auto right) { + auto l0 = thrust::get<0>(left); + auto l1 = thrust::get<1>(left); + auto l2 = thrust::get<2>(left); + auto r0 = thrust::get<0>(right); + auto r1 = thrust::get<1>(right); + auto r2 = thrust::get<2>(right); + if (!((l0 == r0) && (l1 == r1) && (l2 == r2))) + printf("edge mismatch: (%d,%d,%g) != (%d,%d,%g)\n", + (int)l0, + (int)l1, + (float)l2, + (int)r0, + (int)r1, + (float)r2); + return (l0 == r0) && (l1 == r1) && (l2 == r2); + })) + << "Extracted egonet edges do not match with the edges extracted by the reference " + "implementation."; + } else { + thrust::sort( + handle.get_thrust_policy(), + thrust::make_zip_iterator(d_reference_egonet_src.begin(), d_reference_egonet_dst.begin()) + + h_offsets[i], + thrust::make_zip_iterator(d_reference_egonet_src.begin(), d_reference_egonet_dst.begin()) + + h_offsets[i + 1]); + thrust::sort( + handle.get_thrust_policy(), + thrust::make_zip_iterator(d_cugraph_egonet_src.begin(), d_cugraph_egonet_dst.begin()) + + h_offsets[i], + thrust::make_zip_iterator(d_cugraph_egonet_src.begin(), d_cugraph_egonet_dst.begin()) + + h_offsets[i + 1]); + + ASSERT_TRUE(thrust::equal( + handle.get_thrust_policy(), + thrust::make_zip_iterator(d_reference_egonet_src.begin(), d_reference_egonet_dst.begin()) + + h_offsets[i], + thrust::make_zip_iterator(d_reference_egonet_src.begin(), d_reference_egonet_dst.begin()) + + h_offsets[i + 1], + thrust::make_zip_iterator(d_cugraph_egonet_src.begin(), d_cugraph_egonet_dst.begin()) + + h_offsets[i], + [] __device__(auto left, auto right) { + auto l0 = thrust::get<0>(left); + auto l1 = thrust::get<1>(left); + auto r0 = thrust::get<0>(right); + auto r1 = thrust::get<1>(right); + if (!((l0 == r0) && (l1 == r1))) + printf("edge mismatch: (%d,%d) != (%d,%d)\n", (int)l0, (int)l1, (int)r0, (int)r1); + return (l0 == r0) && (l1 == r1); + })) + << "Extracted egonet edges do not match with the edges extracted by the reference " + "implementation."; + } + } +} + +template std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +egonet_reference(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span ego_sources, + int radius); + +template void egonet_validate(raft::handle_t const& handle, + rmm::device_uvector& d_cugraph_egonet_src, + rmm::device_uvector& d_cugraph_egonet_dst, + std::optional>& d_cugraph_egonet_wgt, + rmm::device_uvector& d_cugraph_egonet_offsets, + rmm::device_uvector& d_reference_egonet_src, + rmm::device_uvector& d_reference_egonet_dst, + std::optional>& d_reference_egonet_wgt, + rmm::device_uvector& d_reference_egonet_offsets); + +template void egonet_validate(raft::handle_t const& handle, + rmm::device_uvector& d_cugraph_egonet_src, + rmm::device_uvector& d_cugraph_egonet_dst, + std::optional>& d_cugraph_egonet_wgt, + rmm::device_uvector& d_cugraph_egonet_offsets, + rmm::device_uvector& d_reference_egonet_src, + rmm::device_uvector& d_reference_egonet_dst, + std::optional>& d_reference_egonet_wgt, + rmm::device_uvector& d_reference_egonet_offsets); + +#if 0 +// NOT CURRENTLY USED +template +std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +egonet_reference(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span ego_sources, + int radius); + +template +std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +egonet_reference(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span ego_sources, + int radius); + +template +std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +egonet_reference(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span ego_sources, + int radius); + +template +std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +egonet_reference(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span ego_sources, + int radius); + +template +std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +egonet_reference(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span ego_sources, + int radius); + +template +void egonet_validate( + raft::handle_t const& handle, + rmm::device_uvector& d_cugraph_egonet_src, + rmm::device_uvector& d_cugraph_egonet_dst, + std::optional>& d_cugraph_egonet_wgt, + rmm::device_uvector& d_cugraph_egonet_offsets, + rmm::device_uvector& d_reference_egonet_src, + rmm::device_uvector& d_reference_egonet_dst, + std::optional>& d_reference_egonet_wgt, + rmm::device_uvector& d_reference_egonet_offsets); + +template +void egonet_validate( + raft::handle_t const& handle, + rmm::device_uvector& d_cugraph_egonet_src, + rmm::device_uvector& d_cugraph_egonet_dst, + std::optional>& d_cugraph_egonet_wgt, + rmm::device_uvector& d_cugraph_egonet_offsets, + rmm::device_uvector& d_reference_egonet_src, + rmm::device_uvector& d_reference_egonet_dst, + std::optional>& d_reference_egonet_wgt, + rmm::device_uvector& d_reference_egonet_offsets); + +#endif + +} // namespace test +} // namespace cugraph diff --git a/cpp/tests/community/egonet_validate.hpp b/cpp/tests/community/egonet_validate.hpp new file mode 100644 index 00000000000..86046fb6a2f --- /dev/null +++ b/cpp/tests/community/egonet_validate.hpp @@ -0,0 +1,50 @@ +/* + * Copyright (c) 2022, 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 + +#include +#include + +#include + +#include + +namespace cugraph { +namespace test { + +template +std::tuple, + rmm::device_uvector, + std::optional>, + rmm::device_uvector> +egonet_reference(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span ego_sources, + int radius); + +template +void egonet_validate(raft::handle_t const& handle, + rmm::device_uvector& d_cugraph_egonet_src, + rmm::device_uvector& d_cugraph_egonet_dst, + std::optional>& d_cugraph_egonet_wgt, + rmm::device_uvector& d_cugraph_egonet_offsets, + rmm::device_uvector& d_reference_egonet_src, + rmm::device_uvector& d_reference_egonet_dst, + std::optional>& d_reference_egonet_wgt, + rmm::device_uvector& d_reference_egonet_offsets); + +} // namespace test +} // namespace cugraph diff --git a/cpp/tests/community/mg_egonet_test.cpp b/cpp/tests/community/mg_egonet_test.cpp new file mode 100644 index 00000000000..ed79ac8232b --- /dev/null +++ b/cpp/tests/community/mg_egonet_test.cpp @@ -0,0 +1,338 @@ +/* + * Copyright (c) 2021-2022, 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 "egonet_validate.hpp" + +#include +#include +#include +#include +#include + +#include +#include +#include + +#include +#include +#include +#include + +#include +#include +#include + +#include + +struct Egonet_Usecase { + std::vector ego_sources_{}; + int32_t radius_{1}; + bool test_weighted_{false}; + bool check_correctness_{false}; +}; + +template +class Tests_MGEgonet + : public ::testing::TestWithParam> { + public: + static void SetUpTestCase() { handle_ = cugraph::test::initialize_mg_handle(128); } + static void TearDownTestCase() { handle_.reset(); } + + // Run once for each test instance + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(std::tuple const& param) + { + auto [egonet_usecase, input_usecase] = param; + + HighResClock hr_clock{}; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle_->get_comms().barrier(); + hr_clock.start(); + } + + auto [mg_graph, d_renumber_map_labels] = + cugraph::test::construct_graph( + *handle_, input_usecase, egonet_usecase.test_weighted_, true); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle_->get_comms().barrier(); + double elapsed_time{0.0}; + hr_clock.stop(&elapsed_time); + std::cout << "MG construct_graph took " << elapsed_time * 1e-6 << " s.\n"; + } + + auto mg_graph_view = mg_graph.view(); + + int my_rank = handle_->get_comms().get_rank(); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle_->get_comms().barrier(); + hr_clock.start(); + } + + rmm::device_uvector d_ego_sources(0, handle_->get_stream()); + + if (my_rank == 0) { + d_ego_sources.resize(egonet_usecase.ego_sources_.size(), handle_->get_stream()); + + if constexpr (std::is_same::value) { + raft::update_device(d_ego_sources.data(), + egonet_usecase.ego_sources_.data(), + egonet_usecase.ego_sources_.size(), + handle_->get_stream()); + } else { + std::vector h_ego_sources(d_ego_sources.size()); + std::transform(egonet_usecase.ego_sources_.begin(), + egonet_usecase.ego_sources_.end(), + h_ego_sources.begin(), + [](auto v) { return static_cast(v); }); + raft::update_device( + d_ego_sources.data(), h_ego_sources.data(), h_ego_sources.size(), handle_->get_stream()); + } + } + + cugraph::renumber_ext_vertices( + *handle_, + d_ego_sources.data(), + d_ego_sources.size(), + d_renumber_map_labels->data(), + mg_graph_view.local_vertex_partition_range_first(), + mg_graph_view.local_vertex_partition_range_last()); + + d_ego_sources = cugraph::detail::shuffle_int_vertices_by_gpu_id( + *handle_, std::move(d_ego_sources), mg_graph_view.vertex_partition_range_lasts()); + + auto [d_ego_edgelist_src, d_ego_edgelist_dst, d_ego_edgelist_wgt, d_ego_edgelist_offsets] = + cugraph::extract_ego( + *handle_, + mg_graph_view, + raft::device_span{d_ego_sources.data(), egonet_usecase.ego_sources_.size()}, + static_cast(egonet_usecase.radius_)); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle_->get_comms().barrier(); + double elapsed_time{0.0}; + hr_clock.stop(&elapsed_time); + std::cout << "MG Egonet took " << elapsed_time * 1e-6 << " s.\n"; + } + + if (egonet_usecase.check_correctness_) { + *d_renumber_map_labels = cugraph::test::device_gatherv( + *handle_, + raft::device_span(d_renumber_map_labels->data(), + d_renumber_map_labels->size())); + + d_ego_edgelist_src = cugraph::test::device_gatherv( + *handle_, + raft::device_span(d_ego_edgelist_src.data(), d_ego_edgelist_src.size())); + d_ego_edgelist_dst = cugraph::test::device_gatherv( + *handle_, + raft::device_span(d_ego_edgelist_dst.data(), d_ego_edgelist_dst.size())); + + if (d_ego_edgelist_wgt) { + *d_ego_edgelist_wgt = + cugraph::test::device_gatherv(*handle_, + raft::device_span( + d_ego_edgelist_wgt->data(), d_ego_edgelist_wgt->size())); + } + + d_ego_edgelist_offsets = cugraph::test::device_gatherv( + *handle_, + raft::device_span(d_ego_edgelist_offsets.data(), + d_ego_edgelist_offsets.size())); + + auto [sg_graph, sg_number_map] = + cugraph::test::mg_graph_to_sg_graph(*handle_, mg_graph_view, d_renumber_map_labels, false); + + if (my_rank == 0) { + cugraph::unrenumber_int_vertices( + *handle_, + d_ego_edgelist_src.data(), + d_ego_edgelist_src.size(), + d_renumber_map_labels->data(), + std::vector{mg_graph_view.number_of_vertices()}); + + cugraph::unrenumber_int_vertices( + *handle_, + d_ego_edgelist_dst.data(), + d_ego_edgelist_dst.size(), + d_renumber_map_labels->data(), + std::vector{mg_graph_view.number_of_vertices()}); + + rmm::device_uvector d_sg_ego_sources(egonet_usecase.ego_sources_.size(), + handle_->get_stream()); + + if constexpr (std::is_same::value) { + raft::update_device(d_sg_ego_sources.data(), + egonet_usecase.ego_sources_.data(), + egonet_usecase.ego_sources_.size(), + handle_->get_stream()); + } else { + std::vector h_ego_sources(d_sg_ego_sources.size()); + std::transform(egonet_usecase.ego_sources_.begin(), + egonet_usecase.ego_sources_.end(), + h_ego_sources.begin(), + [](auto v) { return static_cast(v); }); + raft::update_device(d_sg_ego_sources.data(), + h_ego_sources.data(), + h_ego_sources.size(), + handle_->get_stream()); + } + + auto [d_reference_src, d_reference_dst, d_reference_wgt, d_reference_offsets] = + cugraph::extract_ego( + *handle_, + sg_graph.view(), + raft::device_span{d_sg_ego_sources.data(), d_sg_ego_sources.size()}, + static_cast(egonet_usecase.radius_)); + + cugraph::test::egonet_validate(*handle_, + d_ego_edgelist_src, + d_ego_edgelist_dst, + d_ego_edgelist_wgt, + d_ego_edgelist_offsets, + d_reference_src, + d_reference_dst, + d_reference_wgt, + d_reference_offsets); + } + } + } + + private: + static std::unique_ptr handle_; +}; + +template +std::unique_ptr Tests_MGEgonet::handle_ = nullptr; + +using Tests_MGEgonet_File = Tests_MGEgonet; +using Tests_MGEgonet_File64 = Tests_MGEgonet; +using Tests_MGEgonet_Rmat = Tests_MGEgonet; +using Tests_MGEgonet_Rmat64 = Tests_MGEgonet; + +TEST_P(Tests_MGEgonet_File, CheckInt32Int32Float) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGEgonet_File64, CheckInt64Int64Float) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGEgonet_Rmat, CheckInt32Int32Float) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGEgonet_Rmat64, CheckInt64Int64Float) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +INSTANTIATE_TEST_SUITE_P( + simple_file_test, + Tests_MGEgonet_File, + ::testing::Combine( + // enable correctness checks for small graphs + ::testing::Values(Egonet_Usecase{std::vector{0}, 1, false, true}, + Egonet_Usecase{std::vector{0}, 2, false, true}, + Egonet_Usecase{std::vector{0}, 3, false, true}, + Egonet_Usecase{std::vector{10, 0, 5}, 2, false, true}, + Egonet_Usecase{std::vector{9, 3, 10}, 2, false, true}, + Egonet_Usecase{std::vector{5, 9, 3, 10, 12, 13}, 2, true, true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), + cugraph::test::File_Usecase("test/datasets/dolphins.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + simple_rmat_test, + Tests_MGEgonet_Rmat, + ::testing::Combine( + // enable correctness checks for small graphs + ::testing::Values(Egonet_Usecase{std::vector{0}, 1, false, true}, + Egonet_Usecase{std::vector{0}, 2, false, true}, + Egonet_Usecase{std::vector{0}, 3, false, true}), + ::testing::Values(cugraph::test::Rmat_Usecase(10, 16, 0.57, 0.19, 0.19, 0, true, false)))); + +INSTANTIATE_TEST_SUITE_P( + file_benchmark_test, /* note that the test filename can be overridden in benchmarking (with + --gtest_filter to select only the file_benchmark_test with a specific + vertex & edge type combination) by command line arguments and do not + include more than one File_Usecase that differ only in filename + (to avoid running same benchmarks more than once) */ + Tests_MGEgonet_File, + ::testing::Combine( + // disable correctness checks for large graphs + ::testing::Values(Egonet_Usecase{std::vector{5, 9, 3, 10, 12, 13}, 2, true, true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + file64_benchmark_test, /* note that the test filename can be overridden in benchmarking (with + --gtest_filter to select only the file_benchmark_test with a specific + vertex & edge type combination) by command line arguments and do not + include more than one File_Usecase that differ only in filename + (to avoid running same benchmarks more than once) */ + Tests_MGEgonet_File64, + ::testing::Combine( + // disable correctness checks for large graphs + ::testing::Values(Egonet_Usecase{std::vector{5, 9, 3, 10, 12, 13}, 2, true, true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +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_MGEgonet_Rmat, + ::testing::Combine( + // disable correctness checks for large graphs + ::testing::Values(Egonet_Usecase{std::vector{0}, 1, false, true}, + Egonet_Usecase{std::vector{0}, 2, false, true}, + Egonet_Usecase{std::vector{0}, 3, false, true}, + Egonet_Usecase{std::vector{10, 0, 5}, 2, false, true}, + Egonet_Usecase{std::vector{9, 3, 10}, 2, false, true}, + Egonet_Usecase{std::vector{5, 9, 3, 10, 12, 13}, 2, true, true}), + ::testing::Values( + cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, true, false, 0, true)))); + +INSTANTIATE_TEST_SUITE_P( + rmat64_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_MGEgonet_Rmat64, + ::testing::Combine( + // disable correctness checks for large graphs + ::testing::Values(Egonet_Usecase{std::vector{5, 9, 3, 10, 12, 13}, 2, true, true}), + ::testing::Values( + cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, true, false, 0, true)))); + +CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/utilities/mg_utilities.cpp b/cpp/tests/utilities/mg_utilities.cpp index be98c96fc1b..ae1152e16ab 100644 --- a/cpp/tests/utilities/mg_utilities.cpp +++ b/cpp/tests/utilities/mg_utilities.cpp @@ -44,12 +44,10 @@ int query_mpi_comm_world_size() return comm_size; } -std::unique_ptr initialize_mg_handle() +std::unique_ptr initialize_mg_handle(size_t pool_size) { std::unique_ptr handle{nullptr}; - auto constexpr pool_size = 64; // FIXME: tuning parameter - handle = std::make_unique(rmm::cuda_stream_per_thread, std::make_shared(pool_size)); diff --git a/cpp/tests/utilities/mg_utilities.hpp b/cpp/tests/utilities/mg_utilities.hpp index 88212b3420e..26c29be2aa1 100644 --- a/cpp/tests/utilities/mg_utilities.hpp +++ b/cpp/tests/utilities/mg_utilities.hpp @@ -29,7 +29,7 @@ void finalize_mpi(); int query_mpi_comm_world_rank(); int query_mpi_comm_world_size(); -std::unique_ptr initialize_mg_handle(); +std::unique_ptr initialize_mg_handle(size_t pool_size = 64); // NCCL lazily initializes for P2P, and this enforces P2P initialization for better performance // measurements