diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index d211fe9ed5a..108cb0748a8 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -420,6 +420,7 @@ add_library(cugraph SHARED src/components/connectivity.cu src/centrality/katz_centrality.cu src/centrality/betweenness_centrality.cu + src/experimental/generate_rmat_edgelist.cu src/experimental/graph.cu src/experimental/graph_view.cu src/experimental/coarsen_graph.cu diff --git a/cpp/include/experimental/graph_generator.hpp b/cpp/include/experimental/graph_generator.hpp new file mode 100644 index 00000000000..b8495ed7581 --- /dev/null +++ b/cpp/include/experimental/graph_generator.hpp @@ -0,0 +1,84 @@ +/* + * 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 + +#include +#include + +namespace cugraph { +namespace experimental { + +/** + * @brief generate an edge list for an R-mat graph. + * + * This function allows multi-edges and self-loops similar to the Graph 500 reference + * implementation. + * + * @p scramble_vertex_ids needs to be set to `true` to generate a graph conforming to the Graph 500 + * specification (note that scrambling does not affect cuGraph's graph construction performance, so + * this is generally unnecessary). If `edge_factor` is given (e.g. Graph 500), set @p num_edges to + * (size_t{1} << @p scale) * `edge_factor`. To generate an undirected graph, set @p b == @p c and @p + * clip_and_flip = true. All the resulting edges will be placed in the lower triangular part + * (inculding the diagonal) of the graph adjacency matrix. + * + * For multi-GPU generation with `P` GPUs, @p seed should be set to different values in different + * GPUs to avoid every GPU generating the same set of edges. @p num_edges should be adjusted as + * well; e.g. assuming `edge_factor` is given, set @p num_edges = (size_t{1} << @p scale) * + * `edge_factor` / `P` + (rank < (((size_t{1} << @p scale) * `edge_factor`) % P) ? 1 : 0). + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param scale Scale factor to set the number of verties in the graph. Vertex IDs have values in + * [0, V), where V = 1 << @p scale. + * @param num_edges Number of edges to generate. + * @param a a, b, c, d (= 1.0 - (a + b + c)) in the R-mat graph generator (vist https://graph500.org + * for additional details). a, b, c, d should be non-negative and a + b + c should be no larger + * than 1.0. + * @param b a, b, c, d (= 1.0 - (a + b + c)) in the R-mat graph generator (vist https://graph500.org + * for additional details). a, b, c, d should be non-negative and a + b + c should be no larger + * than 1.0. + * @param c a, b, c, d (= 1.0 - (a + b + c)) in the R-mat graph generator (vist https://graph500.org + * for additional details). a, b, c, d should be non-negative and a + b + c should be no larger + * than 1.0. + * @param seed Seed value for the random number generator. + * @param clip_and_flip Flag controlling whether to generate edges only in the lower triangular part + * (including the diagonal) of the graph adjacency matrix (if set to `true`) or not (if set to + * `false`). + * @param scramble_vertex_ids Flag controlling whether to scramble vertex ID bits (if set to `true`) + * or not (if set to `false`); scrambling vertx ID bits breaks correlation between vertex ID values + * and vertex degrees. The scramble code here follows the algorithm in the Graph 500 reference + * implementation version 3.0.0. + * @return std::tuple, rmm::device_uvector> A tuple of + * rmm::device_uvector objects for edge source vertex IDs and edge destination vertex IDs. + */ +template +std::tuple, rmm::device_uvector> generate_rmat_edgelist( + raft::handle_t const& handle, + size_t scale, + size_t edge_factor = 16, + double a = 0.57, + double b = 0.19, + double c = 0.19, + uint64_t seed = 0, + bool clip_and_flip = false, + bool scramble_vertex_ids = false); + +} // namespace experimental +} // namespace cugraph diff --git a/cpp/src/experimental/generate_rmat_edgelist.cu b/cpp/src/experimental/generate_rmat_edgelist.cu new file mode 100644 index 00000000000..0a6d666432f --- /dev/null +++ b/cpp/src/experimental/generate_rmat_edgelist.cu @@ -0,0 +1,149 @@ +/* + * 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 + +#include +#include + +#include +#include +#include +#include + +#include +#include + +#include + +namespace cugraph { +namespace experimental { + +template +std::tuple, rmm::device_uvector> generate_rmat_edgelist( + raft::handle_t const& handle, + size_t scale, + size_t num_edges, + double a, + double b, + double c, + uint64_t seed, + bool clip_and_flip, + bool scramble_vertex_ids) +{ + CUGRAPH_EXPECTS(size_t{1} << scale <= std::numeric_limits::max(), + "Invalid input argument: scale too large for vertex_t."); + CUGRAPH_EXPECTS((a >= 0.0) && (b >= 0.0) && (c >= 0.0) && (a + b + c <= 1.0), + "Invalid input argument: a, b, c should be non-negative and a + b + c should not " + "be larger than 1.0."); + + raft::random::Rng rng(seed + 10); + // to limit memory footprint (1024 is a tuning parameter) + auto max_edges_to_generate_per_iteration = + static_cast(handle.get_device_properties().multiProcessorCount) * 1024; + rmm::device_uvector rands( + std::min(num_edges, max_edges_to_generate_per_iteration) * 2 * scale, handle.get_stream()); + + rmm::device_uvector srcs(num_edges, handle.get_stream()); + rmm::device_uvector dsts(num_edges, handle.get_stream()); + + size_t num_edges_generated{0}; + while (num_edges_generated < num_edges) { + auto num_edges_to_generate = + std::min(num_edges - num_edges_generated, max_edges_to_generate_per_iteration); + auto pair_first = thrust::make_zip_iterator(thrust::make_tuple(srcs.begin(), dsts.begin())) + + num_edges_generated; + rng.uniform( + rands.data(), num_edges_to_generate * 2 * scale, 0.0f, 1.0f, handle.get_stream()); + thrust::transform( + rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(num_edges_to_generate), + pair_first, + // if a + b == 0.0, a_norm is irrelevant, if (1.0 - (a+b)) == 0.0, c_norm is irrelevant + [scale, + clip_and_flip, + rands = rands.data(), + a_plus_b = a + b, + a_norm = (a + b) > 0.0 ? a / (a + b) : 0.0, + c_norm = (1.0 - (a + b)) > 0.0 ? c / (1.0 - (a + b)) : 0.0] __device__(auto i) { + vertex_t src{0}; + vertex_t dst{0}; + for (size_t bit = scale - 1; bit != 0; --bit) { + auto r0 = rands[i * 2 * scale + 2 * bit]; + auto r1 = rands[i * 2 * scale + 2 * bit + 1]; + auto src_bit_set = r0 > a_plus_b; + auto dst_bit_set = r1 > (src_bit_set ? c_norm : a_norm); + if (clip_and_flip) { + if (src == dst) { + if (!src_bit_set && dst_bit_set) { + src_bit_set = !src_bit_set; + dst_bit_set = !dst_bit_set; + } + } + } + src += src_bit_set ? static_cast(1 << bit) : 0; + dst += dst_bit_set ? static_cast(1 << bit) : 0; + } + return thrust::make_tuple(src, dst); + }); + num_edges_generated += num_edges_to_generate; + } + + if (scramble_vertex_ids) { + rands.resize(0, handle.get_stream()); + rands.shrink_to_fit(handle.get_stream()); + + auto pair_first = thrust::make_zip_iterator(thrust::make_tuple(srcs.begin(), dsts.begin())); + thrust::transform(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + pair_first, + pair_first + srcs.size(), + pair_first, + [scale] __device__(auto pair) { + return thrust::make_tuple(detail::scramble(thrust::get<0>(pair), scale), + detail::scramble(thrust::get<1>(pair), scale)); + }); + } + + return std::make_tuple(std::move(srcs), std::move(dsts)); +} + +// explicit instantiation + +template std::tuple, rmm::device_uvector> +generate_rmat_edgelist(raft::handle_t const& handle, + size_t scale, + size_t num_edges, + double a, + double b, + double c, + uint64_t seed, + bool clip_and_flip, + bool scramble_vertex_ids); + +template std::tuple, rmm::device_uvector> +generate_rmat_edgelist(raft::handle_t const& handle, + size_t scale, + size_t num_edges, + double a, + double b, + double c, + uint64_t seed, + bool clip_and_flip, + bool scramble_vertex_ids); + +} // namespace experimental +} // namespace cugraph diff --git a/cpp/src/experimental/scramble.cuh b/cpp/src/experimental/scramble.cuh new file mode 100644 index 00000000000..875bb5feff0 --- /dev/null +++ b/cpp/src/experimental/scramble.cuh @@ -0,0 +1,82 @@ +/* Copyright (C) 2009-2010 The Trustees of Indiana University. */ +/* */ +/* Use, modification and distribution is subject to the Boost Software */ +/* License, Version 1.0. (See accompanying file LICENSE_1_0.txt or copy at */ +/* http://www.boost.org/LICENSE_1_0.txt) */ +/* */ +/* Authors: Jeremiah Willcock */ +/* Andrew Lumsdaine */ + +/* + * 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 experimental { +namespace detail { + +template +__device__ std::enable_if_t bitreversal(uvertex_t value) +{ + return __brevll(value); +} + +template +__device__ std::enable_if_t bitreversal(uvertex_t value) +{ + return __brev(value); +} + +template +__device__ std::enable_if_t bitreversal(uvertex_t value) +{ + return static_cast(__brev(value) >> 16); +} + +/* Apply a permutation to scramble vertex numbers; a randomly generated + * permutation is not used because applying it at scale is too expensive. */ +template +__device__ vertex_t scramble(vertex_t value, size_t lgN) +{ + constexpr size_t number_of_bits = sizeof(vertex_t) * 8; + + static_assert((number_of_bits == 64) || (number_of_bits == 32) || (number_of_bits == 16)); + assert((std::is_unsigned::value && lgN <= number_of_bits) || + (!std::is_unsigned::value && lgN < number_of_bits)); + assert(value >= 0); + + using uvertex_t = typename std::make_unsigned::type; + + constexpr auto scramble_value0 = static_cast( + sizeof(vertex_t) == 8 ? 606610977102444280 : (sizeof(vertex_t) == 4 ? 282475248 : 0)); + constexpr auto scramble_value1 = static_cast( + sizeof(vertex_t) == 8 ? 11680327234415193037 : (sizeof(vertex_t) == 4 ? 2617694917 : 8620)); + + auto v = static_cast(value); + v += scramble_value0 + scramble_value1; + v *= (scramble_value0 | static_cast(0x4519840211493211)); + v = bitreversal(v) >> (number_of_bits - lgN); + v *= (scramble_value1 | static_cast(0x3050852102C843A5)); + v = bitreversal(v) >> (number_of_bits - lgN); + return static_cast(v); +} + +} // namespace detail +} // namespace experimental +} // namespace cugraph diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index a93aa0cfabb..18dfdbc8f63 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -317,6 +317,14 @@ set(MST_TEST_SRC ConfigureTest(MST_TEST "${MST_TEST_SRC}") +################################################################################################### +# - Experimental R-mat graph generation tests ----------------------------------------------------- + +set(EXPERIMENTAL_GENERATE_RMAT_TEST_SRCS + "${CMAKE_CURRENT_SOURCE_DIR}/experimental/generate_rmat_test.cpp") + +ConfigureTest(EXPERIMENTAL_GENERATE_RMAT_TEST "${EXPERIMENTAL_GENERATE_RMAT_TEST_SRCS}" "") + ################################################################################################### # - Experimental Graph tests ---------------------------------------------------------------------- @@ -329,8 +337,6 @@ 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}") @@ -339,8 +345,6 @@ 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}") diff --git a/cpp/tests/experimental/coarsen_graph_test.cpp b/cpp/tests/experimental/coarsen_graph_test.cpp index 941b33e5661..789619f2cd9 100644 --- a/cpp/tests/experimental/coarsen_graph_test.cpp +++ b/cpp/tests/experimental/coarsen_graph_test.cpp @@ -36,20 +36,6 @@ #include #include -template -std::enable_if_t::value, bool> is_valid_vertex(vertex_t num_vertices, - vertex_t v) -{ - return (v >= 0) && (v < num_vertices); -} - -template -std::enable_if_t::value, bool> is_valid_vertex(vertex_t num_vertices, - vertex_t v) -{ - return v < num_vertices; -} - template void check_coarsened_graph_results(edge_t* org_offsets, vertex_t* org_indices, @@ -68,13 +54,13 @@ void check_coarsened_graph_results(edge_t* org_offsets, ASSERT_TRUE(std::count_if(org_indices, org_indices + org_offsets[num_org_vertices], [num_org_vertices](auto nbr) { - return !is_valid_vertex(num_org_vertices, nbr); + return !cugraph::test::is_valid_vertex(num_org_vertices, nbr); }) == 0); ASSERT_TRUE(std::is_sorted(coarse_offsets, coarse_offsets + num_coarse_vertices)); ASSERT_TRUE(std::count_if(coarse_indices, coarse_indices + coarse_offsets[num_coarse_vertices], [num_coarse_vertices](auto nbr) { - return !is_valid_vertex(num_coarse_vertices, nbr); + return !cugraph::test::is_valid_vertex(num_coarse_vertices, nbr); }) == 0); ASSERT_TRUE(num_coarse_vertices <= num_org_vertices); diff --git a/cpp/tests/experimental/degree_test.cpp b/cpp/tests/experimental/degree_test.cpp index 7c7b41cdacc..581b6b29f64 100644 --- a/cpp/tests/experimental/degree_test.cpp +++ b/cpp/tests/experimental/degree_test.cpp @@ -83,9 +83,11 @@ class Tests_Degree : public ::testing::TestWithParam { { raft::handle_t handle{}; - auto graph = cugraph::test:: - read_graph_from_matrix_market_file( - handle, configuration.graph_file_full_path, false); + cugraph::experimental::graph_t graph( + handle); + std::tie(graph, std::ignore) = cugraph::test:: + read_graph_from_matrix_market_file( + handle, configuration.graph_file_full_path, false, false); auto graph_view = graph.view(); std::vector h_offsets(graph_view.get_number_of_vertices() + 1); diff --git a/cpp/tests/experimental/generate_rmat_test.cpp b/cpp/tests/experimental/generate_rmat_test.cpp new file mode 100644 index 00000000000..249a1a3c6c8 --- /dev/null +++ b/cpp/tests/experimental/generate_rmat_test.cpp @@ -0,0 +1,285 @@ +/* + * 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 + +// this function assumes that vertex IDs are not scrambled +template +void validate_rmat_distribution( + std::tuple* edges, + size_t num_edges, + vertex_t src_first, + vertex_t src_last, + vertex_t dst_first, + vertex_t dst_last, + double a, + double b, + double c, + bool clip_and_flip, + size_t min_edges /* stop recursion if # edges < min_edges */, + double error_tolerance /* (computed a|b|c - input a|b|c) shoud be smaller than error_tolerance*/) +{ + // we cannot expect the ratios of the edges in the four quadrants of the graph adjacency matrix to + // converge close to a, b, c, d if num_edges is not large enough. + if (num_edges < min_edges) { return; } + + auto src_threshold = (src_first + src_last) / 2; + auto dst_threshold = (dst_first + dst_last) / 2; + + auto a_plus_b_last = std::partition(edges, edges + num_edges, [src_threshold](auto edge) { + return std::get<0>(edge) < src_threshold; + }); + auto a_last = std::partition( + edges, a_plus_b_last, [dst_threshold](auto edge) { return std::get<1>(edge) < dst_threshold; }); + auto c_last = std::partition(a_plus_b_last, edges + num_edges, [dst_threshold](auto edge) { + return std::get<1>(edge) < dst_threshold; + }); + + ASSERT_TRUE(std::abs((double)std::distance(edges, a_last) / num_edges - a) < error_tolerance) + << "# edges=" << num_edges << " computed a=" << (double)std::distance(edges, a_last) / num_edges + << " iput a=" << a << " error tolerance=" << error_tolerance << "."; + if (clip_and_flip && (src_first == dst_first) && + (src_last == dst_last)) { // if clip_and_flip and in the diagonal + ASSERT_TRUE(std::distance(a_last, a_plus_b_last) == 0); + ASSERT_TRUE(std::abs((double)std::distance(a_plus_b_last, c_last) / num_edges - (b + c)) < + error_tolerance) + << "# edges=" << num_edges + << " computed c=" << (double)std::distance(a_plus_b_last, c_last) / num_edges + << " iput (b + c)=" << (b + c) << " error tolerance=" << error_tolerance << "."; + } else { + ASSERT_TRUE(std::abs((double)std::distance(a_last, a_plus_b_last) / num_edges - b) < + error_tolerance) + << "# edges=" << num_edges + << " computed b=" << (double)std::distance(a_last, a_plus_b_last) / num_edges + << " iput b=" << b << " error tolerance=" << error_tolerance << "."; + ASSERT_TRUE(std::abs((double)std::distance(a_plus_b_last, c_last) / num_edges - c) < + error_tolerance) + << "# edges=" << num_edges + << " computed c=" << (double)std::distance(a_plus_b_last, c_last) / num_edges + << " iput c=" << c << " error tolerance=" << error_tolerance << "."; + } + + validate_rmat_distribution(edges, + std::distance(edges, a_last), + src_first, + src_threshold, + dst_first, + dst_threshold, + a, + b, + c, + clip_and_flip, + min_edges, + error_tolerance); + validate_rmat_distribution(a_last, + std::distance(a_last, a_plus_b_last), + src_first, + (src_first + src_last) / 2, + dst_threshold, + dst_last, + a, + b, + c, + clip_and_flip, + min_edges, + error_tolerance); + validate_rmat_distribution(a_plus_b_last, + std::distance(a_plus_b_last, c_last), + src_threshold, + src_last, + dst_first, + dst_threshold, + a, + b, + c, + clip_and_flip, + min_edges, + error_tolerance); + validate_rmat_distribution(c_last, + std::distance(c_last, edges + num_edges), + src_threshold, + src_last, + dst_threshold, + dst_last, + a, + b, + c, + clip_and_flip, + min_edges, + error_tolerance); + + return; +} + +typedef struct GenerateRmat_Usecase_t { + size_t scale{0}; + size_t edge_factor{0}; + double a{0.0}; + double b{0.0}; + double c{0.0}; + bool clip_and_flip{false}; + + GenerateRmat_Usecase_t( + size_t scale, size_t edge_factor, double a, double b, double c, bool clip_and_flip) + : scale(scale), edge_factor(edge_factor), a(a), b(b), c(c), clip_and_flip(clip_and_flip){}; +} GenerateRmat_Usecase; + +class Tests_GenerateRmat : public ::testing::TestWithParam { + public: + Tests_GenerateRmat() {} + static void SetupTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(GenerateRmat_Usecase const& configuration) + { + raft::handle_t handle{}; + + auto num_vertices = static_cast(size_t{1} << configuration.scale); + std::vector no_scramble_out_degrees(num_vertices, 0); + std::vector no_scramble_in_degrees(num_vertices, 0); + std::vector scramble_out_degrees(num_vertices, 0); + std::vector scramble_in_degrees(num_vertices, 0); + for (size_t scramble = 0; scramble < 2; ++scramble) { + rmm::device_uvector d_srcs(0, handle.get_stream()); + rmm::device_uvector d_dsts(0, handle.get_stream()); + + CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + + std::tie(d_srcs, d_dsts) = cugraph::experimental::generate_rmat_edgelist( + handle, + configuration.scale, + (size_t{1} << configuration.scale) * configuration.edge_factor, + configuration.a, + configuration.b, + configuration.c, + uint64_t{0}, + configuration.clip_and_flip, + static_cast(scramble)); + + CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + + std::vector h_cugraph_srcs(d_srcs.size()); + std::vector h_cugraph_dsts(d_dsts.size()); + + raft::update_host(h_cugraph_srcs.data(), d_srcs.data(), d_srcs.size(), handle.get_stream()); + raft::update_host(h_cugraph_dsts.data(), d_dsts.data(), d_dsts.size(), handle.get_stream()); + CUDA_TRY(cudaStreamSynchronize(handle.get_stream())); + + ASSERT_TRUE( + (h_cugraph_srcs.size() == (size_t{1} << configuration.scale) * configuration.edge_factor) && + (h_cugraph_dsts.size() == (size_t{1} << configuration.scale) * configuration.edge_factor)) + << "Returned an invalid number of R-mat graph edges."; + ASSERT_TRUE( + std::count_if(h_cugraph_srcs.begin(), + h_cugraph_srcs.end(), + [num_vertices = static_cast(size_t{1} << configuration.scale)]( + auto v) { return !cugraph::test::is_valid_vertex(num_vertices, v); }) == 0) + << "Returned R-mat graph edges have invalid source vertex IDs."; + ASSERT_TRUE( + std::count_if(h_cugraph_dsts.begin(), + h_cugraph_dsts.end(), + [num_vertices = static_cast(size_t{1} << configuration.scale)]( + auto v) { return !cugraph::test::is_valid_vertex(num_vertices, v); }) == 0) + << "Returned R-mat graph edges have invalid destination vertex IDs."; + + if (!scramble) { + if (configuration.clip_and_flip) { + for (size_t i = 0; i < h_cugraph_srcs.size(); ++i) { + ASSERT_TRUE(h_cugraph_srcs[i] >= h_cugraph_dsts[i]); + } + } + + std::vector> h_cugraph_edges(h_cugraph_srcs.size()); + for (size_t i = 0; i < h_cugraph_srcs.size(); ++i) { + h_cugraph_edges[i] = std::make_tuple(h_cugraph_srcs[i], h_cugraph_dsts[i]); + } + + validate_rmat_distribution(h_cugraph_edges.data(), + h_cugraph_edges.size(), + vertex_t{0}, + num_vertices, + vertex_t{0}, + num_vertices, + configuration.a, + configuration.b, + configuration.c, + configuration.clip_and_flip, + size_t{100000}, + 0.01); + } + + if (scramble) { + std::for_each(h_cugraph_srcs.begin(), + h_cugraph_srcs.end(), + [&scramble_out_degrees](auto src) { scramble_out_degrees[src]++; }); + std::for_each(h_cugraph_dsts.begin(), + h_cugraph_dsts.end(), + [&scramble_in_degrees](auto dst) { scramble_in_degrees[dst]++; }); + std::sort(scramble_out_degrees.begin(), scramble_out_degrees.end()); + std::sort(scramble_in_degrees.begin(), scramble_in_degrees.end()); + } else { + std::for_each(h_cugraph_srcs.begin(), + h_cugraph_srcs.end(), + [&no_scramble_out_degrees](auto src) { no_scramble_out_degrees[src]++; }); + std::for_each(h_cugraph_dsts.begin(), + h_cugraph_dsts.end(), + [&no_scramble_in_degrees](auto dst) { no_scramble_in_degrees[dst]++; }); + std::sort(no_scramble_out_degrees.begin(), no_scramble_out_degrees.end()); + std::sort(no_scramble_in_degrees.begin(), no_scramble_in_degrees.end()); + } + } + + // this relies on the fact that the edge generator is deterministic. + // ideally, we should test that the two graphs are isomorphic, but this is NP hard; insted, we + // just check out-degree & in-degree distributions + ASSERT_TRUE(std::equal(no_scramble_out_degrees.begin(), + no_scramble_out_degrees.end(), + scramble_out_degrees.begin())); + ASSERT_TRUE(std::equal( + no_scramble_in_degrees.begin(), no_scramble_in_degrees.end(), scramble_in_degrees.begin())); + } +}; + +// FIXME: add tests for type combinations + +TEST_P(Tests_GenerateRmat, CheckInt32) { run_current_test(GetParam()); } + +INSTANTIATE_TEST_CASE_P(simple_test, + Tests_GenerateRmat, + ::testing::Values(GenerateRmat_Usecase(20, 16, 0.57, 0.19, 0.19, true), + GenerateRmat_Usecase(20, 16, 0.57, 0.19, 0.19, false), + GenerateRmat_Usecase(20, 16, 0.45, 0.22, 0.22, true), + GenerateRmat_Usecase(20, 16, 0.45, 0.22, 0.22, false))); + +CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/experimental/weight_sum_test.cpp b/cpp/tests/experimental/weight_sum_test.cpp index aeda7386314..9ab47b69baa 100644 --- a/cpp/tests/experimental/weight_sum_test.cpp +++ b/cpp/tests/experimental/weight_sum_test.cpp @@ -85,9 +85,11 @@ class Tests_WeightSum : public ::testing::TestWithParam { { raft::handle_t handle{}; - auto graph = cugraph::test:: - read_graph_from_matrix_market_file( - handle, configuration.graph_file_full_path, true); + cugraph::experimental::graph_t graph( + handle); + std::tie(graph, std::ignore) = cugraph::test:: + read_graph_from_matrix_market_file( + handle, configuration.graph_file_full_path, true, false); auto graph_view = graph.view(); std::vector h_offsets(graph_view.get_number_of_vertices() + 1); diff --git a/cpp/tests/utilities/test_utilities.hpp b/cpp/tests/utilities/test_utilities.hpp index 4b5517271f5..4682699df2d 100644 --- a/cpp/tests/utilities/test_utilities.hpp +++ b/cpp/tests/utilities/test_utilities.hpp @@ -130,5 +130,19 @@ read_graph_from_matrix_market_file(raft::handle_t const& handle, bool test_weighted, bool renumber); +template +std::enable_if_t::value, bool> is_valid_vertex(vertex_t num_vertices, + vertex_t v) +{ + return (v >= 0) && (v < num_vertices); +} + +template +std::enable_if_t::value, bool> is_valid_vertex(vertex_t num_vertices, + vertex_t v) +{ + return v < num_vertices; +} + } // namespace test } // namespace cugraph