From fb1c319668119ca998cf06f1e2c17ccae5784db3 Mon Sep 17 00:00:00 2001 From: Naim Date: Tue, 12 Mar 2024 02:42:38 +0100 Subject: [PATCH] Add SG test for MIS and Vertex Coloring --- cpp/src/components/vertex_coloring_mg.cu | 3 +- cpp/src/components/vertex_coloring_sg.cu | 2 +- cpp/tests/CMakeLists.txt | 12 +- ...ing_test.cu => mg_vertex_coloring_test.cu} | 11 +- cpp/tests/components/mis_test.cu | 245 +++++++++++++++++ cpp/tests/components/vertex_coloring_test.cu | 249 ++++++++++++++++++ 6 files changed, 512 insertions(+), 10 deletions(-) rename cpp/tests/components/{mg_coloring_test.cu => mg_vertex_coloring_test.cu} (96%) create mode 100644 cpp/tests/components/mis_test.cu create mode 100644 cpp/tests/components/vertex_coloring_test.cu diff --git a/cpp/src/components/vertex_coloring_mg.cu b/cpp/src/components/vertex_coloring_mg.cu index a25f2166ed4..8f87e8bd534 100644 --- a/cpp/src/components/vertex_coloring_mg.cu +++ b/cpp/src/components/vertex_coloring_mg.cu @@ -13,7 +13,8 @@ * See the License for the specific language governing permissions and * limitations under the License. */ -#include "coloring_impl.cuh" +#include "vertex_coloring_impl.cuh" + namespace cugraph { template rmm::device_uvector vertex_coloring( diff --git a/cpp/src/components/vertex_coloring_sg.cu b/cpp/src/components/vertex_coloring_sg.cu index a0de950a081..427bc0b2c81 100644 --- a/cpp/src/components/vertex_coloring_sg.cu +++ b/cpp/src/components/vertex_coloring_sg.cu @@ -13,7 +13,7 @@ * See the License for the specific language governing permissions and * limitations under the License. */ -#include "coloring_impl.cuh" +#include "vertex_coloring_impl.cuh" namespace cugraph { diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 32953e67566..c892ed7ac86 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -372,6 +372,14 @@ ConfigureTest(EDGE_BETWEENNESS_CENTRALITY_TEST centrality/edge_betweenness_centr # - WEAKLY CONNECTED COMPONENTS tests ------------------------------------------------------------- ConfigureTest(WEAKLY_CONNECTED_COMPONENTS_TEST components/weakly_connected_components_test.cpp) +############################################################################################### +# - MIS tests ------------------------------------------------------------------------------ +ConfigureTestMG(MIS_TEST components/mis_test.cu) + +############################################################################################### +# - VERTEX COLORING tests ------------------------------------------------------------------- +ConfigureTestMG(VERTEX_COLORING_TEST components/vertex_coloring_test.cu) + ################################################################################################### # - SIMILARITY tests ------------------------------------------------------------------------------ ConfigureTest(SIMILARITY_TEST link_prediction/similarity_test.cu) @@ -553,8 +561,8 @@ if(BUILD_CUGRAPH_MG_TESTS) ConfigureTestMG(MG_MIS_TEST components/mg_mis_test.cu) ############################################################################################### - # - MG Graph Coloring tests ------------------------------------------------------------------- - ConfigureTestMG(MG_COLORING_TEST components/mg_coloring_test.cu) + # - MG VERTEX COLORING tests ------------------------------------------------------------------- + ConfigureTestMG(MG_VERTEX_COLORING_TEST components/mg_vertex_coloring_test.cu) ############################################################################################### # - MG Core Number tests ---------------------------------------------------------------------- diff --git a/cpp/tests/components/mg_coloring_test.cu b/cpp/tests/components/mg_vertex_coloring_test.cu similarity index 96% rename from cpp/tests/components/mg_coloring_test.cu rename to cpp/tests/components/mg_vertex_coloring_test.cu index 753de4b09be..0b2b55f1be9 100644 --- a/cpp/tests/components/mg_coloring_test.cu +++ b/cpp/tests/components/mg_vertex_coloring_test.cu @@ -63,9 +63,6 @@ class Tests_MGGraphColoring { auto [coloring_usecase, input_usecase] = param; - auto const comm_rank = handle_->get_comms().get_rank(); - auto const comm_size = handle_->get_comms().get_size(); - HighResTimer hr_timer{}; if (cugraph::test::g_perf) { @@ -93,7 +90,7 @@ class Tests_MGGraphColoring raft::random::RngState rng_state(multi_gpu ? handle_->get_comms().get_rank() : 0); auto d_colors = - cugraph::coloring(*handle_, mg_graph_view, rng_state); + cugraph::vertex_coloring(*handle_, mg_graph_view, rng_state); // Test Graph Coloring @@ -160,7 +157,7 @@ class Tests_MGGraphColoring RAFT_CUDA_TRY(cudaDeviceSynchronize()); - weight_t nr_conflicts = cugraph::transform_reduce_e( + size_t nr_conflicts = cugraph::transform_reduce_e( *handle_, mg_graph_view, multi_gpu ? src_color_cache.view() @@ -180,6 +177,8 @@ class Tests_MGGraphColoring }, vertex_t{0}); + ASSERT_TRUE(nr_conflicts == edge_t{0}) << "adjacent vertices can't have same color." << std::endl; + { thrust::for_each( thrust::host, @@ -187,7 +186,7 @@ class Tests_MGGraphColoring h_colors.begin(), h_vertices_in_this_proces.begin(), h_color_conflicts.begin())), thrust::make_zip_iterator(thrust::make_tuple( h_colors.end(), h_vertices_in_this_proces.end(), h_color_conflicts.end())), - [comm_rank](auto color_vetex_and_conflict_flag) { + [](auto color_vetex_and_conflict_flag) { auto color = thrust::get<0>(color_vetex_and_conflict_flag); auto v = thrust::get<1>(color_vetex_and_conflict_flag); auto conflict_flag = thrust::get<2>(color_vetex_and_conflict_flag); diff --git a/cpp/tests/components/mis_test.cu b/cpp/tests/components/mis_test.cu new file mode 100644 index 00000000000..633611d38a8 --- /dev/null +++ b/cpp/tests/components/mis_test.cu @@ -0,0 +1,245 @@ +/* + * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governin_from_mtxg permissions and + * limitations under the License. + */ + +#include "prims/fill_edge_src_dst_property.cuh" +#include "prims/per_v_transform_reduce_incoming_outgoing_e.cuh" +#include "prims/property_generator.cuh" +#include "prims/reduce_op.cuh" +#include "prims/update_edge_src_dst_property.cuh" +#include "utilities/base_fixture.hpp" +#include "utilities/test_graphs.hpp" +#include "utilities/test_utilities.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include + +struct MaximalIndependentSet_Usecase { + bool check_correctness{true}; +}; + +template +class Tests_SGMaximalIndependentSet + : public ::testing::TestWithParam> { + public: + Tests_SGMaximalIndependentSet() {} + + static void SetUpTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(std::tuple const& param) + { + auto [mis_usecase, input_usecase] = param; + + raft::handle_t handle{}; + HighResTimer hr_timer{}; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + hr_timer.start("Construct graph"); + } + + constexpr bool multi_gpu = false; + + auto [sg_graph, sg_edge_weights, sg_renumber_map] = + cugraph::test::construct_graph( + handle, input_usecase, false, true); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + auto sg_graph_view = sg_graph.view(); + auto sg_edge_weight_view = + sg_edge_weights ? std::make_optional((*sg_edge_weights).view()) : std::nullopt; + + raft::random::RngState rng_state(0); + auto d_mis = cugraph::maximal_independent_set( + handle, sg_graph_view, rng_state); + + // Test MIS + if (mis_usecase.check_correctness) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + std::vector h_mis(d_mis.size()); + raft::update_host(h_mis.data(), d_mis.data(), d_mis.size(), handle.get_stream()); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + + auto vertex_first = sg_graph_view.local_vertex_partition_range_first(); + auto vertex_last = sg_graph_view.local_vertex_partition_range_last(); + + std::for_each(h_mis.begin(), h_mis.end(), [vertex_first, vertex_last](vertex_t v) { + ASSERT_TRUE((v >= vertex_first) && (v < vertex_last)); + }); + + // If a vertex is included in MIS, then none of its neighbor should be + + vertex_t local_vtx_partitoin_size = sg_graph_view.local_vertex_partition_range_size(); + rmm::device_uvector d_total_outgoing_nbrs_included_mis(local_vtx_partitoin_size, + handle.get_stream()); + + rmm::device_uvector inclusiong_flags(local_vtx_partitoin_size, handle.get_stream()); + + thrust::uninitialized_fill( + handle.get_thrust_policy(), inclusiong_flags.begin(), inclusiong_flags.end(), vertex_t{0}); + + thrust::for_each( + handle.get_thrust_policy(), + d_mis.begin(), + d_mis.end(), + [inclusiong_flags = + raft::device_span(inclusiong_flags.data(), inclusiong_flags.size()), + v_first = sg_graph_view.local_vertex_partition_range_first()] __device__(auto v) { + auto v_offset = v - v_first; + inclusiong_flags[v_offset] = vertex_t{1}; + }); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + + per_v_transform_reduce_outgoing_e( + handle, + sg_graph_view, + cugraph::detail::edge_major_property_view_t( + inclusiong_flags.data()), + cugraph::detail::edge_minor_property_view_t( + inclusiong_flags.data(), vertex_t{0}), + cugraph::edge_dummy_property_t{}.view(), + [] __device__(auto src, auto dst, auto src_included, auto dst_included, auto wt) { + return (src == dst) ? 0 : dst_included; + }, + vertex_t{0}, + cugraph::reduce_op::plus{}, + d_total_outgoing_nbrs_included_mis.begin()); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + + std::vector h_total_outgoing_nbrs_included_mis( + d_total_outgoing_nbrs_included_mis.size()); + raft::update_host(h_total_outgoing_nbrs_included_mis.data(), + d_total_outgoing_nbrs_included_mis.data(), + d_total_outgoing_nbrs_included_mis.size(), + handle.get_stream()); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + + { + auto vertex_first = sg_graph_view.local_vertex_partition_range_first(); + auto vertex_last = sg_graph_view.local_vertex_partition_range_last(); + + std::for_each(h_mis.begin(), + h_mis.end(), + [vertex_first, vertex_last, &h_total_outgoing_nbrs_included_mis](vertex_t v) { + ASSERT_TRUE((v >= vertex_first) && (v < vertex_last)) + << v << " is not within vertex parition range" << std::endl; + + ASSERT_TRUE(h_total_outgoing_nbrs_included_mis[v - vertex_first] == 0) + << v << "'s neighbor is included in MIS" << std::endl; + }); + } + } + } +}; + +using Tests_SGMaximalIndependentSet_File = + Tests_SGMaximalIndependentSet; +using Tests_SGMaximalIndependentSet_Rmat = + Tests_SGMaximalIndependentSet; + +TEST_P(Tests_SGMaximalIndependentSet_File, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGMaximalIndependentSet_File, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGMaximalIndependentSet_File, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGMaximalIndependentSet_Rmat, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGMaximalIndependentSet_Rmat, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGMaximalIndependentSet_Rmat, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +bool constexpr check_correctness = false; +INSTANTIATE_TEST_SUITE_P( + file_test, + Tests_SGMaximalIndependentSet_File, + ::testing::Combine(::testing::Values(MaximalIndependentSet_Usecase{check_correctness}, + MaximalIndependentSet_Usecase{check_correctness}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + rmat_small_test, + Tests_SGMaximalIndependentSet_Rmat, + ::testing::Combine( + ::testing::Values(MaximalIndependentSet_Usecase{check_correctness}), + ::testing::Values(cugraph::test::Rmat_Usecase(3, 4, 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_SGMaximalIndependentSet_Rmat, + ::testing::Combine( + ::testing::Values(MaximalIndependentSet_Usecase{check_correctness}, + MaximalIndependentSet_Usecase{check_correctness}), + ::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/components/vertex_coloring_test.cu b/cpp/tests/components/vertex_coloring_test.cu new file mode 100644 index 00000000000..c0bad63ac1b --- /dev/null +++ b/cpp/tests/components/vertex_coloring_test.cu @@ -0,0 +1,249 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governin_from_mtxg permissions and + * limitations under the License. + */ + +#include "prims/fill_edge_src_dst_property.cuh" +#include "prims/per_v_transform_reduce_incoming_outgoing_e.cuh" +#include "prims/property_generator.cuh" +#include "prims/reduce_op.cuh" +#include "prims/transform_reduce_e.cuh" +#include "prims/update_edge_src_dst_property.cuh" +#include "utilities/base_fixture.hpp" +#include "utilities/test_graphs.hpp" +#include "utilities/test_utilities.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include + +struct GraphColoring_UseCase { + bool check_correctness{true}; +}; + +template +class Tests_SGGraphColoring + : public ::testing::TestWithParam> { + public: + Tests_SGGraphColoring() {} + + static void SetUpTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(std::tuple const& param) + { + auto [coloring_usecase, input_usecase] = param; + + raft::handle_t handle{}; + HighResTimer hr_timer{}; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + hr_timer.start("Construct graph"); + } + + constexpr bool multi_gpu = false; + + auto [sg_graph, sg_edge_weights, sg_renumber_map] = + cugraph::test::construct_graph( + handle, input_usecase, false, true); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + auto sg_graph_view = sg_graph.view(); + auto sg_edge_weight_view = + sg_edge_weights ? std::make_optional((*sg_edge_weights).view()) : std::nullopt; + + raft::random::RngState rng_state(0); + auto d_colors = + cugraph::vertex_coloring(handle, sg_graph_view, rng_state); + + // Test Graph Coloring + + if (coloring_usecase.check_correctness) { + std::vector h_colors(d_colors.size()); + raft::update_host(h_colors.data(), d_colors.data(), d_colors.size(), handle.get_stream()); + + std::for_each(h_colors.begin(), + h_colors.end(), + [num_vertices = sg_graph_view.number_of_vertices()](vertex_t color_id) { + ASSERT_TRUE(color_id <= num_vertices); + }); + + rmm::device_uvector d_color_conflict_flags( + sg_graph_view.local_vertex_partition_range_size(), handle.get_stream()); + + per_v_transform_reduce_outgoing_e( + handle, + sg_graph_view, + cugraph::detail::edge_major_property_view_t(d_colors.data()), + cugraph::detail::edge_minor_property_view_t(d_colors.data(), + vertex_t{0}), + cugraph::edge_dummy_property_t{}.view(), + [] __device__(auto src, auto dst, auto src_color, auto dst_color, thrust::nullopt_t) { + if ((src != dst) && (src_color == dst_color)) { + return uint8_t{1}; + } else { + return uint8_t{0}; + } + }, + uint8_t{0}, + cugraph::reduce_op::maximum{}, + d_color_conflict_flags.begin()); + + std::vector h_color_conflict_flags(d_color_conflict_flags.size()); + raft::update_host(h_color_conflict_flags.data(), + d_color_conflict_flags.data(), + d_color_conflict_flags.size(), + handle.get_stream()); + + std::vector h_vertices_in_this_proces((*sg_renumber_map).size()); + + raft::update_host(h_vertices_in_this_proces.data(), + (*sg_renumber_map).data(), + (*sg_renumber_map).size(), + handle.get_stream()); + handle.sync_stream(); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + + edge_t nr_conflicts = cugraph::transform_reduce_e( + handle, + sg_graph_view, + cugraph::detail::edge_major_property_view_t(d_colors.begin()), + cugraph::detail::edge_minor_property_view_t(d_colors.begin(), + vertex_t{0}), + cugraph::edge_dummy_property_t{}.view(), + [renumber_map = (*sg_renumber_map).data()] __device__( + auto src, auto dst, auto src_color, auto dst_color, thrust::nullopt_t) { + if ((src != dst) && (src_color == dst_color)) { + return vertex_t{1}; + } else { + return vertex_t{0}; + } + }, + vertex_t{0}); + + ASSERT_TRUE(nr_conflicts == edge_t{0}) + << "adjacent vertices can't have same color." << std::endl; + + if (nr_conflicts >= 0) { + thrust::for_each( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple( + h_colors.begin(), h_vertices_in_this_proces.begin(), h_color_conflict_flags.begin())), + thrust::make_zip_iterator(thrust::make_tuple( + h_colors.end(), h_vertices_in_this_proces.end(), h_color_conflict_flags.end())), + [](auto color_vetex_and_conflict_flag) { + auto color = thrust::get<0>(color_vetex_and_conflict_flag); + auto v = thrust::get<1>(color_vetex_and_conflict_flag); + auto conflict_flag = thrust::get<2>(color_vetex_and_conflict_flag); + ASSERT_TRUE(conflict_flag == 0) + << v << " got same color as one of its neighbor" << std::endl; + }); + } + } + } +}; + +using Tests_SGGraphColoring_File = Tests_SGGraphColoring; +using Tests_SGGraphColoring_Rmat = Tests_SGGraphColoring; + +TEST_P(Tests_SGGraphColoring_File, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGGraphColoring_File, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGGraphColoring_File, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGGraphColoring_Rmat, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGGraphColoring_Rmat, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGGraphColoring_Rmat, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +bool constexpr check_correctness = false; + +INSTANTIATE_TEST_SUITE_P( + file_test, + Tests_SGGraphColoring_File, + ::testing::Combine(::testing::Values(GraphColoring_UseCase{check_correctness}, + GraphColoring_UseCase{check_correctness}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + rmat_small_test, + Tests_SGGraphColoring_Rmat, + ::testing::Combine( + ::testing::Values(GraphColoring_UseCase{check_correctness}), + ::testing::Values(cugraph::test::Rmat_Usecase(3, 4, 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_SGGraphColoring_Rmat, + ::testing::Combine( + ::testing::Values(GraphColoring_UseCase{check_correctness}, + GraphColoring_UseCase{check_correctness}), + ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); + +CUGRAPH_TEST_PROGRAM_MAIN()