diff --git a/.gitignore b/.gitignore index 485b71cac68..8ae12d147f9 100644 --- a/.gitignore +++ b/.gitignore @@ -11,6 +11,7 @@ __pycache__ .lock *.swp *.pytest_cache +*~ DartConfiguration.tcl .DS_Store diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 427188b4c56..6b3308f22e6 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -230,6 +230,8 @@ add_library(cugraph src/link_analysis/pagerank_mg.cu src/centrality/katz_centrality_sg.cu src/centrality/katz_centrality_mg.cu + src/centrality/eigenvector_centrality_sg.cu + src/centrality/eigenvector_centrality_mg.cu src/serialization/serializer.cu src/tree/mst.cu src/components/weakly_connected_components_sg.cu diff --git a/cpp/include/cugraph/algorithms.hpp b/cpp/include/cugraph/algorithms.hpp index c02200b5495..eef2c1c5ed3 100644 --- a/cpp/include/cugraph/algorithms.hpp +++ b/cpp/include/cugraph/algorithms.hpp @@ -1221,6 +1221,40 @@ void pagerank(raft::handle_t const& handle, bool has_initial_guess = false, bool do_expensive_check = false); +/** + * @brief Compute Eigenvector Centrality scores. + * + * This function computes eigenvector centrality scores using the power method. + * + * @throws cugraph::logic_error on erroneous input arguments or if fails to converge before @p + * max_iterations. + * + * @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. + * @param graph_view Graph view object. + * @param initial_centralities Optional device span containing initial values for the eigenvector + * centralities + * @param epsilon Error tolerance to check convergence. Convergence is assumed if the sum of the + * differences in eigenvector centrality values between two consecutive iterations is less than the + * number of vertices in the graph multiplied by @p epsilon. + * @param max_iterations Maximum number of power iterations. + * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`). + * @return device vector containing the centralities. + */ +template +rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + weight_t epsilon, + size_t max_iterations = 500, + bool do_expensive_check = false); + /** * @brief Compute HITS scores. * diff --git a/cpp/include/cugraph/detail/utility_wrappers.hpp b/cpp/include/cugraph/detail/utility_wrappers.hpp index f5ac0c3a3f2..a97be447ec0 100644 --- a/cpp/include/cugraph/detail/utility_wrappers.hpp +++ b/cpp/include/cugraph/detail/utility_wrappers.hpp @@ -45,19 +45,6 @@ void uniform_random_fill(rmm::cuda_stream_view const& stream_view, value_t max_value, uint64_t seed); -/** - * @brief Normalize the values in an array - * - * @tparam value_t type of the value to operate on - * - * @param[in] stream_view stream view - * @param[out] d_value device array to reduce - * @param[in] size number of elements in array - * - */ -template -void normalize(rmm::cuda_stream_view const& stream_view, value_t* d_value, size_t size); - /** * @brief Fill a buffer with a sequence of values * diff --git a/cpp/src/c_api/eigenvector_centrality.cpp b/cpp/src/c_api/eigenvector_centrality.cpp index 3828f42cfd2..c78079a98a8 100644 --- a/cpp/src/c_api/eigenvector_centrality.cpp +++ b/cpp/src/c_api/eigenvector_centrality.cpp @@ -78,22 +78,12 @@ struct eigenvector_centrality_functor : public cugraph::c_api::abstract_functor auto number_map = reinterpret_cast*>(graph_->number_map_); - rmm::device_uvector centralities(graph_view.local_vertex_partition_range_size(), - handle_.get_stream()); - - // FIXME: For now we'll call pagerank which returns a similarly formatted thing - cugraph::pagerank( + auto centralities = cugraph::eigenvector_centrality( handle_, graph_view, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - centralities.data(), - weight_t{0.95}, + std::optional>{}, static_cast(epsilon_), max_iterations_, - false, do_expensive_check_); rmm::device_uvector vertex_ids(graph_view.local_vertex_partition_range_size(), diff --git a/cpp/src/centrality/eigenvector_centrality_impl.cuh b/cpp/src/centrality/eigenvector_centrality_impl.cuh new file mode 100644 index 00000000000..dd5588e6283 --- /dev/null +++ b/cpp/src/centrality/eigenvector_centrality_impl.cuh @@ -0,0 +1,166 @@ +/* + * 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. + */ +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +#include +#include +#include +#include +#include + +namespace cugraph { +namespace detail { + +template +rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& pull_graph_view, + std::optional> initial_centralities, + weight_t epsilon, + size_t max_iterations, + bool do_expensive_check) +{ + using GraphViewType = graph_view_t; + auto const num_vertices = pull_graph_view.number_of_vertices(); + if (num_vertices == 0) { return rmm::device_uvector(0, handle.get_stream()); } + + if (do_expensive_check) { + if (pull_graph_view.is_weighted()) { + auto num_nonpositive_edge_weights = + count_if_e(handle, + pull_graph_view, + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), + [] __device__(vertex_t, vertex_t, weight_t w, auto, auto) { return w <= 0.0; }); + CUGRAPH_EXPECTS(num_nonpositive_edge_weights == 0, + "Invalid input argument: input graph should have postive edge weights."); + } + } + + rmm::device_uvector centralities(pull_graph_view.local_vertex_partition_range_size(), + handle.get_stream()); + if (initial_centralities) { + thrust::copy(handle.get_thrust_policy(), + initial_centralities->begin(), + initial_centralities->end(), + centralities.begin()); + } else { + thrust::fill(handle.get_thrust_policy(), + centralities.begin(), + centralities.end(), + weight_t{1.0} / static_cast(num_vertices)); + } + + // Power iteration + rmm::device_uvector old_centralities(centralities.size(), handle.get_stream()); + + edge_partition_src_property_t edge_partition_src_centralities( + handle, pull_graph_view); + + size_t iter{0}; + while (true) { + thrust::copy(handle.get_thrust_policy(), + centralities.begin(), + centralities.end(), + old_centralities.data()); + + update_edge_partition_src_property( + handle, pull_graph_view, centralities.begin(), edge_partition_src_centralities); + + per_v_transform_reduce_incoming_e( + handle, + pull_graph_view, + edge_partition_src_centralities.device_view(), + dummy_property_t{}.device_view(), + [] __device__(vertex_t, vertex_t, weight_t w, auto src_val, auto) { return src_val * w; }, + weight_t{0}, + centralities.begin()); + + // Normalize the centralities + auto hypotenuse = sqrt(transform_reduce_v( + handle, + pull_graph_view, + centralities.begin(), + [] __device__(auto, auto val) { return val * val; }, + weight_t{0.0})); + + thrust::transform(handle.get_thrust_policy(), + centralities.begin(), + centralities.end(), + centralities.begin(), + [hypotenuse] __device__(auto val) { return val / hypotenuse; }); + + auto diff_sum = transform_reduce_v( + handle, + pull_graph_view, + thrust::make_zip_iterator(thrust::make_tuple(centralities.begin(), old_centralities.data())), + [] __device__(auto, auto val) { return std::abs(thrust::get<0>(val) - thrust::get<1>(val)); }, + weight_t{0.0}); + + iter++; + + if (diff_sum < (pull_graph_view.number_of_vertices() * epsilon)) { + break; + } else if (iter >= max_iterations) { + CUGRAPH_FAIL("Eigenvector Centrality failed to converge."); + } + } + + return centralities; +} + +} // namespace detail + +template +rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + weight_t epsilon, + size_t max_iterations, + bool do_expensive_check) +{ + static_assert(std::is_integral::value, + "GraphViewType::vertex_type should be integral."); + static_assert(std::is_floating_point::value, + "weight_t should be a floating-point type."); + + CUGRAPH_EXPECTS(epsilon >= 0.0, "Invalid input argument: epsilon should be non-negative."); + if (initial_centralities) + CUGRAPH_EXPECTS(initial_centralities->size() == + static_cast(graph_view.local_vertex_partition_range_size()), + "Centralities should be same size as vertex range"); + + return detail::eigenvector_centrality( + handle, graph_view, initial_centralities, epsilon, max_iterations, do_expensive_check); +} + +} // namespace cugraph diff --git a/cpp/src/centrality/eigenvector_centrality_mg.cu b/cpp/src/centrality/eigenvector_centrality_mg.cu new file mode 100644 index 00000000000..09b3d1b3058 --- /dev/null +++ b/cpp/src/centrality/eigenvector_centrality_mg.cu @@ -0,0 +1,69 @@ +/* + * 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 + +namespace cugraph { + +// MG instantiation +template rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + float epsilon, + size_t max_iterations, + bool do_expensive_check); + +template rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + float epsilon, + size_t max_iterations, + bool do_expensive_check); + +template rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + float epsilon, + size_t max_iterations, + bool do_expensive_check); + +template rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + double epsilon, + size_t max_iterations, + bool do_expensive_check); + +template rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + double epsilon, + size_t max_iterations, + bool do_expensive_check); + +template rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + double epsilon, + size_t max_iterations, + bool do_expensive_check); + +} // namespace cugraph diff --git a/cpp/src/centrality/eigenvector_centrality_sg.cu b/cpp/src/centrality/eigenvector_centrality_sg.cu new file mode 100644 index 00000000000..3d025638d6d --- /dev/null +++ b/cpp/src/centrality/eigenvector_centrality_sg.cu @@ -0,0 +1,69 @@ +/* + * 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 + +namespace cugraph { + +// SG instantiation +template rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + float epsilon, + size_t max_iterations, + bool do_expensive_check); + +template rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + float epsilon, + size_t max_iterations, + bool do_expensive_check); + +template rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + float epsilon, + size_t max_iterations, + bool do_expensive_check); + +template rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + double epsilon, + size_t max_iterations, + bool do_expensive_check); + +template rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + double epsilon, + size_t max_iterations, + bool do_expensive_check); + +template rmm::device_uvector eigenvector_centrality( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> initial_centralities, + double epsilon, + size_t max_iterations, + bool do_expensive_check); + +} // namespace cugraph diff --git a/cpp/src/detail/utility_wrappers.cu b/cpp/src/detail/utility_wrappers.cu index b66523d4d59..ba3881dac08 100644 --- a/cpp/src/detail/utility_wrappers.cu +++ b/cpp/src/detail/utility_wrappers.cu @@ -97,20 +97,6 @@ template int64_t compute_maximum_vertex_id(rmm::cuda_stream_view const& stream_v int64_t const* d_edgelist_dsts, size_t num_edges); -template -void normalize(rmm::cuda_stream_view const& stream_view, value_t* d_value, size_t size) -{ - auto sum = thrust::reduce(rmm::exec_policy(stream_view), d_value, d_value + size); - - thrust::transform( - rmm::exec_policy(stream_view), d_value, d_value + size, d_value, [sum] __device__(auto v) { - return v / sum; - }); -} - -template void normalize(rmm::cuda_stream_view const& stream_view, float* d_value, size_t size); -template void normalize(rmm::cuda_stream_view const& stream_view, double* d_value, size_t size); - template std::tuple, rmm::device_uvector> filter_degree_0_vertices( raft::handle_t const& handle, diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 15ff56acfb6..c5e419b0889 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -447,6 +447,10 @@ ConfigureTest(PAGERANK_TEST link_analysis/pagerank_test.cpp) # - KATZ_CENTRALITY tests ------------------------------------------------------------------------- ConfigureTest(KATZ_CENTRALITY_TEST centrality/katz_centrality_test.cpp) +################################################################################################### +# - EIGENVECTOR_CENTRALITY tests ------------------------------------------------------------------------- +ConfigureTest(EIGENVECTOR_CENTRALITY_TEST centrality/eigenvector_centrality_test.cpp) + ################################################################################################### # - WEAKLY CONNECTED COMPONENTS tests ------------------------------------------------------------- ConfigureTest(WEAKLY_CONNECTED_COMPONENTS_TEST components/weakly_connected_components_test.cpp) @@ -546,6 +550,10 @@ if(BUILD_CUGRAPH_MG_TESTS) # - MG KATZ CENTRALITY tests -------------------------------------------------------------- ConfigureTestMG(MG_KATZ_CENTRALITY_TEST centrality/mg_katz_centrality_test.cpp) + ########################################################################################### + # - MG EIGENVECTOR CENTRALITY tests -------------------------------------------------------------- + ConfigureTestMG(MG_EIGENVECTOR_CENTRALITY_TEST centrality/mg_eigenvector_centrality_test.cpp) + ########################################################################################### # - MG BFS tests -------------------------------------------------------------------------- ConfigureTestMG(MG_BFS_TEST traversal/mg_bfs_test.cpp) diff --git a/cpp/tests/c_api/eigenvector_centrality_test.c b/cpp/tests/c_api/eigenvector_centrality_test.c index 7127c949d98..9fd2d2bee6f 100644 --- a/cpp/tests/c_api/eigenvector_centrality_test.c +++ b/cpp/tests/c_api/eigenvector_centrality_test.c @@ -26,23 +26,22 @@ typedef int32_t edge_t; typedef float weight_t; int generic_eigenvector_centrality_test(vertex_t* h_src, - vertex_t* h_dst, - weight_t* h_wgt, - weight_t* h_result, - size_t num_vertices, - size_t num_edges, - bool_t store_transposed, - double alpha, - double epsilon, - size_t max_iterations) + vertex_t* h_dst, + weight_t* h_wgt, + weight_t* h_result, + size_t num_vertices, + size_t num_edges, + bool_t store_transposed, + double epsilon, + size_t max_iterations) { int test_ret_value = 0; cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; cugraph_error_t* ret_error; - cugraph_resource_handle_t* p_handle = NULL; - cugraph_graph_t* p_graph = NULL; + cugraph_resource_handle_t* p_handle = NULL; + cugraph_graph_t* p_graph = NULL; cugraph_centrality_result_t* p_result = NULL; p_handle = cugraph_create_resource_handle(NULL); @@ -56,12 +55,14 @@ int generic_eigenvector_centrality_test(vertex_t* h_src, ret_code = cugraph_eigenvector_centrality( p_handle, p_graph, epsilon, max_iterations, FALSE, &p_result, &ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "cugraph_eigenvector_centrality failed."); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + TEST_ASSERT( + test_ret_value, ret_code == CUGRAPH_SUCCESS, "cugraph_eigenvector_centrality failed."); cugraph_type_erased_device_array_view_t* vertices; cugraph_type_erased_device_array_view_t* centralities; - vertices = cugraph_centrality_result_get_vertices(p_result); + vertices = cugraph_centrality_result_get_vertices(p_result); centralities = cugraph_centrality_result_get_values(p_result); vertex_t h_vertices[num_vertices]; @@ -91,21 +92,21 @@ int generic_eigenvector_centrality_test(vertex_t* h_src, int test_eigenvector_centrality() { - size_t num_edges = 8; + size_t num_edges = 16; size_t num_vertices = 6; - vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 4}; - vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; - weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f}; - weight_t h_result[] = {0.0915528, 0.168382, 0.0656831, 0.191468, 0.120677, 0.362237}; + vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 4, 1, 3, 4, 0, 1, 3, 5, 5}; + vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5, 0, 1, 1, 2, 2, 2, 3, 4}; + weight_t h_wgt[] = { + 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f}; + weight_t h_result[] = {0.236325, 0.292055, 0.458457, 0.60533, 0.190498, 0.495942}; - double alpha = 0.95; - double epsilon = 0.0001; - size_t max_iterations = 20; + double epsilon = 1e-6; + size_t max_iterations = 200; - // Pagerank wants store_transposed = TRUE + // Eigenvector centrality wants store_transposed = TRUE return generic_eigenvector_centrality_test( - h_src, h_dst, h_wgt, h_result, num_vertices, num_edges, TRUE, alpha, epsilon, max_iterations); + h_src, h_dst, h_wgt, h_result, num_vertices, num_edges, TRUE, epsilon, max_iterations); } /******************************************************************************/ diff --git a/cpp/tests/c_api/mg_eigenvector_centrality_test.c b/cpp/tests/c_api/mg_eigenvector_centrality_test.c index 967f3467963..56a5c3911a2 100644 --- a/cpp/tests/c_api/mg_eigenvector_centrality_test.c +++ b/cpp/tests/c_api/mg_eigenvector_centrality_test.c @@ -33,7 +33,6 @@ int generic_eigenvector_centrality_test(const cugraph_resource_handle_t* handle, size_t num_vertices, size_t num_edges, bool_t store_transposed, - double alpha, double epsilon, size_t max_iterations) { @@ -93,19 +92,19 @@ int generic_eigenvector_centrality_test(const cugraph_resource_handle_t* handle, int test_eigenvector_centrality(const cugraph_resource_handle_t* handle) { - size_t num_edges = 8; + size_t num_edges = 16; size_t num_vertices = 6; - vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 4}; - vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; - weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f}; - weight_t h_result[] = {0.0915528, 0.168382, 0.0656831, 0.191468, 0.120677, 0.362237}; + vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 4, 1, 3, 4, 0, 1, 3, 5, 5}; + vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5, 0, 1, 1, 2, 2, 2, 3, 4}; + weight_t h_wgt[] = { + 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f}; + weight_t h_result[] = {0.236374, 0.292046, 0.458369, 0.605472, 0.190544, 0.495814}; - double alpha = 0.95; - double epsilon = 0.0001; - size_t max_iterations = 20; + double epsilon = 1e-6; + size_t max_iterations = 200; - // Pagerank wants store_transposed = TRUE + // Eigenvector centrality wants store_transposed = TRUE return generic_eigenvector_centrality_test(handle, h_src, h_dst, @@ -114,7 +113,6 @@ int test_eigenvector_centrality(const cugraph_resource_handle_t* handle) num_vertices, num_edges, TRUE, - alpha, epsilon, max_iterations); } diff --git a/cpp/tests/centrality/eigenvector_centrality_test.cpp b/cpp/tests/centrality/eigenvector_centrality_test.cpp new file mode 100644 index 00000000000..a212baf8e0f --- /dev/null +++ b/cpp/tests/centrality/eigenvector_centrality_test.cpp @@ -0,0 +1,299 @@ +/* + * 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 governin_from_mtxg permissions and + * limitations under the License. + */ + +#include +#include +#include +#include +#include + +#include +#include +#include +#include + +#include +#include +#include +#include + +#include + +#include +#include +#include +#include +#include + +template +void eigenvector_centrality_reference(vertex_t const* src, + vertex_t const* dst, + std::optional weights, + size_t num_edges, + weight_t* centralities, + vertex_t num_vertices, + weight_t epsilon, + size_t max_iterations) +{ + if (num_vertices == 0) { return; } + + std::vector tmp_centralities(num_vertices, double{0}); + std::vector old_centralities(num_vertices, double{0.0}); + + std::fill(tmp_centralities.begin(), + tmp_centralities.end(), + double{1} / static_cast(num_vertices)); + + size_t iter{0}; + while (true) { + std::copy(tmp_centralities.begin(), tmp_centralities.end(), old_centralities.begin()); + std::fill(tmp_centralities.begin(), tmp_centralities.end(), double{0}); + + for (size_t e = 0; e < num_edges; ++e) { + auto w = weights ? (*weights)[e] : weight_t{1.0}; + tmp_centralities[src[e]] += old_centralities[dst[e]] * w; + } + + auto l2_norm = std::sqrt(std::inner_product( + tmp_centralities.begin(), tmp_centralities.end(), tmp_centralities.begin(), double{0.0})); + + std::transform(tmp_centralities.begin(), + tmp_centralities.end(), + tmp_centralities.begin(), + [l2_norm](auto& val) { return val / l2_norm; }); + + double diff_sum{0.0}; + double diff_max{0}; + for (vertex_t i = 0; i < num_vertices; ++i) { + diff_sum += std::abs(tmp_centralities[i] - old_centralities[i]); + if (std::abs(tmp_centralities[i] - old_centralities[i]) > diff_max) + diff_max = std::abs(tmp_centralities[i] - old_centralities[i]); + } + + if (diff_sum < (num_vertices * epsilon)) { break; } + iter++; + ASSERT_TRUE(iter < max_iterations); + } + + std::transform(tmp_centralities.begin(), tmp_centralities.end(), centralities, [](auto v) { + return static_cast(v); + }); + + return; +} + +struct EigenvectorCentrality_Usecase { + size_t max_iterations{std::numeric_limits::max()}; + bool test_weighted{false}; + bool check_correctness{true}; +}; + +template +class Tests_EigenvectorCentrality + : public ::testing::TestWithParam> { + public: + Tests_EigenvectorCentrality() {} + static void SetupTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(EigenvectorCentrality_Usecase const& eigenvector_usecase, + input_usecase_t const& input_usecase) + { + constexpr bool renumber = true; + + raft::handle_t handle{}; + HighResClock hr_clock{}; + + 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, eigenvector_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(); + + weight_t constexpr epsilon{1e-6}; + + rmm::device_uvector d_centralities(graph_view.number_of_vertices(), + handle.get_stream()); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_clock.start(); + } + + d_centralities = + cugraph::eigenvector_centrality(handle, + graph_view, + std::optional>{}, + epsilon, + eigenvector_usecase.max_iterations, + false); + + 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 << "Eigenvector Centrality took " << elapsed_time * 1e-6 << " s.\n"; + } + + if (eigenvector_usecase.check_correctness) { + auto [dst_v, src_v, opt_wgt_v] = graph_view.decompress_to_edgelist(handle, std::nullopt); + + std::vector h_src(src_v.size()); + std::vector h_dst(dst_v.size()); + auto h_weights = + opt_wgt_v ? std::make_optional>(opt_wgt_v->size()) : std::nullopt; + + raft::update_host(h_src.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(h_dst.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + if (h_weights) { + raft::update_host( + h_weights->data(), opt_wgt_v->data(), opt_wgt_v->size(), handle.get_stream()); + } + + handle.sync_stream(); + + std::vector h_reference_centralities(graph_view.number_of_vertices()); + + eigenvector_centrality_reference( + h_src.data(), + h_dst.data(), + h_weights ? std::make_optional(h_weights->data()) : std::nullopt, + h_src.size(), + h_reference_centralities.data(), + graph_view.number_of_vertices(), + epsilon, + eigenvector_usecase.max_iterations); + + std::vector h_cugraph_centralities(graph_view.number_of_vertices()); + raft::update_host(h_cugraph_centralities.data(), + d_centralities.data(), + d_centralities.size(), + handle.get_stream()); + + handle.sync_stream(); + + auto max_centrality = + *std::max_element(h_cugraph_centralities.begin(), h_cugraph_centralities.end()); + + // skip comparison for low Eigenvector Centrality vertices (lowly ranked vertices) + auto threshold_magnitude = max_centrality * epsilon; + + auto nearly_equal = [epsilon, threshold_magnitude](auto lhs, auto rhs) { + return std::abs(lhs - rhs) < std::max(std::max(lhs, rhs) * epsilon, threshold_magnitude); + }; + + // FIND DIFFERENCES... + size_t count_differences{0}; + for (size_t i = 0; i < h_reference_centralities.size(); ++i) { + if (nearly_equal(h_reference_centralities[i], h_cugraph_centralities[i])) { + } else { + if (count_differences < 10) { + std::cout << "unequal [" << i << "] " << h_reference_centralities[i] + << " != " << h_cugraph_centralities[i] << std::endl; + } + ++count_differences; + } + } + + ASSERT_EQ(count_differences, size_t{0}) + << "Eigenvector centrality values do not match with the reference " + "values."; + } + } +}; + +using Tests_EigenvectorCentrality_File = Tests_EigenvectorCentrality; +using Tests_EigenvectorCentrality_Rmat = Tests_EigenvectorCentrality; + +// FIXME: add tests for type combinations +TEST_P(Tests_EigenvectorCentrality_File, CheckInt32Int32FloatFloat) +{ + auto param = GetParam(); + run_current_test(std::get<0>(param), std::get<1>(param)); +} + +TEST_P(Tests_EigenvectorCentrality_Rmat, CheckInt32Int32FloatFloat) +{ + auto param = GetParam(); + run_current_test( + std::get<0>(param), override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); +} + +TEST_P(Tests_EigenvectorCentrality_Rmat, CheckInt32Int64FloatFloat) +{ + auto param = GetParam(); + run_current_test( + std::get<0>(param), override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); +} + +TEST_P(Tests_EigenvectorCentrality_Rmat, CheckInt64Int64FloatFloat) +{ + auto param = GetParam(); + run_current_test( + std::get<0>(param), override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); +} + +INSTANTIATE_TEST_SUITE_P( + file_test_pass, + Tests_EigenvectorCentrality_File, + ::testing::Combine( + // enable correctness checks + ::testing::Values(EigenvectorCentrality_Usecase{500, false}, + EigenvectorCentrality_Usecase{500, true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), + cugraph::test::File_Usecase("test/datasets/web-Google.mtx"), + cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), + cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + rmat_small_test, + Tests_EigenvectorCentrality_Rmat, + // enable correctness checks + ::testing::Combine( + ::testing::Values(EigenvectorCentrality_Usecase{500, false}, + EigenvectorCentrality_Usecase{500, 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_EigenvectorCentrality_Rmat, + // disable correctness checks for large graphs + ::testing::Combine( + ::testing::Values(EigenvectorCentrality_Usecase{500, false, false}, + EigenvectorCentrality_Usecase{500, 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/centrality/mg_eigenvector_centrality_test.cpp b/cpp/tests/centrality/mg_eigenvector_centrality_test.cpp new file mode 100644 index 00000000000..649120d2ebc --- /dev/null +++ b/cpp/tests/centrality/mg_eigenvector_centrality_test.cpp @@ -0,0 +1,284 @@ +/* + * 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 +#include +#include + +#include + +#include + +struct EigenvectorCentrality_Usecase { + size_t max_iterations{std::numeric_limits::max()}; + bool test_weighted{false}; + bool check_correctness{true}; +}; + +template +class Tests_MGEigenvectorCentrality + : public ::testing::TestWithParam> { + public: + Tests_MGEigenvectorCentrality() {} + static void SetupTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + // Compare the results of running Eigenvector Centrality on multiple GPUs to that of a single-GPU + // run + template + void run_current_test(EigenvectorCentrality_Usecase const& eigenvector_usecase, + input_usecase_t const& input_usecase) + { + // 1. initialize handle + + raft::handle_t handle{}; + HighResClock hr_clock{}; + + raft::comms::initialize_mpi_comms(&handle, MPI_COMM_WORLD); + auto& comm = handle.get_comms(); + auto const comm_size = comm.get_size(); + auto const comm_rank = comm.get_rank(); + + auto row_comm_size = static_cast(sqrt(static_cast(comm_size))); + while (comm_size % row_comm_size != 0) { + --row_comm_size; + } + cugraph::partition_2d::subcomm_factory_t + subcomm_factory(handle, row_comm_size); + + // 2. create MG graph + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle.get_comms().barrier(); + hr_clock.start(); + } + + auto [mg_graph, d_mg_renumber_map_labels] = + cugraph::test::construct_graph( + handle, input_usecase, eigenvector_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(); + + // 3. run MG Eigenvector Centrality + + weight_t constexpr epsilon{1e-6}; + + rmm::device_uvector d_mg_centralities( + mg_graph_view.local_vertex_partition_range_size(), handle.get_stream()); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle.get_comms().barrier(); + hr_clock.start(); + } + + d_mg_centralities = cugraph::eigenvector_centrality( + handle, + mg_graph_view, + std::optional>{}, + // std::make_optional(raft::device_span{d_mg_centralities.data(), d_mg_centralities.size()}), + epsilon, + eigenvector_usecase.max_iterations, + false); + + 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 Eigenvector Centrality took " << elapsed_time * 1e-6 << " s.\n"; + } + + // 4. compare SG & MG results + + if (eigenvector_usecase.check_correctness) { + // 4-1. aggregate MG results + + auto d_mg_aggregate_renumber_map_labels = cugraph::test::device_gatherv( + handle, (*d_mg_renumber_map_labels).data(), (*d_mg_renumber_map_labels).size()); + auto d_mg_aggregate_centralities = + cugraph::test::device_gatherv(handle, d_mg_centralities.data(), d_mg_centralities.size()); + + if (handle.get_comms().get_rank() == int{0}) { + // 4-2. Sort MG results by original vertex id + std::tie(std::ignore, d_mg_aggregate_centralities) = cugraph::test::sort_by_key( + handle, d_mg_aggregate_renumber_map_labels, d_mg_aggregate_centralities); + + // 4-3. create SG graph + auto [sg_graph, d_sg_renumber_map_labels] = + cugraph::test::construct_graph( + handle, input_usecase, eigenvector_usecase.test_weighted, true); + + auto sg_graph_view = sg_graph.view(); + + ASSERT_TRUE(mg_graph_view.number_of_vertices() == sg_graph_view.number_of_vertices()); + + // 4-4. run SG Eigenvector Centrality + rmm::device_uvector d_sg_centralities(sg_graph_view.number_of_vertices(), + handle.get_stream()); + + d_sg_centralities = cugraph::eigenvector_centrality( + handle, + sg_graph_view, + std::optional>{}, + // std::make_optional(raft::device_span{d_sg_centralities.data(), + // d_sg_centralities.size()}), + epsilon, + eigenvector_usecase.max_iterations, + false); + + std::tie(std::ignore, d_sg_centralities) = + cugraph::test::sort_by_key(handle, *d_sg_renumber_map_labels, d_sg_centralities); + + // 4-5. compare + std::vector h_mg_aggregate_centralities(mg_graph_view.number_of_vertices()); + raft::update_host(h_mg_aggregate_centralities.data(), + d_mg_aggregate_centralities.data(), + d_mg_aggregate_centralities.size(), + handle.get_stream()); + + std::vector h_sg_centralities(sg_graph_view.number_of_vertices()); + raft::update_host(h_sg_centralities.data(), + d_sg_centralities.data(), + d_sg_centralities.size(), + handle.get_stream()); + + handle.sync_stream(); + + auto max_centrality = + *std::max_element(h_mg_aggregate_centralities.begin(), h_mg_aggregate_centralities.end()); + + // skip comparison for low Eigenvector Centrality vertices (lowly ranked vertices) + auto threshold_magnitude = max_centrality * epsilon; + + auto nearly_equal = [epsilon, threshold_magnitude](auto lhs, auto rhs) { + return std::abs(lhs - rhs) < std::max(std::max(lhs, rhs) * epsilon, threshold_magnitude); + }; + + // FIND DIFFERENCES... + size_t count_differences{0}; + for (size_t i = 0; i < h_mg_aggregate_centralities.size(); ++i) { + if (nearly_equal(h_mg_aggregate_centralities[i], h_sg_centralities[i])) { + } else { + if (count_differences < 10) { + std::cout << "unequal [" << i << "] " << h_mg_aggregate_centralities[i] + << " != " << h_sg_centralities[i] << std::endl; + } + ++count_differences; + } + } + + ASSERT_EQ(count_differences, size_t{0}) + << "Eigenvector centrality values do not match with the reference " + "values."; + } + } + } +}; + +using Tests_MGEigenvectorCentrality_File = + Tests_MGEigenvectorCentrality; +using Tests_MGEigenvectorCentrality_Rmat = + Tests_MGEigenvectorCentrality; + +TEST_P(Tests_MGEigenvectorCentrality_File, CheckInt32Int32FloatFloat) +{ + auto param = GetParam(); + run_current_test(std::get<0>(param), std::get<1>(param)); +} + +TEST_P(Tests_MGEigenvectorCentrality_Rmat, CheckInt32Int32FloatFloat) +{ + auto param = GetParam(); + run_current_test( + std::get<0>(param), override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); +} + +TEST_P(Tests_MGEigenvectorCentrality_Rmat, CheckInt32Int64FloatFloat) +{ + auto param = GetParam(); + run_current_test( + std::get<0>(param), override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); +} + +TEST_P(Tests_MGEigenvectorCentrality_Rmat, CheckInt64Int64FloatFloat) +{ + auto param = GetParam(); + run_current_test( + std::get<0>(param), override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); +} + +INSTANTIATE_TEST_SUITE_P( + file_test, + Tests_MGEigenvectorCentrality_File, + ::testing::Combine( + // enable correctness checks + ::testing::Values(EigenvectorCentrality_Usecase{500, false}, + EigenvectorCentrality_Usecase{500, true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), + cugraph::test::File_Usecase("test/datasets/web-Google.mtx"), + cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), + cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); + +INSTANTIATE_TEST_SUITE_P(rmat_small_test, + Tests_MGEigenvectorCentrality_Rmat, + ::testing::Combine( + // enable correctness checks + ::testing::Values(EigenvectorCentrality_Usecase{500, false}, + EigenvectorCentrality_Usecase{500, true}), + ::testing::Values(cugraph::test::Rmat_Usecase( + 10, 16, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); + +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_MGEigenvectorCentrality_Rmat, + ::testing::Combine( + // disable correctness checks for large graphs + ::testing::Values(EigenvectorCentrality_Usecase{500, false, false}, + EigenvectorCentrality_Usecase{500, true, false}), + ::testing::Values( + cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); + +CUGRAPH_MG_TEST_PROGRAM_MAIN()