diff --git a/cpp/include/experimental/detail/graph_utils.cuh b/cpp/include/experimental/detail/graph_utils.cuh index 3ac2e2163c6..084d68b8ba4 100644 --- a/cpp/include/experimental/detail/graph_utils.cuh +++ b/cpp/include/experimental/detail/graph_utils.cuh @@ -25,6 +25,7 @@ #include #include +#include #include #include @@ -39,7 +40,7 @@ namespace detail { // compute the numbers of nonzeros in rows (of the graph adjacency matrix, if store_transposed = // false) or columns (of the graph adjacency matrix, if store_transposed = true) template -rmm::device_uvector compute_major_degree( +rmm::device_uvector compute_major_degrees( raft::handle_t const &handle, std::vector const &adj_matrix_partition_offsets, partition_t const &partition) @@ -120,7 +121,7 @@ rmm::device_uvector compute_major_degree( // compute the numbers of nonzeros in rows (of the graph adjacency matrix, if store_transposed = // false) or columns (of the graph adjacency matrix, if store_transposed = true) template -rmm::device_uvector compute_major_degree( +rmm::device_uvector compute_major_degrees( raft::handle_t const &handle, std::vector> const &adj_matrix_partition_offsets, partition_t const &partition) @@ -131,7 +132,22 @@ rmm::device_uvector compute_major_degree( adj_matrix_partition_offsets.end(), tmp_offsets.begin(), [](auto const &offsets) { return offsets.data(); }); - return compute_major_degree(handle, tmp_offsets, partition); + return compute_major_degrees(handle, tmp_offsets, partition); +} + +// compute the numbers of nonzeros in rows (of the graph adjacency matrix, if store_transposed = +// false) or columns (of the graph adjacency matrix, if store_transposed = true) +template +rmm::device_uvector compute_major_degrees(raft::handle_t const &handle, + edge_t const *offsets, + vertex_t number_of_vertices) +{ + rmm::device_uvector degrees(number_of_vertices, handle.get_stream()); + thrust::tabulate(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + degrees.begin(), + degrees.end(), + [offsets] __device__(auto i) { return offsets[i + 1] - offsets[i]; }); + return degrees; } template diff --git a/cpp/include/experimental/graph_view.hpp b/cpp/include/experimental/graph_view.hpp index d2ae1150970..7598841fc1a 100644 --- a/cpp/include/experimental/graph_view.hpp +++ b/cpp/include/experimental/graph_view.hpp @@ -494,6 +494,12 @@ class graph_view_t(nullptr); } + rmm::device_uvector compute_in_degrees(raft::handle_t const& handle) const; + rmm::device_uvector compute_out_degrees(raft::handle_t const& handle) const; + + rmm::device_uvector compute_in_weight_sums(raft::handle_t const& handle) const; + rmm::device_uvector compute_out_weight_sums(raft::handle_t const& handle) const; + private: std::vector adj_matrix_partition_offsets_{}; std::vector adj_matrix_partition_indices_{}; @@ -638,6 +644,12 @@ class graph_view_t compute_in_degrees(raft::handle_t const& handle) const; + rmm::device_uvector compute_out_degrees(raft::handle_t const& handle) const; + + rmm::device_uvector compute_in_weight_sums(raft::handle_t const& handle) const; + rmm::device_uvector compute_out_weight_sums(raft::handle_t const& handle) const; + private: edge_t const* offsets_{nullptr}; vertex_t const* indices_{nullptr}; diff --git a/cpp/include/utilities/shuffle_comm.cuh b/cpp/include/utilities/shuffle_comm.cuh index 05fe51184ca..d1341ee01c6 100644 --- a/cpp/include/utilities/shuffle_comm.cuh +++ b/cpp/include/utilities/shuffle_comm.cuh @@ -69,7 +69,7 @@ rmm::device_uvector sort_and_count(raft::comms::comms_t const &comm, d_tx_value_counts = std::move(d_counts); } - return std::move(d_tx_value_counts); + return d_tx_value_counts; } template @@ -111,7 +111,7 @@ rmm::device_uvector sort_and_count(raft::comms::comms_t const &comm, d_tx_value_counts = std::move(d_counts); } - return std::move(d_tx_value_counts); + return d_tx_value_counts; } // inline to suppress a complaint about ODR violation diff --git a/cpp/src/experimental/graph.cu b/cpp/src/experimental/graph.cu index 5cf393bfce4..498bb4eaefe 100644 --- a/cpp/src/experimental/graph.cu +++ b/cpp/src/experimental/graph.cu @@ -278,7 +278,7 @@ graph_tget_handle_ptr()), adj_matrix_partition_offsets_, partition_); // optional expensive checks (part 2/3) diff --git a/cpp/src/experimental/graph_view.cu b/cpp/src/experimental/graph_view.cu index df92fd94194..f443608e424 100644 --- a/cpp/src/experimental/graph_view.cu +++ b/cpp/src/experimental/graph_view.cu @@ -17,6 +17,7 @@ #include #include #include +#include #include #include @@ -70,6 +71,83 @@ std::vector update_adj_matrix_partition_edge_counts( return adj_matrix_partition_edge_counts; } +template +rmm::device_uvector compute_minor_degrees( + raft::handle_t const& handle, + graph_view_t const& graph_view) +{ + rmm::device_uvector minor_degrees(graph_view.get_number_of_local_vertices(), + handle.get_stream()); + if (store_transposed) { + copy_v_transform_reduce_out_nbr( + handle, + graph_view, + thrust::make_constant_iterator(0) /* dummy */, + thrust::make_constant_iterator(0) /* dummy */, + [] __device__(vertex_t src, vertex_t dst, weight_t w, auto src_val, auto dst_val) { + return edge_t{1}; + }, + edge_t{0}, + minor_degrees.data()); + } else { + copy_v_transform_reduce_in_nbr( + handle, + graph_view, + thrust::make_constant_iterator(0) /* dummy */, + thrust::make_constant_iterator(0) /* dummy */, + [] __device__(vertex_t src, vertex_t dst, weight_t w, auto src_val, auto dst_val) { + return edge_t{1}; + }, + edge_t{0}, + minor_degrees.data()); + } + + return minor_degrees; +} + +template +rmm::device_uvector compute_weight_sums( + raft::handle_t const& handle, + graph_view_t const& graph_view) +{ + rmm::device_uvector weight_sums(graph_view.get_number_of_local_vertices(), + handle.get_stream()); + if (major == store_transposed) { + copy_v_transform_reduce_in_nbr( + handle, + graph_view, + thrust::make_constant_iterator(0) /* dummy */, + thrust::make_constant_iterator(0) /* dummy */, + [] __device__(vertex_t src, vertex_t dst, weight_t w, auto src_val, auto dst_val) { + return w; + }, + weight_t{0.0}, + weight_sums.data()); + } else { + copy_v_transform_reduce_out_nbr( + handle, + graph_view, + thrust::make_constant_iterator(0) /* dummy */, + thrust::make_constant_iterator(0) /* dummy */, + [] __device__(vertex_t src, vertex_t dst, weight_t w, auto src_val, auto dst_val) { + return w; + }, + weight_t{0.0}, + weight_sums.data()); + } + + return weight_sums; +} + } // namespace template on(default_stream), degrees.begin(), @@ -301,6 +379,154 @@ graph_view_t +rmm::device_uvector +graph_view_t>:: + compute_in_degrees(raft::handle_t const& handle) const +{ + if (store_transposed) { + return detail::compute_major_degrees( + handle, this->adj_matrix_partition_offsets_, this->partition_); + } else { + return compute_minor_degrees(handle, *this); + } +} + +template +rmm::device_uvector +graph_view_t>::compute_in_degrees(raft::handle_t const& handle) const +{ + if (store_transposed) { + return detail::compute_major_degrees( + handle, this->offsets_, this->get_number_of_local_vertices()); + } else { + return compute_minor_degrees(handle, *this); + } +} + +template +rmm::device_uvector +graph_view_t>:: + compute_out_degrees(raft::handle_t const& handle) const +{ + if (store_transposed) { + return compute_minor_degrees(handle, *this); + } else { + return detail::compute_major_degrees( + handle, this->adj_matrix_partition_offsets_, this->partition_); + } +} + +template +rmm::device_uvector +graph_view_t>::compute_out_degrees(raft::handle_t const& handle) const +{ + if (store_transposed) { + return compute_minor_degrees(handle, *this); + } else { + return detail::compute_major_degrees( + handle, this->offsets_, this->get_number_of_local_vertices()); + } +} + +template +rmm::device_uvector +graph_view_t>:: + compute_in_weight_sums(raft::handle_t const& handle) const +{ + if (store_transposed) { + return compute_weight_sums(handle, *this); + } else { + return compute_weight_sums(handle, *this); + } +} + +template +rmm::device_uvector graph_view_t< + vertex_t, + edge_t, + weight_t, + store_transposed, + multi_gpu, + std::enable_if_t>::compute_in_weight_sums(raft::handle_t const& handle) const +{ + if (store_transposed) { + return compute_weight_sums(handle, *this); + } else { + return compute_weight_sums(handle, *this); + } +} + +template +rmm::device_uvector +graph_view_t>:: + compute_out_weight_sums(raft::handle_t const& handle) const +{ + if (store_transposed) { + return compute_weight_sums(handle, *this); + } else { + return compute_weight_sums(handle, *this); + } +} + +template +rmm::device_uvector graph_view_t< + vertex_t, + edge_t, + weight_t, + store_transposed, + multi_gpu, + std::enable_if_t>::compute_out_weight_sums(raft::handle_t const& handle) const +{ + if (store_transposed) { + return compute_weight_sums(handle, *this); + } else { + return compute_weight_sums(handle, *this); + } +} + // explicit instantiation template class graph_view_t; diff --git a/cpp/src/experimental/pagerank.cu b/cpp/src/experimental/pagerank.cu index 058cbfe5966..c498d2864b4 100644 --- a/cpp/src/experimental/pagerank.cu +++ b/cpp/src/experimental/pagerank.cu @@ -142,23 +142,9 @@ void pagerank(raft::handle_t const& handle, // 2. compute the sums of the out-going edge weights (if not provided) - rmm::device_uvector tmp_vertex_out_weight_sums(0, handle.get_stream()); - if (precomputed_vertex_out_weight_sums == nullptr) { - tmp_vertex_out_weight_sums.resize(pull_graph_view.get_number_of_local_vertices(), - handle.get_stream()); - // FIXME: better refactor this out (computing out-degree). - copy_v_transform_reduce_out_nbr( - handle, - pull_graph_view, - thrust::make_constant_iterator(0) /* dummy */, - thrust::make_constant_iterator(0) /* dummy */, - [alpha] __device__(vertex_t src, vertex_t dst, weight_t w, auto src_val, auto dst_val) { - return w; - }, - weight_t{0.0}, - tmp_vertex_out_weight_sums.data()); - } - + auto tmp_vertex_out_weight_sums = precomputed_vertex_out_weight_sums == nullptr + ? pull_graph_view.compute_out_weight_sums(handle) + : rmm::device_uvector(0, handle.get_stream()); auto vertex_out_weight_sums = precomputed_vertex_out_weight_sums != nullptr ? precomputed_vertex_out_weight_sums : tmp_vertex_out_weight_sums.data(); diff --git a/cpp/src/experimental/renumber_edgelist.cu b/cpp/src/experimental/renumber_edgelist.cu index 1f9a5a573fa..a87e4b627ed 100644 --- a/cpp/src/experimental/renumber_edgelist.cu +++ b/cpp/src/experimental/renumber_edgelist.cu @@ -224,7 +224,7 @@ rmm::device_uvector compute_renumber_map( labels.begin(), thrust::greater()); - return std::move(labels); + return labels; } template @@ -609,7 +609,7 @@ std::enable_if_t> renumber_edgelist( renumber_map.find( edgelist_minor_vertices, edgelist_minor_vertices + num_edgelist_edges, edgelist_minor_vertices); - return std::move(renumber_map_labels); + return renumber_map_labels; #else return rmm::device_uvector(0, handle.get_stream()); #endif diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 5425c68e896..68b277871b1 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -331,6 +331,26 @@ set(EXPERIMENTAL_GRAPH_TEST_SRCS ConfigureTest(EXPERIMENTAL_GRAPH_TEST "${EXPERIMENTAL_GRAPH_TEST_SRCS}") +################################################################################################### +# - Experimental weight-sum tests ----------------------------------------------------------------- + +set(EXPERIMENTAL_WEIGHT_SUM_TEST_SRCS + "${CMAKE_CURRENT_SOURCE_DIR}/../../thirdparty/mmio/mmio.c" + "${CMAKE_CURRENT_SOURCE_DIR}/utilities/test_utilities.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/experimental/weight_sum_test.cpp") + +ConfigureTest(EXPERIMENTAL_WEIGHT_SUM_TEST "${EXPERIMENTAL_WEIGHT_SUM_TEST_SRCS}") + +################################################################################################### +# - Experimental degree tests --------------------------------------------------------------------- + +set(EXPERIMENTAL_DEGREE_TEST_SRCS + "${CMAKE_CURRENT_SOURCE_DIR}/../../thirdparty/mmio/mmio.c" + "${CMAKE_CURRENT_SOURCE_DIR}/utilities/test_utilities.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/experimental/degree_test.cpp") + +ConfigureTest(EXPERIMENTAL_DEGREE_TEST "${EXPERIMENTAL_DEGREE_TEST_SRCS}") + ################################################################################################### # - Experimental coarsening tests ----------------------------------------------------------------- diff --git a/cpp/tests/experimental/degree_test.cpp b/cpp/tests/experimental/degree_test.cpp new file mode 100644 index 00000000000..7c7b41cdacc --- /dev/null +++ b/cpp/tests/experimental/degree_test.cpp @@ -0,0 +1,165 @@ +/* + * Copyright (c) 2020-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 governin_from_mtxg permissions and + * limitations under the License. + */ + +#include +#include + +#include +#include +#include + +#include +#include +#include +#include + +#include + +#include +#include +#include +#include +#include + +template +void degree_reference(edge_t const* offsets, + vertex_t const* indices, + edge_t* degrees, + vertex_t num_vertices, + bool major) +{ + if (major) { + std::adjacent_difference(offsets + 1, offsets + num_vertices + 1, degrees); + } else { + std::fill(degrees, degrees + num_vertices, edge_t{0}); + for (vertex_t i = 0; i < num_vertices; ++i) { + for (auto j = offsets[i]; j < offsets[i + 1]; ++j) { + auto nbr = indices[j]; + ++degrees[nbr]; + } + } + } + + return; +} + +typedef struct Degree_Usecase_t { + std::string graph_file_full_path{}; + + Degree_Usecase_t(std::string const& graph_file_path) + { + if ((graph_file_path.length() > 0) && (graph_file_path[0] != '/')) { + graph_file_full_path = cugraph::test::get_rapids_dataset_root_dir() + "/" + graph_file_path; + } else { + graph_file_full_path = graph_file_path; + } + }; +} Degree_Usecase; + +class Tests_Degree : public ::testing::TestWithParam { + public: + Tests_Degree() {} + static void SetupTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(Degree_Usecase const& configuration) + { + raft::handle_t handle{}; + + auto graph = cugraph::test:: + read_graph_from_matrix_market_file( + handle, configuration.graph_file_full_path, false); + auto graph_view = graph.view(); + + std::vector h_offsets(graph_view.get_number_of_vertices() + 1); + std::vector h_indices(graph_view.get_number_of_edges()); + raft::update_host(h_offsets.data(), + graph_view.offsets(), + graph_view.get_number_of_vertices() + 1, + handle.get_stream()); + raft::update_host(h_indices.data(), + graph_view.indices(), + graph_view.get_number_of_edges(), + handle.get_stream()); + CUDA_TRY(cudaStreamSynchronize(handle.get_stream())); + + std::vector h_reference_in_degrees(graph_view.get_number_of_vertices()); + std::vector h_reference_out_degrees(graph_view.get_number_of_vertices()); + + degree_reference(h_offsets.data(), + h_indices.data(), + h_reference_in_degrees.data(), + graph_view.get_number_of_vertices(), + store_transposed); + + degree_reference(h_offsets.data(), + h_indices.data(), + h_reference_out_degrees.data(), + graph_view.get_number_of_vertices(), + !store_transposed); + + CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + + auto d_in_degrees = graph_view.compute_in_degrees(handle); + auto d_out_degrees = graph_view.compute_out_degrees(handle); + + CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + + std::vector h_cugraph_in_degrees(graph_view.get_number_of_vertices()); + std::vector h_cugraph_out_degrees(graph_view.get_number_of_vertices()); + + raft::update_host( + h_cugraph_in_degrees.data(), d_in_degrees.data(), d_in_degrees.size(), handle.get_stream()); + raft::update_host(h_cugraph_out_degrees.data(), + d_out_degrees.data(), + d_out_degrees.size(), + handle.get_stream()); + CUDA_TRY(cudaStreamSynchronize(handle.get_stream())); + + ASSERT_TRUE(std::equal( + h_reference_in_degrees.begin(), h_reference_in_degrees.end(), h_cugraph_in_degrees.begin())) + << "In-degree values do not match with the reference values."; + ASSERT_TRUE(std::equal(h_reference_out_degrees.begin(), + h_reference_out_degrees.end(), + h_cugraph_out_degrees.begin())) + << "Out-degree values do not match with the reference values."; + } +}; + +// FIXME: add tests for type combinations + +TEST_P(Tests_Degree, CheckInt32Int32FloatTransposed) +{ + run_current_test(GetParam()); +} + +TEST_P(Tests_Degree, CheckInt32Int32FloatUntransposed) +{ + run_current_test(GetParam()); +} + +INSTANTIATE_TEST_CASE_P(simple_test, + Tests_Degree, + ::testing::Values(Degree_Usecase("test/datasets/karate.mtx"), + Degree_Usecase("test/datasets/web-Google.mtx"), + Degree_Usecase("test/datasets/ljournal-2008.mtx"), + Degree_Usecase("test/datasets/webbase-1M.mtx"))); + +CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/experimental/weight_sum_test.cpp b/cpp/tests/experimental/weight_sum_test.cpp new file mode 100644 index 00000000000..aeda7386314 --- /dev/null +++ b/cpp/tests/experimental/weight_sum_test.cpp @@ -0,0 +1,186 @@ +/* + * Copyright (c) 2020-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 governin_from_mtxg permissions and + * limitations under the License. + */ + +#include +#include + +#include +#include +#include + +#include +#include +#include +#include + +#include + +#include +#include +#include +#include +#include + +template +void weight_sum_reference(edge_t const* offsets, + vertex_t const* indices, + weight_t const* weights, + weight_t* weight_sums, + vertex_t num_vertices, + bool major) +{ + if (!major) { std::fill(weight_sums, weight_sums + num_vertices, weight_t{0.0}); } + for (vertex_t i = 0; i < num_vertices; ++i) { + if (major) { + weight_sums[i] = + std::accumulate(weights + offsets[i], weights + offsets[i + 1], weight_t{0.0}); + } else { + for (auto j = offsets[i]; j < offsets[i + 1]; ++j) { + auto nbr = indices[j]; + weight_sums[nbr] += weights[j]; + } + } + } + + return; +} + +typedef struct WeightSum_Usecase_t { + std::string graph_file_full_path{}; + + WeightSum_Usecase_t(std::string const& graph_file_path) + { + if ((graph_file_path.length() > 0) && (graph_file_path[0] != '/')) { + graph_file_full_path = cugraph::test::get_rapids_dataset_root_dir() + "/" + graph_file_path; + } else { + graph_file_full_path = graph_file_path; + } + }; +} WeightSum_Usecase; + +class Tests_WeightSum : public ::testing::TestWithParam { + public: + Tests_WeightSum() {} + static void SetupTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(WeightSum_Usecase const& configuration) + { + raft::handle_t handle{}; + + auto graph = cugraph::test:: + read_graph_from_matrix_market_file( + handle, configuration.graph_file_full_path, true); + auto graph_view = graph.view(); + + std::vector h_offsets(graph_view.get_number_of_vertices() + 1); + std::vector h_indices(graph_view.get_number_of_edges()); + std::vector h_weights(graph_view.get_number_of_edges()); + raft::update_host(h_offsets.data(), + graph_view.offsets(), + graph_view.get_number_of_vertices() + 1, + handle.get_stream()); + raft::update_host(h_indices.data(), + graph_view.indices(), + graph_view.get_number_of_edges(), + handle.get_stream()); + raft::update_host(h_weights.data(), + graph_view.weights(), + graph_view.get_number_of_edges(), + handle.get_stream()); + CUDA_TRY(cudaStreamSynchronize(handle.get_stream())); + + std::vector h_reference_in_weight_sums(graph_view.get_number_of_vertices()); + std::vector h_reference_out_weight_sums(graph_view.get_number_of_vertices()); + + weight_sum_reference(h_offsets.data(), + h_indices.data(), + h_weights.data(), + h_reference_in_weight_sums.data(), + graph_view.get_number_of_vertices(), + store_transposed); + + weight_sum_reference(h_offsets.data(), + h_indices.data(), + h_weights.data(), + h_reference_out_weight_sums.data(), + graph_view.get_number_of_vertices(), + !store_transposed); + + CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + + auto d_in_weight_sums = graph_view.compute_in_weight_sums(handle); + auto d_out_weight_sums = graph_view.compute_out_weight_sums(handle); + + CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + + std::vector h_cugraph_in_weight_sums(graph_view.get_number_of_vertices()); + std::vector h_cugraph_out_weight_sums(graph_view.get_number_of_vertices()); + + raft::update_host(h_cugraph_in_weight_sums.data(), + d_in_weight_sums.data(), + d_in_weight_sums.size(), + handle.get_stream()); + raft::update_host(h_cugraph_out_weight_sums.data(), + d_out_weight_sums.data(), + d_out_weight_sums.size(), + handle.get_stream()); + CUDA_TRY(cudaStreamSynchronize(handle.get_stream())); + + auto threshold_ratio = weight_t{1e-4}; + auto threshold_magnitude = std::numeric_limits::min(); + 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_reference_in_weight_sums.begin(), + h_reference_in_weight_sums.end(), + h_cugraph_in_weight_sums.begin(), + nearly_equal)) + << "In-weight-sum values do not match with the reference values."; + ASSERT_TRUE(std::equal(h_reference_out_weight_sums.begin(), + h_reference_out_weight_sums.end(), + h_cugraph_out_weight_sums.begin(), + nearly_equal)) + << "Out-weight-sum values do not match with the reference values."; + } +}; + +// FIXME: add tests for type combinations + +TEST_P(Tests_WeightSum, CheckInt32Int32FloatTransposed) +{ + run_current_test(GetParam()); +} + +TEST_P(Tests_WeightSum, CheckInt32Int32FloatUntransposed) +{ + run_current_test(GetParam()); +} + +INSTANTIATE_TEST_CASE_P(simple_test, + Tests_WeightSum, + ::testing::Values(WeightSum_Usecase("test/datasets/karate.mtx"), + WeightSum_Usecase("test/datasets/web-Google.mtx"), + WeightSum_Usecase("test/datasets/ljournal-2008.mtx"), + WeightSum_Usecase("test/datasets/webbase-1M.mtx"))); + +CUGRAPH_TEST_PROGRAM_MAIN()