From 7217cfa36849cf8e46eb8ecc90015ac1f462ffc3 Mon Sep 17 00:00:00 2001 From: Chuck Hastings Date: Wed, 23 Mar 2022 10:53:45 -0400 Subject: [PATCH 1/2] add hits implementation --- .../cugraph/detail/shuffle_wrappers.hpp | 24 +++ cpp/include/cugraph_c/graph.h | 4 +- cpp/src/c_api/hits.cpp | 65 +++--- cpp/src/detail/shuffle_wrappers.cu | 90 ++++++++ cpp/tests/CMakeLists.txt | 1 + cpp/tests/c_api/hits_test.c | 17 +- cpp/tests/c_api/mg_hits_test.c | 202 ++++++++++++++++++ cpp/tests/c_api/mg_pagerank_test.c | 9 +- cpp/tests/c_api/mg_test_utils.cpp | 19 ++ cpp/tests/c_api/mg_test_utils.h | 2 + 10 files changed, 386 insertions(+), 47 deletions(-) create mode 100644 cpp/tests/c_api/mg_hits_test.c diff --git a/cpp/include/cugraph/detail/shuffle_wrappers.hpp b/cpp/include/cugraph/detail/shuffle_wrappers.hpp index db02ab94a5d..c22c98a5672 100644 --- a/cpp/include/cugraph/detail/shuffle_wrappers.hpp +++ b/cpp/include/cugraph/detail/shuffle_wrappers.hpp @@ -93,5 +93,29 @@ rmm::device_uvector groupby_and_count_edgelist_by_local_partition_id( std::optional>& d_edgelist_weights, bool groupby_and_count_local_partition_by_minor = false); +/** + * @brief Collect vertex values (represented as k/v pairs across cluster) and update the + * local value arrays on the GPU responsible for each vertex. + * + * Data will be shuffled and d_local_values[d_vertices[i]] = d_values[i] + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam value_t Type of value associated with the vertex. + * @tparam bool multi_gpu flag + * + * @param[in] handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, + * and handles to various CUDA libraries) to run graph algorithms. + * @param[in/out] d_vertices Vertex IDs for the k/v pair + * @param[in/out] d_values Values for the k/v pair + * @param[out] d_local_values The device vector on each GPU that should be updated + * @param[in] local_vertex_first The first vertex id assigned to the local GPU + */ +template +void collect_vertex_values_to_local(raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + vertex_t local_vertex_first); + } // namespace detail } // namespace cugraph diff --git a/cpp/include/cugraph_c/graph.h b/cpp/include/cugraph_c/graph.h index 092c3c4b91a..a4fde4dde22 100644 --- a/cpp/include/cugraph_c/graph.h +++ b/cpp/include/cugraph_c/graph.h @@ -32,7 +32,7 @@ typedef struct { bool_t is_multigraph; } cugraph_graph_properties_t; -// FIXME: Add support for specifying isloated vertices +// FIXME: Add support for specifying isolated vertices /** * @brief Construct an SG graph * @@ -74,7 +74,7 @@ cugraph_error_code_t cugraph_sg_graph_create(const cugraph_resource_handle_t* ha // but didn't want to confuse with original cugraph_free_graph void cugraph_sg_graph_free(cugraph_graph_t* graph); -// FIXME: Add support for specifying isloated vertices +// FIXME: Add support for specifying isolated vertices /** * @brief Construct an MG graph * diff --git a/cpp/src/c_api/hits.cpp b/cpp/src/c_api/hits.cpp index 7ca5f20e4d0..09f6ac3b7ea 100644 --- a/cpp/src/c_api/hits.cpp +++ b/cpp/src/c_api/hits.cpp @@ -22,6 +22,7 @@ #include #include +#include #include #include @@ -113,40 +114,46 @@ struct hits_functor : public cugraph::c_api::abstract_functor { weight_t hub_score_differences{0}; size_t number_of_iterations{0}; -#if 0 - // FIXME: Implementation will look something like this. - - if (initial_hubs_guess_ != nullptr) { - // - // Need to renumber initial_hubs_guess_vertices - // Need to shuffle and populate hubs - // - // This is the original pagerank code, it will be sort of like this - renumber_ext_vertices(handle_, - personalization_vertices_->as_type(), - personalization_vertices_->size_, - number_map->data(), - graph_view.get_local_vertex_first(), - graph_view.get_local_vertex_last(), - do_expensive_check_); + if (initial_hubs_guess_vertices_ != nullptr) { + rmm::device_uvector guess_vertices(initial_hubs_guess_vertices_->size_, + handle_.get_stream()); + rmm::device_uvector guess_values(initial_hubs_guess_values_->size_, + handle_.get_stream()); + + raft::copy(guess_vertices.data(), + initial_hubs_guess_vertices_->as_type(), + guess_vertices.size(), + handle_.get_stream()); + raft::copy(guess_values.data(), + initial_hubs_guess_values_->as_type(), + guess_values.size(), + handle_.get_stream()); + + cugraph::renumber_ext_vertices(handle_, + guess_vertices.data(), + guess_vertices.size(), + number_map->data(), + graph_view.get_local_vertex_first(), + graph_view.get_local_vertex_last(), + do_expensive_check_); + + cugraph::detail::collect_vertex_values_to_local( + handle_, guess_vertices, guess_values, hubs, graph_view.get_local_vertex_first()); } - // TODO: Add these to the result std::tie(hub_score_differences, number_of_iterations) = - cugraph::hits(handle_, - graph_view, - hubs.data(), - authorities.data(), - epsilon_, - max_iterations_, - has_initial_hubs_guess, - normalize_, - do_expensive_check_); + cugraph::hits( + handle_, + graph_view, + hubs.data(), + authorities.data(), + epsilon_, + max_iterations_, + (initial_hubs_guess_vertices_ != nullptr), + normalize_, + do_expensive_check_); raft::copy(vertex_ids.data(), number_map->data(), vertex_ids.size(), handle_.get_stream()); -#else - unsupported(); -#endif result_ = new cugraph::c_api::cugraph_hits_result_t{ new cugraph::c_api::cugraph_type_erased_device_array_t(vertex_ids, graph_->vertex_type_), diff --git a/cpp/src/detail/shuffle_wrappers.cu b/cpp/src/detail/shuffle_wrappers.cu index 26bdd21a1f9..bb69b0dd6ca 100644 --- a/cpp/src/detail/shuffle_wrappers.cu +++ b/cpp/src/detail/shuffle_wrappers.cu @@ -335,5 +335,95 @@ template rmm::device_uvector groupby_and_count_edgelist_by_local_partiti std::optional>& d_edgelist_weights, bool groupby_and_counts_local_partition); +template +void collect_vertex_values_to_local(raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + vertex_t local_vertex_first) +{ + if constexpr (multi_gpu) { + auto& comm = handle.get_comms(); + auto const comm_size = comm.get_size(); + + rmm::device_uvector d_rx_vertices(0, handle.get_stream()); + rmm::device_uvector d_rx_values(0, handle.get_stream()); + + std::tie(d_rx_vertices, d_rx_values, std::ignore) = + cugraph::groupby_gpu_id_and_shuffle_kv_pairs( + comm, + d_vertices.begin(), + d_vertices.end(), + d_values.begin(), + [key_func = cugraph::detail::compute_gpu_id_from_vertex_t{comm_size}] __device__( + auto val) { return key_func(val); }, + handle.get_stream()); + + auto vertex_iterator = thrust::make_transform_iterator( + d_rx_vertices.begin(), + [local_vertex_first] __device__(vertex_t v) { return v - local_vertex_first; }); + + thrust::scatter(handle.get_thrust_policy(), + d_rx_values.begin(), + d_rx_values.end(), + vertex_iterator, + d_local_values.begin()); + } else { + thrust::scatter(handle.get_thrust_policy(), + d_values.begin(), + d_values.end(), + d_vertices.begin(), + d_local_values.begin()); + } +} + +template void collect_vertex_values_to_local(raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int32_t local_vertex_first); + +template void collect_vertex_values_to_local(raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int64_t local_vertex_first); + +template void collect_vertex_values_to_local(raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int32_t local_vertex_first); + +template void collect_vertex_values_to_local(raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int64_t local_vertex_first); + +template void collect_vertex_values_to_local(raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int32_t local_vertex_first); + +template void collect_vertex_values_to_local(raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int64_t local_vertex_first); + +template void collect_vertex_values_to_local(raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int32_t local_vertex_first); + +template void collect_vertex_values_to_local(raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int64_t local_vertex_first); + } // namespace detail } // namespace cugraph diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index b7826bceeec..1482f4f810f 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -620,6 +620,7 @@ if(BUILD_CUGRAPH_MG_TESTS) # - MG C API tests ------------------------------------------------------------------------ ConfigureCTestMG(MG_CAPI_CREATE_GRAPH c_api/mg_create_graph_test.c c_api/mg_test_utils.cpp) ConfigureCTestMG(MG_CAPI_PAGERANK c_api/mg_pagerank_test.c c_api/mg_test_utils.cpp) + ConfigureCTestMG(MG_CAPI_HITS c_api/mg_hits_test.c c_api/mg_test_utils.cpp) else() message(FATAL_ERROR "OpenMPI NOT found, cannot build MG tests.") endif() diff --git a/cpp/tests/c_api/hits_test.c b/cpp/tests/c_api/hits_test.c index caa2bb292e5..7548d7148f0 100644 --- a/cpp/tests/c_api/hits_test.c +++ b/cpp/tests/c_api/hits_test.c @@ -33,7 +33,6 @@ int generic_hits_test(vertex_t* h_src, size_t num_vertices, size_t num_edges, bool_t store_transposed, - double alpha, double epsilon, size_t max_iterations) { @@ -58,9 +57,6 @@ int generic_hits_test(vertex_t* h_src, ret_code = cugraph_hits( p_handle, p_graph, epsilon, max_iterations, NULL, NULL, FALSE, FALSE, &p_result, &ret_error); - TEST_ASSERT(test_ret_value, ret_code != CUGRAPH_SUCCESS, "cugraph_hits worked, but it's not implemented!!!"); - -#if 0 TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "cugraph_hits failed."); cugraph_type_erased_device_array_view_t* vertices; @@ -102,7 +98,6 @@ int generic_hits_test(vertex_t* h_src, cugraph_sg_graph_free(p_graph); cugraph_free_resource_handle(p_handle); cugraph_error_free(ret_error); -#endif return test_ret_value; } @@ -115,10 +110,9 @@ int test_hits() 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_hubs[] = {0.0915528, 0.168382, 0.0656831, 0.191468, 0.120677, 0.362237}; - weight_t h_authorities[] = {0.0915528, 0.168382, 0.0656831, 0.191468, 0.120677, 0.362237}; + weight_t h_hubs[] = {0.347296, 0.532089, 1, 0.00003608, 0.00003608, 0}; + weight_t h_authorities[] = {0.652703, 0.879385, 0, 1, 0.347296, 0.00009136}; - double alpha = 0.95; double epsilon = 0.0001; size_t max_iterations = 20; @@ -131,7 +125,6 @@ int test_hits() num_vertices, num_edges, TRUE, - alpha, epsilon, max_iterations); } @@ -144,10 +137,9 @@ int test_hits_with_transpose() 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_hubs[] = {0.0915528, 0.168382, 0.0656831, 0.191468, 0.120677, 0.362237}; - weight_t h_authorities[] = {0.0915528, 0.168382, 0.0656831, 0.191468, 0.120677, 0.362237}; + weight_t h_hubs[] = {0.347296, 0.532089, 1, 0.00003608, 0.00003608, 0}; + weight_t h_authorities[] = {0.652703, 0.879385, 0, 1, 0.347296, 0.00009136}; - double alpha = 0.95; double epsilon = 0.0001; size_t max_iterations = 20; @@ -162,7 +154,6 @@ int test_hits_with_transpose() num_vertices, num_edges, FALSE, - alpha, epsilon, max_iterations); } diff --git a/cpp/tests/c_api/mg_hits_test.c b/cpp/tests/c_api/mg_hits_test.c new file mode 100644 index 00000000000..914a0da0594 --- /dev/null +++ b/cpp/tests/c_api/mg_hits_test.c @@ -0,0 +1,202 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "mg_test_utils.h" /* RUN_TEST */ + +#include +#include + +#include + +typedef int32_t vertex_t; +typedef int32_t edge_t; +typedef float weight_t; + +int generic_hits_test(const cugraph_resource_handle_t *p_handle, + vertex_t* h_src, + vertex_t* h_dst, + weight_t* h_wgt, + weight_t* h_result_hubs, + weight_t* h_result_authorities, + 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_graph_t* p_graph = NULL; + cugraph_hits_result_t* p_result = NULL; + + ret_code = create_mg_test_graph( + p_handle, h_src, h_dst, h_wgt, num_edges, store_transposed, &p_graph, &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "create_mg_test_graph failed."); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + ret_code = cugraph_hits( + p_handle, p_graph, epsilon, max_iterations, NULL, NULL, FALSE, FALSE, &p_result, &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "cugraph_hits failed."); + + // NOTE: Because we get back vertex ids, hubs and authorities, we can + // simply compare the returned values with the expected results + // for the entire graph. Each GPU will have a subset of the + // total vertices, so they will do a subset of the comparisons. + cugraph_type_erased_device_array_view_t* vertices; + cugraph_type_erased_device_array_view_t* hubs; + cugraph_type_erased_device_array_view_t* authorities; + + vertices = cugraph_hits_result_get_vertices(p_result); + hubs = cugraph_hits_result_get_hubs(p_result); + authorities = cugraph_hits_result_get_authorities(p_result); + double score_differences = cugraph_hits_result_get_hub_score_differences(p_result); + size_t num_iterations = cugraph_hits_result_get_number_of_iterations(p_result); + + vertex_t h_vertices[num_vertices]; + weight_t h_hubs[num_vertices]; + weight_t h_authorities[num_vertices]; + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + p_handle, (byte_t*)h_vertices, vertices, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = + cugraph_type_erased_device_array_view_copy_to_host(p_handle, (byte_t*)h_hubs, hubs, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + p_handle, (byte_t*)h_authorities, authorities, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + size_t num_local_vertices = cugraph_type_erased_device_array_view_size(vertices); + + for (int i = 0; (i < num_local_vertices) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + nearlyEqual(h_result_hubs[h_vertices[i]], h_hubs[i], 0.001), + "hubs results don't match"); + TEST_ASSERT(test_ret_value, + nearlyEqual(h_result_authorities[h_vertices[i]], h_authorities[i], 0.001), + "authorities results don't match"); + } + + cugraph_hits_result_free(p_result); + cugraph_mg_graph_free(p_graph); + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int test_hits(const cugraph_resource_handle_t* handle) +{ + size_t num_edges = 8; + 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_hubs[] = {0.347296, 0.532089, 1, 0.00003608, 0.00003608, 0}; + weight_t h_authorities[] = {0.652703, 0.879385, 0, 1, 0.347296, 0.00009136}; + + double epsilon = 0.0001; + size_t max_iterations = 20; + + // hits wants store_transposed = TRUE + return generic_hits_test(handle, + h_src, + h_dst, + h_wgt, + h_hubs, + h_authorities, + num_vertices, + num_edges, + TRUE, + epsilon, + max_iterations); +} + +int test_hits_with_transpose(const cugraph_resource_handle_t* handle) +{ + size_t num_edges = 8; + 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_hubs[] = {0.347296, 0.532089, 1, 0.00003608, 0.00003608, 0}; + weight_t h_authorities[] = {0.652703, 0.879385, 0, 1, 0.347296, 0.00009136}; + + double epsilon = 0.0001; + size_t max_iterations = 20; + + // Hits wants store_transposed = TRUE + // This call will force cugraph_hits to transpose the graph + // But we're passing src/dst backwards so the results will be the same + return generic_hits_test(handle, + h_src, + h_dst, + h_wgt, + h_hubs, + h_authorities, + num_vertices, + num_edges, + FALSE, + epsilon, + max_iterations); +} + +/******************************************************************************/ + +int main(int argc, char** argv) +{ + // Set up MPI: + int comm_rank; + int comm_size; + int num_gpus_per_node; + cudaError_t status; + int mpi_status; + int result = 0; + cugraph_resource_handle_t* handle = NULL; + cugraph_error_t* ret_error; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + int prows = 1; + + C_MPI_TRY(MPI_Init(&argc, &argv)); + C_MPI_TRY(MPI_Comm_rank(MPI_COMM_WORLD, &comm_rank)); + C_MPI_TRY(MPI_Comm_size(MPI_COMM_WORLD, &comm_size)); + C_CUDA_TRY(cudaGetDeviceCount(&num_gpus_per_node)); + C_CUDA_TRY(cudaSetDevice(comm_rank % num_gpus_per_node)); + + void* raft_handle = create_raft_handle(prows); + handle = cugraph_create_resource_handle(raft_handle); + + if (result == 0) { + result |= RUN_MG_TEST(test_hits, handle); + result |= RUN_MG_TEST(test_hits_with_transpose, handle); + + cugraph_free_resource_handle(handle); + } + + free_raft_handle(raft_handle); + + C_MPI_TRY(MPI_Finalize()); + + return result; +} diff --git a/cpp/tests/c_api/mg_pagerank_test.c b/cpp/tests/c_api/mg_pagerank_test.c index 453b493b484..941e5e09c73 100644 --- a/cpp/tests/c_api/mg_pagerank_test.c +++ b/cpp/tests/c_api/mg_pagerank_test.c @@ -14,7 +14,6 @@ * limitations under the License. */ -#include "c_test_utils.h" /* RUN_TEST */ #include "mg_test_utils.h" /* RUN_TEST */ #include @@ -49,12 +48,16 @@ int generic_pagerank_test(const cugraph_resource_handle_t* handle, ret_code = create_mg_test_graph( handle, h_src, h_dst, h_wgt, num_edges, store_transposed, &p_graph, &ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "create_test_graph failed."); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "create_mg_test_graph failed."); ret_code = cugraph_pagerank( handle, p_graph, NULL, alpha, epsilon, max_iterations, FALSE, FALSE, &p_result, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "cugraph_pagerank failed."); + // NOTE: Because we get back vertex ids and pageranks, we can simply compare + // the returned values with the expected results for the entire + // graph. Each GPU will have a subset of the total vertices, so + // they will do a subset of the comparisons. cugraph_type_erased_device_array_view_t* vertices; cugraph_type_erased_device_array_view_t* pageranks; @@ -83,7 +86,7 @@ int generic_pagerank_test(const cugraph_resource_handle_t* handle, cugraph_type_erased_device_array_view_free(pageranks); cugraph_type_erased_device_array_view_free(vertices); cugraph_pagerank_result_free(p_result); - cugraph_sg_graph_free(p_graph); + cugraph_mg_graph_free(p_graph); cugraph_error_free(ret_error); return test_ret_value; diff --git a/cpp/tests/c_api/mg_test_utils.cpp b/cpp/tests/c_api/mg_test_utils.cpp index 3b687df21de..6fa64c78b73 100644 --- a/cpp/tests/c_api/mg_test_utils.cpp +++ b/cpp/tests/c_api/mg_test_utils.cpp @@ -18,9 +18,12 @@ #include #include +#include #include +#include + extern "C" int run_mg_test(int (*test)(const cugraph_resource_handle_t*), const char* test_name, const cugraph_resource_handle_t* handle) @@ -29,6 +32,10 @@ extern "C" int run_mg_test(int (*test)(const cugraph_resource_handle_t*), time_t start_time, end_time; int rank = 0; + auto raft_handle = + reinterpret_cast(handle)->handle_; + auto &comm = raft_handle->get_comms(); + rank = cugraph_resource_handle_get_rank(handle); if (rank == 0) { @@ -40,6 +47,18 @@ extern "C" int run_mg_test(int (*test)(const cugraph_resource_handle_t*), ret_val = test(handle); + // FIXME: This is copied from host_scalar_allreduce + // which is in a file of thrust enabled code which can't + // be inclined in a cpp file. Either make this file a .cu + // or refactor host_scalar_comm.cuh to separate the thrust + // code from the non-thrust code + rmm::device_uvector d_input(1, raft_handle->get_stream()); + raft::update_device(d_input.data(), &ret_val, 1, raft_handle->get_stream()); + comm.allreduce(d_input.data(), d_input.data(), 1, raft::comms::op_t::SUM, raft_handle->get_stream()); + raft::update_host(&ret_val, d_input.data(), 1, raft_handle->get_stream()); + auto status = comm.sync_stream(raft_handle->get_stream()); + CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); + if (rank == 0) { time(&end_time); diff --git a/cpp/tests/c_api/mg_test_utils.h b/cpp/tests/c_api/mg_test_utils.h index ec6507c0a69..0425d8cf74b 100644 --- a/cpp/tests/c_api/mg_test_utils.h +++ b/cpp/tests/c_api/mg_test_utils.h @@ -15,6 +15,8 @@ */ #pragma once +#include "c_test_utils.h" + #include #include From 38d63f8a10ebe2289200a0b4ba888991a4237c66 Mon Sep 17 00:00:00 2001 From: Chuck Hastings Date: Wed, 23 Mar 2022 10:55:50 -0400 Subject: [PATCH 2/2] fix clang-format issues --- cpp/src/detail/shuffle_wrappers.cu | 102 ++++++++++++++++------------- cpp/tests/c_api/mg_test_utils.cpp | 5 +- 2 files changed, 58 insertions(+), 49 deletions(-) diff --git a/cpp/src/detail/shuffle_wrappers.cu b/cpp/src/detail/shuffle_wrappers.cu index bb69b0dd6ca..593748e22cd 100644 --- a/cpp/src/detail/shuffle_wrappers.cu +++ b/cpp/src/detail/shuffle_wrappers.cu @@ -377,53 +377,61 @@ void collect_vertex_values_to_local(raft::handle_t const& handle, } } -template void collect_vertex_values_to_local(raft::handle_t const& handle, - rmm::device_uvector& d_vertices, - rmm::device_uvector& d_values, - rmm::device_uvector& d_local_values, - int32_t local_vertex_first); - -template void collect_vertex_values_to_local(raft::handle_t const& handle, - rmm::device_uvector& d_vertices, - rmm::device_uvector& d_values, - rmm::device_uvector& d_local_values, - int64_t local_vertex_first); - -template void collect_vertex_values_to_local(raft::handle_t const& handle, - rmm::device_uvector& d_vertices, - rmm::device_uvector& d_values, - rmm::device_uvector& d_local_values, - int32_t local_vertex_first); - -template void collect_vertex_values_to_local(raft::handle_t const& handle, - rmm::device_uvector& d_vertices, - rmm::device_uvector& d_values, - rmm::device_uvector& d_local_values, - int64_t local_vertex_first); - -template void collect_vertex_values_to_local(raft::handle_t const& handle, - rmm::device_uvector& d_vertices, - rmm::device_uvector& d_values, - rmm::device_uvector& d_local_values, - int32_t local_vertex_first); - -template void collect_vertex_values_to_local(raft::handle_t const& handle, - rmm::device_uvector& d_vertices, - rmm::device_uvector& d_values, - rmm::device_uvector& d_local_values, - int64_t local_vertex_first); - -template void collect_vertex_values_to_local(raft::handle_t const& handle, - rmm::device_uvector& d_vertices, - rmm::device_uvector& d_values, - rmm::device_uvector& d_local_values, - int32_t local_vertex_first); - -template void collect_vertex_values_to_local(raft::handle_t const& handle, - rmm::device_uvector& d_vertices, - rmm::device_uvector& d_values, - rmm::device_uvector& d_local_values, - int64_t local_vertex_first); +template void collect_vertex_values_to_local( + raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int32_t local_vertex_first); + +template void collect_vertex_values_to_local( + raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int64_t local_vertex_first); + +template void collect_vertex_values_to_local( + raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int32_t local_vertex_first); + +template void collect_vertex_values_to_local( + raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int64_t local_vertex_first); + +template void collect_vertex_values_to_local( + raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int32_t local_vertex_first); + +template void collect_vertex_values_to_local( + raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int64_t local_vertex_first); + +template void collect_vertex_values_to_local( + raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int32_t local_vertex_first); + +template void collect_vertex_values_to_local( + raft::handle_t const& handle, + rmm::device_uvector& d_vertices, + rmm::device_uvector& d_values, + rmm::device_uvector& d_local_values, + int64_t local_vertex_first); } // namespace detail } // namespace cugraph diff --git a/cpp/tests/c_api/mg_test_utils.cpp b/cpp/tests/c_api/mg_test_utils.cpp index 6fa64c78b73..9539e73cf18 100644 --- a/cpp/tests/c_api/mg_test_utils.cpp +++ b/cpp/tests/c_api/mg_test_utils.cpp @@ -34,7 +34,7 @@ extern "C" int run_mg_test(int (*test)(const cugraph_resource_handle_t*), auto raft_handle = reinterpret_cast(handle)->handle_; - auto &comm = raft_handle->get_comms(); + auto& comm = raft_handle->get_comms(); rank = cugraph_resource_handle_get_rank(handle); @@ -54,7 +54,8 @@ extern "C" int run_mg_test(int (*test)(const cugraph_resource_handle_t*), // code from the non-thrust code rmm::device_uvector d_input(1, raft_handle->get_stream()); raft::update_device(d_input.data(), &ret_val, 1, raft_handle->get_stream()); - comm.allreduce(d_input.data(), d_input.data(), 1, raft::comms::op_t::SUM, raft_handle->get_stream()); + comm.allreduce( + d_input.data(), d_input.data(), 1, raft::comms::op_t::SUM, raft_handle->get_stream()); raft::update_host(&ret_val, d_input.data(), 1, raft_handle->get_stream()); auto status = comm.sync_stream(raft_handle->get_stream()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure.");