From 390bc7c2dbc442d7e652477440c1653363408b4c Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Wed, 12 May 2021 16:29:18 -0400 Subject: [PATCH 1/9] add missing include statement --- cpp/include/cugraph/utilities/host_scalar_comm.cuh | 1 + 1 file changed, 1 insertion(+) diff --git a/cpp/include/cugraph/utilities/host_scalar_comm.cuh b/cpp/include/cugraph/utilities/host_scalar_comm.cuh index 4505d35e011..85994ed22bf 100644 --- a/cpp/include/cugraph/utilities/host_scalar_comm.cuh +++ b/cpp/include/cugraph/utilities/host_scalar_comm.cuh @@ -15,6 +15,7 @@ */ #pragma once +#include #include #include From c7ab994c1e650eb54575b2292ffa6ba9a60bb7e6 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Wed, 12 May 2021 17:40:43 -0400 Subject: [PATCH 2/9] move gather_distributed_vector from tests/community to tests/utilities and renumber to device_gatherv --- cpp/tests/CMakeLists.txt | 34 +++++++++++ cpp/tests/community/mg_louvain_helper.cu | 34 ----------- cpp/tests/community/mg_louvain_helper.hpp | 5 -- cpp/tests/community/mg_louvain_test.cpp | 5 +- cpp/tests/utilities/device_comm_wrapper.cu | 64 +++++++++++++++++++++ cpp/tests/utilities/device_comm_wrapper.hpp | 29 ++++++++++ 6 files changed, 130 insertions(+), 41 deletions(-) create mode 100644 cpp/tests/utilities/device_comm_wrapper.cu create mode 100644 cpp/tests/utilities/device_comm_wrapper.hpp diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 80484fdfad6..15a90410711 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -58,6 +58,39 @@ target_link_libraries(cugraphtestutil cugraph) set_target_properties(cugraphtestutil PROPERTIES CUDA_ARCHITECTURES OFF) +add_library(cugraphmgtestutil STATIC + "${CMAKE_CURRENT_SOURCE_DIR}/utilities/device_comm_wrapper.cu") + +set_property(TARGET cugraphmgtestutil PROPERTY POSITION_INDEPENDENT_CODE ON) + +target_include_directories(cugraphmgtestutil + PRIVATE + "${CUB_INCLUDE_DIR}" + "${THRUST_INCLUDE_DIR}" + "${CUCO_INCLUDE_DIR}" + "${LIBCUDACXX_INCLUDE_DIR}" + "${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES}" + "${RMM_INCLUDE}" + "${NCCL_INCLUDE_DIRS}" + "${CMAKE_CURRENT_SOURCE_DIR}/../../thirdparty/mmio" + "${CMAKE_CURRENT_SOURCE_DIR}/../include" + "${CMAKE_CURRENT_SOURCE_DIR}" + "${RAFT_DIR}/cpp/include" +) + +target_link_libraries(cugraphmgtestutil cugraph) + +# CUDA_ARCHITECTURES=OFF implies cmake will not pass arch flags to the +# compiler. CUDA_ARCHITECTURES must be set to a non-empty value to prevent +# cmake warnings about policy CMP0104. With this setting, arch flags must be +# manually set! ("evaluate_gpu_archs(GPU_ARCHS)" is the current mechanism +# used in cpp/CMakeLists.txt for setting arch options). +# Run "cmake --help-policy CMP0104" for policy details. +# NOTE: the CUDA_ARCHITECTURES=OFF setting may be removed after migrating to +# the findcudatoolkit features in cmake 3.17+ +set_target_properties(cugraphmgtestutil PROPERTIES + CUDA_ARCHITECTURES OFF) + ################################################################################################### # - compiler function ----------------------------------------------------------------------------- @@ -194,6 +227,7 @@ function(ConfigureTestMG CMAKE_TEST_NAME CMAKE_TEST_SRC) target_link_libraries(${CMAKE_TEST_NAME} PRIVATE + cugraphmgtestutil cugraphtestutil cugraph GTest::GTest diff --git a/cpp/tests/community/mg_louvain_helper.cu b/cpp/tests/community/mg_louvain_helper.cu index 1311970292a..2b1b5ade41e 100644 --- a/cpp/tests/community/mg_louvain_helper.cu +++ b/cpp/tests/community/mg_louvain_helper.cu @@ -31,36 +31,6 @@ namespace cugraph { namespace test { -template -rmm::device_uvector gather_distributed_vector(raft::handle_t const &handle, - T const *d_input, - size_t size) -{ - auto rx_sizes = - cugraph::experimental::host_scalar_gather(handle.get_comms(), size, 0, handle.get_stream()); - std::vector rx_displs(static_cast(handle.get_comms().get_rank()) == 0 - ? handle.get_comms().get_size() - : int{0}, - size_t{0}); - if (static_cast(handle.get_comms().get_rank()) == 0) { - std::partial_sum(rx_sizes.begin(), rx_sizes.end() - 1, rx_displs.begin() + 1); - } - - auto total_size = thrust::reduce(thrust::host, rx_sizes.begin(), rx_sizes.end()); - rmm::device_uvector gathered_v(total_size, handle.get_stream()); - - cugraph::experimental::device_gatherv(handle.get_comms(), - d_input, - gathered_v.data(), - size, - rx_sizes, - rx_displs, - 0, - handle.get_stream()); - - return gathered_v; -} - template bool compare_renumbered_vectors(raft::handle_t const &handle, rmm::device_uvector const &v1, @@ -336,10 +306,6 @@ template void single_gpu_renumber_edgelist_given_number_map( rmm::device_uvector &d_edgelist_cols, rmm::device_uvector &d_renumber_map_gathered_v); -template rmm::device_uvector gather_distributed_vector(raft::handle_t const &handle, - int const *d_input, - size_t size); - template bool compare_renumbered_vectors(raft::handle_t const &handle, rmm::device_uvector const &v1, rmm::device_uvector const &v2); diff --git a/cpp/tests/community/mg_louvain_helper.hpp b/cpp/tests/community/mg_louvain_helper.hpp index 456301f4d7b..6d074e2d5e9 100644 --- a/cpp/tests/community/mg_louvain_helper.hpp +++ b/cpp/tests/community/mg_louvain_helper.hpp @@ -24,11 +24,6 @@ namespace cugraph { namespace test { -template -rmm::device_uvector gather_distributed_vector(raft::handle_t const &handle, - T const *d_input, - size_t size); - template bool compare_renumbered_vectors(raft::handle_t const &handle, rmm::device_uvector const &v1, diff --git a/cpp/tests/community/mg_louvain_test.cpp b/cpp/tests/community/mg_louvain_test.cpp index e8cc94edf99..7851fac76a2 100644 --- a/cpp/tests/community/mg_louvain_test.cpp +++ b/cpp/tests/community/mg_louvain_test.cpp @@ -17,6 +17,7 @@ #include "mg_louvain_helper.hpp" #include +#include #include #include @@ -144,7 +145,7 @@ class Louvain_MG_Testfixture : public ::testing::TestWithParam thrust::make_counting_iterator(dendrogram.num_levels()), [&dendrogram, &sg_graph, &d_clustering_v, &sg_modularity, &handle, resolution, rank]( size_t i) { - auto d_dendrogram_gathered_v = cugraph::test::gather_distributed_vector( + auto d_dendrogram_gathered_v = cugraph::test::device_gatherv( handle, dendrogram.get_level_ptr_nocheck(i), dendrogram.get_level_size_nocheck(i)); if (rank == 0) { @@ -207,7 +208,7 @@ class Louvain_MG_Testfixture : public ::testing::TestWithParam SCOPED_TRACE("compare modularity input: " + param.graph_file_full_path); - auto d_renumber_map_gathered_v = cugraph::test::gather_distributed_vector( + auto d_renumber_map_gathered_v = cugraph::test::device_gatherv( handle, d_renumber_map_labels.data(), d_renumber_map_labels.size()); compare_sg_results(handle, diff --git a/cpp/tests/utilities/device_comm_wrapper.cu b/cpp/tests/utilities/device_comm_wrapper.cu new file mode 100644 index 00000000000..4da442863eb --- /dev/null +++ b/cpp/tests/utilities/device_comm_wrapper.cu @@ -0,0 +1,64 @@ +/* + * Copyright (c) 2021, 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 "device_comm_wrapper.hpp" + +#include +#include + +#include +#include + +namespace cugraph { +namespace test { + +template +rmm::device_uvector device_gatherv(raft::handle_t const &handle, T const *d_input, size_t size) +{ + bool is_root = handle.get_comms().get_rank() == int{0}; + auto rx_sizes = cugraph::experimental::host_scalar_gather( + handle.get_comms(), size, int{0}, handle.get_stream()); + std::vector rx_displs(is_root ? static_cast(handle.get_comms().get_size()) + : size_t{0}); + if (is_root) { std::partial_sum(rx_sizes.begin(), rx_sizes.end() - 1, rx_displs.begin() + 1); } + + rmm::device_uvector gathered_v( + is_root ? std::reduce(rx_sizes.begin(), rx_sizes.end()) : size_t{0}, handle.get_stream()); + + cugraph::experimental::device_gatherv(handle.get_comms(), + d_input, + gathered_v.data(), + size, + rx_sizes, + rx_displs, + int{0}, + handle.get_stream()); + + return gathered_v; +} + +// explicit instantiation + +template rmm::device_uvector device_gatherv(raft::handle_t const &handle, + int32_t const *d_input, + size_t size); + +template rmm::device_uvector device_gatherv(raft::handle_t const &handle, + int64_t const *d_input, + size_t size); + +} // namespace test +} // namespace cugraph diff --git a/cpp/tests/utilities/device_comm_wrapper.hpp b/cpp/tests/utilities/device_comm_wrapper.hpp new file mode 100644 index 00000000000..f56f24248d6 --- /dev/null +++ b/cpp/tests/utilities/device_comm_wrapper.hpp @@ -0,0 +1,29 @@ +/* + * Copyright (c) 2021, 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 + +namespace cugraph { +namespace test { + +template +rmm::device_uvector device_gatherv(raft::handle_t const &handle, T const *d_input, size_t size); + +} // namespace test +} // namespace cugraph From a04c39bcf49cafae742ac50f2be9dab9d6c8f9ae Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 13 May 2021 11:02:10 -0400 Subject: [PATCH 3/9] update thrust wrapper sort_by_key to return both sorted keys & values to better reflect the behavior of thrust sort_by_key --- cpp/tests/utilities/thrust_wrapper.cu | 62 ++++++++++++++------------ cpp/tests/utilities/thrust_wrapper.hpp | 8 ++-- 2 files changed, 37 insertions(+), 33 deletions(-) diff --git a/cpp/tests/utilities/thrust_wrapper.cu b/cpp/tests/utilities/thrust_wrapper.cu index 5d32fb8a5d1..dfd420b1e2d 100644 --- a/cpp/tests/utilities/thrust_wrapper.cu +++ b/cpp/tests/utilities/thrust_wrapper.cu @@ -26,10 +26,8 @@ namespace cugraph { namespace test { template -rmm::device_uvector sort_by_key(raft::handle_t const& handle, - vertex_t const* keys, - value_t const* values, - size_t num_pairs) +std::tuple, rmm::device_uvector> sort_by_key( + raft::handle_t const& handle, vertex_t const* keys, value_t const* values, size_t num_pairs) { rmm::device_uvector sorted_keys(num_pairs, handle.get_stream_view()); rmm::device_uvector sorted_values(num_pairs, handle.get_stream_view()); @@ -44,38 +42,44 @@ rmm::device_uvector sort_by_key(raft::handle_t const& handle, sorted_keys.end(), sorted_values.begin()); - return sorted_values; + return std::make_tuple(std::move(sorted_keys), std::move(sorted_values)); } -template rmm::device_uvector sort_by_key(raft::handle_t const& handle, - int32_t const* keys, - float const* values, - size_t num_pairs); +template std::tuple, rmm::device_uvector> +sort_by_key(raft::handle_t const& handle, + int32_t const* keys, + float const* values, + size_t num_pairs); -template rmm::device_uvector sort_by_key(raft::handle_t const& handle, - int32_t const* keys, - double const* values, - size_t num_pairs); +template std::tuple, rmm::device_uvector> +sort_by_key(raft::handle_t const& handle, + int32_t const* keys, + double const* values, + size_t num_pairs); -template rmm::device_uvector sort_by_key(raft::handle_t const& handle, - int32_t const* keys, - int32_t const* values, - size_t num_pairs); +template std::tuple, rmm::device_uvector> +sort_by_key(raft::handle_t const& handle, + int32_t const* keys, + int32_t const* values, + size_t num_pairs); -template rmm::device_uvector sort_by_key(raft::handle_t const& handle, - int64_t const* keys, - float const* values, - size_t num_pairs); +template std::tuple, rmm::device_uvector> +sort_by_key(raft::handle_t const& handle, + int64_t const* keys, + float const* values, + size_t num_pairs); -template rmm::device_uvector sort_by_key(raft::handle_t const& handle, - int64_t const* keys, - double const* values, - size_t num_pairs); +template std::tuple, rmm::device_uvector> +sort_by_key(raft::handle_t const& handle, + int64_t const* keys, + double const* values, + size_t num_pairs); -template rmm::device_uvector sort_by_key(raft::handle_t const& handle, - int64_t const* keys, - int64_t const* values, - size_t num_pairs); +template std::tuple, rmm::device_uvector> +sort_by_key(raft::handle_t const& handle, + int64_t const* keys, + int64_t const* values, + size_t num_pairs); } // namespace test } // namespace cugraph diff --git a/cpp/tests/utilities/thrust_wrapper.hpp b/cpp/tests/utilities/thrust_wrapper.hpp index 579dc3c550f..96f370f884c 100644 --- a/cpp/tests/utilities/thrust_wrapper.hpp +++ b/cpp/tests/utilities/thrust_wrapper.hpp @@ -17,14 +17,14 @@ #include #include +#include + namespace cugraph { namespace test { template -rmm::device_uvector sort_by_key(raft::handle_t const& handle, - vertex_t const* keys, - value_t const* values, - size_t num_pairs); +std::tuple, rmm::device_uvector> sort_by_key( + raft::handle_t const& handle, vertex_t const* keys, value_t const* values, size_t num_pairs); } // namespace test } // namespace cugraph From 2aa2a8951d374add80e02de0c75d2331f2ea3912 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 13 May 2021 11:04:07 -0400 Subject: [PATCH 4/9] add additional explicit instantiations of device_gatherv --- cpp/tests/utilities/device_comm_wrapper.cu | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/cpp/tests/utilities/device_comm_wrapper.cu b/cpp/tests/utilities/device_comm_wrapper.cu index 4da442863eb..cda355db826 100644 --- a/cpp/tests/utilities/device_comm_wrapper.cu +++ b/cpp/tests/utilities/device_comm_wrapper.cu @@ -60,5 +60,13 @@ template rmm::device_uvector device_gatherv(raft::handle_t const &handl int64_t const *d_input, size_t size); +template rmm::device_uvector device_gatherv(raft::handle_t const &handle, + float const *d_input, + size_t size); + +template rmm::device_uvector device_gatherv(raft::handle_t const &handle, + double const *d_input, + size_t size); + } // namespace test } // namespace cugraph From 93b5164a91aab899aef195fa228128cdef9c0b7a Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 13 May 2021 11:08:07 -0400 Subject: [PATCH 5/9] add const to a can-be-const input parameter of renumber_edgelist --- .../cugraph/experimental/graph_functions.hpp | 2 +- cpp/src/experimental/renumber_utils.cu | 70 ++++++++++--------- 2 files changed, 38 insertions(+), 34 deletions(-) diff --git a/cpp/include/cugraph/experimental/graph_functions.hpp b/cpp/include/cugraph/experimental/graph_functions.hpp index cb1b90a6e8e..ea7d0df3ba2 100644 --- a/cpp/include/cugraph/experimental/graph_functions.hpp +++ b/cpp/include/cugraph/experimental/graph_functions.hpp @@ -284,7 +284,7 @@ void unrenumber_int_vertices(raft::handle_t const& handle, vertex_t const* renumber_map_labels, vertex_t local_int_vertex_first, vertex_t local_int_vertex_last, - std::vector& vertex_partition_lasts, + std::vector const& vertex_partition_lasts, bool do_expensive_check = false); /** diff --git a/cpp/src/experimental/renumber_utils.cu b/cpp/src/experimental/renumber_utils.cu index 5e5e88ef8f7..765dbb19886 100644 --- a/cpp/src/experimental/renumber_utils.cu +++ b/cpp/src/experimental/renumber_utils.cu @@ -214,7 +214,7 @@ void unrenumber_int_vertices(raft::handle_t const& handle, vertex_t const* renumber_map_labels, vertex_t local_int_vertex_first, vertex_t local_int_vertex_last, - std::vector& vertex_partition_lasts, + std::vector const& vertex_partition_lasts, bool do_expensive_check) { double constexpr load_factor = 0.7; @@ -385,41 +385,45 @@ template void unrenumber_local_int_vertices(raft::handle_t const& handl int64_t local_int_vertex_last, bool do_expensive_check); -template void unrenumber_int_vertices(raft::handle_t const& handle, - int32_t* vertices, - size_t num_vertices, - int32_t const* renumber_map_labels, - int32_t local_int_vertex_first, - int32_t local_int_vertex_last, - std::vector& vertex_partition_lasts, - bool do_expensive_check); +template void unrenumber_int_vertices( + raft::handle_t const& handle, + int32_t* vertices, + size_t num_vertices, + int32_t const* renumber_map_labels, + int32_t local_int_vertex_first, + int32_t local_int_vertex_last, + std::vector const& vertex_partition_lasts, + bool do_expensive_check); -template void unrenumber_int_vertices(raft::handle_t const& handle, - int32_t* vertices, - size_t num_vertices, - int32_t const* renumber_map_labels, - int32_t local_int_vertex_first, - int32_t local_int_vertex_last, - std::vector& vertex_partition_lasts, - bool do_expensive_check); +template void unrenumber_int_vertices( + raft::handle_t const& handle, + int32_t* vertices, + size_t num_vertices, + int32_t const* renumber_map_labels, + int32_t local_int_vertex_first, + int32_t local_int_vertex_last, + std::vector const& vertex_partition_lasts, + bool do_expensive_check); -template void unrenumber_int_vertices(raft::handle_t const& handle, - int64_t* vertices, - size_t num_vertices, - int64_t const* renumber_map_labels, - int64_t local_int_vertex_first, - int64_t local_int_vertex_last, - std::vector& vertex_partition_lasts, - bool do_expensive_check); +template void unrenumber_int_vertices( + raft::handle_t const& handle, + int64_t* vertices, + size_t num_vertices, + int64_t const* renumber_map_labels, + int64_t local_int_vertex_first, + int64_t local_int_vertex_last, + std::vector const& vertex_partition_lasts, + bool do_expensive_check); -template void unrenumber_int_vertices(raft::handle_t const& handle, - int64_t* vertices, - size_t num_vertices, - int64_t const* renumber_map_labels, - int64_t local_int_vertex_first, - int64_t local_int_vertex_last, - std::vector& vertex_partition_lasts, - bool do_expensive_check); +template void unrenumber_int_vertices( + raft::handle_t const& handle, + int64_t* vertices, + size_t num_vertices, + int64_t const* renumber_map_labels, + int64_t local_int_vertex_first, + int64_t local_int_vertex_last, + std::vector const& vertex_partition_lasts, + bool do_expensive_check); } // namespace experimental } // namespace cugraph From 4a0af52b750aa0c932a76270f4f6a7aff1c029ac Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 13 May 2021 11:37:46 -0400 Subject: [PATCH 6/9] update C++ tests for the thrust wrapper API change --- cpp/tests/experimental/bfs_test.cpp | 13 ++++++++----- cpp/tests/experimental/katz_centrality_test.cpp | 4 +++- cpp/tests/experimental/pagerank_test.cpp | 12 +++++++----- cpp/tests/experimental/sssp_test.cpp | 13 ++++++++----- cpp/tests/pagerank/mg_pagerank_test.cpp | 9 +++++---- 5 files changed, 31 insertions(+), 20 deletions(-) diff --git a/cpp/tests/experimental/bfs_test.cpp b/cpp/tests/experimental/bfs_test.cpp index da4ef2f5dfb..2c8ab894096 100644 --- a/cpp/tests/experimental/bfs_test.cpp +++ b/cpp/tests/experimental/bfs_test.cpp @@ -212,12 +212,15 @@ class Tests_BFS : public ::testing::TestWithParam d_unrenumbered_distances(size_t{0}, handle.get_stream()); + std::tie(std::ignore, d_unrenumbered_distances) = cugraph::test::sort_by_key( handle, d_renumber_map_labels.data(), d_distances.data(), d_renumber_map_labels.size()); - auto d_unrenumbered_predecessors = cugraph::test::sort_by_key(handle, - d_renumber_map_labels.data(), - d_predecessors.data(), - d_renumber_map_labels.size()); + rmm::device_uvector d_unrenumbered_predecessors(size_t{0}, handle.get_stream()); + std::tie(std::ignore, d_unrenumbered_predecessors) = + cugraph::test::sort_by_key(handle, + d_renumber_map_labels.data(), + d_predecessors.data(), + d_renumber_map_labels.size()); raft::update_host(h_cugraph_distances.data(), d_unrenumbered_distances.data(), d_unrenumbered_distances.size(), diff --git a/cpp/tests/experimental/katz_centrality_test.cpp b/cpp/tests/experimental/katz_centrality_test.cpp index 35773073757..d0fc558c89f 100644 --- a/cpp/tests/experimental/katz_centrality_test.cpp +++ b/cpp/tests/experimental/katz_centrality_test.cpp @@ -225,7 +225,9 @@ class Tests_KatzCentrality std::vector h_cugraph_katz_centralities(graph_view.get_number_of_vertices()); if (renumber) { - auto d_unrenumbered_katz_centralities = + rmm::device_uvector d_unrenumbered_katz_centralities(size_t{0}, + handle.get_stream()); + std::tie(std::ignore, d_unrenumbered_katz_centralities) = cugraph::test::sort_by_key(handle, d_renumber_map_labels.data(), d_katz_centralities.data(), diff --git a/cpp/tests/experimental/pagerank_test.cpp b/cpp/tests/experimental/pagerank_test.cpp index e1b7b121b1f..9b07059d2da 100644 --- a/cpp/tests/experimental/pagerank_test.cpp +++ b/cpp/tests/experimental/pagerank_test.cpp @@ -303,10 +303,11 @@ class Tests_PageRank d_renumber_map_labels.data(), vertex_t{0}, graph_view.get_number_of_vertices()); - cugraph::test::sort_by_key(handle, - d_unrenumbered_personalization_vertices.data(), - d_unrenumbered_personalization_values.data(), - d_unrenumbered_personalization_vertices.size()); + std::tie(d_unrenumbered_personalization_vertices, d_unrenumbered_personalization_values) = + cugraph::test::sort_by_key(handle, + d_unrenumbered_personalization_vertices.data(), + d_unrenumbered_personalization_values.data(), + d_unrenumbered_personalization_vertices.size()); raft::update_host(h_unrenumbered_personalization_vertices.data(), d_unrenumbered_personalization_vertices.data(), @@ -346,7 +347,8 @@ class Tests_PageRank std::vector h_cugraph_pageranks(graph_view.get_number_of_vertices()); if (renumber) { - auto d_unrenumbered_pageranks = cugraph::test::sort_by_key( + rmm::device_uvector d_unrenumbered_pageranks(size_t{0}, handle.get_stream()); + std::tie(std::ignore, d_unrenumbered_pageranks) = cugraph::test::sort_by_key( handle, d_renumber_map_labels.data(), d_pageranks.data(), d_renumber_map_labels.size()); raft::update_host(h_cugraph_pageranks.data(), d_unrenumbered_pageranks.data(), diff --git a/cpp/tests/experimental/sssp_test.cpp b/cpp/tests/experimental/sssp_test.cpp index 3095afad8fc..e12df163551 100644 --- a/cpp/tests/experimental/sssp_test.cpp +++ b/cpp/tests/experimental/sssp_test.cpp @@ -220,12 +220,15 @@ class Tests_SSSP : public ::testing::TestWithParam d_unrenumbered_distances(size_t{0}, handle.get_stream()); + std::tie(std::ignore, d_unrenumbered_distances) = cugraph::test::sort_by_key( handle, d_renumber_map_labels.data(), d_distances.data(), d_renumber_map_labels.size()); - auto d_unrenumbered_predecessors = cugraph::test::sort_by_key(handle, - d_renumber_map_labels.data(), - d_predecessors.data(), - d_renumber_map_labels.size()); + rmm::device_uvector d_unrenumbered_predecessors(size_t{0}, handle.get_stream()); + std::tie(std::ignore, d_unrenumbered_predecessors) = + cugraph::test::sort_by_key(handle, + d_renumber_map_labels.data(), + d_predecessors.data(), + d_renumber_map_labels.size()); raft::update_host(h_cugraph_distances.data(), d_unrenumbered_distances.data(), diff --git a/cpp/tests/pagerank/mg_pagerank_test.cpp b/cpp/tests/pagerank/mg_pagerank_test.cpp index adedfa2e3bc..8ccf967d95a 100644 --- a/cpp/tests/pagerank/mg_pagerank_test.cpp +++ b/cpp/tests/pagerank/mg_pagerank_test.cpp @@ -242,10 +242,11 @@ class Tests_MGPageRank displacements.data(), handle.get_stream()); - cugraph::test::sort_by_key(handle, - d_unrenumbered_personalization_vertices.data(), - d_unrenumbered_personalization_values.data(), - d_unrenumbered_personalization_vertices.size()); + std::tie(d_unrenumbered_personalization_vertices, d_unrenumbered_personalization_values) = + cugraph::test::sort_by_key(handle, + d_unrenumbered_personalization_vertices.data(), + d_unrenumbered_personalization_values.data(), + d_unrenumbered_personalization_vertices.size()); } // 5-3. run SG PageRank From 707c2c3f7413be517ee8a372872c7f599b942dbd Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 13 May 2021 14:13:09 -0400 Subject: [PATCH 7/9] update MG tests to gather the MG results in root and compare with the SG results --- cpp/tests/experimental/mg_bfs_test.cpp | 235 ++++++++-------- .../experimental/mg_katz_centrality_test.cpp | 138 +++++----- cpp/tests/experimental/mg_sssp_test.cpp | 260 +++++++++--------- cpp/tests/pagerank/mg_pagerank_test.cpp | 214 ++++++-------- 4 files changed, 424 insertions(+), 423 deletions(-) diff --git a/cpp/tests/experimental/mg_bfs_test.cpp b/cpp/tests/experimental/mg_bfs_test.cpp index e498e403334..ef980a2448e 100644 --- a/cpp/tests/experimental/mg_bfs_test.cpp +++ b/cpp/tests/experimental/mg_bfs_test.cpp @@ -16,6 +16,7 @@ #include #include +#include #include #include @@ -82,6 +83,7 @@ class Tests_MGBFS : public ::testing::TestWithParam mg_graph(handle); rmm::device_uvector d_mg_renumber_map_labels(0, handle.get_stream()); std::tie(mg_graph, d_mg_renumber_map_labels) = @@ -128,121 +130,132 @@ class Tests_MGBFS : public ::testing::TestWithParam sg_graph(handle); - std::tie(sg_graph, std::ignore) = - input_usecase.template construct_graph( - handle, false, false); - - auto sg_graph_view = sg_graph.view(); - - std::vector vertex_partition_lasts(comm_size); - for (size_t i = 0; i < vertex_partition_lasts.size(); ++i) { - vertex_partition_lasts[i] = mg_graph_view.get_vertex_partition_last(i); - } - - rmm::device_scalar d_source(static_cast(bfs_usecase.source), - handle.get_stream()); - cugraph::experimental::unrenumber_int_vertices( - handle, - d_source.data(), - size_t{1}, - d_mg_renumber_map_labels.data(), - mg_graph_view.get_local_vertex_first(), - mg_graph_view.get_local_vertex_last(), - vertex_partition_lasts, - true); - auto unrenumbered_source = d_source.value(handle.get_stream()); - - // 5-2. run SG BFS - - rmm::device_uvector d_sg_distances(sg_graph_view.get_number_of_local_vertices(), - handle.get_stream()); - rmm::device_uvector d_sg_predecessors(sg_graph_view.get_number_of_local_vertices(), - handle.get_stream()); - - cugraph::experimental::bfs(handle, - sg_graph_view, - d_sg_distances.data(), - d_sg_predecessors.data(), - unrenumbered_source, - false, - std::numeric_limits::max()); - - // 5-3. compare - - std::vector h_sg_offsets(sg_graph_view.get_number_of_vertices() + 1); - std::vector h_sg_indices(sg_graph_view.get_number_of_edges()); - raft::update_host(h_sg_offsets.data(), - sg_graph_view.offsets(), - sg_graph_view.get_number_of_vertices() + 1, - handle.get_stream()); - raft::update_host(h_sg_indices.data(), - sg_graph_view.indices(), - sg_graph_view.get_number_of_edges(), - handle.get_stream()); - - std::vector h_sg_distances(sg_graph_view.get_number_of_vertices()); - std::vector h_sg_predecessors(sg_graph_view.get_number_of_vertices()); - raft::update_host( - h_sg_distances.data(), d_sg_distances.data(), d_sg_distances.size(), handle.get_stream()); - raft::update_host(h_sg_predecessors.data(), - d_sg_predecessors.data(), - d_sg_predecessors.size(), - handle.get_stream()); - - std::vector h_mg_distances(mg_graph_view.get_number_of_local_vertices()); - std::vector h_mg_predecessors(mg_graph_view.get_number_of_local_vertices()); - raft::update_host( - h_mg_distances.data(), d_mg_distances.data(), d_mg_distances.size(), handle.get_stream()); - cugraph::experimental::unrenumber_int_vertices( - handle, - d_mg_predecessors.data(), - d_mg_predecessors.size(), - d_mg_renumber_map_labels.data(), - mg_graph_view.get_local_vertex_first(), - mg_graph_view.get_local_vertex_last(), - vertex_partition_lasts, - true); - raft::update_host(h_mg_predecessors.data(), - d_mg_predecessors.data(), - d_mg_predecessors.size(), - handle.get_stream()); - - std::vector h_mg_renumber_map_labels(d_mg_renumber_map_labels.size()); - raft::update_host(h_mg_renumber_map_labels.data(), - d_mg_renumber_map_labels.data(), - d_mg_renumber_map_labels.size(), - handle.get_stream()); - - handle.get_stream_view().synchronize(); - - for (vertex_t i = 0; i < mg_graph_view.get_number_of_local_vertices(); ++i) { - auto mapped_vertex = h_mg_renumber_map_labels[i]; - ASSERT_TRUE(h_mg_distances[i] == h_sg_distances[mapped_vertex]) - << "MG BFS distance for vertex: " << mapped_vertex << " in rank: " << comm_rank - << " has value: " << h_mg_distances[i] - << " different from the corresponding SG value: " << h_sg_distances[mapped_vertex]; - if (h_mg_predecessors[i] == cugraph::invalid_vertex_id::value) { - ASSERT_TRUE(h_sg_predecessors[mapped_vertex] == h_mg_predecessors[i]) - << "vertex reachability does not match with the SG result."; - } else { - ASSERT_TRUE(h_sg_distances[h_mg_predecessors[i]] + 1 == h_sg_distances[mapped_vertex]) - << "distances to this vertex != distances to the predecessor vertex + 1."; - bool found{false}; - for (auto j = h_sg_offsets[h_mg_predecessors[i]]; - j < h_sg_offsets[h_mg_predecessors[i] + 1]; - ++j) { - if (h_sg_indices[j] == mapped_vertex) { - found = true; - break; + // 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_distances = + cugraph::test::device_gatherv(handle, d_mg_distances.data(), d_mg_distances.size()); + auto d_mg_aggregate_predecessors = + cugraph::test::device_gatherv(handle, d_mg_predecessors.data(), d_mg_predecessors.size()); + + if (handle.get_comms().get_rank() == int{0}) { + // 4-2. unrenumbr MG results + + cugraph::experimental::unrenumber_int_vertices( + handle, + d_mg_aggregate_predecessors.data(), + d_mg_aggregate_predecessors.size(), + d_mg_aggregate_renumber_map_labels.data(), + vertex_t{0}, + mg_graph_view.get_number_of_vertices(), + std::vector{mg_graph_view.get_number_of_vertices()}); + + std::tie(std::ignore, d_mg_aggregate_distances) = + cugraph::test::sort_by_key(handle, + d_mg_aggregate_renumber_map_labels.data(), + d_mg_aggregate_distances.data(), + d_mg_aggregate_renumber_map_labels.size()); + std::tie(std::ignore, d_mg_aggregate_predecessors) = + cugraph::test::sort_by_key(handle, + d_mg_aggregate_renumber_map_labels.data(), + d_mg_aggregate_predecessors.data(), + d_mg_aggregate_renumber_map_labels.size()); + + // 4-3. create SG graph + + cugraph::experimental::graph_t sg_graph(handle); + std::tie(sg_graph, std::ignore) = + input_usecase.template construct_graph( + handle, false, false); + + auto sg_graph_view = sg_graph.view(); + + ASSERT_TRUE(mg_graph_view.get_number_of_vertices() == + sg_graph_view.get_number_of_vertices()); + + // 4-4. run SG BFS + + rmm::device_uvector d_sg_distances(sg_graph_view.get_number_of_vertices(), + handle.get_stream()); + rmm::device_uvector d_sg_predecessors( + sg_graph_view.get_number_of_local_vertices(), handle.get_stream()); + + vertex_t unrenumbered_source{}; + raft::update_host(&unrenumbered_source, + d_mg_aggregate_renumber_map_labels.data() + bfs_usecase.source, + size_t{1}, + handle.get_stream()); + handle.get_stream_view().synchronize(); + + cugraph::experimental::bfs(handle, + sg_graph_view, + d_sg_distances.data(), + d_sg_predecessors.data(), + unrenumbered_source, + false, + std::numeric_limits::max()); + // 4-5. compare + + std::vector h_sg_offsets(sg_graph_view.get_number_of_vertices() + 1); + std::vector h_sg_indices(sg_graph_view.get_number_of_edges()); + raft::update_host(h_sg_offsets.data(), + sg_graph_view.offsets(), + sg_graph_view.get_number_of_vertices() + 1, + handle.get_stream()); + raft::update_host(h_sg_indices.data(), + sg_graph_view.indices(), + sg_graph_view.get_number_of_edges(), + handle.get_stream()); + + std::vector h_mg_aggregate_distances(mg_graph_view.get_number_of_vertices()); + std::vector h_mg_aggregate_predecessors(mg_graph_view.get_number_of_vertices()); + + raft::update_host(h_mg_aggregate_distances.data(), + d_mg_aggregate_distances.data(), + d_mg_aggregate_distances.size(), + handle.get_stream()); + raft::update_host(h_mg_aggregate_predecessors.data(), + d_mg_aggregate_predecessors.data(), + d_mg_aggregate_predecessors.size(), + handle.get_stream()); + + std::vector h_sg_distances(sg_graph_view.get_number_of_vertices()); + std::vector h_sg_predecessors(sg_graph_view.get_number_of_vertices()); + + raft::update_host( + h_sg_distances.data(), d_sg_distances.data(), d_sg_distances.size(), handle.get_stream()); + raft::update_host(h_sg_predecessors.data(), + d_sg_predecessors.data(), + d_sg_predecessors.size(), + handle.get_stream()); + + handle.get_stream_view().synchronize(); + + ASSERT_TRUE(std::equal(h_mg_aggregate_distances.begin(), + h_mg_aggregate_distances.end(), + h_sg_distances.begin())); + for (size_t i = 0; i < h_mg_aggregate_predecessors.size(); ++i) { + if (h_mg_aggregate_predecessors[i] == cugraph::invalid_vertex_id::value) { + ASSERT_TRUE(h_sg_predecessors[i] == h_mg_aggregate_predecessors[i]) + << "vertex reachability does not match with the SG result."; + } else { + ASSERT_TRUE(h_sg_distances[h_mg_aggregate_predecessors[i]] + 1 == h_sg_distances[i]) + << "distances to this vertex != distances to the predecessor vertex + 1."; + bool found{false}; + for (auto j = h_sg_offsets[h_mg_aggregate_predecessors[i]]; + j < h_sg_offsets[h_mg_aggregate_predecessors[i] + 1]; + ++j) { + if (h_sg_indices[j] == i) { + found = true; + break; + } } + ASSERT_TRUE(found) << "no edge from the predecessor vertex to this vertex."; } - ASSERT_TRUE(found) << "no edge from the predecessor vertex to this vertex."; } } } diff --git a/cpp/tests/experimental/mg_katz_centrality_test.cpp b/cpp/tests/experimental/mg_katz_centrality_test.cpp index eca04fb3241..e32bb45b5b7 100644 --- a/cpp/tests/experimental/mg_katz_centrality_test.cpp +++ b/cpp/tests/experimental/mg_katz_centrality_test.cpp @@ -16,6 +16,7 @@ #include #include +#include #include #include @@ -132,68 +133,81 @@ class Tests_MGKatzCentrality // 5. copmare SG & MG results if (katz_usecase.check_correctness) { - // 5-1. create SG graph - - cugraph::experimental::graph_t sg_graph(handle); - std::tie(sg_graph, std::ignore) = - input_usecase.template construct_graph( - handle, true, false); - - auto sg_graph_view = sg_graph.view(); - - // 5-3. run SG Katz Centrality - - rmm::device_uvector d_sg_katz_centralities(sg_graph_view.get_number_of_vertices(), - handle.get_stream()); - - cugraph::experimental::katz_centrality(handle, - sg_graph_view, - static_cast(nullptr), - d_sg_katz_centralities.data(), - alpha, - beta, - epsilon, - std::numeric_limits::max(), // max_iterations - false); - - // 5-4. compare - - std::vector h_sg_katz_centralities(sg_graph_view.get_number_of_vertices()); - raft::update_host(h_sg_katz_centralities.data(), - d_sg_katz_centralities.data(), - d_sg_katz_centralities.size(), - handle.get_stream()); - - std::vector h_mg_katz_centralities(mg_graph_view.get_number_of_local_vertices()); - raft::update_host(h_mg_katz_centralities.data(), - d_mg_katz_centralities.data(), - d_mg_katz_centralities.size(), - handle.get_stream()); - - std::vector h_mg_renumber_map_labels(d_mg_renumber_map_labels.size()); - raft::update_host(h_mg_renumber_map_labels.data(), - d_mg_renumber_map_labels.data(), - d_mg_renumber_map_labels.size(), - handle.get_stream()); - - handle.get_stream_view().synchronize(); - - auto threshold_ratio = 1e-3; - auto threshold_magnitude = - (1.0 / static_cast(mg_graph_view.get_number_of_vertices())) * - threshold_ratio; // skip comparison for low KatzCentrality verties (lowly ranked vertices) - auto nearly_equal = [threshold_ratio, threshold_magnitude](auto lhs, auto rhs) { - return std::abs(lhs - rhs) < - std::max(std::max(lhs, rhs) * threshold_ratio, threshold_magnitude); - }; - - for (vertex_t i = 0; i < mg_graph_view.get_number_of_local_vertices(); ++i) { - auto mapped_vertex = h_mg_renumber_map_labels[i]; - ASSERT_TRUE(nearly_equal(h_mg_katz_centralities[i], h_sg_katz_centralities[mapped_vertex])) - << "MG KatzCentrality value for vertex: " << mapped_vertex << " in rank: " << comm_rank - << " has value: " << h_mg_katz_centralities[i] - << " which exceeds the error margin for comparing to SG value: " - << h_sg_katz_centralities[mapped_vertex]; + // 5-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_katz_centralities = cugraph::test::device_gatherv( + handle, d_mg_katz_centralities.data(), d_mg_katz_centralities.size()); + + if (handle.get_comms().get_rank() == int{0}) { + // 5-2. unrenumbr MG results + + std::tie(std::ignore, d_mg_aggregate_katz_centralities) = + cugraph::test::sort_by_key(handle, + d_mg_aggregate_renumber_map_labels.data(), + d_mg_aggregate_katz_centralities.data(), + d_mg_aggregate_renumber_map_labels.size()); + + // 5-3. create SG graph + + cugraph::experimental::graph_t sg_graph(handle); + std::tie(sg_graph, std::ignore) = + input_usecase.template construct_graph( + handle, true, false); + + auto sg_graph_view = sg_graph.view(); + + ASSERT_TRUE(mg_graph_view.get_number_of_vertices() == + sg_graph_view.get_number_of_vertices()); + + // 5-4. run SG Katz Centrality + + rmm::device_uvector d_sg_katz_centralities(sg_graph_view.get_number_of_vertices(), + handle.get_stream()); + + cugraph::experimental::katz_centrality( + handle, + sg_graph_view, + static_cast(nullptr), + d_sg_katz_centralities.data(), + alpha, + beta, + epsilon, + std::numeric_limits::max(), // max_iterations + false); + + // 5-5. compare + + std::vector h_mg_aggregate_katz_centralities( + mg_graph_view.get_number_of_vertices()); + raft::update_host(h_mg_aggregate_katz_centralities.data(), + d_mg_aggregate_katz_centralities.data(), + d_mg_aggregate_katz_centralities.size(), + handle.get_stream()); + + std::vector h_sg_katz_centralities(sg_graph_view.get_number_of_vertices()); + raft::update_host(h_sg_katz_centralities.data(), + d_sg_katz_centralities.data(), + d_sg_katz_centralities.size(), + handle.get_stream()); + + handle.get_stream_view().synchronize(); + + auto threshold_ratio = 1e-3; + auto threshold_magnitude = + (1.0 / static_cast(mg_graph_view.get_number_of_vertices())) * + threshold_ratio; // skip comparison for low KatzCentrality verties (lowly ranked + // vertices) + auto nearly_equal = [threshold_ratio, threshold_magnitude](auto lhs, auto rhs) { + return std::abs(lhs - rhs) < + std::max(std::max(lhs, rhs) * threshold_ratio, threshold_magnitude); + }; + + ASSERT_TRUE(std::equal(h_mg_aggregate_katz_centralities.begin(), + h_mg_aggregate_katz_centralities.end(), + h_sg_katz_centralities.begin(), + nearly_equal)); } } } diff --git a/cpp/tests/experimental/mg_sssp_test.cpp b/cpp/tests/experimental/mg_sssp_test.cpp index d3da904afc9..190e831841f 100644 --- a/cpp/tests/experimental/mg_sssp_test.cpp +++ b/cpp/tests/experimental/mg_sssp_test.cpp @@ -16,6 +16,7 @@ #include #include +#include #include #include @@ -110,7 +111,6 @@ class Tests_MGSSSP : public ::testing::TestWithParam sg_graph(handle); - std::tie(sg_graph, std::ignore) = - input_usecase.template construct_graph( - handle, true, false); - - auto sg_graph_view = sg_graph.view(); - - std::vector vertex_partition_lasts(comm_size); - for (size_t i = 0; i < vertex_partition_lasts.size(); ++i) { - vertex_partition_lasts[i] = mg_graph_view.get_vertex_partition_last(i); - } - - rmm::device_scalar d_source(static_cast(sssp_usecase.source), - handle.get_stream()); - cugraph::experimental::unrenumber_int_vertices( - handle, - d_source.data(), - size_t{1}, - d_mg_renumber_map_labels.data(), - mg_graph_view.get_local_vertex_first(), - mg_graph_view.get_local_vertex_last(), - vertex_partition_lasts, - true); - auto unrenumbered_source = d_source.value(handle.get_stream()); - - // 5-2. run SG SSSP - - rmm::device_uvector d_sg_distances(sg_graph_view.get_number_of_local_vertices(), - handle.get_stream()); - rmm::device_uvector d_sg_predecessors(sg_graph_view.get_number_of_local_vertices(), - handle.get_stream()); - - // FIXME: disable do_expensive_check - cugraph::experimental::sssp(handle, - sg_graph_view, - d_sg_distances.data(), - d_sg_predecessors.data(), - unrenumbered_source, - std::numeric_limits::max()); - - // 5-3. compare - - std::vector h_sg_offsets(sg_graph_view.get_number_of_vertices() + 1); - std::vector h_sg_indices(sg_graph_view.get_number_of_edges()); - std::vector h_sg_weights(sg_graph_view.get_number_of_edges()); - raft::update_host(h_sg_offsets.data(), - sg_graph_view.offsets(), - sg_graph_view.get_number_of_vertices() + 1, - handle.get_stream()); - raft::update_host(h_sg_indices.data(), - sg_graph_view.indices(), - sg_graph_view.get_number_of_edges(), - handle.get_stream()); - raft::update_host(h_sg_weights.data(), - sg_graph_view.weights(), - sg_graph_view.get_number_of_edges(), - handle.get_stream()); - - std::vector h_sg_distances(sg_graph_view.get_number_of_vertices()); - std::vector h_sg_predecessors(sg_graph_view.get_number_of_vertices()); - raft::update_host( - h_sg_distances.data(), d_sg_distances.data(), d_sg_distances.size(), handle.get_stream()); - raft::update_host(h_sg_predecessors.data(), - d_sg_predecessors.data(), - d_sg_predecessors.size(), - handle.get_stream()); - - std::vector h_mg_distances(mg_graph_view.get_number_of_local_vertices()); - std::vector h_mg_predecessors(mg_graph_view.get_number_of_local_vertices()); - raft::update_host( - h_mg_distances.data(), d_mg_distances.data(), d_mg_distances.size(), handle.get_stream()); - cugraph::experimental::unrenumber_int_vertices( - handle, - d_mg_predecessors.data(), - d_mg_predecessors.size(), - d_mg_renumber_map_labels.data(), - mg_graph_view.get_local_vertex_first(), - mg_graph_view.get_local_vertex_last(), - vertex_partition_lasts, - true); - raft::update_host(h_mg_predecessors.data(), - d_mg_predecessors.data(), - d_mg_predecessors.size(), - handle.get_stream()); - - std::vector h_mg_renumber_map_labels(d_mg_renumber_map_labels.size()); - raft::update_host(h_mg_renumber_map_labels.data(), - d_mg_renumber_map_labels.data(), - d_mg_renumber_map_labels.size(), - handle.get_stream()); - - handle.get_stream_view().synchronize(); - - auto max_weight_element = std::max_element(h_sg_weights.begin(), h_sg_weights.end()); - auto epsilon = *max_weight_element * weight_t{1e-6}; - auto nearly_equal = [epsilon](auto lhs, auto rhs) { return std::fabs(lhs - rhs) < epsilon; }; - - for (vertex_t i = 0; i < mg_graph_view.get_number_of_local_vertices(); ++i) { - auto mapped_vertex = h_mg_renumber_map_labels[i]; - ASSERT_TRUE(nearly_equal(h_mg_distances[i], h_sg_distances[mapped_vertex])) - << "MG SSSP distance for vertex: " << mapped_vertex << " in rank: " << comm_rank - << " has value: " << h_mg_distances[i] - << " different from the corresponding SG value: " << h_sg_distances[mapped_vertex]; - if (h_mg_predecessors[i] == cugraph::invalid_vertex_id::value) { - ASSERT_TRUE(h_sg_predecessors[mapped_vertex] == h_mg_predecessors[i]) - << "vertex reachability does not match with the SG result."; - } else { - auto pred_distance = h_sg_distances[h_mg_predecessors[i]]; - bool found{false}; - for (auto j = h_sg_offsets[h_mg_predecessors[i]]; - j < h_sg_offsets[h_mg_predecessors[i] + 1]; - ++j) { - if (h_sg_indices[j] == mapped_vertex) { - if (nearly_equal(pred_distance + h_sg_weights[j], h_sg_distances[mapped_vertex])) { - found = true; - break; + // 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_distances = + cugraph::test::device_gatherv(handle, d_mg_distances.data(), d_mg_distances.size()); + auto d_mg_aggregate_predecessors = + cugraph::test::device_gatherv(handle, d_mg_predecessors.data(), d_mg_predecessors.size()); + + if (handle.get_comms().get_rank() == int{0}) { + // 4-2. unrenumber MG results + + cugraph::experimental::unrenumber_int_vertices( + handle, + d_mg_aggregate_predecessors.data(), + d_mg_aggregate_predecessors.size(), + d_mg_aggregate_renumber_map_labels.data(), + vertex_t{0}, + mg_graph_view.get_number_of_vertices(), + std::vector{mg_graph_view.get_number_of_vertices()}); + + std::tie(std::ignore, d_mg_aggregate_distances) = + cugraph::test::sort_by_key(handle, + d_mg_aggregate_renumber_map_labels.data(), + d_mg_aggregate_distances.data(), + d_mg_aggregate_renumber_map_labels.size()); + std::tie(std::ignore, d_mg_aggregate_predecessors) = + cugraph::test::sort_by_key(handle, + d_mg_aggregate_renumber_map_labels.data(), + d_mg_aggregate_predecessors.data(), + d_mg_aggregate_renumber_map_labels.size()); + + // 4-3. create SG graph + + cugraph::experimental::graph_t sg_graph(handle); + std::tie(sg_graph, std::ignore) = + input_usecase.template construct_graph( + handle, true, false); + + auto sg_graph_view = sg_graph.view(); + + ASSERT_TRUE(mg_graph_view.get_number_of_vertices() == + sg_graph_view.get_number_of_vertices()); + + // 4-4. run SG SSSP + + rmm::device_uvector d_sg_distances(sg_graph_view.get_number_of_local_vertices(), + handle.get_stream()); + rmm::device_uvector d_sg_predecessors( + sg_graph_view.get_number_of_local_vertices(), handle.get_stream()); + vertex_t unrenumbered_source{}; + raft::update_host(&unrenumbered_source, + d_mg_aggregate_renumber_map_labels.data() + sssp_usecase.source, + size_t{1}, + handle.get_stream()); + handle.get_stream_view().synchronize(); + + cugraph::experimental::sssp(handle, + sg_graph_view, + d_sg_distances.data(), + d_sg_predecessors.data(), + unrenumbered_source, + std::numeric_limits::max()); + + // 4-5. compare + + std::vector h_sg_offsets(sg_graph_view.get_number_of_vertices() + 1); + std::vector h_sg_indices(sg_graph_view.get_number_of_edges()); + std::vector h_sg_weights(sg_graph_view.get_number_of_edges()); + raft::update_host(h_sg_offsets.data(), + sg_graph_view.offsets(), + sg_graph_view.get_number_of_vertices() + 1, + handle.get_stream()); + raft::update_host(h_sg_indices.data(), + sg_graph_view.indices(), + sg_graph_view.get_number_of_edges(), + handle.get_stream()); + raft::update_host(h_sg_weights.data(), + sg_graph_view.weights(), + sg_graph_view.get_number_of_edges(), + handle.get_stream()); + + std::vector h_mg_aggregate_distances(mg_graph_view.get_number_of_vertices()); + std::vector h_mg_aggregate_predecessors(mg_graph_view.get_number_of_vertices()); + raft::update_host(h_mg_aggregate_distances.data(), + d_mg_aggregate_distances.data(), + d_mg_aggregate_distances.size(), + handle.get_stream()); + raft::update_host(h_mg_aggregate_predecessors.data(), + d_mg_aggregate_predecessors.data(), + d_mg_aggregate_predecessors.size(), + handle.get_stream()); + + std::vector h_sg_distances(sg_graph_view.get_number_of_vertices()); + std::vector h_sg_predecessors(sg_graph_view.get_number_of_vertices()); + raft::update_host( + h_sg_distances.data(), d_sg_distances.data(), d_sg_distances.size(), handle.get_stream()); + raft::update_host(h_sg_predecessors.data(), + d_sg_predecessors.data(), + d_sg_predecessors.size(), + handle.get_stream()); + + handle.get_stream_view().synchronize(); + + auto max_weight_element = std::max_element(h_sg_weights.begin(), h_sg_weights.end()); + auto epsilon = *max_weight_element * weight_t{1e-6}; + auto nearly_equal = [epsilon](auto lhs, auto rhs) { + return std::fabs(lhs - rhs) < epsilon; + }; + + ASSERT_TRUE(std::equal(h_mg_aggregate_distances.begin(), + h_mg_aggregate_distances.end(), + h_sg_distances.begin(), + nearly_equal)); + + for (size_t i = 0; i < h_mg_aggregate_predecessors.size(); ++i) { + if (h_mg_aggregate_predecessors[i] == cugraph::invalid_vertex_id::value) { + ASSERT_TRUE(h_sg_predecessors[i] == h_mg_aggregate_predecessors[i]) + << "vertex reachability does not match with the SG result."; + } else { + auto pred_distance = h_sg_distances[h_mg_aggregate_predecessors[i]]; + bool found{false}; + for (auto j = h_sg_offsets[h_mg_aggregate_predecessors[i]]; + j < h_sg_offsets[h_mg_aggregate_predecessors[i] + 1]; + ++j) { + if (h_sg_indices[j] == i) { + if (nearly_equal(pred_distance + h_sg_weights[j], h_sg_distances[i])) { + found = true; + break; + } } } + ASSERT_TRUE(found) + << "no edge from the predecessor vertex to this vertex with the matching weight."; } - ASSERT_TRUE(found) - << "no edge from the predecessor vertex to this vertex with the matching weight."; } } } @@ -288,7 +300,7 @@ INSTANTIATE_TEST_SUITE_P( ::testing::Values( // enable correctness checks std::make_tuple(SSSP_Usecase{0}, - cugraph::test::Rmat_Usecase(10, 16, 0.57, 0.19, 0.19, 0, false, false)))); + cugraph::test::Rmat_Usecase(10, 16, 0.57, 0.19, 0.19, 0, false, false, true)))); INSTANTIATE_TEST_SUITE_P( rmat_large_test, @@ -296,6 +308,6 @@ INSTANTIATE_TEST_SUITE_P( ::testing::Values( // disable correctness checks for large graphs std::make_tuple(SSSP_Usecase{0, false}, - cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); + cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false, true)))); CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/pagerank/mg_pagerank_test.cpp b/cpp/tests/pagerank/mg_pagerank_test.cpp index 8ccf967d95a..9c0a56daf29 100644 --- a/cpp/tests/pagerank/mg_pagerank_test.cpp +++ b/cpp/tests/pagerank/mg_pagerank_test.cpp @@ -16,6 +16,7 @@ #include #include +#include #include #include @@ -171,135 +172,96 @@ class Tests_MGPageRank // 5. copmare SG & MG results if (pagerank_usecase.check_correctness) { - // 5-1. create SG graph - - cugraph::experimental::graph_t sg_graph(handle); - std::tie(sg_graph, std::ignore) = - input_usecase.template construct_graph( - handle, true, false); - - auto sg_graph_view = sg_graph.view(); - - // 5-2. collect personalization vertex/value pairs - - rmm::device_uvector d_sg_personalization_vertices(0, handle.get_stream()); - rmm::device_uvector d_sg_personalization_values(0, handle.get_stream()); - if (pagerank_usecase.personalization_ratio > 0.0) { - rmm::device_uvector d_unrenumbered_personalization_vertices( - d_mg_personalization_vertices.size(), handle.get_stream()); - rmm::device_uvector d_unrenumbered_personalization_values( - d_unrenumbered_personalization_vertices.size(), handle.get_stream()); - raft::copy_async(d_unrenumbered_personalization_vertices.data(), - d_mg_personalization_vertices.data(), - d_mg_personalization_vertices.size(), - handle.get_stream()); - raft::copy_async(d_unrenumbered_personalization_values.data(), - d_mg_personalization_values.data(), - d_mg_personalization_values.size(), - handle.get_stream()); - - std::vector vertex_partition_lasts(comm_size); - for (size_t i = 0; i < vertex_partition_lasts.size(); ++i) { - vertex_partition_lasts[i] = mg_graph_view.get_vertex_partition_last(i); - } - cugraph::experimental::unrenumber_int_vertices( + // 5-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_personalization_vertices = cugraph::test::device_gatherv( + handle, d_mg_personalization_vertices.data(), d_mg_personalization_vertices.size()); + auto d_mg_aggregate_personalization_values = cugraph::test::device_gatherv( + handle, d_mg_personalization_values.data(), d_mg_personalization_values.size()); + auto d_mg_aggregate_pageranks = + cugraph::test::device_gatherv(handle, d_mg_pageranks.data(), d_mg_pageranks.size()); + + if (handle.get_comms().get_rank() == int{0}) { + // 5-2. unrenumbr MG results + + cugraph::experimental::unrenumber_int_vertices( handle, - d_unrenumbered_personalization_vertices.data(), - d_unrenumbered_personalization_vertices.size(), - d_mg_renumber_map_labels.data(), - mg_graph_view.get_local_vertex_first(), - mg_graph_view.get_local_vertex_last(), - vertex_partition_lasts, - handle.get_stream()); - - rmm::device_scalar d_local_personalization_vector_size( - d_unrenumbered_personalization_vertices.size(), handle.get_stream()); - rmm::device_uvector d_recvcounts(comm_size, handle.get_stream()); - comm.allgather( - d_local_personalization_vector_size.data(), d_recvcounts.data(), 1, handle.get_stream()); - std::vector recvcounts(d_recvcounts.size()); - raft::update_host( - recvcounts.data(), d_recvcounts.data(), d_recvcounts.size(), handle.get_stream()); - auto status = comm.sync_stream(handle.get_stream()); - ASSERT_EQ(status, raft::comms::status_t::SUCCESS); - - std::vector displacements(recvcounts.size(), size_t{0}); - std::partial_sum(recvcounts.begin(), recvcounts.end() - 1, displacements.begin() + 1); - - d_sg_personalization_vertices.resize(displacements.back() + recvcounts.back(), - handle.get_stream()); - d_sg_personalization_values.resize(d_sg_personalization_vertices.size(), - handle.get_stream()); - - comm.allgatherv(d_unrenumbered_personalization_vertices.data(), - d_sg_personalization_vertices.data(), - recvcounts.data(), - displacements.data(), - handle.get_stream()); - comm.allgatherv(d_unrenumbered_personalization_values.data(), - d_sg_personalization_values.data(), - recvcounts.data(), - displacements.data(), - handle.get_stream()); - - std::tie(d_unrenumbered_personalization_vertices, d_unrenumbered_personalization_values) = + d_mg_aggregate_personalization_vertices.data(), + d_mg_aggregate_personalization_vertices.size(), + d_mg_aggregate_renumber_map_labels.data(), + vertex_t{0}, + mg_graph_view.get_number_of_vertices(), + std::vector{mg_graph_view.get_number_of_vertices()}); + std::tie(d_mg_aggregate_personalization_vertices, d_mg_aggregate_personalization_values) = cugraph::test::sort_by_key(handle, - d_unrenumbered_personalization_vertices.data(), - d_unrenumbered_personalization_values.data(), - d_unrenumbered_personalization_vertices.size()); - } + d_mg_aggregate_personalization_vertices.data(), + d_mg_aggregate_personalization_values.data(), + d_mg_aggregate_personalization_vertices.size()); + std::tie(std::ignore, d_mg_aggregate_pageranks) = + cugraph::test::sort_by_key(handle, + d_mg_aggregate_renumber_map_labels.data(), + d_mg_aggregate_pageranks.data(), + d_mg_aggregate_renumber_map_labels.size()); + + // 5-3. create SG graph + + cugraph::experimental::graph_t sg_graph(handle); + std::tie(sg_graph, std::ignore) = + input_usecase.template construct_graph( + handle, true, false); - // 5-3. run SG PageRank - - rmm::device_uvector d_sg_pageranks(sg_graph_view.get_number_of_vertices(), - handle.get_stream()); - - cugraph::experimental::pagerank(handle, - sg_graph_view, - static_cast(nullptr), - d_sg_personalization_vertices.data(), - d_sg_personalization_values.data(), - static_cast(d_sg_personalization_vertices.size()), - d_sg_pageranks.data(), - alpha, - epsilon, - std::numeric_limits::max(), // max_iterations - false); - - // 5-4. compare - - std::vector h_sg_pageranks(sg_graph_view.get_number_of_vertices()); - raft::update_host( - h_sg_pageranks.data(), d_sg_pageranks.data(), d_sg_pageranks.size(), handle.get_stream()); - - std::vector h_mg_pageranks(mg_graph_view.get_number_of_local_vertices()); - raft::update_host( - h_mg_pageranks.data(), d_mg_pageranks.data(), d_mg_pageranks.size(), handle.get_stream()); - - std::vector h_mg_renumber_map_labels(d_mg_renumber_map_labels.size()); - raft::update_host(h_mg_renumber_map_labels.data(), - d_mg_renumber_map_labels.data(), - d_mg_renumber_map_labels.size(), - handle.get_stream()); - - handle.get_stream_view().synchronize(); - - auto threshold_ratio = 1e-3; - auto threshold_magnitude = - (1.0 / static_cast(mg_graph_view.get_number_of_vertices())) * - threshold_ratio; // skip comparison for low PageRank verties (lowly ranked vertices) - auto nearly_equal = [threshold_ratio, threshold_magnitude](auto lhs, auto rhs) { - return std::abs(lhs - rhs) < - std::max(std::max(lhs, rhs) * threshold_ratio, threshold_magnitude); - }; - - for (vertex_t i = 0; i < mg_graph_view.get_number_of_local_vertices(); ++i) { - auto mapped_vertex = h_mg_renumber_map_labels[i]; - ASSERT_TRUE(nearly_equal(h_mg_pageranks[i], h_sg_pageranks[mapped_vertex])) - << "MG PageRank value for vertex: " << mapped_vertex << " in rank: " << comm_rank - << " has value: " << h_mg_pageranks[i] - << " which exceeds the error margin for comparing to SG value: " - << h_sg_pageranks[mapped_vertex]; + auto sg_graph_view = sg_graph.view(); + + ASSERT_TRUE(mg_graph_view.get_number_of_vertices() == + sg_graph_view.get_number_of_vertices()); + + // 5-4. run SG PageRank + + rmm::device_uvector d_sg_pageranks(sg_graph_view.get_number_of_vertices(), + handle.get_stream()); + + cugraph::experimental::pagerank( + handle, + sg_graph_view, + static_cast(nullptr), + d_mg_aggregate_personalization_vertices.data(), + d_mg_aggregate_personalization_values.data(), + static_cast(d_mg_aggregate_personalization_vertices.size()), + d_sg_pageranks.data(), + alpha, + epsilon, + std::numeric_limits::max(), // max_iterations + false); + + // 5-4. compare + + std::vector h_mg_aggregate_pageranks(mg_graph_view.get_number_of_vertices()); + raft::update_host(h_mg_aggregate_pageranks.data(), + d_mg_aggregate_pageranks.data(), + d_mg_aggregate_pageranks.size(), + handle.get_stream()); + + std::vector h_sg_pageranks(sg_graph_view.get_number_of_vertices()); + raft::update_host( + h_sg_pageranks.data(), d_sg_pageranks.data(), d_sg_pageranks.size(), handle.get_stream()); + + handle.get_stream_view().synchronize(); + + auto threshold_ratio = 1e-3; + auto threshold_magnitude = + (1.0 / static_cast(mg_graph_view.get_number_of_vertices())) * + threshold_ratio; // skip comparison for low PageRank verties (lowly ranked vertices) + auto nearly_equal = [threshold_ratio, threshold_magnitude](auto lhs, auto rhs) { + return std::abs(lhs - rhs) < + std::max(std::max(lhs, rhs) * threshold_ratio, threshold_magnitude); + }; + + ASSERT_TRUE(std::equal(h_mg_aggregate_pageranks.begin(), + h_mg_aggregate_pageranks.end(), + h_sg_pageranks.begin(), + nearly_equal)); } } } From 42c7c5f49789c6b386fcc96ef3445a3da74603fb Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 13 May 2021 14:16:32 -0400 Subject: [PATCH 8/9] add barrier in MG performance measurements (PERF=1) for more accurate performance measurement --- cpp/tests/experimental/mg_bfs_test.cpp | 4 ++++ cpp/tests/experimental/mg_katz_centrality_test.cpp | 4 ++++ cpp/tests/experimental/mg_sssp_test.cpp | 4 ++++ cpp/tests/pagerank/mg_pagerank_test.cpp | 4 ++++ 4 files changed, 16 insertions(+) diff --git a/cpp/tests/experimental/mg_bfs_test.cpp b/cpp/tests/experimental/mg_bfs_test.cpp index ef980a2448e..a832e0f99ac 100644 --- a/cpp/tests/experimental/mg_bfs_test.cpp +++ b/cpp/tests/experimental/mg_bfs_test.cpp @@ -81,6 +81,7 @@ class Tests_MGBFS : public ::testing::TestWithParam mg_graph(handle); @@ -88,6 +89,7 @@ class Tests_MGKatzCentrality if (PERF) { 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"; @@ -110,6 +112,7 @@ class Tests_MGKatzCentrality if (PERF) { CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle.get_comms().barrier(); hr_clock.start(); } @@ -125,6 +128,7 @@ class Tests_MGKatzCentrality if (PERF) { CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle.get_comms().barrier(); double elapsed_time{0.0}; hr_clock.stop(&elapsed_time); std::cout << "MG Katz Centrality took " << elapsed_time * 1e-6 << " s.\n"; diff --git a/cpp/tests/experimental/mg_sssp_test.cpp b/cpp/tests/experimental/mg_sssp_test.cpp index 190e831841f..8568545cbd6 100644 --- a/cpp/tests/experimental/mg_sssp_test.cpp +++ b/cpp/tests/experimental/mg_sssp_test.cpp @@ -78,6 +78,7 @@ class Tests_MGSSSP : public ::testing::TestWithParam mg_graph(handle); @@ -88,6 +89,7 @@ class Tests_MGSSSP : public ::testing::TestWithParam mg_graph(handle); @@ -90,6 +91,7 @@ class Tests_MGPageRank if (PERF) { 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"; @@ -147,6 +149,7 @@ class Tests_MGPageRank if (PERF) { CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle.get_comms().barrier(); hr_clock.start(); } @@ -164,6 +167,7 @@ class Tests_MGPageRank if (PERF) { CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle.get_comms().barrier(); double elapsed_time{0.0}; hr_clock.stop(&elapsed_time); std::cout << "MG PageRank took " << elapsed_time * 1e-6 << " s.\n"; From 175971446acedd6b9ac99b3437b253f8d86eaec8 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 13 May 2021 14:24:25 -0400 Subject: [PATCH 9/9] clang-format --- cpp/tests/utilities/device_comm_wrapper.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/tests/utilities/device_comm_wrapper.cu b/cpp/tests/utilities/device_comm_wrapper.cu index cda355db826..2d66e05c59f 100644 --- a/cpp/tests/utilities/device_comm_wrapper.cu +++ b/cpp/tests/utilities/device_comm_wrapper.cu @@ -61,12 +61,12 @@ template rmm::device_uvector device_gatherv(raft::handle_t const &handl size_t size); template rmm::device_uvector device_gatherv(raft::handle_t const &handle, - float const *d_input, - size_t size); + float const *d_input, + size_t size); template rmm::device_uvector device_gatherv(raft::handle_t const &handle, - double const *d_input, - size_t size); + double const *d_input, + size_t size); } // namespace test } // namespace cugraph