From 4e20f7338711e108263274c265e40ca8dc7298fc Mon Sep 17 00:00:00 2001 From: Chuck Hastings <45364586+ChuckHastings@users.noreply.github.com> Date: Thu, 3 Jun 2021 12:04:06 -0400 Subject: [PATCH] adding test graphs - part 2 (#1603) This effort was originally targeted toward the WCC effort, but has been expanded a bit. This supersedes #1545 which I will close. The goal here is to create a means for constructing test graphs in an easier fashion. Testing the capabilities of different graph algorithms might require a variety of graphs. The objective of this PR is to better organize the graph generation components and to introduce some components to help in composing graphs out of multiple components. This PR introduces the following capabilities: * Create an ER graph * Create a collection of Complete Graphs * Create a collection of 2D mesh graphs * Create a collection of 3D mesh graphs * Create a random path graph (connect all vertices with a single randomly ordered path) * Translate vertex ids of a graph * Combine multiple edge lists into a single graph Closes #1543 Authors: - Chuck Hastings (https://github.com/ChuckHastings) Approvers: - Andrei Schaffer (https://github.com/aschaffer) - Brad Rees (https://github.com/BradReesWork) URL: https://github.com/rapidsai/cugraph/pull/1603 --- cpp/CMakeLists.txt | 5 +- .../cugraph/experimental/graph_generator.hpp | 139 ---- cpp/include/cugraph/graph_generators.hpp | 364 +++++++++ cpp/include/cugraph/utilities/cython.hpp | 6 +- cpp/src/generators/erdos_renyi_generator.cu | 122 ++++ .../generate_rmat_edgelist.cu | 43 +- cpp/src/generators/generator_tools.cu | 301 ++++++++ cpp/src/generators/simple_generators.cu | 343 +++++++++ cpp/src/utilities/cython.cu | 85 ++- cpp/tests/CMakeLists.txt | 16 + cpp/tests/community/mg_louvain_helper.cu | 42 -- cpp/tests/community/mg_louvain_helper.hpp | 5 + cpp/tests/community/mg_louvain_test.cpp | 2 +- .../mg_weakly_connected_components_test.cpp | 32 +- cpp/tests/components/wcc_test.cpp | 26 +- .../weakly_connected_components_test.cpp | 2 +- cpp/tests/experimental/bfs_test.cpp | 1 + cpp/tests/experimental/generate_rmat_test.cpp | 82 +-- .../experimental/katz_centrality_test.cpp | 1 + cpp/tests/experimental/mg_bfs_test.cpp | 31 +- .../experimental/mg_katz_centrality_test.cpp | 5 +- cpp/tests/experimental/mg_sssp_test.cpp | 31 +- cpp/tests/experimental/pagerank_test.cpp | 1 + cpp/tests/experimental/sssp_test.cpp | 1 + cpp/tests/generators/erdos_renyi_test.cpp | 98 +++ cpp/tests/generators/generators_test.cpp | 689 ++++++++++++++++++ cpp/tests/pagerank/mg_pagerank_test.cpp | 40 +- cpp/tests/utilities/rmat_utilities.cu | 19 +- cpp/tests/utilities/test_graphs.hpp | 558 ++++++++++++++ cpp/tests/utilities/test_utilities.hpp | 162 ++-- cpp/tests/utilities/thrust_wrapper.cu | 48 ++ cpp/tests/utilities/thrust_wrapper.hpp | 11 + python/cugraph/generators/rmat.pxd | 6 +- 33 files changed, 2811 insertions(+), 506 deletions(-) delete mode 100644 cpp/include/cugraph/experimental/graph_generator.hpp create mode 100644 cpp/include/cugraph/graph_generators.hpp create mode 100644 cpp/src/generators/erdos_renyi_generator.cu rename cpp/src/{experimental => generators}/generate_rmat_edgelist.cu (84%) create mode 100644 cpp/src/generators/generator_tools.cu create mode 100644 cpp/src/generators/simple_generators.cu create mode 100644 cpp/tests/generators/erdos_renyi_test.cpp create mode 100644 cpp/tests/generators/generators_test.cpp create mode 100644 cpp/tests/utilities/test_graphs.hpp diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index d1f50c42b14..8c85e18fd2f 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -185,7 +185,10 @@ 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/generators/generate_rmat_edgelist.cu + src/generators/generator_tools.cu + src/generators/simple_generators.cu + src/generators/erdos_renyi_generator.cu src/experimental/graph.cu src/experimental/graph_view.cu src/experimental/coarsen_graph.cu diff --git a/cpp/include/cugraph/experimental/graph_generator.hpp b/cpp/include/cugraph/experimental/graph_generator.hpp deleted file mode 100644 index 78a73e2fe05..00000000000 --- a/cpp/include/cugraph/experimental/graph_generator.hpp +++ /dev/null @@ -1,139 +0,0 @@ -/* - * 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 -#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 num_edges, - 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); - -enum class generator_distribution_t { POWER_LAW = 0, UNIFORM }; - -/** - * @brief generate multiple edge lists using the R-mat graph generator. - * - * 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. - * - * - * @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 n_edgelists Number of edge lists (graphs) to generate - * @param min_scale Scale factor to set the minimum number of verties in the graph. - * @param max_scale Scale factor to set the maximum number of verties in the graph. - * @param edge_factor Average number of edges per vertex to generate. - * @param size_distribution Distribution of the graph sizes, impacts the scale parameter of the - * R-MAT generator - * @param edge_distribution Edges distribution for each graph, impacts how R-MAT parameters a,b,c,d, - * are set. - * @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 A vector of std::tuple, rmm::device_uvector> of - *size @p n_edgelists, each vector element being a tuple of rmm::device_uvector objects for edge - *source vertex IDs and edge destination vertex IDs. - */ -template -std::vector, rmm::device_uvector>> -generate_rmat_edgelists( - raft::handle_t const& handle, - size_t n_edgelists, - size_t min_scale, - size_t max_scale, - size_t edge_factor = 16, - generator_distribution_t size_distribution = generator_distribution_t::POWER_LAW, - generator_distribution_t edge_distribution = generator_distribution_t::POWER_LAW, - uint64_t seed = 0, - bool clip_and_flip = false, - bool scramble_vertex_ids = false); - -} // namespace experimental -} // namespace cugraph diff --git a/cpp/include/cugraph/graph_generators.hpp b/cpp/include/cugraph/graph_generators.hpp new file mode 100644 index 00000000000..9bd002b4299 --- /dev/null +++ b/cpp/include/cugraph/graph_generators.hpp @@ -0,0 +1,364 @@ +/* + * 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 +#include + +namespace cugraph { + +/** + * @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. + * + * NOTE: The scramble_vertex_ids function needs to be called in order 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 + * (including 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`). + * @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 num_edges, + double a = 0.57, + double b = 0.19, + double c = 0.19, + uint64_t seed = 0, + bool clip_and_flip = false); + +enum class generator_distribution_t { POWER_LAW = 0, UNIFORM }; + +/** + * @brief generate multiple edge lists using the R-mat graph generator. + * + * This function allows multi-edges and self-loops similar to the Graph 500 reference + * implementation. + * + * NOTE: The scramble_vertex_ids function needs to be called in order 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 + * (including the diagonal) of the graph adjacency matrix. + * + * @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 n_edgelists Number of edge lists (graphs) to generate + * @param min_scale Scale factor to set the minimum number of verties in the graph. + * @param max_scale Scale factor to set the maximum number of verties in the graph. + * @param edge_factor Average number of edges per vertex to generate. + * @param size_distribution Distribution of the graph sizes, impacts the scale parameter of the + * R-MAT generator + * @param edge_distribution Edges distribution for each graph, impacts how R-MAT parameters a,b,c,d, + * are set. + * @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`). + * @return A vector of std::tuple, rmm::device_uvector> of + *size @p n_edgelists, each vector element being a tuple of rmm::device_uvector objects for edge + *source vertex IDs and edge destination vertex IDs. + */ +template +std::vector, rmm::device_uvector>> +generate_rmat_edgelists( + raft::handle_t const &handle, + size_t n_edgelists, + size_t min_scale, + size_t max_scale, + size_t edge_factor = 16, + generator_distribution_t size_distribution = generator_distribution_t::POWER_LAW, + generator_distribution_t edge_distribution = generator_distribution_t::POWER_LAW, + uint64_t seed = 0, + bool clip_and_flip = false); + +/** + * @brief generate an edge list for path graph + * + * A path graph of size n connects the vertices from 0 to (n - 1) + * in a single long path: ((0,1), (1,2), ..., (n - 2, n - 1) + * + * If executed in a multi-gpu context (handle comms has been initialized) + * the path will span all GPUs including an edge from the last vertex on + * GPU i to the first vertex on GPU (i+1) + * + * This function will generate a collection of path graphs. @p component_parameters_v + * defines the parameters for generating each component. Each element of + * @p component_parameters_v defines a tuple consisting of the number of vertices + * and the base vertex id for the component. + * + * @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 component_parameters_v A vector containing tuples consisting of the number of vertices and + * base vertex id for each component to generate. + * @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_path_graph_edgelist( + raft::handle_t const &handle, + std::vector> const &component_parameters_v); + +/** + * @brief generate an edge list for a 2D Mesh Graph + * + * A sequence of 2D mesh graphs will be constructed according to the + * component specifications. Each 2D mesh graph is configured with a tuple + * containing (x, y, base_vertex_id). @p component_parameters_v will contain + * a tuple for each component. + * + * If executed in a multi-gpu context (handle comms has been initialized) + * each GPU will generate disjoint 2D mesh constructs of equal size. + * + * @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 component_parameters_v Vector containing tuple defining the configuration of each + * component + * @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_2d_mesh_graph_edgelist( + raft::handle_t const &handle, + std::vector> const &component_parameters_v); + +/** + * @brief generate an edge list for a 3D Mesh Graph + * + * A sequence of 3D mesh graphs will be constructed according to the + * component specifications. Each 3D mesh graph is configured with a tuple + * containing (x, y, z, base_vertex_id). @p component_parameters_v will contain + * a tuple for each component. + * + * If executed in a multi-gpu context (handle comms has been initialized) + * each GPU will generate disjoint 3D mesh constructs of equal size. + * + * @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 component_parameters_v Vector containing tuple defining the configuration of each + * component + * @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_3d_mesh_graph_edgelist( + raft::handle_t const &handle, + std::vector> const &component_parameters_v); + +/** + * @brief generate an edge lists for some complete graphs + * + * A sequence of complete graphs will be constructed according to the + * component specifications. Each complete graph is configured with a tuple + * containing (n, base_vertex_id). @p component_parameters_v will contain + * a tuple for each component. + * + * If executed in a multi-gpu context (handle comms has been initialized) + * each GPU will generate disjoint complete graph constructs of equal size. + * + * @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 component_parameters_v Vector containing tuple defining the configuration of each + * component + * @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_complete_graph_edgelist( + raft::handle_t const &handle, + std::vector> const &component_parameters_v); + +/** + * @brief generate an edge lists for an Erdos-Renyi graph + * + * This API supports the G(n,p) model which requires O(n^2) work. + * + * If executed in a multi-gpu context (handle comms has been initialized) + * each GPU will generate Erdos-Renyi edges for its portion of the 2D + * partitioning of the adjacency matrix. + * + * @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 num_vertices Number of vertices to use in the generated graph + * @param p Probability for edge creation + * @param base_vertex_id Starting vertex id for the generated graph + * @param seed Seed value for the random number generator. + * @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_erdos_renyi_graph_edgelist_gnp(raft::handle_t const &handle, + vertex_t num_vertices, + float p, + vertex_t base_vertex_id, + uint64_t seed = 0); + +/** + * @brief generate an edge lists for an Erdos-Renyi graph + * + * This API supports the G(n,m) model + * + * If executed in a multi-gpu context (handle comms has been initialized) + * each GPU will generate Erdos-Renyi edges for its portion of the 2D + * partitioning of the adjacency matrix. + * + * @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 num_vertices Number of vertices to use in each complete graph + * @param m Number of edges to generate + * @param base_vertex_id Starting vertex id for the generated graph + * @param seed Seed value for the random number generator. + * @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_erdos_renyi_graph_edgelist_gnm(raft::handle_t const &handle, + vertex_t num_vertices, + size_t m, + vertex_t base_vertex_id, + uint64_t seed = 0); + +/** + * @brief symmetrize an edgelist + * + * Given an edgelist for a graph, symmetrize and deduplicate edges. + * + * If a duplicate edge exists in a weighted graph, one of the weights is arbitrarily + * returned. + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam weight_t Type of weights. + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param d_src_v Vector of source vertices + * @param d_dst_v Vector of destination vertices + * @param d_weights_v Optional vector of edge weights + * @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, + std::optional>> +symmetrize_edgelist(raft::handle_t const &handle, + rmm::device_uvector &&d_src_v, + rmm::device_uvector &&d_dst_v, + std::optional> &&optional_d_weights_v); + +/** + * @brief scramble vertex ids in a graph + * + * Given an edgelist for a graph, scramble all vertex ids by the given offset. + * This translation is done in place. + * + * The scramble code here follows the algorithm in the Graph 500 reference + * implementation version 3.0.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 d_src_v Vector of source vertices + * @param d_dst_v Vector of destination vertices + * @param vertex_id_offset Offset to add to each vertex id + * @param seed Used to initialize random number generator + */ +template +void scramble_vertex_ids(raft::handle_t const &handle, + rmm::device_uvector &d_src_v, + rmm::device_uvector &d_dst_v, + vertex_t vertex_id_offset, + uint64_t seed = 0); + +/** + * @brief Combine edgelists from multiple sources into a single edgelist + * + * If executed in a multi-gpu context (handle comms has been initialized) + * each GPU will operate only on its subset of data. Any shuffling to get + * edges onto the same GPU should be done prior to calling this function. + * + * @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 sources The source vertex ids to combine + * @param dests The destination vertex ids to combine + * @param weights Optional vector of weights to combine + * @param remove_multi_edges If true (the default) then remove multi edges, if false leave them in + * @return std::tuple, rmm::device_uvector, + * rmm::device_uvector> A tuple of rmm::device_uvector objects for edge source vertex IDs + * and edge destination vertex IDs and edge weights. + */ +template +std::tuple, + rmm::device_uvector, + std::optional>> +combine_edgelists(raft::handle_t const &handle, + std::vector> &&d_sources, + std::vector> &&d_dests, + std::optional>> &&optional_d_weights, + bool remove_multi_edges = true); + +} // namespace cugraph diff --git a/cpp/include/cugraph/utilities/cython.hpp b/cpp/include/cugraph/utilities/cython.hpp index f8284a16ae3..f187a985108 100644 --- a/cpp/include/cugraph/utilities/cython.hpp +++ b/cpp/include/cugraph/utilities/cython.hpp @@ -16,8 +16,8 @@ #pragma once #include -#include #include +#include #include #include @@ -523,8 +523,8 @@ call_generate_rmat_edgelists(raft::handle_t const& handle, size_t min_scale, size_t max_scale, size_t edge_factor, - cugraph::experimental::generator_distribution_t size_distribution, - cugraph::experimental::generator_distribution_t edge_distribution, + cugraph::generator_distribution_t size_distribution, + cugraph::generator_distribution_t edge_distribution, uint64_t seed, bool clip_and_flip, bool scramble_vertex_ids); diff --git a/cpp/src/generators/erdos_renyi_generator.cu b/cpp/src/generators/erdos_renyi_generator.cu new file mode 100644 index 00000000000..8452a613174 --- /dev/null +++ b/cpp/src/generators/erdos_renyi_generator.cu @@ -0,0 +1,122 @@ +/* + * 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 + +namespace cugraph { + +template +std::tuple, rmm::device_uvector> +generate_erdos_renyi_graph_edgelist_gnp(raft::handle_t const& handle, + vertex_t num_vertices, + float p, + vertex_t base_vertex_id, + uint64_t seed) +{ + CUGRAPH_EXPECTS(num_vertices < std::numeric_limits::max(), + "Implementation cannot support specified value"); + + auto random_iterator = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), [seed] __device__(size_t index) { + thrust::default_random_engine rng(seed); + thrust::uniform_real_distribution dist(0.0, 1.0); + rng.discard(index); + return dist(rng); + }); + + size_t count = thrust::count_if(rmm::exec_policy(handle.get_stream()), + random_iterator, + random_iterator + num_vertices * num_vertices, + [p] __device__(float prob) { return prob < p; }); + + rmm::device_uvector indices_v(count, handle.get_stream()); + + thrust::copy_if(rmm::exec_policy(handle.get_stream()), + random_iterator, + random_iterator + num_vertices * num_vertices, + indices_v.begin(), + [p] __device__(float prob) { return prob < p; }); + + rmm::device_uvector src_v(count, handle.get_stream()); + rmm::device_uvector dst_v(count, handle.get_stream()); + + thrust::transform(rmm::exec_policy(handle.get_stream()), + indices_v.begin(), + indices_v.end(), + thrust::make_zip_iterator(thrust::make_tuple(src_v.begin(), src_v.end())), + [num_vertices] __device__(size_t index) { + size_t src = index / num_vertices; + size_t dst = index % num_vertices; + + return thrust::make_tuple(static_cast(src), + static_cast(dst)); + }); + + handle.get_stream_view().synchronize(); + + return std::make_tuple(std::move(src_v), std::move(dst_v)); +} + +template +std::tuple, rmm::device_uvector> +generate_erdos_renyi_graph_edgelist_gnm(raft::handle_t const& handle, + vertex_t num_vertices, + size_t m, + vertex_t base_vertex_id, + uint64_t seed) +{ + CUGRAPH_FAIL("Not implemented"); +} + +template std::tuple, rmm::device_uvector> +generate_erdos_renyi_graph_edgelist_gnp(raft::handle_t const& handle, + int32_t num_vertices, + float p, + int32_t base_vertex_id, + uint64_t seed); + +template std::tuple, rmm::device_uvector> +generate_erdos_renyi_graph_edgelist_gnp(raft::handle_t const& handle, + int64_t num_vertices, + float p, + int64_t base_vertex_id, + uint64_t seed); + +template std::tuple, rmm::device_uvector> +generate_erdos_renyi_graph_edgelist_gnm(raft::handle_t const& handle, + int32_t num_vertices, + size_t m, + int32_t base_vertex_id, + uint64_t seed); + +template std::tuple, rmm::device_uvector> +generate_erdos_renyi_graph_edgelist_gnm(raft::handle_t const& handle, + int64_t num_vertices, + size_t m, + int64_t base_vertex_id, + uint64_t seed); + +} // namespace cugraph diff --git a/cpp/src/experimental/generate_rmat_edgelist.cu b/cpp/src/generators/generate_rmat_edgelist.cu similarity index 84% rename from cpp/src/experimental/generate_rmat_edgelist.cu rename to cpp/src/generators/generate_rmat_edgelist.cu index e0cccd70071..638d18b1831 100644 --- a/cpp/src/experimental/generate_rmat_edgelist.cu +++ b/cpp/src/generators/generate_rmat_edgelist.cu @@ -14,9 +14,7 @@ * limitations under the License. */ -#include - -#include +#include #include #include @@ -28,11 +26,10 @@ #include #include +#include #include -#include "rmm/detail/error.hpp" namespace cugraph { -namespace experimental { template std::tuple, rmm::device_uvector> generate_rmat_edgelist( @@ -43,8 +40,7 @@ std::tuple, rmm::device_uvector> generat double b, double c, uint64_t seed, - bool clip_and_flip, - bool scramble_vertex_ids) + bool clip_and_flip) { CUGRAPH_EXPECTS((size_t{1} << scale) <= static_cast(std::numeric_limits::max()), "Invalid input argument: scale too large for vertex_t."); @@ -105,21 +101,6 @@ std::tuple, rmm::device_uvector> generat 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)); } @@ -133,8 +114,7 @@ generate_rmat_edgelists(raft::handle_t const& handle, generator_distribution_t component_distribution, generator_distribution_t edge_distribution, uint64_t seed, - bool clip_and_flip, - bool scramble_vertex_ids) + bool clip_and_flip) { CUGRAPH_EXPECTS(min_scale > 0, "minimum graph scale is 1."); CUGRAPH_EXPECTS( @@ -171,7 +151,7 @@ generate_rmat_edgelists(raft::handle_t const& handle, for (size_t i = 0; i < n_edgelists; i++) { output.push_back(generate_rmat_edgelist( - handle, scale[i], scale[i] * edge_factor, a, b, c, i, clip_and_flip, scramble_vertex_ids)); + handle, scale[i], scale[i] * edge_factor, a, b, c, i, clip_and_flip)); } return output; } @@ -184,8 +164,7 @@ generate_rmat_edgelist(raft::handle_t const& handle, double b, double c, uint64_t seed, - bool clip_and_flip, - bool scramble_vertex_ids); + bool clip_and_flip); template std::tuple, rmm::device_uvector> generate_rmat_edgelist(raft::handle_t const& handle, @@ -195,8 +174,7 @@ generate_rmat_edgelist(raft::handle_t const& handle, double b, double c, uint64_t seed, - bool clip_and_flip, - bool scramble_vertex_ids); + bool clip_and_flip); template std::vector, rmm::device_uvector>> generate_rmat_edgelists(raft::handle_t const& handle, @@ -207,8 +185,7 @@ generate_rmat_edgelists(raft::handle_t const& handle, generator_distribution_t component_distribution, generator_distribution_t edge_distribution, uint64_t seed, - bool clip_and_flip, - bool scramble_vertex_ids); + bool clip_and_flip); template std::vector, rmm::device_uvector>> generate_rmat_edgelists(raft::handle_t const& handle, @@ -219,8 +196,6 @@ generate_rmat_edgelists(raft::handle_t const& handle, generator_distribution_t component_distribution, generator_distribution_t edge_distribution, uint64_t seed, - bool clip_and_flip, - bool scramble_vertex_ids); + bool clip_and_flip); -} // namespace experimental } // namespace cugraph diff --git a/cpp/src/generators/generator_tools.cu b/cpp/src/generators/generator_tools.cu new file mode 100644 index 00000000000..3ebef13f3b1 --- /dev/null +++ b/cpp/src/generators/generator_tools.cu @@ -0,0 +1,301 @@ +/* + * 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 + +namespace cugraph { + +namespace detail { + +template +rmm::device_uvector append_all(raft::handle_t const &handle, + std::vector> &&input) +{ + size_t size{0}; + // for (size_t i = 0; i < input.size(); ++i) size += input[i].size(); + for (auto &element : input) size += element.size(); + + rmm::device_uvector output(size, handle.get_stream()); + auto output_iter = output.begin(); + + for (auto &element : input) { + raft::copy(output_iter, element.begin(), element.size(), handle.get_stream()); + output_iter += element.size(); + } + + /* +for (size_t i = 0; i < input.size(); ++i) { + raft::copy(output_iter, input[i].begin(), input[i].size(), handle.get_stream()); + output_iter += input[i].size(); +} + */ + + return output; +} + +} // namespace detail + +template +void scramble_vertex_ids(raft::handle_t const &handle, + rmm::device_uvector &d_src_v, + rmm::device_uvector &d_dst_v, + vertex_t vertex_id_offset, + uint64_t seed) +{ + vertex_t scale = 1 + raft::log2(d_src_v.size()); + + auto pair_first = thrust::make_zip_iterator(thrust::make_tuple(d_src_v.begin(), d_dst_v.begin())); + thrust::transform(rmm::exec_policy(handle.get_stream()), + pair_first, + pair_first + d_src_v.size(), + pair_first, + [scale] __device__(auto pair) { + return thrust::make_tuple( + experimental::detail::scramble(thrust::get<0>(pair), scale), + experimental::detail::scramble(thrust::get<1>(pair), scale)); + }); +} + +template +std::tuple, + rmm::device_uvector, + std::optional>> +combine_edgelists(raft::handle_t const &handle, + std::vector> &&sources, + std::vector> &&dests, + std::optional>> &&optional_d_weights, + bool remove_multi_edges) +{ + CUGRAPH_EXPECTS(sources.size() == dests.size(), + "sources and dests vertex lists must be the same size"); + + if (optional_d_weights) { + CUGRAPH_EXPECTS(sources.size() == optional_d_weights.value().size(), + "has_weights is specified, sources and weights must be the same size"); + + thrust::for_each_n( + thrust::host, + thrust::make_zip_iterator( + thrust::make_tuple(sources.begin(), dests.begin(), optional_d_weights.value().begin())), + sources.size(), + [](auto tuple) { + CUGRAPH_EXPECTS(thrust::get<0>(tuple).size() != thrust::get<1>(tuple).size(), + "source vertex and dest vertex uvectors must be same size"); + CUGRAPH_EXPECTS(thrust::get<0>(tuple).size() != thrust::get<2>(tuple).size(), + "source vertex and weights uvectors must be same size"); + }); + } else { + thrust::for_each_n( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(sources.begin(), dests.begin())), + sources.size(), + [](auto tuple) { + CUGRAPH_EXPECTS(thrust::get<0>(tuple).size() == thrust::get<1>(tuple).size(), + "source vertex and dest vertex uvectors must be same size"); + }); + } + + std::vector> d_weights; + + rmm::device_uvector srcs_v(0, handle.get_stream()); + rmm::device_uvector dsts_v(0, handle.get_stream()); + rmm::device_uvector weights_v(0, handle.get_stream()); + + srcs_v = detail::append_all(handle, std::move(sources)); + dsts_v = detail::append_all(handle, std::move(dests)); + + if (optional_d_weights) { + weights_v = detail::append_all(handle, std::move(optional_d_weights.value())); + } + + if (remove_multi_edges) { + size_t number_of_edges{srcs_v.size()}; + + if (optional_d_weights) { + thrust::sort( + rmm::exec_policy(handle.get_stream()), + thrust::make_zip_iterator( + thrust::make_tuple(srcs_v.begin(), dsts_v.begin(), weights_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(srcs_v.end(), dsts_v.end(), weights_v.end()))); + + auto pair_first = + thrust::make_zip_iterator(thrust::make_tuple(srcs_v.begin(), dsts_v.begin())); + auto end_iter = thrust::unique_by_key(rmm::exec_policy(handle.get_stream()), + pair_first, + pair_first + srcs_v.size(), + weights_v.begin()); + + number_of_edges = thrust::distance(pair_first, thrust::get<0>(end_iter)); + } else { + thrust::sort(rmm::exec_policy(handle.get_stream()), + thrust::make_zip_iterator(thrust::make_tuple(srcs_v.begin(), dsts_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(srcs_v.end(), dsts_v.end()))); + + auto pair_first = + thrust::make_zip_iterator(thrust::make_tuple(srcs_v.begin(), dsts_v.begin())); + + auto end_iter = thrust::unique( + rmm::exec_policy(handle.get_stream()), + thrust::make_zip_iterator(thrust::make_tuple(srcs_v.begin(), dsts_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(srcs_v.end(), dsts_v.end()))); + + number_of_edges = thrust::distance(pair_first, end_iter); + } + + srcs_v.resize(number_of_edges, handle.get_stream()); + srcs_v.shrink_to_fit(handle.get_stream()); + dsts_v.resize(number_of_edges, handle.get_stream()); + dsts_v.shrink_to_fit(handle.get_stream()); + + if (optional_d_weights) { + weights_v.resize(number_of_edges, handle.get_stream()); + weights_v.shrink_to_fit(handle.get_stream()); + } + } + + return std::make_tuple( + std::move(srcs_v), + std::move(dsts_v), + optional_d_weights + ? std::move(std::optional>(std::move(weights_v))) + : std::nullopt); +} + +template +std::tuple, + rmm::device_uvector, + std::optional>> +symmetrize_edgelist(raft::handle_t const &handle, + rmm::device_uvector &&d_src_v, + rmm::device_uvector &&d_dst_v, + std::optional> &&optional_d_weights_v) +{ + auto offset = d_src_v.size(); + d_src_v.resize(offset * 2, handle.get_stream_view()); + d_dst_v.resize(offset * 2, handle.get_stream_view()); + + thrust::copy(rmm::exec_policy(handle.get_stream_view()), + d_dst_v.begin(), + d_dst_v.begin() + offset, + d_src_v.begin() + offset); + thrust::copy(rmm::exec_policy(handle.get_stream_view()), + d_src_v.begin(), + d_src_v.begin() + offset, + d_dst_v.begin() + offset); + if (optional_d_weights_v) { + optional_d_weights_v->resize(d_src_v.size(), handle.get_stream_view()); + thrust::copy(rmm::exec_policy(handle.get_stream_view()), + optional_d_weights_v->begin(), + optional_d_weights_v->begin() + offset, + optional_d_weights_v->begin() + offset); + } + + return std::make_tuple(std::move(d_src_v), + std::move(d_dst_v), + optional_d_weights_v ? std::move(optional_d_weights_v) : std::nullopt); +} + +template void scramble_vertex_ids(raft::handle_t const &handle, + rmm::device_uvector &d_src_v, + rmm::device_uvector &d_dst_v, + int32_t vertex_id_offset, + uint64_t seed); + +template void scramble_vertex_ids(raft::handle_t const &handle, + rmm::device_uvector &d_src_v, + rmm::device_uvector &d_dst_v, + int64_t vertex_id_offset, + uint64_t seed); + +template std::tuple, + rmm::device_uvector, + std::optional>> +combine_edgelists(raft::handle_t const &handle, + std::vector> &&sources, + std::vector> &&dests, + std::optional>> &&optional_d_weights, + bool remove_multi_edges); + +template std::tuple, + rmm::device_uvector, + std::optional>> +combine_edgelists(raft::handle_t const &handle, + std::vector> &&sources, + std::vector> &&dests, + std::optional>> &&optional_d_weights, + bool remove_multi_edges); + +template std::tuple, + rmm::device_uvector, + std::optional>> +combine_edgelists(raft::handle_t const &handle, + std::vector> &&sources, + std::vector> &&dests, + std::optional>> &&optional_d_weights, + bool remove_multi_edges); + +template std::tuple, + rmm::device_uvector, + std::optional>> +combine_edgelists(raft::handle_t const &handle, + std::vector> &&sources, + std::vector> &&dests, + std::optional>> &&optional_d_weights, + bool remove_multi_edges); + +template std::tuple, + rmm::device_uvector, + std::optional>> +symmetrize_edgelist(raft::handle_t const &handle, + rmm::device_uvector &&d_src_v, + rmm::device_uvector &&d_dst_v, + std::optional> &&optional_d_weights_v); +template std::tuple, + rmm::device_uvector, + std::optional>> +symmetrize_edgelist(raft::handle_t const &handle, + rmm::device_uvector &&d_src_v, + rmm::device_uvector &&d_dst_v, + std::optional> &&optional_d_weights_v); + +template std::tuple, + rmm::device_uvector, + std::optional>> +symmetrize_edgelist(raft::handle_t const &handle, + rmm::device_uvector &&d_src_v, + rmm::device_uvector &&d_dst_v, + std::optional> &&optional_d_weights_v); +template std::tuple, + rmm::device_uvector, + std::optional>> +symmetrize_edgelist(raft::handle_t const &handle, + rmm::device_uvector &&d_src_v, + rmm::device_uvector &&d_dst_v, + std::optional> &&optional_d_weights_v); + +} // namespace cugraph diff --git a/cpp/src/generators/simple_generators.cu b/cpp/src/generators/simple_generators.cu new file mode 100644 index 00000000000..413e08962e7 --- /dev/null +++ b/cpp/src/generators/simple_generators.cu @@ -0,0 +1,343 @@ +/* + * 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 + +namespace cugraph { + +template +std::tuple, rmm::device_uvector> +generate_path_graph_edgelist(raft::handle_t const& handle, + std::vector> const& component_parms_v) +{ + size_t num_edges = thrust::transform_reduce( + thrust::host, + component_parms_v.begin(), + component_parms_v.end(), + [](auto tuple) { return (std::get<0>(tuple) - 1); }, + size_t{0}, + std::plus()); + + bool edge_off_end{false}; + + if (handle.comms_initialized()) { + auto& comm = handle.get_comms(); + auto const comm_size = comm.get_size(); + auto const comm_rank = comm.get_rank(); + + if (comm_size > 1) { + if (comm_rank < comm_size) { + num_edges += component_parms_v.size(); + edge_off_end = true; + } + } + } + + rmm::device_uvector d_src_v(num_edges, handle.get_stream()); + rmm::device_uvector d_dst_v(num_edges, handle.get_stream()); + + auto src_iterator = d_src_v.begin(); + auto dst_iterator = d_dst_v.begin(); + + for (auto tuple : component_parms_v) { + vertex_t num_vertices, base_vertex_id; + std::tie(num_vertices, base_vertex_id) = tuple; + + vertex_t num_edges{num_vertices - 1}; + + if (edge_off_end) ++num_edges; + + thrust::sequence(rmm::exec_policy(handle.get_stream()), + src_iterator, + src_iterator + num_edges, + base_vertex_id); + + thrust::sequence(rmm::exec_policy(handle.get_stream()), + dst_iterator, + dst_iterator + num_edges, + base_vertex_id + 1); + + src_iterator += num_edges; + dst_iterator += num_edges; + } + + handle.get_stream_view().synchronize(); + + return std::make_tuple(std::move(d_src_v), std::move(d_dst_v)); +} + +template +std::tuple, rmm::device_uvector> +generate_2d_mesh_graph_edgelist( + raft::handle_t const& handle, + std::vector> const& component_parms_v) +{ + size_t num_edges = thrust::transform_reduce( + thrust::host, + component_parms_v.begin(), + component_parms_v.end(), + [](auto tuple) { + vertex_t x, y; + std::tie(x, y, std::ignore) = tuple; + + return ((x - 1) * y) + (x * (y - 1)); + }, + size_t{0}, + std::plus()); + + rmm::device_uvector d_src_v(num_edges, handle.get_stream()); + rmm::device_uvector d_dst_v(num_edges, handle.get_stream()); + + auto output_iterator = + thrust::make_zip_iterator(thrust::make_tuple(d_src_v.begin(), d_dst_v.begin())); + + for (auto tuple : component_parms_v) { + vertex_t x, y, base_vertex_id; + std::tie(x, y, base_vertex_id) = tuple; + + vertex_t num_vertices = x * y; + + auto x_iterator = thrust::make_zip_iterator( + thrust::make_tuple(thrust::make_counting_iterator(base_vertex_id), + thrust::make_counting_iterator(base_vertex_id + 1))); + + output_iterator = thrust::copy_if(rmm::exec_policy(handle.get_stream()), + x_iterator, + x_iterator + num_vertices - 1, + output_iterator, + [base_vertex_id, x] __device__(auto pair) { + vertex_t dst = thrust::get<1>(pair); + // Want to skip if dst is in the last column of a graph + return ((dst - base_vertex_id) % x) != 0; + }); + + auto y_iterator = thrust::make_zip_iterator( + thrust::make_tuple(thrust::make_counting_iterator(base_vertex_id), + thrust::make_counting_iterator(base_vertex_id + x))); + + output_iterator = thrust::copy_if(rmm::exec_policy(handle.get_stream()), + y_iterator, + y_iterator + num_vertices - x, + output_iterator, + [base_vertex_id, x, y] __device__(auto pair) { + vertex_t dst = thrust::get<1>(pair); + + // Want to skip if dst is in the first row of a new graph + return ((dst - base_vertex_id) % (x * y)) >= x; + }); + } + + handle.get_stream_view().synchronize(); + + return std::make_tuple(std::move(d_src_v), std::move(d_dst_v)); +} + +template +std::tuple, rmm::device_uvector> +generate_3d_mesh_graph_edgelist( + raft::handle_t const& handle, + std::vector> const& component_parms_v) +{ + size_t num_edges = thrust::transform_reduce( + thrust::host, + component_parms_v.begin(), + component_parms_v.end(), + [](auto tuple) { + vertex_t x, y, z; + std::tie(x, y, z, std::ignore) = tuple; + + return ((x - 1) * y * z) + (x * (y - 1) * z) + (x * y * (z - 1)); + }, + size_t{0}, + std::plus()); + + rmm::device_uvector d_src_v(num_edges, handle.get_stream()); + rmm::device_uvector d_dst_v(num_edges, handle.get_stream()); + + auto output_iterator = + thrust::make_zip_iterator(thrust::make_tuple(d_src_v.begin(), d_dst_v.begin())); + + for (auto tuple : component_parms_v) { + vertex_t x, y, z, base_vertex_id; + std::tie(x, y, z, base_vertex_id) = tuple; + + vertex_t num_vertices = x * y * z; + + auto x_iterator = thrust::make_zip_iterator( + thrust::make_tuple(thrust::make_counting_iterator(base_vertex_id), + thrust::make_counting_iterator(base_vertex_id + 1))); + + output_iterator = thrust::copy_if(rmm::exec_policy(handle.get_stream()), + x_iterator, + x_iterator + num_vertices - 1, + output_iterator, + [base_vertex_id, x] __device__(auto pair) { + vertex_t dst = thrust::get<1>(pair); + // Want to skip if dst is in the last column of a graph + return ((dst - base_vertex_id) % x) != 0; + }); + + auto y_iterator = thrust::make_zip_iterator( + thrust::make_tuple(thrust::make_counting_iterator(base_vertex_id), + thrust::make_counting_iterator(base_vertex_id + x))); + + output_iterator = thrust::copy_if(rmm::exec_policy(handle.get_stream()), + y_iterator, + y_iterator + num_vertices - x, + output_iterator, + [base_vertex_id, x, y] __device__(auto pair) { + vertex_t dst = thrust::get<1>(pair); + // Want to skip if dst is in the first row of a new graph + return ((dst - base_vertex_id) % (x * y)) >= x; + }); + + auto z_iterator = thrust::make_zip_iterator( + thrust::make_tuple(thrust::make_counting_iterator(base_vertex_id), + thrust::make_counting_iterator(base_vertex_id + x * y))); + + output_iterator = thrust::copy_if(rmm::exec_policy(handle.get_stream()), + z_iterator, + z_iterator + num_vertices - x * y, + output_iterator, + [base_vertex_id, x, y, z] __device__(auto pair) { + vertex_t dst = thrust::get<1>(pair); + // Want to skip if dst is in the first row of a new graph + return ((dst - base_vertex_id) % (x * y * z)) >= (x * y); + }); + } + + handle.get_stream_view().synchronize(); + + return std::make_tuple(std::move(d_src_v), std::move(d_dst_v)); +} + +template +std::tuple, rmm::device_uvector> +generate_complete_graph_edgelist( + raft::handle_t const& handle, + std::vector> const& component_parms_v) +{ + std::for_each(component_parms_v.begin(), component_parms_v.end(), [](auto tuple) { + vertex_t num_vertices = std::get<0>(tuple); + CUGRAPH_EXPECTS(num_vertices < std::numeric_limits::max(), + "Implementation cannot support specified value"); + }); + + size_t num_edges = thrust::transform_reduce( + thrust::host, + component_parms_v.begin(), + component_parms_v.end(), + [](auto tuple) { + vertex_t num_vertices = std::get<0>(tuple); + return num_vertices * (num_vertices - 1) / 2; + }, + size_t{0}, + std::plus()); + + vertex_t invalid_vertex{std::numeric_limits::max()}; + + rmm::device_uvector d_src_v(num_edges, handle.get_stream()); + rmm::device_uvector d_dst_v(num_edges, handle.get_stream()); + + auto output_iterator = + thrust::make_zip_iterator(thrust::make_tuple(d_src_v.begin(), d_dst_v.begin())); + + for (auto tuple : component_parms_v) { + vertex_t num_vertices, base_vertex_id; + std::tie(num_vertices, base_vertex_id) = tuple; + + auto transform_iter = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), + [base_vertex_id, num_vertices, invalid_vertex] __device__(size_t index) { + size_t graph_index = index / (num_vertices * num_vertices); + size_t local_index = index % (num_vertices * num_vertices); + + vertex_t src = base_vertex_id + static_cast(local_index / num_vertices); + vertex_t dst = base_vertex_id + static_cast(local_index % num_vertices); + + if (src == dst) { + src = invalid_vertex; + dst = invalid_vertex; + } else { + src += (graph_index * num_vertices); + dst += (graph_index * num_vertices); + } + + return thrust::make_tuple(src, dst); + }); + + output_iterator = thrust::copy_if(rmm::exec_policy(handle.get_stream()), + transform_iter, + transform_iter + num_vertices * num_vertices, + output_iterator, + [invalid_vertex] __device__(auto tuple) { + auto src = thrust::get<0>(tuple); + auto dst = thrust::get<1>(tuple); + + return (src != invalid_vertex) && (src < dst); + }); + } + + handle.get_stream_view().synchronize(); + + return std::make_tuple(std::move(d_src_v), std::move(d_dst_v)); +} + +template std::tuple, rmm::device_uvector> +generate_path_graph_edgelist(raft::handle_t const& handle, + std::vector> const& component_parms_v); + +template std::tuple, rmm::device_uvector> +generate_path_graph_edgelist(raft::handle_t const& handle, + std::vector> const& component_parms_v); + +template std::tuple, rmm::device_uvector> +generate_2d_mesh_graph_edgelist( + raft::handle_t const& handle, + std::vector> const& component_parms_v); + +template std::tuple, rmm::device_uvector> +generate_2d_mesh_graph_edgelist( + raft::handle_t const& handle, + std::vector> const& component_parms_v); + +template std::tuple, rmm::device_uvector> +generate_3d_mesh_graph_edgelist( + raft::handle_t const& handle, + std::vector> const& component_parms_v); + +template std::tuple, rmm::device_uvector> +generate_3d_mesh_graph_edgelist( + raft::handle_t const& handle, + std::vector> const& component_parms_v); + +template std::tuple, rmm::device_uvector> +generate_complete_graph_edgelist( + raft::handle_t const& handle, std::vector> const& component_parms_v); + +template std::tuple, rmm::device_uvector> +generate_complete_graph_edgelist( + raft::handle_t const& handle, std::vector> const& component_parms_v); + +} // namespace cugraph diff --git a/cpp/src/utilities/cython.cu b/cpp/src/utilities/cython.cu index 41b1b406063..a95e4eb5421 100644 --- a/cpp/src/utilities/cython.cu +++ b/cpp/src/utilities/cython.cu @@ -17,9 +17,9 @@ #include #include #include -#include #include #include +#include #include #include #include @@ -807,8 +807,13 @@ std::unique_ptr call_generate_rmat_edgelist(raft::handle_t co bool clip_and_flip, bool scramble_vertex_ids) { - auto src_dst_tuple = cugraph::experimental::generate_rmat_edgelist( - handle, scale, num_edges, a, b, c, seed, clip_and_flip, scramble_vertex_ids); + auto src_dst_tuple = cugraph::generate_rmat_edgelist( + handle, scale, num_edges, a, b, c, seed, clip_and_flip); + + if (scramble_vertex_ids) { + cugraph::scramble_vertex_ids( + handle, std::get<0>(src_dst_tuple), std::get<1>(src_dst_tuple), vertex_t{0}, seed); + } graph_generator_t gg_vals{ std::make_unique(std::get<0>(src_dst_tuple).release()), @@ -824,23 +829,29 @@ call_generate_rmat_edgelists(raft::handle_t const& handle, size_t min_scale, size_t max_scale, size_t edge_factor, - cugraph::experimental::generator_distribution_t size_distribution, - cugraph::experimental::generator_distribution_t edge_distribution, + cugraph::generator_distribution_t size_distribution, + cugraph::generator_distribution_t edge_distribution, uint64_t seed, bool clip_and_flip, bool scramble_vertex_ids) { - auto src_dst_vec_tuple = - cugraph::experimental::generate_rmat_edgelists(handle, - n_edgelists, - min_scale, - max_scale, - edge_factor, - size_distribution, - edge_distribution, - seed, - clip_and_flip, - scramble_vertex_ids); + auto src_dst_vec_tuple = cugraph::generate_rmat_edgelists(handle, + n_edgelists, + min_scale, + max_scale, + edge_factor, + size_distribution, + edge_distribution, + seed, + clip_and_flip); + + if (scramble_vertex_ids) { + std::for_each( + src_dst_vec_tuple.begin(), src_dst_vec_tuple.end(), [&handle, seed](auto& src_dst_tuple) { + cugraph::scramble_vertex_ids( + handle, std::get<0>(src_dst_tuple), std::get<1>(src_dst_tuple), vertex_t{0}, seed); + }); + } std::vector, std::unique_ptr>> gg_vec; @@ -1504,31 +1515,29 @@ template std::unique_ptr call_generate_rmat_edgelist template std::vector< std::pair, std::unique_ptr>> -call_generate_rmat_edgelists( - raft::handle_t const& handle, - size_t n_edgelists, - size_t min_scale, - size_t max_scale, - size_t edge_factor, - cugraph::experimental::generator_distribution_t size_distribution, - cugraph::experimental::generator_distribution_t edge_distribution, - uint64_t seed, - bool clip_and_flip, - bool scramble_vertex_ids); +call_generate_rmat_edgelists(raft::handle_t const& handle, + size_t n_edgelists, + size_t min_scale, + size_t max_scale, + size_t edge_factor, + cugraph::generator_distribution_t size_distribution, + cugraph::generator_distribution_t edge_distribution, + uint64_t seed, + bool clip_and_flip, + bool scramble_vertex_ids); template std::vector< std::pair, std::unique_ptr>> -call_generate_rmat_edgelists( - raft::handle_t const& handle, - size_t n_edgelists, - size_t min_scale, - size_t max_scale, - size_t edge_factor, - cugraph::experimental::generator_distribution_t size_distribution, - cugraph::experimental::generator_distribution_t edge_distribution, - uint64_t seed, - bool clip_and_flip, - bool scramble_vertex_ids); +call_generate_rmat_edgelists(raft::handle_t const& handle, + size_t n_edgelists, + size_t min_scale, + size_t max_scale, + size_t edge_factor, + cugraph::generator_distribution_t size_distribution, + cugraph::generator_distribution_t edge_distribution, + uint64_t seed, + bool clip_and_flip, + bool scramble_vertex_ids); } // namespace cython } // namespace cugraph diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 0608744400d..ec18640bc11 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -240,6 +240,22 @@ endif(RAPIDS_DATASET_ROOT_DIR) ### test sources ################################################################################## ################################################################################################### +################################################################################################### +# - graph generator tests ------------------------------------------------------------------------- + +set(GRAPH_GENERATORS_TEST_SRC + "${CMAKE_CURRENT_SOURCE_DIR}/generators/generators_test.cpp") + + ConfigureTest(GRAPH_GENERATORS_TEST "${GRAPH_GENERATORS_TEST_SRC}") + +################################################################################################### +# - erdos renyi graph generator tests ------------------------------------------------------------- + +set(ERDOS_RENYI_GENERATOR_TEST_SRC + "${CMAKE_CURRENT_SOURCE_DIR}/generators/erdos_renyi_test.cpp") + + ConfigureTest(ERDOS_RENYI_GENERATOR_TEST "${ERDOS_RENYI_GENERATOR_TEST_SRC}") + ################################################################################################### # - katz centrality tests ------------------------------------------------------------------------- ConfigureTest(KATZ_TEST centrality/katz_centrality_test.cu) diff --git a/cpp/tests/community/mg_louvain_helper.cu b/cpp/tests/community/mg_louvain_helper.cu index 2b1b5ade41e..935c36c9232 100644 --- a/cpp/tests/community/mg_louvain_helper.cu +++ b/cpp/tests/community/mg_louvain_helper.cu @@ -31,44 +31,6 @@ namespace cugraph { namespace test { -template -bool compare_renumbered_vectors(raft::handle_t const &handle, - rmm::device_uvector const &v1, - rmm::device_uvector const &v2) -{ - vertex_t max = 1 + thrust::reduce(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), - v1.begin(), - v1.end(), - vertex_t{0}); - - rmm::device_uvector map(max, handle.get_stream()); - - auto iter = thrust::make_zip_iterator(thrust::make_tuple(v1.begin(), v2.begin())); - - thrust::for_each(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), - iter, - iter + v1.size(), - [d_map = map.data()] __device__(auto pair) { - vertex_t e1 = thrust::get<0>(pair); - vertex_t e2 = thrust::get<1>(pair); - - d_map[e1] = e2; - }); - - auto error_count = - thrust::count_if(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), - iter, - iter + v1.size(), - [d_map = map.data()] __device__(auto pair) { - vertex_t e1 = thrust::get<0>(pair); - vertex_t e2 = thrust::get<1>(pair); - - return (d_map[e1] != e2); - }); - - return (error_count == 0); -} - template void single_gpu_renumber_edgelist_given_number_map(raft::handle_t const &handle, rmm::device_uvector &edgelist_rows_v, @@ -306,10 +268,6 @@ template void single_gpu_renumber_edgelist_given_number_map( rmm::device_uvector &d_edgelist_cols, rmm::device_uvector &d_renumber_map_gathered_v); -template bool compare_renumbered_vectors(raft::handle_t const &handle, - rmm::device_uvector const &v1, - rmm::device_uvector const &v2); - template std::unique_ptr> coarsen_graph( raft::handle_t const &handle, diff --git a/cpp/tests/community/mg_louvain_helper.hpp b/cpp/tests/community/mg_louvain_helper.hpp index 6d074e2d5e9..5ed710b7417 100644 --- a/cpp/tests/community/mg_louvain_helper.hpp +++ b/cpp/tests/community/mg_louvain_helper.hpp @@ -24,6 +24,11 @@ namespace cugraph { namespace test { +template +bool compare_renumbered_vectors(raft::handle_t const &handle, + std::vector const &v1, + std::vector const &v2); + template bool compare_renumbered_vectors(raft::handle_t const &handle, rmm::device_uvector const &v1, diff --git a/cpp/tests/community/mg_louvain_test.cpp b/cpp/tests/community/mg_louvain_test.cpp index 7e085919fd7..9c6d7bb4491 100644 --- a/cpp/tests/community/mg_louvain_test.cpp +++ b/cpp/tests/community/mg_louvain_test.cpp @@ -157,7 +157,7 @@ class Louvain_MG_Testfixture : public ::testing::TestWithParam std::tie(std::ignore, sg_modularity) = cugraph::louvain(handle, graph_view, d_clustering_v.data(), size_t{1}, resolution); - EXPECT_TRUE(cugraph::test::compare_renumbered_vectors( + EXPECT_TRUE(cugraph::test::renumbered_vectors_same( handle, d_clustering_v, d_dendrogram_gathered_v)); sg_graph = diff --git a/cpp/tests/components/mg_weakly_connected_components_test.cpp b/cpp/tests/components/mg_weakly_connected_components_test.cpp index dd61dafc682..a64919c4f92 100644 --- a/cpp/tests/components/mg_weakly_connected_components_test.cpp +++ b/cpp/tests/components/mg_weakly_connected_components_test.cpp @@ -17,7 +17,7 @@ #include #include #include -#include +#include #include #include @@ -224,20 +224,20 @@ INSTANTIATE_TEST_SUITE_P( cugraph::test::File_Usecase("test/datasets/polbooks.mtx"), cugraph::test::File_Usecase("test/datasets/netscience.mtx")))); -INSTANTIATE_TEST_SUITE_P( - rmat_small_test, - Tests_MGWeaklyConnectedComponents_Rmat, - ::testing::Values( - // enable correctness checks - std::make_tuple(WeaklyConnectedComponents_Usecase{}, - cugraph::test::Rmat_Usecase(10, 16, 0.57, 0.19, 0.19, 0, true, false, true)))); - -INSTANTIATE_TEST_SUITE_P( - rmat_large_test, - Tests_MGWeaklyConnectedComponents_Rmat, - ::testing::Values( - // disable correctness checks - std::make_tuple(WeaklyConnectedComponents_Usecase{false}, - cugraph::test::Rmat_Usecase(20, 16, 0.57, 0.19, 0.19, 0, true, false, true)))); +INSTANTIATE_TEST_SUITE_P(rmat_small_test, + Tests_MGWeaklyConnectedComponents_Rmat, + ::testing::Values( + // enable correctness checks + std::make_tuple(WeaklyConnectedComponents_Usecase{}, + cugraph::test::Rmat_Usecase( + 10, 16, 0.57, 0.19, 0.19, 0, true, false, 0, true)))); + +INSTANTIATE_TEST_SUITE_P(rmat_large_test, + Tests_MGWeaklyConnectedComponents_Rmat, + ::testing::Values( + // disable correctness checks + std::make_tuple(WeaklyConnectedComponents_Usecase{false}, + cugraph::test::Rmat_Usecase( + 20, 16, 0.57, 0.19, 0.19, 0, true, false, 0, true)))); CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/components/wcc_test.cpp b/cpp/tests/components/wcc_test.cpp index 381757bc977..9f6254d445f 100644 --- a/cpp/tests/components/wcc_test.cpp +++ b/cpp/tests/components/wcc_test.cpp @@ -9,8 +9,8 @@ * */ -#include #include +#include #include #include @@ -42,9 +42,13 @@ class Tests_WCC : public ::testing::TestWithParam graph(handle); + std::cout << "calling construct_graph" << std::endl; + std::tie(graph, std::ignore) = - input_usecase.template construct_graph( - handle, false, false); + cugraph::test::construct_graph( + handle, input_usecase, false, false); + + std::cout << "back from construct_graph" << std::endl; auto graph_view = graph.view(); @@ -59,7 +63,7 @@ class Tests_WCC : public ::testing::TestWithParam; using Tests_WCC_Rmat = Tests_WCC; -using Tests_WCC_LineGraph = Tests_WCC; +using Tests_WCC_PathGraph = Tests_WCC; TEST_P(Tests_WCC_File, WCC) { @@ -71,7 +75,7 @@ TEST_P(Tests_WCC_Rmat, WCC) auto param = GetParam(); run_current_test(std::get<0>(param), std::get<1>(param)); } -TEST_P(Tests_WCC_LineGraph, WCC) +TEST_P(Tests_WCC_PathGraph, WCC) { auto param = GetParam(); run_current_test(std::get<0>(param), std::get<1>(param)); @@ -89,9 +93,13 @@ INSTANTIATE_TEST_SUITE_P( std::make_tuple(WCC_Usecase{}, cugraph::test::File_Usecase("test/datasets/hollywood.mtx")))); INSTANTIATE_TEST_SUITE_P( - line_graph_test, - Tests_WCC_LineGraph, - ::testing::Values(std::make_tuple(WCC_Usecase{}, cugraph::test::LineGraph_Usecase(1000)), - std::make_tuple(WCC_Usecase{}, cugraph::test::LineGraph_Usecase(100000)))); + path_graph_test, + Tests_WCC_PathGraph, + ::testing::Values(std::make_tuple(WCC_Usecase{}, + cugraph::test::PathGraph_Usecase( + std::vector>({{1000, 0}}))), + std::make_tuple(WCC_Usecase{}, + cugraph::test::PathGraph_Usecase( + std::vector>({{100000, 0}}))))); CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/components/weakly_connected_components_test.cpp b/cpp/tests/components/weakly_connected_components_test.cpp index 2332aaff261..6523b6a280a 100644 --- a/cpp/tests/components/weakly_connected_components_test.cpp +++ b/cpp/tests/components/weakly_connected_components_test.cpp @@ -16,7 +16,7 @@ #include #include -#include +#include #include #include diff --git a/cpp/tests/experimental/bfs_test.cpp b/cpp/tests/experimental/bfs_test.cpp index 2c8ab894096..3fea9f371e0 100644 --- a/cpp/tests/experimental/bfs_test.cpp +++ b/cpp/tests/experimental/bfs_test.cpp @@ -16,6 +16,7 @@ #include #include +#include #include #include diff --git a/cpp/tests/experimental/generate_rmat_test.cpp b/cpp/tests/experimental/generate_rmat_test.cpp index 6d97628e83d..7c2dbb3911a 100644 --- a/cpp/tests/experimental/generate_rmat_test.cpp +++ b/cpp/tests/experimental/generate_rmat_test.cpp @@ -19,7 +19,7 @@ #include #include -#include +#include #include #include @@ -178,7 +178,7 @@ class Tests_GenerateRmat : public ::testing::TestWithParam CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement - std::tie(d_srcs, d_dsts) = cugraph::experimental::generate_rmat_edgelist( + std::tie(d_srcs, d_dsts) = cugraph::generate_rmat_edgelist( handle, configuration.scale, (size_t{1} << configuration.scale) * configuration.edge_factor, @@ -186,8 +186,8 @@ class Tests_GenerateRmat : public ::testing::TestWithParam configuration.b, configuration.c, uint64_t{0}, - configuration.clip_and_flip, - static_cast(scramble)); + configuration.clip_and_flip); + // static_cast(scramble)); CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement @@ -290,15 +290,15 @@ typedef struct GenerateRmats_Usecase_t { size_t min_scale{0}; size_t max_scale{0}; size_t edge_factor{0}; - cugraph::experimental::generator_distribution_t component_distribution; - cugraph::experimental::generator_distribution_t edge_distribution; + cugraph::generator_distribution_t component_distribution; + cugraph::generator_distribution_t edge_distribution; GenerateRmats_Usecase_t(size_t n_edgelists, size_t min_scale, size_t max_scale, size_t edge_factor, - cugraph::experimental::generator_distribution_t component_distribution, - cugraph::experimental::generator_distribution_t edge_distribution) + cugraph::generator_distribution_t component_distribution, + cugraph::generator_distribution_t edge_distribution) : n_edgelists(n_edgelists), min_scale(min_scale), max_scale(max_scale), @@ -322,15 +322,14 @@ class Tests_GenerateRmats : public ::testing::TestWithParam(handle, - configuration.n_edgelists, - configuration.min_scale, - configuration.max_scale, - configuration.edge_factor, - configuration.component_distribution, - configuration.edge_distribution, - uint64_t{0}); + auto outputs = cugraph::generate_rmat_edgelists(handle, + configuration.n_edgelists, + configuration.min_scale, + configuration.max_scale, + configuration.edge_factor, + configuration.component_distribution, + configuration.edge_distribution, + uint64_t{0}); CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement ASSERT_EQ(configuration.n_edgelists, outputs.size()); @@ -346,29 +345,28 @@ TEST_P(Tests_GenerateRmats, CheckInt32) { run_current_test(GetParam()); INSTANTIATE_TEST_SUITE_P( simple_test, Tests_GenerateRmats, - ::testing::Values( - GenerateRmats_Usecase(8, - 1, - 16, - 32, - cugraph::experimental::generator_distribution_t::UNIFORM, - cugraph::experimental::generator_distribution_t::UNIFORM), - GenerateRmats_Usecase(8, - 1, - 16, - 32, - cugraph::experimental::generator_distribution_t::UNIFORM, - cugraph::experimental::generator_distribution_t::POWER_LAW), - GenerateRmats_Usecase(8, - 3, - 16, - 32, - cugraph::experimental::generator_distribution_t::POWER_LAW, - cugraph::experimental::generator_distribution_t::UNIFORM), - GenerateRmats_Usecase(8, - 3, - 16, - 32, - cugraph::experimental::generator_distribution_t::POWER_LAW, - cugraph::experimental::generator_distribution_t::POWER_LAW))); + ::testing::Values(GenerateRmats_Usecase(8, + 1, + 16, + 32, + cugraph::generator_distribution_t::UNIFORM, + cugraph::generator_distribution_t::UNIFORM), + GenerateRmats_Usecase(8, + 1, + 16, + 32, + cugraph::generator_distribution_t::UNIFORM, + cugraph::generator_distribution_t::POWER_LAW), + GenerateRmats_Usecase(8, + 3, + 16, + 32, + cugraph::generator_distribution_t::POWER_LAW, + cugraph::generator_distribution_t::UNIFORM), + GenerateRmats_Usecase(8, + 3, + 16, + 32, + cugraph::generator_distribution_t::POWER_LAW, + cugraph::generator_distribution_t::POWER_LAW))); CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/experimental/katz_centrality_test.cpp b/cpp/tests/experimental/katz_centrality_test.cpp index d0fc558c89f..aa66e69d4f7 100644 --- a/cpp/tests/experimental/katz_centrality_test.cpp +++ b/cpp/tests/experimental/katz_centrality_test.cpp @@ -16,6 +16,7 @@ #include #include +#include #include #include diff --git a/cpp/tests/experimental/mg_bfs_test.cpp b/cpp/tests/experimental/mg_bfs_test.cpp index a832e0f99ac..04eb1bf7b43 100644 --- a/cpp/tests/experimental/mg_bfs_test.cpp +++ b/cpp/tests/experimental/mg_bfs_test.cpp @@ -17,6 +17,7 @@ #include #include #include +#include #include #include @@ -292,20 +293,20 @@ INSTANTIATE_TEST_SUITE_P( cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); -INSTANTIATE_TEST_SUITE_P( - rmat_small_test, - Tests_MGBFS_Rmat, - ::testing::Values( - // enable correctness checks - std::make_tuple(BFS_Usecase{0}, - cugraph::test::Rmat_Usecase(10, 16, 0.57, 0.19, 0.19, 0, false, false, true)))); - -INSTANTIATE_TEST_SUITE_P( - rmat_large_test, - Tests_MGBFS_Rmat, - ::testing::Values( - // disable correctness checks for large graphs - std::make_tuple(BFS_Usecase{0, false}, - cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false, true)))); +INSTANTIATE_TEST_SUITE_P(rmat_small_test, + Tests_MGBFS_Rmat, + ::testing::Values( + // enable correctness checks + std::make_tuple(BFS_Usecase{0}, + cugraph::test::Rmat_Usecase( + 10, 16, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); + +INSTANTIATE_TEST_SUITE_P(rmat_large_test, + Tests_MGBFS_Rmat, + ::testing::Values( + // disable correctness checks for large graphs + std::make_tuple(BFS_Usecase{0, false}, + cugraph::test::Rmat_Usecase( + 20, 32, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/experimental/mg_katz_centrality_test.cpp b/cpp/tests/experimental/mg_katz_centrality_test.cpp index d67cd9090b8..27ef64d124e 100644 --- a/cpp/tests/experimental/mg_katz_centrality_test.cpp +++ b/cpp/tests/experimental/mg_katz_centrality_test.cpp @@ -17,6 +17,7 @@ #include #include #include +#include #include #include @@ -250,7 +251,7 @@ INSTANTIATE_TEST_SUITE_P(rmat_small_test, ::testing::Values(KatzCentrality_Usecase{false}, KatzCentrality_Usecase{true}), ::testing::Values(cugraph::test::Rmat_Usecase( - 10, 16, 0.57, 0.19, 0.19, 0, false, false, true)))); + 10, 16, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); INSTANTIATE_TEST_SUITE_P(rmat_large_test, Tests_MGKatzCentrality_Rmat, @@ -259,6 +260,6 @@ INSTANTIATE_TEST_SUITE_P(rmat_large_test, ::testing::Values(KatzCentrality_Usecase{false, false}, KatzCentrality_Usecase{true, false}), ::testing::Values(cugraph::test::Rmat_Usecase( - 20, 32, 0.57, 0.19, 0.19, 0, false, false, true)))); + 20, 32, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/experimental/mg_sssp_test.cpp b/cpp/tests/experimental/mg_sssp_test.cpp index 8568545cbd6..da5120163df 100644 --- a/cpp/tests/experimental/mg_sssp_test.cpp +++ b/cpp/tests/experimental/mg_sssp_test.cpp @@ -17,6 +17,7 @@ #include #include #include +#include #include #include @@ -298,20 +299,20 @@ INSTANTIATE_TEST_SUITE_P( std::make_tuple(SSSP_Usecase{1000}, cugraph::test::File_Usecase("test/datasets/wiki2003.mtx")))); -INSTANTIATE_TEST_SUITE_P( - rmat_small_test, - Tests_MGSSSP_Rmat, - ::testing::Values( - // enable correctness checks - std::make_tuple(SSSP_Usecase{0}, - cugraph::test::Rmat_Usecase(10, 16, 0.57, 0.19, 0.19, 0, false, false, true)))); - -INSTANTIATE_TEST_SUITE_P( - rmat_large_test, - Tests_MGSSSP_Rmat, - ::testing::Values( - // disable correctness checks for large graphs - std::make_tuple(SSSP_Usecase{0, false}, - cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false, true)))); +INSTANTIATE_TEST_SUITE_P(rmat_small_test, + Tests_MGSSSP_Rmat, + ::testing::Values( + // enable correctness checks + std::make_tuple(SSSP_Usecase{0}, + cugraph::test::Rmat_Usecase( + 10, 16, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); + +INSTANTIATE_TEST_SUITE_P(rmat_large_test, + Tests_MGSSSP_Rmat, + ::testing::Values( + // disable correctness checks for large graphs + std::make_tuple(SSSP_Usecase{0, false}, + cugraph::test::Rmat_Usecase( + 20, 32, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/experimental/pagerank_test.cpp b/cpp/tests/experimental/pagerank_test.cpp index 9b07059d2da..5c0b0f288d4 100644 --- a/cpp/tests/experimental/pagerank_test.cpp +++ b/cpp/tests/experimental/pagerank_test.cpp @@ -16,6 +16,7 @@ #include #include +#include #include #include diff --git a/cpp/tests/experimental/sssp_test.cpp b/cpp/tests/experimental/sssp_test.cpp index e12df163551..9a50553a114 100644 --- a/cpp/tests/experimental/sssp_test.cpp +++ b/cpp/tests/experimental/sssp_test.cpp @@ -16,6 +16,7 @@ #include #include +#include #include #include diff --git a/cpp/tests/generators/erdos_renyi_test.cpp b/cpp/tests/generators/erdos_renyi_test.cpp new file mode 100644 index 00000000000..c91a9af7c41 --- /dev/null +++ b/cpp/tests/generators/erdos_renyi_test.cpp @@ -0,0 +1,98 @@ +/* + * 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 governin_from_mtxg permissions and + * limitations under the License. + */ + +#include +#include + +#include +#include + +#include + +#include + +struct GenerateErdosRenyiTest : public ::testing::Test { +}; + +template +void test_symmetric(std::vector &h_src_v, std::vector &h_dst_v) +{ + std::vector reverse_src_v(h_src_v.size()); + std::vector reverse_dst_v(h_dst_v.size()); + + std::copy(h_src_v.begin(), h_src_v.end(), reverse_dst_v.begin()); + std::copy(h_dst_v.begin(), h_dst_v.end(), reverse_src_v.begin()); + + thrust::sort(thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(h_src_v.begin(), h_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(h_src_v.end(), h_dst_v.end()))); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(reverse_src_v.begin(), reverse_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(reverse_src_v.end(), reverse_dst_v.end()))); + + EXPECT_EQ(reverse_src_v, h_src_v); + EXPECT_EQ(reverse_dst_v, h_dst_v); +} + +template +void er_test(size_t num_vertices, float p) +{ + raft::handle_t handle; + rmm::device_uvector d_src_v(0, handle.get_stream()); + rmm::device_uvector d_dst_v(0, handle.get_stream()); + + std::tie(d_src_v, d_dst_v) = + cugraph::generate_erdos_renyi_graph_edgelist_gnp(handle, num_vertices, p, 0); + + handle.get_stream_view().synchronize(); + + std::vector h_src_v(d_src_v.size()); + std::vector h_dst_v(d_dst_v.size()); + + raft::update_host(h_src_v.data(), d_src_v.data(), d_src_v.size(), handle.get_stream()); + raft::update_host(h_dst_v.data(), d_dst_v.data(), d_dst_v.size(), handle.get_stream()); + + handle.get_stream_view().synchronize(); + + float expected_edge_count = p * num_vertices * num_vertices; + + ASSERT_GE(h_src_v.size(), static_cast(expected_edge_count * 0.8)); + ASSERT_LE(h_src_v.size(), static_cast(expected_edge_count * 1.2)); + ASSERT_EQ(std::count_if(h_src_v.begin(), + h_src_v.end(), + [n = static_cast(num_vertices)](auto v) { + return !cugraph::experimental::is_valid_vertex(n, v); + }), + 0); + ASSERT_EQ(std::count_if(h_dst_v.begin(), + h_dst_v.end(), + [n = static_cast(num_vertices)](auto v) { + return !cugraph::experimental::is_valid_vertex(n, v); + }), + 0); +} + +TEST_F(GenerateErdosRenyiTest, ERTest) +{ + er_test(size_t{10}, float{0.1}); + er_test(size_t{20}, float{0.1}); + er_test(size_t{50}, float{0.1}); + er_test(size_t{10000}, float{0.1}); +} + +CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/generators/generators_test.cpp b/cpp/tests/generators/generators_test.cpp new file mode 100644 index 00000000000..11e63d81f36 --- /dev/null +++ b/cpp/tests/generators/generators_test.cpp @@ -0,0 +1,689 @@ +/* + * 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 + +struct GeneratorsTest : public ::testing::Test { +}; + +TEST_F(GeneratorsTest, PathGraphTest) +{ + using vertex_t = int32_t; + + std::vector expected_src_v({0, 1, 2, 3}); + std::vector expected_dst_v({1, 2, 3, 4}); + std::vector actual_src_v; + std::vector actual_dst_v; + + std::vector> parameters({{5, 0}}); + + raft::handle_t handle; + + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::tie(src_v, dst_v) = cugraph::generate_path_graph_edgelist(handle, parameters); + + actual_src_v.resize(src_v.size()); + actual_dst_v.resize(src_v.size()); + + raft::update_host(actual_src_v.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(actual_dst_v.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + + EXPECT_EQ(expected_src_v, actual_src_v); + EXPECT_EQ(expected_dst_v, actual_dst_v); +} + +TEST_F(GeneratorsTest, Mesh2DGraphTest) +{ + using vertex_t = int32_t; + + std::vector expected_src_v({0, 1, 2, 4, 5, 6, 8, 9, 10, 12, 13, 14, 16, 17, 18, + 20, 21, 22, 0, 1, 2, 3, 8, 9, 10, 11, 16, 17, 18, 19}); + std::vector expected_dst_v({1, 2, 3, 5, 6, 7, 9, 10, 11, 13, 14, 15, 17, 18, 19, + 21, 22, 23, 4, 5, 6, 7, 12, 13, 14, 15, 20, 21, 22, 23}); + std::vector actual_src_v; + std::vector actual_dst_v; + + raft::handle_t handle; + + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::vector> parameters( + {{4, 2, 0}, {4, 2, 8}, {4, 2, 16}}); + + std::tie(src_v, dst_v) = cugraph::generate_2d_mesh_graph_edgelist(handle, parameters); + + actual_src_v.resize(src_v.size()); + actual_dst_v.resize(src_v.size()); + + raft::update_host(actual_src_v.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(actual_dst_v.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())) + + expected_src_v.size()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())) + + actual_src_v.size()); + + EXPECT_EQ(expected_src_v, actual_src_v); + EXPECT_EQ(expected_dst_v, actual_dst_v); +} + +TEST_F(GeneratorsTest, Mesh3DGraphTest) +{ + using vertex_t = int32_t; + + std::vector expected_src_v( + {0, 1, 3, 4, 6, 7, 9, 10, 12, 13, 15, 16, 18, 19, 21, 22, 24, 25, 27, 28, 30, 31, 33, 34, + 36, 37, 39, 40, 42, 43, 45, 46, 48, 49, 51, 52, 54, 55, 57, 58, 60, 61, 63, 64, 66, 67, 69, 70, + 72, 73, 75, 76, 78, 79, 0, 1, 2, 3, 4, 5, 9, 10, 11, 12, 13, 14, 18, 19, 20, 21, 22, 23, + 27, 28, 29, 30, 31, 32, 36, 37, 38, 39, 40, 41, 45, 46, 47, 48, 49, 50, 54, 55, 56, 57, 58, 59, + 63, 64, 65, 66, 67, 68, 72, 73, 74, 75, 76, 77, 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, + 12, 13, 14, 15, 16, 17, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, + 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71}); + + std::vector expected_dst_v( + {1, 2, 4, 5, 7, 8, 10, 11, 13, 14, 16, 17, 19, 20, 22, 23, 25, 26, 28, 29, 31, 32, 34, 35, + 37, 38, 40, 41, 43, 44, 46, 47, 49, 50, 52, 53, 55, 56, 58, 59, 61, 62, 64, 65, 67, 68, 70, 71, + 73, 74, 76, 77, 79, 80, 3, 4, 5, 6, 7, 8, 12, 13, 14, 15, 16, 17, 21, 22, 23, 24, 25, 26, + 30, 31, 32, 33, 34, 35, 39, 40, 41, 42, 43, 44, 48, 49, 50, 51, 52, 53, 57, 58, 59, 60, 61, 62, + 66, 67, 68, 69, 70, 71, 75, 76, 77, 78, 79, 80, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, + 21, 22, 23, 24, 25, 26, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, + 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80}); + + std::vector actual_src_v; + std::vector actual_dst_v; + + raft::handle_t handle; + + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::vector> parameters( + {{3, 3, 3, 0}, {3, 3, 3, 27}, {3, 3, 3, 54}}); + + std::tie(src_v, dst_v) = cugraph::generate_3d_mesh_graph_edgelist(handle, parameters); + + actual_src_v.resize(src_v.size()); + actual_dst_v.resize(src_v.size()); + + raft::update_host(actual_src_v.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(actual_dst_v.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())) + + expected_src_v.size()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())) + + actual_src_v.size()); + + EXPECT_EQ(expected_src_v, actual_src_v); + EXPECT_EQ(expected_dst_v, actual_dst_v); +} + +TEST_F(GeneratorsTest, CompleteGraphTestTriangles) +{ + using vertex_t = int32_t; + + std::vector expected_src_v({0, 0, 1, 3, 3, 4, 6, 6, 7}); + std::vector expected_dst_v({1, 2, 2, 4, 5, 5, 7, 8, 8}); + std::vector actual_src_v; + std::vector actual_dst_v; + + raft::handle_t handle; + + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::vector> parameters({{3, 0}, {3, 3}, {3, 6}}); + + std::tie(src_v, dst_v) = cugraph::generate_complete_graph_edgelist(handle, parameters); + + actual_src_v.resize(src_v.size()); + actual_dst_v.resize(src_v.size()); + + raft::update_host(actual_src_v.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(actual_dst_v.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())) + + expected_src_v.size()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())) + + actual_src_v.size()); + + EXPECT_EQ(expected_src_v, actual_src_v); + EXPECT_EQ(expected_dst_v, actual_dst_v); +} + +TEST_F(GeneratorsTest, CompleteGraphTest5) +{ + using vertex_t = int32_t; + + size_t num_vertices{5}; + size_t num_graphs{3}; + + std::vector expected_src_v({0, 0, 0, 0, 1, 1, 1, 2, 2, 3, 5, 5, 5, 5, 6, + 6, 6, 7, 7, 8, 10, 10, 10, 10, 11, 11, 11, 12, 12, 13}); + std::vector expected_dst_v({1, 2, 3, 4, 2, 3, 4, 3, 4, 4, 6, 7, 8, 9, 7, + 8, 9, 8, 9, 9, 11, 12, 13, 14, 12, 13, 14, 13, 14, 14}); + std::vector actual_src_v; + std::vector actual_dst_v; + + raft::handle_t handle; + + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::vector> parameters({{5, 0}, {5, 5}, {5, 10}}); + + std::tie(src_v, dst_v) = cugraph::generate_complete_graph_edgelist(handle, parameters); + + actual_src_v.resize(src_v.size()); + actual_dst_v.resize(src_v.size()); + + raft::update_host(actual_src_v.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(actual_dst_v.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())) + + expected_src_v.size()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())) + + actual_src_v.size()); + + EXPECT_EQ(expected_src_v, actual_src_v); + EXPECT_EQ(expected_dst_v, actual_dst_v); +} + +TEST_F(GeneratorsTest, LineGraphTestSymmetric) +{ + using vertex_t = int32_t; + + size_t num_vertices{5}; + std::vector expected_src_v({0, 1, 2, 3, 1, 2, 3, 4}); + std::vector expected_dst_v({1, 2, 3, 4, 0, 1, 2, 3}); + std::vector actual_src_v; + std::vector actual_dst_v; + + raft::handle_t handle; + + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::vector> parameters({{5, 0}}); + + std::tie(src_v, dst_v) = cugraph::generate_path_graph_edgelist(handle, parameters); + std::tie(src_v, dst_v, std::ignore) = cugraph::symmetrize_edgelist( + handle, std::move(src_v), std::move(dst_v), std::nullopt); + + actual_src_v.resize(src_v.size()); + actual_dst_v.resize(src_v.size()); + + raft::update_host(actual_src_v.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(actual_dst_v.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())) + + expected_src_v.size()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())) + + actual_src_v.size()); + + EXPECT_EQ(expected_src_v, actual_src_v); + EXPECT_EQ(expected_dst_v, actual_dst_v); +} + +TEST_F(GeneratorsTest, Mesh2DGraphTestSymmetric) +{ + using vertex_t = int32_t; + + size_t x{4}; + size_t y{2}; + size_t num_graphs{3}; + + std::vector expected_src_v({0, 1, 2, 4, 5, 6, 8, 9, 10, 12, 13, 14, 16, 17, 18, + 20, 21, 22, 0, 1, 2, 3, 8, 9, 10, 11, 16, 17, 18, 19, + 1, 2, 3, 5, 6, 7, 9, 10, 11, 13, 14, 15, 17, 18, 19, + 21, 22, 23, 4, 5, 6, 7, 12, 13, 14, 15, 20, 21, 22, 23}); + std::vector expected_dst_v({1, 2, 3, 5, 6, 7, 9, 10, 11, 13, 14, 15, 17, 18, 19, + 21, 22, 23, 4, 5, 6, 7, 12, 13, 14, 15, 20, 21, 22, 23, + 0, 1, 2, 4, 5, 6, 8, 9, 10, 12, 13, 14, 16, 17, 18, + 20, 21, 22, 0, 1, 2, 3, 8, 9, 10, 11, 16, 17, 18, 19}); + std::vector actual_src_v; + std::vector actual_dst_v; + + raft::handle_t handle; + + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::vector> parameters( + {{4, 2, 0}, {4, 2, 8}, {4, 2, 16}}); + + std::tie(src_v, dst_v) = cugraph::generate_2d_mesh_graph_edgelist(handle, parameters); + std::tie(src_v, dst_v, std::ignore) = cugraph::symmetrize_edgelist( + handle, std::move(src_v), std::move(dst_v), std::nullopt); + + actual_src_v.resize(src_v.size()); + actual_dst_v.resize(src_v.size()); + + raft::update_host(actual_src_v.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(actual_dst_v.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())) + + expected_src_v.size()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())) + + actual_src_v.size()); + + EXPECT_EQ(expected_src_v, actual_src_v); + EXPECT_EQ(expected_dst_v, actual_dst_v); +} + +TEST_F(GeneratorsTest, Mesh3DGraphTestSymmetric) +{ + using vertex_t = int32_t; + + size_t x{3}; + size_t y{3}; + size_t z{3}; + size_t num_graphs{3}; + + std::vector expected_src_v( + {0, 1, 3, 4, 6, 7, 9, 10, 12, 13, 15, 16, 18, 19, 21, 22, 24, 25, 27, 28, 30, 31, 33, 34, + 36, 37, 39, 40, 42, 43, 45, 46, 48, 49, 51, 52, 54, 55, 57, 58, 60, 61, 63, 64, 66, 67, 69, 70, + 72, 73, 75, 76, 78, 79, 0, 1, 2, 3, 4, 5, 9, 10, 11, 12, 13, 14, 18, 19, 20, 21, 22, 23, + 27, 28, 29, 30, 31, 32, 36, 37, 38, 39, 40, 41, 45, 46, 47, 48, 49, 50, 54, 55, 56, 57, 58, 59, + 63, 64, 65, 66, 67, 68, 72, 73, 74, 75, 76, 77, 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, + 12, 13, 14, 15, 16, 17, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, + 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 1, 2, 4, 5, 7, 8, + 10, 11, 13, 14, 16, 17, 19, 20, 22, 23, 25, 26, 28, 29, 31, 32, 34, 35, 37, 38, 40, 41, 43, 44, + 46, 47, 49, 50, 52, 53, 55, 56, 58, 59, 61, 62, 64, 65, 67, 68, 70, 71, 73, 74, 76, 77, 79, 80, + 3, 4, 5, 6, 7, 8, 12, 13, 14, 15, 16, 17, 21, 22, 23, 24, 25, 26, 30, 31, 32, 33, 34, 35, + 39, 40, 41, 42, 43, 44, 48, 49, 50, 51, 52, 53, 57, 58, 59, 60, 61, 62, 66, 67, 68, 69, 70, 71, + 75, 76, 77, 78, 79, 80, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, + 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 63, 64, 65, 66, 67, 68, + 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80}); + + std::vector expected_dst_v( + {1, 2, 4, 5, 7, 8, 10, 11, 13, 14, 16, 17, 19, 20, 22, 23, 25, 26, 28, 29, 31, 32, 34, 35, + 37, 38, 40, 41, 43, 44, 46, 47, 49, 50, 52, 53, 55, 56, 58, 59, 61, 62, 64, 65, 67, 68, 70, 71, + 73, 74, 76, 77, 79, 80, 3, 4, 5, 6, 7, 8, 12, 13, 14, 15, 16, 17, 21, 22, 23, 24, 25, 26, + 30, 31, 32, 33, 34, 35, 39, 40, 41, 42, 43, 44, 48, 49, 50, 51, 52, 53, 57, 58, 59, 60, 61, 62, + 66, 67, 68, 69, 70, 71, 75, 76, 77, 78, 79, 80, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, + 21, 22, 23, 24, 25, 26, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, + 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 0, 1, 3, 4, 6, 7, + 9, 10, 12, 13, 15, 16, 18, 19, 21, 22, 24, 25, 27, 28, 30, 31, 33, 34, 36, 37, 39, 40, 42, 43, + 45, 46, 48, 49, 51, 52, 54, 55, 57, 58, 60, 61, 63, 64, 66, 67, 69, 70, 72, 73, 75, 76, 78, 79, + 0, 1, 2, 3, 4, 5, 9, 10, 11, 12, 13, 14, 18, 19, 20, 21, 22, 23, 27, 28, 29, 30, 31, 32, + 36, 37, 38, 39, 40, 41, 45, 46, 47, 48, 49, 50, 54, 55, 56, 57, 58, 59, 63, 64, 65, 66, 67, 68, + 72, 73, 74, 75, 76, 77, 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, + 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 54, 55, 56, 57, 58, 59, + 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71}); + + std::vector actual_src_v; + std::vector actual_dst_v; + + raft::handle_t handle; + + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::vector> parameters( + {{3, 3, 3, 0}, {3, 3, 3, 27}, {3, 3, 3, 54}}); + + std::tie(src_v, dst_v) = cugraph::generate_3d_mesh_graph_edgelist(handle, parameters); + std::tie(src_v, dst_v, std::ignore) = cugraph::symmetrize_edgelist( + handle, std::move(src_v), std::move(dst_v), std::nullopt); + + actual_src_v.resize(src_v.size()); + actual_dst_v.resize(src_v.size()); + + raft::update_host(actual_src_v.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(actual_dst_v.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())) + + expected_src_v.size()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())) + + actual_src_v.size()); + + EXPECT_EQ(expected_src_v, actual_src_v); + EXPECT_EQ(expected_dst_v, actual_dst_v); +} + +TEST_F(GeneratorsTest, CompleteGraphTestTrianglesSymmetric) +{ + using vertex_t = int32_t; + + size_t num_vertices{3}; + size_t num_graphs{3}; + + std::vector expected_src_v({0, 0, 1, 3, 3, 4, 6, 6, 7, 1, 2, 2, 4, 5, 5, 7, 8, 8}); + std::vector expected_dst_v({1, 2, 2, 4, 5, 5, 7, 8, 8, 0, 0, 1, 3, 3, 4, 6, 6, 7}); + std::vector actual_src_v; + std::vector actual_dst_v; + + raft::handle_t handle; + + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::vector> parameters({{3, 0}, {3, 3}, {3, 6}}); + + std::tie(src_v, dst_v) = cugraph::generate_complete_graph_edgelist(handle, parameters); + std::tie(src_v, dst_v, std::ignore) = cugraph::symmetrize_edgelist( + handle, std::move(src_v), std::move(dst_v), std::nullopt); + + actual_src_v.resize(src_v.size()); + actual_dst_v.resize(src_v.size()); + + raft::update_host(actual_src_v.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(actual_dst_v.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())) + + expected_src_v.size()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())) + + actual_src_v.size()); + + EXPECT_EQ(expected_src_v, actual_src_v); + EXPECT_EQ(expected_dst_v, actual_dst_v); +} + +TEST_F(GeneratorsTest, CompleteGraphTest5Symmetric) +{ + using vertex_t = int32_t; + + size_t num_vertices{5}; + size_t num_graphs{3}; + + std::vector expected_src_v({0, 0, 0, 0, 1, 1, 1, 2, 2, 3, 5, 5, 5, 5, 6, + 6, 6, 7, 7, 8, 10, 10, 10, 10, 11, 11, 11, 12, 12, 13, + 1, 2, 3, 4, 2, 3, 4, 3, 4, 4, 6, 7, 8, 9, 7, + 8, 9, 8, 9, 9, 11, 12, 13, 14, 12, 13, 14, 13, 14, 14}); + std::vector expected_dst_v({1, 2, 3, 4, 2, 3, 4, 3, 4, 4, 6, 7, 8, 9, 7, + 8, 9, 8, 9, 9, 11, 12, 13, 14, 12, 13, 14, 13, 14, 14, + 0, 0, 0, 0, 1, 1, 1, 2, 2, 3, 5, 5, 5, 5, 6, + 6, 6, 7, 7, 8, 10, 10, 10, 10, 11, 11, 11, 12, 12, 13}); + std::vector actual_src_v; + std::vector actual_dst_v; + + raft::handle_t handle; + + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::vector> parameters({{5, 0}, {5, 5}, {5, 10}}); + + std::tie(src_v, dst_v) = cugraph::generate_complete_graph_edgelist(handle, parameters); + std::tie(src_v, dst_v, std::ignore) = cugraph::symmetrize_edgelist( + handle, std::move(src_v), std::move(dst_v), std::nullopt); + + actual_src_v.resize(src_v.size()); + actual_dst_v.resize(src_v.size()); + + raft::update_host(actual_src_v.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(actual_dst_v.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())) + + expected_src_v.size()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())) + + actual_src_v.size()); + + EXPECT_EQ(expected_src_v, actual_src_v); + EXPECT_EQ(expected_dst_v, actual_dst_v); +} + +TEST_F(GeneratorsTest, CombineGraphsTest) +{ + using vertex_t = int32_t; + using weight_t = float; + + raft::handle_t handle; + + size_t num_vertices{8}; + + std::vector expected_src_v({0, 1, 2, 3, 4, 5, 6, 8, 9, 10, 12, 13, 14, 16, 17, 18, + 20, 21, 22, 0, 1, 2, 3, 8, 9, 10, 11, 16, 17, 18, 19}); + std::vector expected_dst_v({1, 2, 3, 4, 5, 6, 7, 9, 10, 11, 13, 14, 15, 17, 18, 19, + 21, 22, 23, 4, 5, 6, 7, 12, 13, 14, 15, 20, 21, 22, 23}); + + rmm::device_uvector src_graph_1_v(0, handle.get_stream()); + rmm::device_uvector dst_graph_1_v(0, handle.get_stream()); + rmm::device_uvector src_graph_2_v(0, handle.get_stream()); + rmm::device_uvector dst_graph_2_v(0, handle.get_stream()); + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::vector> parameters1({{num_vertices, 0}}); + std::vector> parameters2( + {{4, 2, 0}, {4, 2, 8}, {4, 2, 16}}); + + std::tie(src_graph_1_v, dst_graph_1_v) = + cugraph::generate_path_graph_edgelist(handle, parameters1); + std::tie(src_graph_2_v, dst_graph_2_v) = + cugraph::generate_2d_mesh_graph_edgelist(handle, parameters2); + + std::vector> sources; + sources.push_back(std::move(src_graph_1_v)); + sources.push_back(std::move(src_graph_2_v)); + + std::vector> dests; + dests.push_back(std::move(dst_graph_1_v)); + dests.push_back(std::move(dst_graph_2_v)); + + std::tie(src_v, dst_v, std::ignore) = cugraph::combine_edgelists( + handle, std::move(sources), std::move(dests), std::nullopt); + + std::vector actual_src_v; + std::vector actual_dst_v; + + actual_src_v.resize(src_v.size()); + actual_dst_v.resize(dst_v.size()); + + raft::update_host(actual_src_v.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(actual_dst_v.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())) + + expected_src_v.size()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())) + + actual_src_v.size()); + + EXPECT_EQ(expected_src_v, actual_src_v); + EXPECT_EQ(expected_dst_v, actual_dst_v); +} + +TEST_F(GeneratorsTest, CombineGraphsOffsetsTest) +{ + using vertex_t = int32_t; + using weight_t = float; + + raft::handle_t handle; + + size_t num_vertices{8}; + vertex_t offset{10}; + + std::vector expected_src_v({0, 1, 2, 3, 4, 5, 6, 10, 11, 12, 14, 15, 16, + 18, 19, 20, 22, 23, 24, 26, 27, 28, 30, 31, 32, 10, + 11, 12, 13, 18, 19, 20, 21, 26, 27, 28, 29}); + std::vector expected_dst_v({1, 2, 3, 4, 5, 6, 7, 11, 12, 13, 15, 16, 17, + 19, 20, 21, 23, 24, 25, 27, 28, 29, 31, 32, 33, 14, + 15, 16, 17, 22, 23, 24, 25, 30, 31, 32, 33}); + + rmm::device_uvector src_graph_1_v(0, handle.get_stream()); + rmm::device_uvector dst_graph_1_v(0, handle.get_stream()); + rmm::device_uvector src_graph_2_v(0, handle.get_stream()); + rmm::device_uvector dst_graph_2_v(0, handle.get_stream()); + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::vector> parameters1({{num_vertices, 0}}); + std::vector> parameters2( + {{4, 2, 10}, {4, 2, 18}, {4, 2, 26}}); + + std::tie(src_graph_1_v, dst_graph_1_v) = + cugraph::generate_path_graph_edgelist(handle, parameters1); + std::tie(src_graph_2_v, dst_graph_2_v) = + cugraph::generate_2d_mesh_graph_edgelist(handle, parameters2); + + std::vector> sources; + sources.push_back(std::move(src_graph_1_v)); + sources.push_back(std::move(src_graph_2_v)); + + std::vector> dests; + dests.push_back(std::move(dst_graph_1_v)); + dests.push_back(std::move(dst_graph_2_v)); + + std::tie(src_v, dst_v, std::ignore) = cugraph::combine_edgelists( + handle, std::move(sources), std::move(dests), std::nullopt); + + std::vector actual_src_v; + std::vector actual_dst_v; + + actual_src_v.resize(src_v.size()); + actual_dst_v.resize(dst_v.size()); + + raft::update_host(actual_src_v.data(), src_v.data(), src_v.size(), handle.get_stream()); + raft::update_host(actual_dst_v.data(), dst_v.data(), dst_v.size(), handle.get_stream()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(expected_src_v.begin(), expected_dst_v.begin())) + + expected_src_v.size()); + + thrust::sort( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())), + thrust::make_zip_iterator(thrust::make_tuple(actual_src_v.begin(), actual_dst_v.begin())) + + actual_src_v.size()); + + EXPECT_EQ(expected_src_v, actual_src_v); + EXPECT_EQ(expected_dst_v, actual_dst_v); +} + +TEST_F(GeneratorsTest, ScrambleTest) +{ + using vertex_t = int32_t; + using edge_t = int32_t; + + edge_t num_vertices{30}; + edge_t num_edges{100}; + + raft::handle_t handle; + + std::vector input_src_v(num_edges); + std::vector input_dst_v(num_edges); + + std::default_random_engine generator{}; + std::uniform_int_distribution distribution{0, num_vertices - 1}; + + std::generate(input_src_v.begin(), input_src_v.end(), [&distribution, &generator]() { + return distribution(generator); + }); + std::generate(input_dst_v.begin(), input_dst_v.end(), [&distribution, &generator]() { + return distribution(generator); + }); + + rmm::device_uvector d_src_v(input_src_v.size(), handle.get_stream()); + rmm::device_uvector d_dst_v(input_src_v.size(), handle.get_stream()); + std::vector output_src_v(input_src_v.size()); + std::vector output_dst_v(input_src_v.size()); + + raft::update_device(d_src_v.data(), input_src_v.data(), input_src_v.size(), handle.get_stream()); + raft::update_device(d_dst_v.data(), input_dst_v.data(), input_dst_v.size(), handle.get_stream()); + + cugraph::scramble_vertex_ids(handle, d_src_v, d_dst_v, 5, 0); + + raft::update_host(output_src_v.data(), d_src_v.data(), d_src_v.size(), handle.get_stream()); + raft::update_host(output_dst_v.data(), d_dst_v.data(), d_dst_v.size(), handle.get_stream()); + + EXPECT_TRUE(cugraph::test::renumbered_vectors_same(handle, input_src_v, output_src_v)); + EXPECT_TRUE(cugraph::test::renumbered_vectors_same(handle, input_dst_v, output_dst_v)); +} + +CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/pagerank/mg_pagerank_test.cpp b/cpp/tests/pagerank/mg_pagerank_test.cpp index 21a2c11f4a7..6370c7b7758 100644 --- a/cpp/tests/pagerank/mg_pagerank_test.cpp +++ b/cpp/tests/pagerank/mg_pagerank_test.cpp @@ -17,6 +17,7 @@ #include #include #include +#include #include #include @@ -218,8 +219,7 @@ class Tests_MGPageRank auto sg_graph_view = sg_graph.view(); - ASSERT_TRUE(mg_graph_view.get_number_of_vertices() == - sg_graph_view.get_number_of_vertices()); + ASSERT_EQ(mg_graph_view.get_number_of_vertices(), sg_graph_view.get_number_of_vertices()); // 5-4. run SG PageRank @@ -300,22 +300,24 @@ INSTANTIATE_TEST_SUITE_P( cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); -INSTANTIATE_TEST_SUITE_P(rmat_small_tests, - Tests_MGPageRank_Rmat, - ::testing::Combine(::testing::Values(PageRank_Usecase{0.0, false}, - PageRank_Usecase{0.5, false}, - PageRank_Usecase{0.0, true}, - PageRank_Usecase{0.5, true}), - ::testing::Values(cugraph::test::Rmat_Usecase( - 10, 16, 0.57, 0.19, 0.19, 0, false, false, true)))); - -INSTANTIATE_TEST_SUITE_P(rmat_large_tests, - Tests_MGPageRank_Rmat, - ::testing::Combine(::testing::Values(PageRank_Usecase{0.0, false, false}, - PageRank_Usecase{0.5, false, false}, - PageRank_Usecase{0.0, true, false}, - PageRank_Usecase{0.5, true, false}), - ::testing::Values(cugraph::test::Rmat_Usecase( - 20, 32, 0.57, 0.19, 0.19, 0, false, false, true)))); +INSTANTIATE_TEST_SUITE_P( + rmat_small_tests, + Tests_MGPageRank_Rmat, + ::testing::Combine(::testing::Values(PageRank_Usecase{0.0, false}, + PageRank_Usecase{0.5, false}, + PageRank_Usecase{0.0, true}, + PageRank_Usecase{0.5, true}), + ::testing::Values(cugraph::test::Rmat_Usecase( + 10, 16, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); + +INSTANTIATE_TEST_SUITE_P( + rmat_large_tests, + Tests_MGPageRank_Rmat, + ::testing::Combine(::testing::Values(PageRank_Usecase{0.0, false, false}, + PageRank_Usecase{0.5, false, false}, + PageRank_Usecase{0.0, true, false}, + PageRank_Usecase{0.5, true, false}), + ::testing::Values(cugraph::test::Rmat_Usecase( + 20, 32, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/utilities/rmat_utilities.cu b/cpp/tests/utilities/rmat_utilities.cu index 8f6a6cf499a..fda72fc9054 100644 --- a/cpp/tests/utilities/rmat_utilities.cu +++ b/cpp/tests/utilities/rmat_utilities.cu @@ -18,7 +18,7 @@ #include #include -#include +#include #include #include #include @@ -94,15 +94,14 @@ generate_graph_from_rmat_params(raft::handle_t const& handle, rmm::device_uvector d_tmp_rows(0, handle.get_stream()); rmm::device_uvector d_tmp_cols(0, handle.get_stream()); std::tie(i == 0 ? d_edgelist_rows : d_tmp_rows, i == 0 ? d_edgelist_cols : d_tmp_cols) = - cugraph::experimental::generate_rmat_edgelist(handle, - scale, - partition_edge_counts[i], - a, - b, - c, - base_seed + id, - undirected ? true : false, - scramble_vertex_ids); + cugraph::generate_rmat_edgelist(handle, + scale, + partition_edge_counts[i], + a, + b, + c, + base_seed + id, + undirected ? true : false); rmm::device_uvector d_tmp_weights(0, handle.get_stream()); if (test_weighted) { diff --git a/cpp/tests/utilities/test_graphs.hpp b/cpp/tests/utilities/test_graphs.hpp new file mode 100644 index 00000000000..b8ee8f024b0 --- /dev/null +++ b/cpp/tests/utilities/test_graphs.hpp @@ -0,0 +1,558 @@ +/* + * 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 test { + +namespace detail { + +class TranslateGraph_Usecase { + public: + TranslateGraph_Usecase() = delete; + TranslateGraph_Usecase(size_t base_vertex_id = 0) : base_vertex_id_(base_vertex_id) {} + + template + void translate(raft::handle_t const& handle, + rmm::device_uvector& d_src, + rmm::device_uvector& d_dst) const + { + if (base_vertex_id_ > 0) + cugraph::test::translate_vertex_ids( + handle, d_src, d_dst, static_cast(base_vertex_id_)); + } + + size_t base_vertex_id_{}; +}; + +} // namespace detail + +class File_Usecase : public detail::TranslateGraph_Usecase { + public: + File_Usecase() = delete; + + File_Usecase(std::string const& graph_file_path, size_t base_vertex_id = 0) + : detail::TranslateGraph_Usecase(base_vertex_id) + { + 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; + } + } + + template + std::tuple, + rmm::device_uvector, + rmm::device_uvector, + vertex_t, + bool> + construct_edgelist(raft::handle_t const& handle, bool test_weighted) const + { + rmm::device_uvector d_src_v(0, handle.get_stream()); + rmm::device_uvector d_dst_v(0, handle.get_stream()); + rmm::device_uvector d_weights_v(0, handle.get_stream()); + vertex_t num_vertices; + bool is_symmetric; + + std::tie(d_src_v, d_dst_v, d_weights_v, num_vertices, is_symmetric) = + read_edgelist_from_matrix_market_file( + handle, graph_file_full_path_, test_weighted); + + translate(handle, d_src_v, d_dst_v); + +#if 0 + if (multi_gpu) { + std::tie(d_src_v, d_dst_v) = filter_edgelist_by_gpu(handle, d_src_v, d_dst_v); + } +#endif + + return std::make_tuple( + std::move(d_src_v), + std::move(d_dst_v), + std::move(d_weights_v), + static_cast(detail::TranslateGraph_Usecase::base_vertex_id_) + num_vertices, + is_symmetric); + } + + template + std::tuple< + cugraph::experimental::graph_t, + rmm::device_uvector> + construct_graph(raft::handle_t const& handle, bool test_weighted, bool renumber = true) const + { + rmm::device_uvector d_src_v(0, handle.get_stream()); + rmm::device_uvector d_dst_v(0, handle.get_stream()); + rmm::device_uvector d_weights_v(0, handle.get_stream()); + vertex_t num_vertices; + bool is_symmetric; + + std::tie(d_src_v, d_dst_v, d_weights_v, num_vertices, is_symmetric) = + this->template construct_edgelist( + handle, test_weighted); + + // TODO: Consider calling construct_edgelist and creating + // a generic test function to take the edgelist and + // do the graph construction. + // + // Would be more reusable across tests + // + return read_graph_from_matrix_market_file( + handle, graph_file_full_path_, test_weighted, renumber); + } + + private: + std::string graph_file_full_path_{}; +}; + +class Rmat_Usecase : public detail::TranslateGraph_Usecase { + public: + Rmat_Usecase() = delete; + + Rmat_Usecase(size_t scale, + size_t edge_factor, + double a, + double b, + double c, + uint64_t seed, + bool undirected, + bool scramble_vertex_ids, + size_t base_vertex_id = 0, + bool multi_gpu_usecase = false) + : detail::TranslateGraph_Usecase(base_vertex_id), + scale_(scale), + edge_factor_(edge_factor), + a_(a), + b_(b), + c_(c), + seed_(seed), + undirected_(undirected), + scramble_vertex_ids_(scramble_vertex_ids), + multi_gpu_usecase_(multi_gpu_usecase) + { + } + + template + std::tuple, + rmm::device_uvector, + rmm::device_uvector, + vertex_t, + bool> + construct_edgelist(raft::handle_t const& handle, bool test_weighted) const + { + // TODO: Tease through generate_graph_from_rmat_params + // to extract the edgelist part + // Call cugraph::translate_vertex_ids(handle, d_src_v, d_dst_v, base_vertex_id_); + + CUGRAPH_FAIL("Not implemented"); + } + + template + std::tuple< + cugraph::experimental::graph_t, + rmm::device_uvector> + construct_graph(raft::handle_t const& handle, bool test_weighted, bool renumber = true) const + { + std::vector partition_ids(1); + size_t comm_size; + + if (multi_gpu_usecase_) { + auto& comm = handle.get_comms(); + comm_size = comm.get_size(); + auto const comm_rank = comm.get_rank(); + + partition_ids.resize(multi_gpu ? size_t{1} : static_cast(comm_size)); + + std::iota(partition_ids.begin(), + partition_ids.end(), + multi_gpu ? static_cast(comm_rank) : size_t{0}); + } else { + comm_size = 1; + partition_ids[0] = size_t{0}; + } + + // TODO: Need to offset by base_vertex_id_ + // static_cast(base_vertex_id_)); + // Consider using construct_edgelist like other options + return generate_graph_from_rmat_params( + handle, + scale_, + edge_factor_, + a_, + b_, + c_, + seed_, + undirected_, + scramble_vertex_ids_, + test_weighted, + renumber, + partition_ids, + comm_size); + } + + private: + size_t scale_{}; + size_t edge_factor_{}; + double a_{}; + double b_{}; + double c_{}; + uint64_t seed_{}; + bool undirected_{}; + bool scramble_vertex_ids_{}; + bool multi_gpu_usecase_{}; +}; + +class PathGraph_Usecase { + public: + PathGraph_Usecase() = delete; + + PathGraph_Usecase(std::vector> parms, + bool weighted = false, + bool scramble = false) + : parms_(parms), weighted_(weighted) + { + } + + template + std::tuple, + rmm::device_uvector, + rmm::device_uvector, + vertex_t, + bool> + construct_edgelist(raft::handle_t const& handle, bool test_weighted) const + { + rmm::device_uvector weights_v(0, handle.get_stream()); + + constexpr bool symmetric{true}; + + std::vector> converted_parms(parms_.size()); + + std::transform(parms_.begin(), parms_.end(), converted_parms.begin(), [](auto p) { + return std::make_tuple(static_cast(std::get<0>(p)), + static_cast(std::get<1>(p))); + }); + + rmm::device_uvector src_v(0, handle.get_stream()); + rmm::device_uvector dst_v(0, handle.get_stream()); + + std::tie(src_v, dst_v) = + cugraph::generate_path_graph_edgelist(handle, converted_parms); + std::tie(src_v, dst_v, std::ignore) = cugraph::symmetrize_edgelist( + handle, std::move(src_v), std::move(dst_v), std::nullopt); + + if (test_weighted) { + auto length = src_v.size(); + weights_v.resize(length, handle.get_stream()); + } + + return std::make_tuple( + std::move(src_v), std::move(dst_v), std::move(weights_v), num_vertices_, symmetric); + } + + template + std::tuple< + cugraph::experimental::graph_t, + rmm::device_uvector> + construct_graph(raft::handle_t const& handle, bool test_weighted, bool renumber = true) const + { + CUGRAPH_FAIL("not implemented"); + } + + private: + std::vector> parms_{}; + size_t num_vertices_{0}; + bool weighted_{false}; +}; + +class Mesh2DGraph_Usecase { + public: + Mesh2DGraph_Usecase() = delete; + + Mesh2DGraph_Usecase(std::vector> const& parms, bool weighted) + : parms_(parms), weighted_(weighted) + { + } + + template + std::tuple, + rmm::device_uvector, + rmm::device_uvector, + vertex_t, + bool> + construct_edgelist(raft::handle_t const& handle, bool test_weighted) const + { + } + + template + std::tuple< + cugraph::experimental::graph_t, + rmm::device_uvector> + construct_graph(raft::handle_t const& handle, bool test_weighted, bool renumber = true) const; + + private: + std::vector> parms_{}; + bool weighted_{false}; +}; + +class Mesh3DGraph_Usecase { + public: + Mesh3DGraph_Usecase() = delete; + + Mesh3DGraph_Usecase(std::vector> const& parms, + bool weighted) + : parms_(parms), weighted_(weighted) + { + } + + template + std::tuple, + rmm::device_uvector, + rmm::device_uvector, + vertex_t, + bool> + construct_edgelist(raft::handle_t const& handle, bool test_weighted) const; + + template + std::tuple< + cugraph::experimental::graph_t, + rmm::device_uvector> + construct_graph(raft::handle_t const& handle, bool test_weighted, bool renumber = true) const; + + private: + std::vector> parms_{}; + bool weighted_{false}; +}; + +class CompleteGraph_Usecase { + public: + CompleteGraph_Usecase() = delete; + + CompleteGraph_Usecase(std::vector> const& parms, bool weighted) + : parms_(parms), weighted_(weighted) + { + } + + template + std::tuple, + rmm::device_uvector, + rmm::device_uvector, + vertex_t, + bool> + construct_edgelist(raft::handle_t const& handle, bool test_weighted) const; + + template + std::tuple< + cugraph::experimental::graph_t, + rmm::device_uvector> + construct_graph(raft::handle_t const& handle, bool test_weighted, bool renumber = true) const; + + private: + std::vector> parms_{}; + bool weighted_{false}; +}; + +namespace detail { + +template +struct combined_construct_graph_tuple_impl { + template + std::vector, + rmm::device_uvector, + rmm::device_uvector, + vertex_t, + bool>> + construct_edges(raft::handle_t const& handle, + bool test_weighted, + generator_tuple_t const& generator_tuple) const + { + return combined_construct_graph_tuple_impl() + .construct_edges(generator_tuple) + .push_back(std::get(generator_tuple).construct_edges(handle, test_weighted)); + } +}; + +template +struct combined_construct_graph_tuple_impl { + template + std::vector, + rmm::device_uvector, + rmm::device_uvector, + vertex_t, + bool>> + construct_edges(raft::handle_t const& handle, + bool test_weighted, + generator_tuple_t const& generator_tuple) const + { + return std::vector, + rmm::device_uvector, + rmm::device_uvector, + vertex_t, + bool>>(); + } +}; + +} // namespace detail + +template +class CombinedGenerator_Usecase { + CombinedGenerator_Usecase() = delete; + + CombinedGenerator_Usecase(generator_tuple_t const& tuple) : generator_tuple_(tuple) {} + + template + std::tuple, + rmm::device_uvector, + rmm::device_uvector, + vertex_t, + bool> + construct_edgelist(raft::handle_t const& handle, bool test_weighted) const + { + size_t constexpr tuple_size{std::tuple_size::value}; + + auto edge_tuple_vector = + detail::combined_construct_graph_tuple_impl() + .construct_edges(handle, test_weighted, generator_tuple_); + + // Need to combine + CUGRAPH_FAIL("not implemented"); + } + + template + std::tuple< + cugraph::experimental::graph_t, + rmm::device_uvector> + construct_graph(raft::handle_t const& handle, bool test_weighted, bool renumber = true) const + { + // Call construct_edgelist to get tuple of edge lists + // return generate_graph_from_edgelist<...>(...) + CUGRAPH_FAIL("not implemented"); + } + + private: + generator_tuple_t const& generator_tuple_; +}; + +template +std::tuple, + rmm::device_uvector> +construct_graph(raft::handle_t const& handle, + input_usecase_t const& input_usecase, + bool test_weighted, + bool renumber = true) +{ + rmm::device_uvector d_src_v(0, handle.get_stream()); + rmm::device_uvector d_dst_v(0, handle.get_stream()); + rmm::device_uvector d_weights_v(0, handle.get_stream()); + vertex_t num_vertices{0}; + bool is_symmetric{false}; + + std::tie(d_src_v, d_dst_v, d_weights_v, num_vertices, is_symmetric) = + input_usecase + .template construct_edgelist( + handle, test_weighted); + + return cugraph::experimental:: + create_graph_from_edgelist( + handle, + std::nullopt, + std::move(d_src_v), + std::move(d_dst_v), + std::move(d_weights_v), + cugraph::experimental::graph_properties_t{is_symmetric, false, test_weighted}, + renumber); +} + +} // namespace test +} // namespace cugraph diff --git a/cpp/tests/utilities/test_utilities.hpp b/cpp/tests/utilities/test_utilities.hpp index 075db1906e1..09da0556e44 100644 --- a/cpp/tests/utilities/test_utilities.hpp +++ b/cpp/tests/utilities/test_utilities.hpp @@ -18,6 +18,8 @@ #include #include +#include +#include #include #include @@ -153,122 +155,6 @@ generate_graph_from_rmat_params(raft::handle_t const& handle, std::vector const& partition_ids, size_t num_partitions); -class File_Usecase { - public: - File_Usecase() = delete; - - File_Usecase(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; - } - } - - template - std::tuple< - cugraph::experimental::graph_t, - rmm::device_uvector> - construct_graph(raft::handle_t const& handle, bool test_weighted, bool renumber = true) const - { - return read_graph_from_matrix_market_file( - handle, graph_file_full_path_, test_weighted, renumber); - } - - private: - std::string graph_file_full_path_{}; -}; - -class Rmat_Usecase { - public: - Rmat_Usecase() = delete; - - Rmat_Usecase(size_t scale, - size_t edge_factor, - double a, - double b, - double c, - uint64_t seed, - bool undirected, - bool scramble_vertex_ids, - bool multi_gpu_usecase = false) - : scale_(scale), - edge_factor_(edge_factor), - a_(a), - b_(b), - c_(c), - seed_(seed), - undirected_(undirected), - scramble_vertex_ids_(scramble_vertex_ids), - multi_gpu_usecase_(multi_gpu_usecase) - { - } - - template - std::tuple< - cugraph::experimental::graph_t, - rmm::device_uvector> - construct_graph(raft::handle_t const& handle, bool test_weighted, bool renumber = true) const - { - std::vector partition_ids(1); - size_t comm_size; - - if (multi_gpu_usecase_) { - auto& comm = handle.get_comms(); - comm_size = comm.get_size(); - auto const comm_rank = comm.get_rank(); - - partition_ids.resize(multi_gpu ? size_t{1} : static_cast(comm_size)); - - std::iota(partition_ids.begin(), - partition_ids.end(), - multi_gpu ? static_cast(comm_rank) : size_t{0}); - } else { - comm_size = 1; - partition_ids[0] = size_t{0}; - } - - return generate_graph_from_rmat_params( - handle, - scale_, - edge_factor_, - a_, - b_, - c_, - seed_, - undirected_, - scramble_vertex_ids_, - test_weighted, - renumber, - partition_ids, - comm_size); - } - - private: - size_t scale_{}; - size_t edge_factor_{}; - double a_{}; - double b_{}; - double c_{}; - uint64_t seed_{}; - bool undirected_{}; - bool scramble_vertex_ids_{}; - bool multi_gpu_usecase_{}; -}; - // alias for easy customization for debug purposes: // template @@ -391,5 +277,49 @@ std::pair compare_graphs(raft::handle_t const& handle, } } +template +bool renumbered_vectors_same(raft::handle_t const& handle, + std::vector const& v1, + std::vector const& v2) +{ + if (v1.size() != v2.size()) return false; + + std::map map; + + auto iter = thrust::make_zip_iterator(thrust::make_tuple(v1.begin(), v2.begin())); + + std::for_each(iter, iter + v1.size(), [&map](auto pair) { + vertex_t e1 = thrust::get<0>(pair); + vertex_t e2 = thrust::get<1>(pair); + + map[e1] = e2; + }); + + auto error_count = std::count_if(iter, iter + v1.size(), [&map](auto pair) { + vertex_t e1 = thrust::get<0>(pair); + vertex_t e2 = thrust::get<1>(pair); + + return (map[e1] != e2); + }); + + return (error_count == 0); +} + +template +bool renumbered_vectors_same(raft::handle_t const& handle, + rmm::device_uvector const& v1, + rmm::device_uvector const& v2) +{ + if (v1.size() != v2.size()) return false; + + std::vector h_v1(v1.size()); + std::vector h_v2(v1.size()); + + raft::update_host(h_v1.data(), v1.data(), v1.size(), handle.get_stream()); + raft::update_host(h_v2.data(), v2.data(), v2.size(), handle.get_stream()); + + return renumbered_vectors_same(handle, h_v1, h_v2); +} + } // namespace test } // namespace cugraph diff --git a/cpp/tests/utilities/thrust_wrapper.cu b/cpp/tests/utilities/thrust_wrapper.cu index dfd420b1e2d..ae36582d18d 100644 --- a/cpp/tests/utilities/thrust_wrapper.cu +++ b/cpp/tests/utilities/thrust_wrapper.cu @@ -81,5 +81,53 @@ sort_by_key(raft::handle_t const& handle, int64_t const* values, size_t num_pairs); +template +void translate_vertex_ids(raft::handle_t const& handle, + rmm::device_uvector& d_src_v, + rmm::device_uvector& d_dst_v, + vertex_t vertex_id_offset) +{ + thrust::transform(rmm::exec_policy(handle.get_stream()), + d_src_v.begin(), + d_src_v.end(), + d_src_v.begin(), + [offset = vertex_id_offset] __device__(vertex_t v) { return offset + v; }); + + thrust::transform(rmm::exec_policy(handle.get_stream()), + d_dst_v.begin(), + d_dst_v.end(), + d_dst_v.begin(), + [offset = vertex_id_offset] __device__(vertex_t v) { return offset + v; }); +} + +template +void populate_vertex_ids(raft::handle_t const& handle, + rmm::device_uvector& d_vertices_v, + vertex_t vertex_id_offset) +{ + thrust::sequence(rmm::exec_policy(handle.get_stream()), + d_vertices_v.begin(), + d_vertices_v.end(), + vertex_id_offset); +} + +template void translate_vertex_ids(raft::handle_t const& handle, + rmm::device_uvector& d_src_v, + rmm::device_uvector& d_dst_v, + int32_t vertex_id_offset); + +template void translate_vertex_ids(raft::handle_t const& handle, + rmm::device_uvector& d_src_v, + rmm::device_uvector& d_dst_v, + int64_t vertex_id_offset); + +template void populate_vertex_ids(raft::handle_t const& handle, + rmm::device_uvector& d_vertices_v, + int32_t vertex_id_offset); + +template void populate_vertex_ids(raft::handle_t const& handle, + rmm::device_uvector& d_vertices_v, + int64_t vertex_id_offset); + } // namespace test } // namespace cugraph diff --git a/cpp/tests/utilities/thrust_wrapper.hpp b/cpp/tests/utilities/thrust_wrapper.hpp index 96f370f884c..45208a6b921 100644 --- a/cpp/tests/utilities/thrust_wrapper.hpp +++ b/cpp/tests/utilities/thrust_wrapper.hpp @@ -26,5 +26,16 @@ template std::tuple, rmm::device_uvector> sort_by_key( raft::handle_t const& handle, vertex_t const* keys, value_t const* values, size_t num_pairs); +template +void translate_vertex_ids(raft::handle_t const& handle, + rmm::device_uvector& d_src_v, + rmm::device_uvector& d_dst_v, + vertex_t vertex_id_offset); + +template +void populate_vertex_ids(raft::handle_t const& handle, + rmm::device_uvector& d_vertices_v, + vertex_t vertex_id_offset); + } // namespace test } // namespace cugraph diff --git a/python/cugraph/generators/rmat.pxd b/python/cugraph/generators/rmat.pxd index 16606b59d0f..3c51108c778 100644 --- a/python/cugraph/generators/rmat.pxd +++ b/python/cugraph/generators/rmat.pxd @@ -14,10 +14,10 @@ from libcpp cimport bool from cugraph.structure.graph_utilities cimport * from libcpp.vector cimport vector -cdef extern from "cugraph/experimental/graph_generator.hpp" namespace "cugraph::experimental": +cdef extern from "cugraph/graph_generators.hpp" namespace "cugraph": ctypedef enum generator_distribution_t: - POWER_LAW "cugraph::experimental::generator_distribution_t::POWER_LAW" - UNIFORM "cugraph::experimental::generator_distribution_t::UNIFORM" + POWER_LAW "cugraph::generator_distribution_t::POWER_LAW" + UNIFORM "cugraph::generator_distribution_t::UNIFORM" cdef extern from "cugraph/utilities/cython.hpp" namespace "cugraph::cython":