diff --git a/.gitignore b/.gitignore index 30bcd5a845d..c2498b35cf1 100644 --- a/.gitignore +++ b/.gitignore @@ -63,6 +63,7 @@ datasets/* !datasets/cyber.csv !datasets/karate-data.csv !datasets/karate_undirected.csv +!datasets/karate-disjoint.csv !datasets/netscience.csv diff --git a/CHANGELOG.md b/CHANGELOG.md index c7594c03f6c..44b9c097774 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,12 +1,15 @@ # cuGraph 0.16.0 (Date TBD) ## New Features +- PR 1098 Add new graph classes to support 2D partitioning ## Improvements - PR 1081 MNMG Renumbering - sort partitions by degree - PR 1115 Replace deprecated rmm::mr::get_default_resource with rmm::mr::get_current_device_resource +- PR #1129 Refactored test to use common dataset and added additional doc pages ## Bug Fixes +- PR #1131 Show style checker errors with set +e # cuGraph 0.15.0 (26 Aug 2020) @@ -100,6 +103,8 @@ - PR #1087 Updated benchmarks README to better describe how to get plugin, added rapids-pytest-benchmark plugin to conda dev environments - PR #1101 Removed unnecessary device-to-host copy which caused a performance regression - PR #1106 Added new release.ipynb to notebook test skip list +- PR #1125 Patch Thrust to workaround `CUDA_CUB_RET_IF_FAIL` macro clearing CUDA errors + # cuGraph 0.14.0 (03 Jun 2020) diff --git a/README.md b/README.md index 45405d902bf..a51b9fb4e0c 100644 --- a/README.md +++ b/README.md @@ -6,7 +6,7 @@ The [RAPIDS](https://rapids.ai) cuGraph library is a collection of GPU accelerat For more project details, see [rapids.ai](https://rapids.ai/). -**NOTE:** For the latest stable [README.md](https://github.com/rapidsai/cudf/blob/main/README.md) ensure you are on the latest branch. +**NOTE:** For the latest stable [README.md](https://github.com/rapidsai/cugraph/blob/main/README.md) ensure you are on the latest branch. diff --git a/ci/checks/style.sh b/ci/checks/style.sh index 696f566a96a..978ac03d85b 100755 --- a/ci/checks/style.sh +++ b/ci/checks/style.sh @@ -11,6 +11,9 @@ # captured for returning a final status code. This allows all style check to # take place to provide a more comprehensive list of style violations. set -o pipefail +# CI does `set -e` then sources this file, so we override that so we can output +# the results from the various style checkers +set +e ERRORCODE=0 PATH=/conda/bin:$PATH diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index a40fc493f2b..03049a21f00 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -214,6 +214,7 @@ FetchContent_Declare( GIT_REPOSITORY https://github.com/thrust/thrust.git GIT_TAG 1.9.10 GIT_SHALLOW true + PATCH_COMMAND COMMAND patch -p1 < "${CMAKE_CURRENT_SOURCE_DIR}/cmake/thrust-ret-if-fail.patch" ) FetchContent_GetProperties(thrust) @@ -356,6 +357,8 @@ add_library(cugraph SHARED src/components/connectivity.cu src/centrality/katz_centrality.cu src/centrality/betweenness_centrality.cu + src/experimental/graph.cu + src/experimental/graph_view.cu ) # diff --git a/cpp/cmake/thrust-ret-if-fail.patch b/cpp/cmake/thrust-ret-if-fail.patch new file mode 100644 index 00000000000..990b3f993be --- /dev/null +++ b/cpp/cmake/thrust-ret-if-fail.patch @@ -0,0 +1,16 @@ +diff --git a/thrust/system/cuda/detail/core/util.h b/thrust/system/cuda/detail/core/util.h +index a2c87772..ea4ed640 100644 +--- a/thrust/system/cuda/detail/core/util.h ++++ b/thrust/system/cuda/detail/core/util.h +@@ -652,7 +652,10 @@ namespace core { + } + + #define CUDA_CUB_RET_IF_FAIL(e) \ +- if (cub::Debug((e), __FILE__, __LINE__)) return e; ++ { \ ++ auto const error = (e); \ ++ if (cub::Debug(error, __FILE__, __LINE__)) return error; \ ++ } + + // uninitialized + // ------- diff --git a/cpp/include/experimental/detail/graph_utils.cuh b/cpp/include/experimental/detail/graph_utils.cuh new file mode 100644 index 00000000000..fe092342f80 --- /dev/null +++ b/cpp/include/experimental/detail/graph_utils.cuh @@ -0,0 +1,129 @@ +/* + * Copyright (c) 2020, 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 +#include + +#include +#include + +namespace cugraph { +namespace experimental { +namespace detail { + +// compute the numbers of nonzeros in rows (of the graph adjacency matrix, if store_transposed = +// false) or columns (of the graph adjacency matrix, if store_transposed = true) +template +rmm::device_uvector compute_major_degree( + raft::handle_t const &handle, + std::vector const &adj_matrix_partition_offsets, + partition_t const &partition) +{ + auto &comm_p_row = handle.get_subcomm(comm_p_row_key); + auto const comm_p_row_rank = comm_p_row.get_rank(); + auto const comm_p_row_size = comm_p_row.get_size(); + auto &comm_p_col = handle.get_subcomm(comm_p_col_key); + auto const comm_p_col_rank = comm_p_col.get_rank(); + auto const comm_p_col_size = comm_p_col.get_size(); + + rmm::device_uvector local_degrees(0, handle.get_stream()); + rmm::device_uvector degrees(0, handle.get_stream()); + + vertex_t max_num_local_degrees{0}; + for (int i = 0; i < comm_p_col_size; ++i) { + auto vertex_partition_idx = + partition.is_hypergraph_partitioned() + ? static_cast(comm_p_row_size) * static_cast(i) + + static_cast(comm_p_row_rank) + : static_cast(comm_p_col_size) * static_cast(comm_p_row_rank) + + static_cast(i); + vertex_t major_first{}; + vertex_t major_last{}; + std::tie(major_first, major_last) = partition.get_vertex_partition_range(vertex_partition_idx); + max_num_local_degrees = std::max(max_num_local_degrees, major_last - major_first); + if (i == comm_p_col_rank) { degrees.resize(major_last - major_first, handle.get_stream()); } + } + local_degrees.resize(max_num_local_degrees, handle.get_stream()); + for (int i = 0; i < comm_p_col_size; ++i) { + auto vertex_partition_idx = + partition.is_hypergraph_partitioned() + ? static_cast(comm_p_row_size) * static_cast(i) + + static_cast(comm_p_row_rank) + : static_cast(comm_p_col_size) * static_cast(comm_p_row_rank) + + static_cast(i); + vertex_t major_first{}; + vertex_t major_last{}; + std::tie(major_first, major_last) = partition.get_vertex_partition_range(vertex_partition_idx); + auto p_offsets = partition.is_hypergraph_partitioned() + ? adj_matrix_partition_offsets[i] + : adj_matrix_partition_offsets[0] + + (major_first - partition.get_vertex_partition_range_first( + comm_p_col_size * comm_p_row_rank)); + thrust::transform(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + thrust::make_counting_iterator(vertex_t{0}), + thrust::make_counting_iterator(major_last - major_first), + local_degrees.data(), + [p_offsets] __device__(auto i) { return p_offsets[i + 1] - p_offsets[i]; }); + comm_p_row.reduce(local_degrees.data(), + i == comm_p_col_rank ? degrees.data() : static_cast(nullptr), + degrees.size(), + raft::comms::op_t::SUM, + comm_p_col_rank, + handle.get_stream()); + } + + auto status = handle.get_comms().sync_stream( + handle.get_stream()); // this is neessary as local_degrees will become out-of-scope once this + // function returns. + CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); + + return degrees; +} + +// compute the numbers of nonzeros in rows (of the graph adjacency matrix, if store_transposed = +// false) or columns (of the graph adjacency matrix, if store_transposed = true) +template +rmm::device_uvector compute_major_degree( + raft::handle_t const &handle, + std::vector> const &adj_matrix_partition_offsets, + partition_t const &partition) +{ + // we can avoid creating this temporary with "if constexpr" supported from C++17 + std::vector tmp_offsets(adj_matrix_partition_offsets.size(), nullptr); + std::transform(adj_matrix_partition_offsets.begin(), + adj_matrix_partition_offsets.end(), + tmp_offsets.begin(), + [](auto const &offsets) { return offsets.data(); }); + return compute_major_degree(handle, tmp_offsets, partition); +} + +template +struct degree_from_offsets_t { + edge_t const *offsets{nullptr}; + + __device__ edge_t operator()(vertex_t v) { return offsets[v + 1] - offsets[v]; } +}; + +} // namespace detail +} // namespace experimental +} // namespace cugraph diff --git a/cpp/include/experimental/graph.hpp b/cpp/include/experimental/graph.hpp new file mode 100644 index 00000000000..ea4a7882363 --- /dev/null +++ b/cpp/include/experimental/graph.hpp @@ -0,0 +1,190 @@ +/* + * Copyright (c) 2020, 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 +#include +#include +#include + +namespace cugraph { +namespace experimental { + +template +struct edgelist_t { + vertex_t const *p_src_vertices{nullptr}; + vertex_t const *p_dst_vertices{nullptr}; + weight_t const *p_edge_weights{nullptr}; + edge_t number_of_edges{0}; +}; + +// graph_t is an owning graph class (note that graph_view_t is a non-owning graph class) +template +class graph_t; + +// multi-GPU version +template +class graph_t> + : public detail::graph_base_t { + public: + using vertex_type = vertex_t; + using edge_type = edge_t; + using weight_type = weight_t; + static constexpr bool is_adj_matrix_transposed = store_transposed; + static constexpr bool is_multi_gpu = multi_gpu; + + graph_t(raft::handle_t const &handle, + std::vector> const &edge_lists, + partition_t const &partition, + vertex_t number_of_vertices, + edge_t number_of_edges, + graph_properties_t properties, + bool sorted_by_global_degree_within_vertex_partition, + bool do_expensive_check = false); + + vertex_t get_number_of_local_vertices() const + { + return partition_.get_vertex_partition_range_last() - + partition_.get_vertex_partition_range_first(); + } + + graph_view_t view() + { + std::vector offsets(adj_matrix_partition_offsets_.size(), nullptr); + std::vector indices(adj_matrix_partition_indices_.size(), nullptr); + std::vector weights(adj_matrix_partition_weights_.size(), nullptr); + for (size_t i = 0; i < offsets.size(); ++i) { + offsets[i] = adj_matrix_partition_offsets_[i].data(); + indices[i] = adj_matrix_partition_indices_[i].data(); + if (weights.size() > 0) { weights[i] = adj_matrix_partition_weights_[i].data(); } + } + + return graph_view_t( + *(this->get_handle_ptr()), + offsets, + indices, + weights, + vertex_partition_segment_offsets_, + partition_, + this->get_number_of_vertices(), + this->get_number_of_edges(), + this->get_graph_properties(), + vertex_partition_segment_offsets_.size() > 0, + false); + } + + private: + std::vector> adj_matrix_partition_offsets_{}; + std::vector> adj_matrix_partition_indices_{}; + std::vector> adj_matrix_partition_weights_{}; + + partition_t partition_{}; + + std::vector + vertex_partition_segment_offsets_{}; // segment offsets within the vertex partition based on + // vertex degree, relevant only if + // sorted_by_global_degree_within_vertex_partition is true +}; + +// single-GPU version +template +class graph_t> + : public detail::graph_base_t { + public: + using vertex_type = vertex_t; + using edge_type = edge_t; + using weight_type = weight_t; + static constexpr bool is_adj_matrix_transposed = store_transposed; + static constexpr bool is_multi_gpu = multi_gpu; + + graph_t(raft::handle_t const &handle, + edgelist_t const &edge_list, + vertex_t number_of_vertices, + graph_properties_t properties, + bool sorted_by_degree, + bool do_expensive_check = false); + + vertex_t get_number_of_local_vertices() const { return this->get_number_of_vertices(); } + + graph_view_t view() + { + return graph_view_t( + *(this->get_handle_ptr()), + offsets_.data(), + indices_.data(), + weights_.data(), + segment_offsets_, + this->get_number_of_vertices(), + this->get_number_of_edges(), + this->get_graph_properties(), + segment_offsets_.size() > 0, + false); + } + + private: + rmm::device_uvector offsets_; + rmm::device_uvector indices_; + rmm::device_uvector weights_; + std::vector segment_offsets_{}; // segment offsets based on vertex degree, relevant + // only if sorted_by_global_degree is true +}; + +template +struct invalid_idx; + +template +struct invalid_idx< + T, + typename std::enable_if_t::value && std::is_signed::value>> + : std::integral_constant { +}; + +template +struct invalid_idx< + T, + typename std::enable_if_t::value && std::is_unsigned::value>> + : std::integral_constant::max()> { +}; + +template +struct invalid_vertex_id : invalid_idx { +}; + +template +struct invalid_edge_id : invalid_idx { +}; + +} // namespace experimental +} // namespace cugraph \ No newline at end of file diff --git a/cpp/include/experimental/graph_view.hpp b/cpp/include/experimental/graph_view.hpp new file mode 100644 index 00000000000..b3b899a5068 --- /dev/null +++ b/cpp/include/experimental/graph_view.hpp @@ -0,0 +1,367 @@ +/* + * Copyright (c) 2020, 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 +#include +#include +#include +#include + +namespace cugraph { +namespace experimental { + +// FIXME: these should better be defined somewhere else. +std::string const comm_p_row_key = "comm_p_row"; +std::string const comm_p_col_key = "comm_p_key"; + +/** + * @brief store vertex partitioning map + * + * Say P = P_row * P_col GPUs. We need to partition 1D vertex arrays (storing per vertex values) and + * the 2D graph adjacency matrix (or transposed 2D graph adjacency matrix) of G. An 1D vertex array + * of size V is divided to P linear partitions; each partition has the size close to V / P. We + * consider two different strategies to partition the 2D matrix: the default strategy and the + * hypergraph partitioning based strategy (the latter is for future extension). + * + * In the default case, one GPU will be responsible for 1 rectangular partition. The matrix will be + * horizontally partitioned first to P_row slabs. Each slab will be further vertically partitioned + * to P_col rectangles. Each rectangular partition will have the size close to V / P_row by V / + * P_col. + * + * To be more specific, a GPU with (row_rank, col_rank) will be responsible for one rectangular + * partition [a,b) by [c,d) where a = vertex_partition_offsets[P_col * row_rank], b = + * vertex_partition_offsets[p_col * (row_rank + 1)], c = vertex_partition_offsets[P_row * col_rank], + * and d = vertex_partition_offsets[p_row * (col_rank + 1)] + * + * In the future, we may apply hyper-graph partitioning to divide V vertices to P groups minimizing + * edge cuts across groups while balancing the number of vertices in each group. We will also + * renumber vertices so the vertices in each group are mapped to consecutive integers. Then, there + * will be more non-zeros in the diagonal partitions of the 2D graph adjacency matrix (or the + * transposed 2D graph adjacency matrix) than the off-diagonal partitions. The default strategy does + * not balance the number of nonzeros if hyper-graph partitioning is applied. To solve this problem, + * the matrix is first horizontally partitioned to P (instead of P_row) slabs, then each slab will + * be further vertically partitioned to P_col rectangles. One GPU will be responsible P_col + * rectangular partitions in this case. + * + * To be more specific, a GPU with (row_rank, col_rank) will be responsible for P_col rectangular + * partitions [a_i,b_i) by [c,d) where a_i = vertex_partition_offsets[P_row * i + row_rank] and b_i + * = vertex_partition_offsets[P_row * i + row_rank + 1]. c and d are same to 1) and i = [0, P_col). + * + * See E. G. Boman et. al., “Scalable matrix computations on large scale-free graphs using 2D graph + * partitioning”, 2013 for additional detail. + * + * @tparam vertex_t Type of vertex ID + */ +template +class partition_t { + public: + partition_t(std::vector const& vertex_partition_offsets, + bool hypergraph_partitioned, + int comm_p_row_size, + int comm_p_col_size, + int comm_p_row_rank, + int comm_p_col_rank) + : vertex_partition_offsets_(vertex_partition_offsets), + hypergraph_partitioned_(hypergraph_partitioned), + comm_p_rank_(comm_p_col_size * comm_p_row_rank + comm_p_col_rank), + comm_p_row_size_(comm_p_row_size), + comm_p_col_size_(comm_p_col_size), + comm_p_row_rank_(comm_p_row_rank), + comm_p_col_rank_(comm_p_col_rank) + { + CUGRAPH_EXPECTS( + vertex_partition_offsets.size() == static_cast(comm_p_row_size * comm_p_col_size), + "Invalid API parameter: erroneous vertex_partition_offsets.size()."); + + CUGRAPH_EXPECTS( + std::is_sorted(vertex_partition_offsets_.begin(), vertex_partition_offsets_.end()), + "Invalid API parameter: partition.vertex_partition_offsets values should be non-descending."); + CUGRAPH_EXPECTS(vertex_partition_offsets_[0] == vertex_t{0}, + "Invalid API parameter: partition.vertex_partition_offsets[0] should be 0."); + } + + std::tuple get_vertex_partition_range() const + { + return std::make_tuple(vertex_partition_offsets_[comm_p_rank_], + vertex_partition_offsets_[comm_p_rank_ + 1]); + } + + vertex_t get_vertex_partition_range_first() const + { + return vertex_partition_offsets_[comm_p_rank_]; + } + + vertex_t get_vertex_partition_range_last() const + { + return vertex_partition_offsets_[comm_p_rank_ + 1]; + } + + std::tuple get_vertex_partition_range(size_t vertex_partition_idx) const + { + return std::make_tuple(vertex_partition_offsets_[vertex_partition_idx], + vertex_partition_offsets_[vertex_partition_idx + 1]); + } + + vertex_t get_vertex_partition_range_first(size_t vertex_partition_idx) const + { + return vertex_partition_offsets_[vertex_partition_idx]; + } + + vertex_t get_vertex_partition_range_last(size_t vertex_partition_idx) const + { + return vertex_partition_offsets_[vertex_partition_idx + 1]; + } + + std::tuple get_matrix_partition_major_range(size_t partition_idx) const + { + auto major_first = + hypergraph_partitioned_ + ? vertex_partition_offsets_[comm_p_row_size_ * partition_idx + comm_p_row_rank_] + : vertex_partition_offsets_[comm_p_row_rank_ * comm_p_col_size_]; + auto major_last = + hypergraph_partitioned_ + ? vertex_partition_offsets_[comm_p_row_size_ * partition_idx + comm_p_row_rank_ + 1] + : vertex_partition_offsets_[(comm_p_row_rank_ + 1) * comm_p_col_size_]; + + return std::make_tuple(major_first, major_last); + } + + std::tuple get_matrix_partition_minor_range() const + { + auto minor_first = vertex_partition_offsets_[comm_p_col_rank_ * comm_p_row_size_]; + auto minor_last = vertex_partition_offsets_[(comm_p_col_rank_ + 1) * comm_p_row_size_]; + + return std::make_tuple(minor_first, minor_last); + } + + bool is_hypergraph_partitioned() const { return hypergraph_partitioned_; } + + private: + std::vector vertex_partition_offsets_{}; // size = P + 1 + bool hypergraph_partitioned_{false}; + + int comm_p_rank_{0}; + int comm_p_row_size_{0}; + int comm_p_col_size_{0}; + int comm_p_row_rank_{0}; + int comm_p_col_rank_{0}; +}; + +struct graph_properties_t { + bool is_symmetric{false}; + bool is_multigraph{false}; +}; + +namespace detail { + +// FIXME: threshold values require tuning +size_t constexpr low_degree_threshold{raft::warp_size()}; +size_t constexpr mid_degree_threshold{1024}; +size_t constexpr num_segments_per_vertex_partition{3}; + +// Common for both graph_view_t & graph_t and both single-GPU & multi-GPU versions +template +class graph_base_t { + public: + graph_base_t(raft::handle_t const& handle, + vertex_t number_of_vertices, + edge_t number_of_edges, + graph_properties_t properties) + : handle_ptr_(&handle), + number_of_vertices_(number_of_vertices), + number_of_edges_(number_of_edges), + properties_(properties){}; + + vertex_t get_number_of_vertices() const { return number_of_vertices_; } + edge_t get_number_of_edges() const { return number_of_edges_; } + + bool is_symmetric() const { return properties_.is_symmetric; } + bool is_multigraph() const { return properties_.is_multigraph; } + + protected: + raft::handle_t const* get_handle_ptr() const { return handle_ptr_; }; + graph_properties_t get_graph_properties() const { return properties_; } + + private: + raft::handle_t const* handle_ptr_{nullptr}; + + vertex_t number_of_vertices_{0}; + edge_t number_of_edges_{0}; + + graph_properties_t properties_{}; +}; + +} // namespace detail + +// graph_view_t is a non-owning graph class (note that graph_t is an owning graph class) +template +class graph_view_t; + +// multi-GPU version +template +class graph_view_t> + : public detail::graph_base_t { + public: + using vertex_type = vertex_t; + using edge_type = edge_t; + using weight_type = weight_t; + static constexpr bool is_adj_matrix_transposed = store_transposed; + static constexpr bool is_multi_gpu = multi_gpu; + + graph_view_t(raft::handle_t const& handle, + std::vector const& adj_matrix_partition_offsets, + std::vector const& adj_matrix_partition_indices, + std::vector const& adj_matrix_partition_weights, + std::vector const& vertex_partition_segment_offsets, + partition_t const& partition, + vertex_t number_of_vertices, + edge_t number_of_edges, + graph_properties_t properties, + bool sorted_by_global_degree_within_vertex_partition, + bool do_expensive_check = false); + + vertex_t get_number_of_local_vertices() const + { + return partition_.get_vertex_partition_range_last() - + partition_.get_vertex_partition_range_first(); + } + + size_t get_number_of_adj_matrix_partitions() { return adj_matrix_partition_offsets_.size(); } + + // FIXME: this function is not part of the public stable API.This function is mainly for pattern + // accelerator implementation. This function is currently public to support the legacy + // implementations directly accessing CSR/CSC data, but this function will eventually become + // private or even disappear if we switch to CSR + DCSR (or CSC + DCSC). + edge_t const* offsets(size_t adj_matrix_partition_idx) const + { + return adj_matrix_partition_offsets_[adj_matrix_partition_idx]; + } + + // FIXME: this function is not part of the public stable API.This function is mainly for pattern + // accelerator implementation. This function is currently public to support the legacy + // implementations directly accessing CSR/CSC data, but this function will eventually become + // private or even disappear if we switch to CSR + DCSR (or CSC + DCSC). + vertex_t const* indices(size_t adj_matrix_partition_idx) const + { + return adj_matrix_partition_indices_[adj_matrix_partition_idx]; + } + + // FIXME: this function is not part of the public stable API.This function is mainly for pattern + // accelerator implementation. This function is currently public to support the legacy + // implementations directly accessing CSR/CSC data, but this function will eventually become + // private or even disappear if we switch to CSR + DCSR (or CSC + DCSC). + weight_t const* weights(size_t adj_matrix_partition_idx) const + { + return adj_matrix_partition_weights_.size() > 0 + ? adj_matrix_partition_weights_[adj_matrix_partition_idx] + : static_cast(nullptr); + } + + private: + std::vector adj_matrix_partition_offsets_{}; + std::vector adj_matrix_partition_indices_{}; + std::vector adj_matrix_partition_weights_{}; + + partition_t partition_{}; + + std::vector + vertex_partition_segment_offsets_{}; // segment offsets within the vertex partition based on + // vertex degree, relevant only if + // sorted_by_global_degree_within_vertex_partition is true +}; + +// single-GPU version +template +class graph_view_t> + : public detail::graph_base_t { + public: + using vertex_type = vertex_t; + using edge_type = edge_t; + using weight_type = weight_t; + static constexpr bool is_adj_matrix_transposed = store_transposed; + static constexpr bool is_multi_gpu = multi_gpu; + + graph_view_t(raft::handle_t const& handle, + edge_t const* offsets, + vertex_t const* indices, + weight_t const* weights, + std::vector const& segment_offsets, + vertex_t number_of_vertices, + edge_t number_of_edges, + graph_properties_t properties, + bool sorted_by_degree, + bool do_expensive_check = false); + + vertex_t get_number_of_local_vertices() const { return this->get_number_of_vertices(); } + + // FIXME: this function is not part of the public stable API.This function is mainly for pattern + // accelerator implementation. This function is currently public to support the legacy + // implementations directly accessing CSR/CSC data, but this function will eventually become + // private. + edge_t const* offsets() const { return offsets_; } + + // FIXME: this function is not part of the public stable API.This function is mainly for pattern + // accelerator implementation. This function is currently public to support the legacy + // implementations directly accessing CSR/CSC data, but this function will eventually become + // private. + vertex_t const* indices() const { return indices_; } + + // FIXME: this function is not part of the public stable API.This function is mainly for pattern + // accelerator implementation. This function is currently public to support the legacy + // implementations directly accessing CSR/CSC data, but this function will eventually become + // private. + weight_t const* weights() const { return weights_; } + + private: + edge_t const* offsets_{nullptr}; + vertex_t const* indices_{nullptr}; + weight_t const* weights_{nullptr}; + std::vector segment_offsets_{}; // segment offsets based on vertex degree, relevant + // only if sorted_by_global_degree is true +}; + +} // namespace experimental +} // namespace cugraph \ No newline at end of file diff --git a/cpp/src/experimental/graph.cu b/cpp/src/experimental/graph.cu new file mode 100644 index 00000000000..7b7625fd911 --- /dev/null +++ b/cpp/src/experimental/graph.cu @@ -0,0 +1,539 @@ +/* + * Copyright (c) 2020, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include + +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +namespace cugraph { +namespace experimental { + +namespace { + +// can't use lambda due to nvcc limitations (The enclosing parent function ("graph_view_t") for an +// extended __device__ lambda must allow its address to be taken) +template +struct out_of_range_t { + vertex_t major_first{}; + vertex_t major_last{}; + vertex_t minor_first{}; + vertex_t minor_last{}; + + __device__ bool operator()(thrust::tuple t) + { + auto major = thrust::get<0>(t); + auto minor = thrust::get<1>(t); + return (major < major_first) || (major >= major_last) || (minor < minor_first) || + (minor >= minor_last); + } +}; + +template +std:: + tuple, rmm::device_uvector, rmm::device_uvector> + edge_list_to_compressed_sparse(raft::handle_t const &handle, + edgelist_t const &edgelist, + vertex_t major_first, + vertex_t major_last, + vertex_t minor_first, + vertex_t minor_last) +{ + rmm::device_uvector offsets((major_last - major_first) + 1, handle.get_stream()); + rmm::device_uvector indices(edgelist.number_of_edges, handle.get_stream()); + rmm::device_uvector weights( + edgelist.p_edge_weights != nullptr ? edgelist.number_of_edges : 0, handle.get_stream()); + thrust::fill(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + offsets.begin(), + offsets.end(), + edge_t{0}); + thrust::fill(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + indices.begin(), + indices.end(), + vertex_t{0}); + + // FIXME: need to performance test this code with R-mat graphs having highly-skewed degree + // distribution. If there is a small number of vertices with very large degrees, atomicAdd can + // sequentialize execution. CUDA9+ & Kepler+ provide complier/architectural optimizations to + // mitigate this impact + // (https://developer.nvidia.com/blog/cuda-pro-tip-optimized-filtering-warp-aggregated-atomics/), + // and we need to check this thrust::for_each based approach delivers the expected performance. + + // FIXME: also need to verify this approach is at least not significantly slower than the sorting + // based approach (this approach does not use extra memory, so better stick to this approach + // unless performance is significantly worse). + + auto p_offsets = offsets.data(); + auto p_indices = indices.data(); + auto p_weights = + edgelist.p_edge_weights != nullptr ? weights.data() : static_cast(nullptr); + + thrust::for_each(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + store_transposed ? edgelist.p_dst_vertices : edgelist.p_src_vertices, + store_transposed ? edgelist.p_dst_vertices + edgelist.number_of_edges + : edgelist.p_src_vertices + edgelist.number_of_edges, + [p_offsets, major_first] __device__(auto v) { + atomicAdd(p_offsets + (v - major_first), edge_t{1}); + }); + + thrust::exclusive_scan(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + offsets.begin(), + offsets.end(), + offsets.begin()); + + if (edgelist.p_edge_weights != nullptr) { + auto edge_first = thrust::make_zip_iterator(thrust::make_tuple( + edgelist.p_src_vertices, edgelist.p_dst_vertices, edgelist.p_edge_weights)); + thrust::for_each(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + edge_first, + edge_first + edgelist.number_of_edges, + [p_offsets, p_indices, p_weights, major_first] __device__(auto e) { + auto s = thrust::get<0>(e); + auto d = thrust::get<1>(e); + auto w = thrust::get<2>(e); + auto major = store_transposed ? d : s; + auto minor = store_transposed ? s : d; + auto start = p_offsets[major - major_first]; + auto degree = p_offsets[(major - major_first) + 1] - start; + auto idx = atomicAdd(p_indices + (start + degree - 1), + vertex_t{1}); // use the last element as a counter + // FIXME: we can actually store minor - minor_first instead of minor to save + // memory if minor can be larger than 32 bit but minor - minor_first fits + // within 32 bit + p_indices[start + idx] = + minor; // overwrite the counter only if idx == degree - 1 (no race) + p_weights[start + idx] = w; + }); + } else { + auto edge_first = thrust::make_zip_iterator( + thrust::make_tuple(edgelist.p_src_vertices, edgelist.p_dst_vertices)); + thrust::for_each(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + edge_first, + edge_first + edgelist.number_of_edges, + [p_offsets, p_indices, p_weights, major_first] __device__(auto e) { + auto s = thrust::get<0>(e); + auto d = thrust::get<1>(e); + auto major = store_transposed ? d : s; + auto minor = store_transposed ? s : d; + auto start = p_offsets[major - major_first]; + auto degree = p_offsets[(major - major_first) + 1] - start; + auto idx = atomicAdd(p_indices + (start + degree - 1), + vertex_t{1}); // use the last element as a counter + // FIXME: we can actually store minor - minor_first instead of minor to save + // memory if minor can be larger than 32 bit but minor - minor_first fits + // within 32 bit + p_indices[start + idx] = + minor; // overwrite the counter only if idx == degree - 1 (no race) + }); + } + + // FIXME: need to add an option to sort neighbor lists + + return std::make_tuple(std::move(offsets), std::move(indices), std::move(weights)); +} + +template +std::vector segment_degree_sorted_vertex_partition(raft::handle_t const &handle, + DegreeIterator degree_first, + DegreeIterator degree_last, + ThresholdIterator threshold_first, + ThresholdIterator threshold_last) +{ + auto num_elements = thrust::distance(degree_first, degree_last); + auto num_segments = thrust::distance(threshold_first, threshold_last) + 1; + + std::vector h_segment_offsets(num_segments + 1); + h_segment_offsets[0] = 0; + h_segment_offsets.back() = num_elements; + + rmm::device_uvector d_segment_offsets(num_segments - 1, handle.get_stream()); + + thrust::upper_bound(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + degree_first, + degree_last, + threshold_first, + threshold_last, + d_segment_offsets.begin()); + + raft::update_host(h_segment_offsets.begin() + 1, + d_segment_offsets.begin(), + d_segment_offsets.size(), + handle.get_stream()); + + CUDA_TRY(cudaStreamSynchronize( + handle.get_stream())); // this is necessary as d_segment_offsets will become out-of-scope once + // this functions and returning a host variable which can be used right + // after return. + + return h_segment_offsets; +} + +} // namespace + +template +graph_t>:: + graph_t(raft::handle_t const &handle, + std::vector> const &edgelists, + partition_t const &partition, + vertex_t number_of_vertices, + edge_t number_of_edges, + graph_properties_t properties, + bool sorted_by_global_degree_within_vertex_partition, + bool do_expensive_check) + : detail::graph_base_t( + handle, number_of_vertices, number_of_edges, properties), + partition_(partition) +{ + // cheap error checks + + auto &comm_p = this->get_handle_ptr()->get_comms(); + auto const comm_p_size = comm_p.get_size(); + auto &comm_p_row = this->get_handle_ptr()->get_subcomm(comm_p_row_key); + auto const comm_p_row_rank = comm_p_row.get_rank(); + auto const comm_p_row_size = comm_p_row.get_size(); + auto &comm_p_col = this->get_handle_ptr()->get_subcomm(comm_p_col_key); + auto const comm_p_col_rank = comm_p_col.get_rank(); + auto const comm_p_col_size = comm_p_col.get_size(); + auto default_stream = this->get_handle_ptr()->get_stream(); + + CUGRAPH_EXPECTS(edgelists.size() > 0, + "Invalid API parameter: edgelists.size() should be non-zero."); + + bool is_weighted = edgelists[0].p_edge_weights != nullptr; + + CUGRAPH_EXPECTS( + std::any_of(edgelists.begin() + 1, + edgelists.end(), + [is_weighted](auto edgelist) { + return (edgelist.p_src_vertices == nullptr) || + (edgelist.p_dst_vertices == nullptr) || + (is_weighted && (edgelist.p_edge_weights == nullptr)) || + (!is_weighted && (edgelist.p_edge_weights != nullptr)); + }) == false, + "Invalid API parameter: edgelists[].p_src_vertices and edgelists[].p_dst_vertices should not " + "be nullptr and edgelists[].p_edge_weights should be nullptr (if edgelists[0].p_edge_weights " + "is nullptr) or should not be nullptr (otherwise)."); + + CUGRAPH_EXPECTS((partition.is_hypergraph_partitioned() && + (edgelists.size() == static_cast(comm_p_row_size))) || + (!(partition.is_hypergraph_partitioned()) && (edgelists.size() == 1)), + "Invalid API parameter: errneous edgelists.size()."); + + // optional expensive checks (part 1/3) + + if (do_expensive_check) { + edge_t number_of_local_edges_sum{}; + for (size_t i = 0; i < edgelists.size(); ++i) { + vertex_t major_first{}; + vertex_t major_last{}; + vertex_t minor_first{}; + vertex_t minor_last{}; + std::tie(major_first, major_last) = partition.get_matrix_partition_major_range(i); + std::tie(minor_first, minor_last) = partition.get_matrix_partition_minor_range(); + + number_of_local_edges_sum += edgelists[i].number_of_edges; + + auto edge_first = thrust::make_zip_iterator(thrust::make_tuple( + store_transposed ? edgelists[i].p_dst_vertices : edgelists[i].p_src_vertices, + store_transposed ? edgelists[i].p_src_vertices : edgelists[i].p_dst_vertices)); + // better use thrust::any_of once https://github.com/thrust/thrust/issues/1016 is resolved + CUGRAPH_EXPECTS(thrust::count_if(rmm::exec_policy(default_stream)->on(default_stream), + edge_first, + edge_first + edgelists[i].number_of_edges, + out_of_range_t{ + major_first, major_last, minor_first, minor_last}) == 0, + "Invalid API parameter: edgelists[] have out-of-range values."); + } + this->get_handle_ptr()->get_comms().allreduce(&number_of_local_edges_sum, + &number_of_local_edges_sum, + 1, + raft::comms::op_t::SUM, + default_stream); + CUGRAPH_EXPECTS(number_of_local_edges_sum == this->get_number_of_edges(), + "Invalid API parameter: the sum of local edges doe counts not match with " + "number_of_local_edges."); + + CUGRAPH_EXPECTS( + partition.get_vertex_partition_range_last(comm_p_size - 1) == number_of_vertices, + "Invalid API parameter: vertex partition should cover [0, number_of_vertices)."); + } + + // convert edge list (COO) to compressed sparse format (CSR or CSC) + + adj_matrix_partition_offsets_.reserve(edgelists.size()); + adj_matrix_partition_indices_.reserve(edgelists.size()); + adj_matrix_partition_weights_.reserve(is_weighted ? edgelists.size() : 0); + for (size_t i = 0; i < edgelists.size(); ++i) { + vertex_t major_first{}; + vertex_t major_last{}; + vertex_t minor_first{}; + vertex_t minor_last{}; + std::tie(major_first, major_last) = partition.get_matrix_partition_major_range(i); + std::tie(minor_first, minor_last) = partition.get_matrix_partition_minor_range(); + + rmm::device_uvector offsets(0, default_stream); + rmm::device_uvector indices(0, default_stream); + rmm::device_uvector weights(0, default_stream); + std::tie(offsets, indices, weights) = edge_list_to_compressed_sparse( + *(this->get_handle_ptr()), edgelists[i], major_first, major_last, minor_first, minor_last); + adj_matrix_partition_offsets_.push_back(std::move(offsets)); + adj_matrix_partition_indices_.push_back(std::move(indices)); + if (adj_matrix_partition_weights_.size() > 0) { + adj_matrix_partition_weights_.push_back(std::move(weights)); + } + } + + // update degree-based segment offsets (to be used for graph analytics kernel optimization) + + if (sorted_by_global_degree_within_vertex_partition) { + auto degrees = detail::compute_major_degree( + *(this->get_handle_ptr()), adj_matrix_partition_offsets_, partition_); + + // optional expensive checks (part 2/3) + + if (do_expensive_check) { + CUGRAPH_EXPECTS(thrust::is_sorted(rmm::exec_policy(default_stream)->on(default_stream), + degrees.begin(), + degrees.end(), + thrust::greater{}), + "Invalid API parameter: sorted_by_global_degree_within_vertex_partition is " + "set to true, but degrees are not non-ascending."); + } + + static_assert(detail::num_segments_per_vertex_partition == 3); + static_assert((detail::low_degree_threshold <= detail::mid_degree_threshold) && + (detail::mid_degree_threshold <= std::numeric_limits::max())); + rmm::device_uvector d_thresholds(detail::num_segments_per_vertex_partition - 1, + default_stream); + std::vector h_thresholds = {static_cast(detail::low_degree_threshold), + static_cast(detail::mid_degree_threshold)}; + raft::update_device( + d_thresholds.data(), h_thresholds.data(), h_thresholds.size(), default_stream); + + rmm::device_uvector segment_offsets(detail::num_segments_per_vertex_partition + 1, + default_stream); + segment_offsets.set_element_async(0, 0, default_stream); + segment_offsets.set_element_async( + detail::num_segments_per_vertex_partition, degrees.size(), default_stream); + + thrust::upper_bound(rmm::exec_policy(default_stream)->on(default_stream), + degrees.begin(), + degrees.end(), + d_thresholds.begin(), + d_thresholds.end(), + segment_offsets.begin() + 1); + + rmm::device_uvector aggregate_segment_offsets( + comm_p_row_size * segment_offsets.size(), default_stream); + comm_p_row.allgather(segment_offsets.data(), + aggregate_segment_offsets.data(), + segment_offsets.size(), + default_stream); + + vertex_partition_segment_offsets_.resize(comm_p_row_size * (segment_offsets.size())); + raft::update_host(vertex_partition_segment_offsets_.data(), + aggregate_segment_offsets.data(), + aggregate_segment_offsets.size(), + default_stream); + + auto status = handle.get_comms().sync_stream( + default_stream); // this is necessary as degrees, d_thresholds, and segment_offsets will + // become out-of-scope once control flow exits this block and + // vertex_partition_segment_offsets_ can be used right after return. + CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); + } + + // optional expensive checks (part 3/3) + + if (do_expensive_check) { + // FIXME: check for symmetricity may better be implemetned with transpose(). + if (this->is_symmetric()) {} + // FIXME: check for duplicate edges may better be implemented after deciding whether to sort + // neighbor list or not. + if (!this->is_multigraph()) {} + } +} + +template +graph_t>:: + graph_t(raft::handle_t const &handle, + edgelist_t const &edgelist, + vertex_t number_of_vertices, + graph_properties_t properties, + bool sorted_by_degree, + bool do_expensive_check) + : detail::graph_base_t( + handle, number_of_vertices, edgelist.number_of_edges, properties), + offsets_(rmm::device_uvector(0, handle.get_stream())), + indices_(rmm::device_uvector(0, handle.get_stream())), + weights_(rmm::device_uvector(0, handle.get_stream())) +{ + // cheap error checks + + auto default_stream = this->get_handle_ptr()->get_stream(); + + CUGRAPH_EXPECTS( + (edgelist.p_src_vertices != nullptr) && (edgelist.p_dst_vertices != nullptr), + "Invalid API parameter: edgelist.p_src_vertices and edgelist.p_dst_vertices should " + "not be nullptr."); + + // optional expensive checks (part 1/2) + + if (do_expensive_check) { + auto edge_first = thrust::make_zip_iterator( + thrust::make_tuple(store_transposed ? edgelist.p_dst_vertices : edgelist.p_src_vertices, + store_transposed ? edgelist.p_src_vertices : edgelist.p_dst_vertices)); + // better use thrust::any_of once https://github.com/thrust/thrust/issues/1016 is resolved + CUGRAPH_EXPECTS(thrust::count_if( + rmm::exec_policy(default_stream)->on(default_stream), + edge_first, + edge_first + edgelist.number_of_edges, + out_of_range_t{ + 0, this->get_number_of_vertices(), 0, this->get_number_of_vertices()}) == 0, + "Invalid API parameter: edgelist have out-of-range values."); + + // FIXME: check for symmetricity may better be implemetned with transpose(). + if (this->is_symmetric()) {} + // FIXME: check for duplicate edges may better be implemented after deciding whether to sort + // neighbor list or not. + if (!this->is_multigraph()) {} + } + + // convert edge list (COO) to compressed sparse format (CSR or CSC) + + std::tie(offsets_, indices_, weights_) = + edge_list_to_compressed_sparse(*(this->get_handle_ptr()), + edgelist, + vertex_t{0}, + this->get_number_of_vertices(), + vertex_t{0}, + this->get_number_of_vertices()); + + // update degree-based segment offsets (to be used for graph analytics kernel optimization) + + if (sorted_by_degree) { + auto degree_first = thrust::make_transform_iterator( + thrust::make_counting_iterator(vertex_t{0}), + detail::degree_from_offsets_t{offsets_.data()}); + + // optional expensive checks (part 2/2) + + if (do_expensive_check) { + CUGRAPH_EXPECTS(thrust::is_sorted(rmm::exec_policy(default_stream)->on(default_stream), + degree_first, + degree_first + this->get_number_of_vertices(), + thrust::greater{}), + "Invalid API parameter: sorted_by_degree is set to true, but degrees are not " + "non-ascending."); + } + + static_assert(detail::num_segments_per_vertex_partition == 3); + static_assert((detail::low_degree_threshold <= detail::mid_degree_threshold) && + (detail::mid_degree_threshold <= std::numeric_limits::max())); + rmm::device_uvector d_thresholds(detail::num_segments_per_vertex_partition - 1, + default_stream); + std::vector h_thresholds = {static_cast(detail::low_degree_threshold), + static_cast(detail::mid_degree_threshold)}; + raft::update_device( + d_thresholds.data(), h_thresholds.data(), h_thresholds.size(), default_stream); + + rmm::device_uvector segment_offsets(detail::num_segments_per_vertex_partition + 1, + default_stream); + segment_offsets.set_element_async(0, 0, default_stream); + segment_offsets.set_element_async( + detail::num_segments_per_vertex_partition, this->get_number_of_vertices(), default_stream); + + thrust::upper_bound(rmm::exec_policy(default_stream)->on(default_stream), + degree_first, + degree_first + this->get_number_of_vertices(), + d_thresholds.begin(), + d_thresholds.end(), + segment_offsets.begin() + 1); + + segment_offsets_.resize(segment_offsets.size()); + raft::update_host( + segment_offsets_.data(), segment_offsets.data(), segment_offsets.size(), default_stream); + + CUDA_TRY(cudaStreamSynchronize( + default_stream)); // this is necessary as d_thresholds and segment_offsets will become + // out-of-scpe once control flow exits this block and segment_offsets_ can + // be used right after return. + } + + // optional expensive checks (part 3/3) + + if (do_expensive_check) { + // FIXME: check for symmetricity may better be implemetned with transpose(). + if (this->is_symmetric()) {} + // FIXME: check for duplicate edges may better be implemented after deciding whether to sort + // neighbor list or not. + if (!this->is_multigraph()) {} + } +} + +// explicit instantiation + +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; + +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; +template class graph_t; + +} // namespace experimental +} // namespace cugraph diff --git a/cpp/src/experimental/graph_view.cu b/cpp/src/experimental/graph_view.cu new file mode 100644 index 00000000000..b297a825a01 --- /dev/null +++ b/cpp/src/experimental/graph_view.cu @@ -0,0 +1,304 @@ +/* + * Copyright (c) 2020, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include + +#include +#include +#include + +#include +#include +#include +#include + +#include +#include +#include +#include + +namespace cugraph { +namespace experimental { + +namespace { + +// can't use lambda due to nvcc limitations (The enclosing parent function ("graph_view_t") for an +// extended __device__ lambda must allow its address to be taken) +template +struct out_of_range_t { + vertex_t min{}; + vertex_t max{}; + + __device__ bool operator()(vertex_t v) { return (v < min) || (v >= max); } +}; + +} // namespace + +template +graph_view_t>:: + graph_view_t(raft::handle_t const& handle, + std::vector const& adj_matrix_partition_offsets, + std::vector const& adj_matrix_partition_indices, + std::vector const& adj_matrix_partition_weights, + std::vector const& vertex_partition_segment_offsets, + partition_t const& partition, + vertex_t number_of_vertices, + edge_t number_of_edges, + graph_properties_t properties, + bool sorted_by_global_degree_within_vertex_partition, + bool do_expensive_check) + : detail::graph_base_t( + handle, number_of_vertices, number_of_edges, properties), + adj_matrix_partition_offsets_(adj_matrix_partition_offsets), + adj_matrix_partition_indices_(adj_matrix_partition_indices), + adj_matrix_partition_weights_(adj_matrix_partition_weights), + partition_(partition), + vertex_partition_segment_offsets_(vertex_partition_segment_offsets) +{ + // cheap error checks + + auto const comm_p_size = this->get_handle_ptr()->get_comms().get_size(); + auto const comm_p_row_size = this->get_handle_ptr()->get_subcomm(comm_p_row_key).get_size(); + auto const comm_p_col_size = this->get_handle_ptr()->get_subcomm(comm_p_col_key).get_size(); + + CUGRAPH_EXPECTS(adj_matrix_partition_offsets.size() == adj_matrix_partition_indices.size(), + "Invalid API parameter: adj_matrix_partition_offsets.size() and " + "adj_matrix_partition_indices.size() should coincide."); + CUGRAPH_EXPECTS((adj_matrix_partition_weights.size() == adj_matrix_partition_offsets.size()) || + (adj_matrix_partition_weights.size() == 0), + "Invalid API parameter: adj_matrix_partition_weights.size() should coincide with " + "adj_matrix_partition_offsets.size() (if weighted) or 0 (if unweighted)."); + + CUGRAPH_EXPECTS( + (partition.is_hypergraph_partitioned() && + (adj_matrix_partition_offsets.size() == static_cast(comm_p_row_size))) || + (!(partition.is_hypergraph_partitioned()) && (adj_matrix_partition_offsets.size() == 1)), + "Invalid API parameter: errneous adj_matrix_partition_offsets.size()."); + + CUGRAPH_EXPECTS((sorted_by_global_degree_within_vertex_partition && + (vertex_partition_segment_offsets.size() == + comm_p_col_size * (detail::num_segments_per_vertex_partition + 1))) || + (!sorted_by_global_degree_within_vertex_partition && + (vertex_partition_segment_offsets.size() == 0)), + "Invalid API parameter: vertex_partition_segment_offsets.size() does not match " + "with sorted_by_global_degree_within_vertex_partition."); + + // optional expensive checks + + if (do_expensive_check) { + auto default_stream = this->get_handle_ptr()->get_stream(); + + auto const comm_p_row_rank = this->get_handle_ptr()->get_subcomm(comm_p_row_key).get_rank(); + auto const comm_p_col_rank = this->get_handle_ptr()->get_subcomm(comm_p_col_key).get_rank(); + + edge_t number_of_local_edges_sum{}; + for (size_t i = 0; i < adj_matrix_partition_offsets.size(); ++i) { + vertex_t major_first{}; + vertex_t major_last{}; + vertex_t minor_first{}; + vertex_t minor_last{}; + std::tie(major_first, major_last) = partition.get_matrix_partition_major_range(i); + std::tie(minor_first, minor_last) = partition.get_matrix_partition_minor_range(); + CUGRAPH_EXPECTS( + thrust::is_sorted(rmm::exec_policy(default_stream)->on(default_stream), + adj_matrix_partition_offsets[i], + adj_matrix_partition_offsets[i] + (major_last - major_first + 1)), + "Invalid API parameter: adj_matrix_partition_offsets[] is not sorted."); + edge_t number_of_local_edges{}; + raft::update_host(&number_of_local_edges, + adj_matrix_partition_offsets[i] + (major_last - major_first), + 1, + default_stream); + CUDA_TRY(cudaStreamSynchronize(default_stream)); + number_of_local_edges_sum += number_of_local_edges; + + // better use thrust::any_of once https://github.com/thrust/thrust/issues/1016 is resolved + CUGRAPH_EXPECTS( + thrust::count_if(rmm::exec_policy(default_stream)->on(default_stream), + adj_matrix_partition_indices[i], + adj_matrix_partition_indices[i] + number_of_local_edges, + out_of_range_t{minor_first, minor_last}) == 0, + "Invalid API parameter: adj_matrix_partition_indices[] have out-of-range vertex IDs."); + } + this->get_handle_ptr()->get_comms().allreduce(&number_of_local_edges_sum, + &number_of_local_edges_sum, + 1, + raft::comms::op_t::SUM, + default_stream); + auto status = handle.get_comms().sync_stream(default_stream); + CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); + CUGRAPH_EXPECTS(number_of_local_edges_sum == this->get_number_of_edges(), + "Invalid API parameter: the sum of local edges doe counts not match with " + "number_of_local_edges."); + + if (sorted_by_global_degree_within_vertex_partition) { + auto degrees = detail::compute_major_degree(handle, adj_matrix_partition_offsets, partition); + CUGRAPH_EXPECTS(thrust::is_sorted(rmm::exec_policy(default_stream)->on(default_stream), + degrees.begin(), + degrees.end(), + thrust::greater{}), + "Invalid API parameter: sorted_by_global_degree_within_vertex_partition is " + "set to true, but degrees are not non-ascending."); + + for (int i = 0; i < comm_p_col_size; ++i) { + CUGRAPH_EXPECTS(std::is_sorted(vertex_partition_segment_offsets.begin() + + (detail::num_segments_per_vertex_partition + 1) * i, + vertex_partition_segment_offsets.begin() + + (detail::num_segments_per_vertex_partition + 1) * (i + 1)), + "Invalid API parameter: erroneous vertex_partition_segment_offsets."); + CUGRAPH_EXPECTS( + vertex_partition_segment_offsets[(detail::num_segments_per_vertex_partition + 1) * i] == + 0, + "Invalid API parameter: erroneous vertex_partition_segment_offsets."); + auto vertex_partition_idx = partition.is_hypergraph_partitioned() + ? comm_p_row_size * i + comm_p_row_rank + : comm_p_col_size * comm_p_row_rank + i; + CUGRAPH_EXPECTS( + vertex_partition_segment_offsets[(detail::num_segments_per_vertex_partition + 1) * i + + detail::num_segments_per_vertex_partition] == + partition.get_vertex_partition_range_first(vertex_partition_idx), + "Invalid API parameter: erroneous vertex_partition_segment_offsets."); + } + } + + CUGRAPH_EXPECTS( + partition.get_vertex_partition_range_last(comm_p_size - 1) == number_of_vertices, + "Invalid API parameter: vertex partition should cover [0, number_of_vertices)."); + + // FIXME: check for symmetricity may better be implemetned with transpose(). + if (this->is_symmetric()) {} + // FIXME: check for duplicate edges may better be implemented after deciding whether to sort + // neighbor list or not. + if (!this->is_multigraph()) {} + } +} + +template +graph_view_t>::graph_view_t(raft::handle_t const& handle, + edge_t const* offsets, + vertex_t const* indices, + weight_t const* weights, + std::vector const& + segment_offsets, + vertex_t number_of_vertices, + edge_t number_of_edges, + graph_properties_t properties, + bool sorted_by_degree, + bool do_expensive_check) + : detail::graph_base_t( + handle, number_of_vertices, number_of_edges, properties), + offsets_(offsets), + indices_(indices), + weights_(weights), + segment_offsets_(segment_offsets) +{ + // cheap error checks + + CUGRAPH_EXPECTS( + (sorted_by_degree && + (segment_offsets.size() == (detail::num_segments_per_vertex_partition + 1))) || + (!sorted_by_degree && (segment_offsets.size() == 0)), + "Invalid API parameter: segment_offsets.size() does not match with sorted_by_degree."); + + // optional expensive checks + + if (do_expensive_check) { + auto default_stream = this->get_handle_ptr()->get_stream(); + + CUGRAPH_EXPECTS(thrust::is_sorted(rmm::exec_policy(default_stream)->on(default_stream), + offsets, + offsets + (this->get_number_of_vertices() + 1)), + "Invalid API parameter: offsets is not sorted."); + + // better use thrust::any_of once https://github.com/thrust/thrust/issues/1016 is resolved + CUGRAPH_EXPECTS( + thrust::count_if(rmm::exec_policy(default_stream)->on(default_stream), + indices, + indices + this->get_number_of_edges(), + out_of_range_t{0, this->get_number_of_vertices()}) == 0, + "Invalid API parameter: adj_matrix_partition_indices[] have out-of-range vertex IDs."); + + if (sorted_by_degree) { + auto degree_first = + thrust::make_transform_iterator(thrust::make_counting_iterator(vertex_t{0}), + detail::degree_from_offsets_t{offsets}); + CUGRAPH_EXPECTS(thrust::is_sorted(rmm::exec_policy(default_stream)->on(default_stream), + degree_first, + degree_first + this->get_number_of_vertices(), + thrust::greater{}), + "Invalid API parameter: sorted_by_degree is set to true, but degrees are not " + "non-ascending."); + + CUGRAPH_EXPECTS(std::is_sorted(segment_offsets.begin(), segment_offsets.end()), + "Invalid API parameter: erroneous segment_offsets."); + CUGRAPH_EXPECTS(segment_offsets[0] == 0, "Invalid API parameter: segment_offsets."); + CUGRAPH_EXPECTS(segment_offsets.back() == this->get_number_of_vertices(), + "Invalid API parameter: segment_offsets."); + } + + // FIXME: check for symmetricity may better be implemetned with transpose(). + if (this->is_symmetric()) {} + // FIXME: check for duplicate edges may better be implemented after deciding whether to sort + // neighbor list or not. + if (!this->is_multigraph()) {} + } +} + +// explicit instantiation + +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; + +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; +template class graph_view_t; + +} // namespace experimental +} // namespace cugraph diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index e0f945639ca..eb10790f328 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -254,19 +254,13 @@ set(FIND_MATCHES_TEST_SRC ConfigureTest(FIND_MATCHES_TEST "${FIND_MATCHES_TEST_SRC}" "") ################################################################################################### -#-NCCL tests --------------------------------------------------------------------- +# - Experimental Graph tests ---------------------------------------------------------------------- -if (BUILD_MPI) - set(NCCL_TEST_SRC - "${CMAKE_CURRENT_SOURCE_DIR}/nccl/nccl_test.cu") - - ConfigureTest(NCCL_TEST "${NCCL_TEST_SRC}" "") - - set(NCCL_DEGREE_TEST_SRC - "${CMAKE_CURRENT_SOURCE_DIR}/nccl/degree_test.cu") +set(EXPERIMENTAL_GRAPH_TEST_SRCS + "${CMAKE_SOURCE_DIR}/../thirdparty/mmio/mmio.c" + "${CMAKE_CURRENT_SOURCE_DIR}/experimental/graph_test.cpp") - ConfigureTest(NCCL_DEGREE_TEST "${NCCL_DEGREE_TEST_SRC}" "") -endif(BUILD_MPI) +ConfigureTest(EXPERIMENTAL_GRAPH_TEST "${EXPERIMENTAL_GRAPH_TEST_SRCS}" "") ################################################################################################### ### enable testing ################################################################################ diff --git a/cpp/tests/experimental/graph_test.cpp b/cpp/tests/experimental/graph_test.cpp new file mode 100644 index 00000000000..b80de68f95c --- /dev/null +++ b/cpp/tests/experimental/graph_test.cpp @@ -0,0 +1,242 @@ +/* + * Copyright (c) 2020, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governin_from_mtxg permissions and + * limitations under the License. + */ + +#include +#include + +#include +#include + +#include +#include +#include +#include + +#include + +#include +#include +#include +#include +#include + +template +std::tuple, std::vector, std::vector> graph_reference( + vertex_t const* p_src_vertices, + vertex_t const* p_dst_vertices, + weight_t const* p_edge_weights, + vertex_t number_of_vertices, + edge_t number_of_edges) +{ + std::vector offsets(number_of_vertices + 1, edge_t{0}); + std::vector indices(number_of_edges, vertex_t{0}); + std::vector weights(p_edge_weights != nullptr ? number_of_edges : 0, weight_t{0.0}); + + for (size_t i = 0; i < number_of_edges; ++i) { + auto major = store_transposed ? p_dst_vertices[i] : p_src_vertices[i]; + offsets[1 + major]++; + } + std::partial_sum(offsets.begin() + 1, offsets.end(), offsets.begin() + 1); + + for (size_t i = 0; i < number_of_edges; ++i) { + auto major = store_transposed ? p_dst_vertices[i] : p_src_vertices[i]; + auto minor = store_transposed ? p_src_vertices[i] : p_dst_vertices[i]; + auto start = offsets[major]; + auto degree = offsets[major + 1] - start; + auto idx = indices[start + degree - 1]++; + indices[start + idx] = minor; + if (p_edge_weights != nullptr) { weights[start + idx] = p_edge_weights[i]; } + } + + return std::make_tuple(std::move(offsets), std::move(indices), std::move(weights)); +} + +typedef struct Graph_Usecase_t { + std::string graph_file_full_path{}; + bool test_weighted{false}; + + Graph_Usecase_t(std::string const& graph_file_path, bool test_weighted) + : test_weighted(test_weighted) + { + 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; + } + }; +} Graph_Usecase; + +class Tests_Graph : public ::testing::TestWithParam { + public: + Tests_Graph() {} + static void SetupTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(Graph_Usecase const& configuration) + { + auto mm_graph = + cugraph::test::read_edgelist_from_matrix_market_file( + configuration.graph_file_full_path); + edge_t number_of_edges = static_cast(mm_graph.h_rows.size()); + + std::vector h_reference_offsets{}; + std::vector h_reference_indices{}; + std::vector h_reference_weights{}; + + std::tie(h_reference_offsets, h_reference_indices, h_reference_weights) = + graph_reference( + mm_graph.h_rows.data(), + mm_graph.h_cols.data(), + configuration.test_weighted ? mm_graph.h_weights.data() : nullptr, + mm_graph.number_of_vertices, + number_of_edges); + + raft::handle_t handle{}; + + rmm::device_uvector d_rows(number_of_edges, handle.get_stream()); + rmm::device_uvector d_cols(number_of_edges, handle.get_stream()); + rmm::device_uvector d_weights(configuration.test_weighted ? number_of_edges : 0, + handle.get_stream()); + + raft::update_device( + d_rows.data(), mm_graph.h_rows.data(), number_of_edges, handle.get_stream()); + raft::update_device( + d_cols.data(), mm_graph.h_cols.data(), number_of_edges, handle.get_stream()); + if (configuration.test_weighted) { + raft::update_device( + d_weights.data(), mm_graph.h_weights.data(), number_of_edges, handle.get_stream()); + } + + cugraph::experimental::edgelist_t edgelist{ + d_rows.data(), + d_cols.data(), + configuration.test_weighted ? d_weights.data() : nullptr, + number_of_edges}; + + CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + + auto graph = + cugraph::experimental::graph_t( + handle, + edgelist, + mm_graph.number_of_vertices, + cugraph::experimental::graph_properties_t{mm_graph.is_symmetric, false}, + false, + true); + + auto graph_view = graph.view(); + + CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + + ASSERT_EQ(graph_view.get_number_of_vertices(), mm_graph.number_of_vertices); + ASSERT_EQ(graph_view.get_number_of_edges(), number_of_edges); + + std::vector h_cugraph_offsets(graph_view.get_number_of_vertices() + 1); + std::vector h_cugraph_indices(graph_view.get_number_of_edges()); + std::vector h_cugraph_weights( + configuration.test_weighted ? graph_view.get_number_of_edges() : 0); + + raft::update_host(h_cugraph_offsets.data(), + graph_view.offsets(), + graph_view.get_number_of_vertices() + 1, + handle.get_stream()); + raft::update_host(h_cugraph_indices.data(), + graph_view.indices(), + graph_view.get_number_of_edges(), + handle.get_stream()); + if (configuration.test_weighted) { + raft::update_host(h_cugraph_weights.data(), + graph_view.weights(), + graph_view.get_number_of_edges(), + handle.get_stream()); + } + + CUDA_TRY(cudaStreamSynchronize(handle.get_stream())); + + ASSERT_TRUE( + std::equal(h_reference_offsets.begin(), h_reference_offsets.end(), h_cugraph_offsets.begin())) + << "Graph compressed sparse format offsets do not match with the reference values."; + ASSERT_EQ(h_reference_weights.size(), h_cugraph_weights.size()); + for (vertex_t i = 0; i < mm_graph.number_of_vertices; ++i) { + auto start = h_reference_offsets[i]; + auto degree = h_reference_offsets[i + 1] - start; + if (configuration.test_weighted) { + std::vector> reference_pairs(degree); + std::vector> cugraph_pairs(degree); + for (edge_t j = 0; j < degree; ++j) { + reference_pairs[j] = + std::make_tuple(h_reference_indices[start + j], h_reference_weights[start + j]); + cugraph_pairs[j] = + std::make_tuple(h_cugraph_indices[start + j], h_cugraph_weights[start + j]); + } + std::sort(reference_pairs.begin(), reference_pairs.end()); + std::sort(cugraph_pairs.begin(), cugraph_pairs.end()); + ASSERT_TRUE( + std::equal(reference_pairs.begin(), reference_pairs.end(), cugraph_pairs.begin())) + << "Graph compressed sparse format indices & weights for vertex " << i + << " do not match with the reference values."; + } else { + std::vector reference_indices(h_reference_indices.begin() + start, + h_reference_indices.begin() + (start + degree)); + std::vector cugraph_indices(h_cugraph_indices.begin() + start, + h_cugraph_indices.begin() + (start + degree)); + std::sort(reference_indices.begin(), reference_indices.end()); + std::sort(cugraph_indices.begin(), cugraph_indices.end()); + ASSERT_TRUE( + std::equal(reference_indices.begin(), reference_indices.end(), cugraph_indices.begin())) + << "Graph compressed sparse format indices for vertex " << i + << " do not match with the reference values."; + } + } + } +}; + +TEST_P(Tests_Graph, CheckStoreTransposedFalse) +{ + run_current_test(GetParam()); + run_current_test(GetParam()); + run_current_test(GetParam()); + run_current_test(GetParam()); + run_current_test(GetParam()); + run_current_test(GetParam()); +} + +TEST_P(Tests_Graph, CheckStoreTransposedTrue) +{ + run_current_test(GetParam()); + run_current_test(GetParam()); + run_current_test(GetParam()); + run_current_test(GetParam()); + run_current_test(GetParam()); + run_current_test(GetParam()); +} + +INSTANTIATE_TEST_CASE_P(simple_test, + Tests_Graph, + ::testing::Values(Graph_Usecase("test/datasets/karate.mtx", false), + Graph_Usecase("test/datasets/karate.mtx", true), + Graph_Usecase("test/datasets/web-Google.mtx", false), + Graph_Usecase("test/datasets/web-Google.mtx", true), + Graph_Usecase("test/datasets/ljournal-2008.mtx", false), + Graph_Usecase("test/datasets/ljournal-2008.mtx", true), + Graph_Usecase("test/datasets/webbase-1M.mtx", false), + Graph_Usecase("test/datasets/webbase-1M.mtx", true))); + +CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/utilities/test_utilities.hpp b/cpp/tests/utilities/test_utilities.hpp index b9aace72f5d..65703e9541d 100644 --- a/cpp/tests/utilities/test_utilities.hpp +++ b/cpp/tests/utilities/test_utilities.hpp @@ -15,6 +15,8 @@ */ #pragma once +#include + #include #include @@ -25,15 +27,7 @@ extern "C" { #include #include - -#define MPICHECK(cmd) \ - { \ - int e = cmd; \ - if (e != MPI_SUCCESS) { \ - printf("Failed: MPI error %s:%d '%d'\n", __FILE__, __LINE__, e); \ - FAIL(); \ - } \ - } +#include namespace cugraph { namespace test { @@ -334,5 +328,53 @@ static const std::string& get_rapids_dataset_root_dir() return rdrd; } +template +struct edgelist_from_market_matrix_file_t { + std::vector h_rows{}; + std::vector h_cols{}; + std::vector h_weights{}; + vertex_t number_of_vertices{}; + bool is_symmetric{}; +}; + +template +edgelist_from_market_matrix_file_t read_edgelist_from_matrix_market_file( + std::string const& graph_file_full_path) +{ + edgelist_from_market_matrix_file_t ret{}; + + MM_typecode mc{}; + vertex_t m{}; + vertex_t k{}; + edge_t nnz{}; + + FILE* file = fopen(graph_file_full_path.c_str(), "r"); + CUGRAPH_EXPECTS(file != nullptr, "fopen failure."); + + edge_t tmp_m{}; + edge_t tmp_k{}; + auto mm_ret = cugraph::test::mm_properties(file, 1, &mc, &tmp_m, &tmp_k, &nnz); + CUGRAPH_EXPECTS(mm_ret == 0, "could not read Matrix Market file properties."); + m = static_cast(tmp_m); + k = static_cast(tmp_k); + CUGRAPH_EXPECTS(mm_is_matrix(mc) && mm_is_coordinate(mc) && !mm_is_complex(mc) && !mm_is_skew(mc), + "invalid Matrix Market file properties."); + + ret.h_rows.assign(nnz, vertex_t{0}); + ret.h_cols.assign(nnz, vertex_t{0}); + ret.h_weights.assign(nnz, weight_t{0.0}); + ret.number_of_vertices = m; + ret.is_symmetric = mm_is_symmetric(mc); + + mm_ret = cugraph::test::mm_to_coo( + file, 1, nnz, ret.h_rows.data(), ret.h_cols.data(), ret.h_weights.data(), nullptr); + CUGRAPH_EXPECTS(mm_ret == 0, "could not read matrix data"); + + auto file_ret = fclose(file); + CUGRAPH_EXPECTS(file_ret == 0, "fclose failure."); + + return std::move(ret); +} + } // namespace test } // namespace cugraph diff --git a/datasets/karate-disjoint.csv b/datasets/karate-disjoint.csv new file mode 100644 index 00000000000..d4f03bb35b9 --- /dev/null +++ b/datasets/karate-disjoint.csv @@ -0,0 +1,312 @@ +1 0 1.0 +2 0 1.0 +3 0 1.0 +4 0 1.0 +5 0 1.0 +6 0 1.0 +7 0 1.0 +8 0 1.0 +10 0 1.0 +11 0 1.0 +12 0 1.0 +13 0 1.0 +17 0 1.0 +19 0 1.0 +21 0 1.0 +31 0 1.0 +2 1 1.0 +3 1 1.0 +7 1 1.0 +13 1 1.0 +17 1 1.0 +19 1 1.0 +21 1 1.0 +30 1 1.0 +3 2 1.0 +7 2 1.0 +8 2 1.0 +9 2 1.0 +13 2 1.0 +27 2 1.0 +28 2 1.0 +32 2 1.0 +7 3 1.0 +12 3 1.0 +13 3 1.0 +6 4 1.0 +10 4 1.0 +6 5 1.0 +10 5 1.0 +16 5 1.0 +16 6 1.0 +30 8 1.0 +32 8 1.0 +33 8 1.0 +33 9 1.0 +33 13 1.0 +32 14 1.0 +33 14 1.0 +32 15 1.0 +33 15 1.0 +32 18 1.0 +33 18 1.0 +33 19 1.0 +32 20 1.0 +33 20 1.0 +32 22 1.0 +33 22 1.0 +25 23 1.0 +27 23 1.0 +29 23 1.0 +32 23 1.0 +33 23 1.0 +25 24 1.0 +27 24 1.0 +31 24 1.0 +31 25 1.0 +29 26 1.0 +33 26 1.0 +33 27 1.0 +31 28 1.0 +33 28 1.0 +32 29 1.0 +33 29 1.0 +32 30 1.0 +33 30 1.0 +32 31 1.0 +33 31 1.0 +33 32 1.0 +0 1 1.0 +0 2 1.0 +0 3 1.0 +0 4 1.0 +0 5 1.0 +0 6 1.0 +0 7 1.0 +0 8 1.0 +0 10 1.0 +0 11 1.0 +0 12 1.0 +0 13 1.0 +0 17 1.0 +0 19 1.0 +0 21 1.0 +0 31 1.0 +1 2 1.0 +1 3 1.0 +1 7 1.0 +1 13 1.0 +1 17 1.0 +1 19 1.0 +1 21 1.0 +1 30 1.0 +2 3 1.0 +2 7 1.0 +2 8 1.0 +2 9 1.0 +2 13 1.0 +2 27 1.0 +2 28 1.0 +2 32 1.0 +3 7 1.0 +3 12 1.0 +3 13 1.0 +4 6 1.0 +4 10 1.0 +5 6 1.0 +5 10 1.0 +5 16 1.0 +6 16 1.0 +8 30 1.0 +8 32 1.0 +8 33 1.0 +9 33 1.0 +13 33 1.0 +14 32 1.0 +14 33 1.0 +15 32 1.0 +15 33 1.0 +18 32 1.0 +18 33 1.0 +19 33 1.0 +20 32 1.0 +20 33 1.0 +22 32 1.0 +22 33 1.0 +23 25 1.0 +23 27 1.0 +23 29 1.0 +23 32 1.0 +23 33 1.0 +24 25 1.0 +24 27 1.0 +24 31 1.0 +25 31 1.0 +26 29 1.0 +26 33 1.0 +27 33 1.0 +28 31 1.0 +28 33 1.0 +29 32 1.0 +29 33 1.0 +30 32 1.0 +30 33 1.0 +31 32 1.0 +31 33 1.0 +32 33 1.0 +101 100 1.0 +102 100 1.0 +103 100 1.0 +104 100 1.0 +105 100 1.0 +106 100 1.0 +107 100 1.0 +108 100 1.0 +110 100 1.0 +111 100 1.0 +112 100 1.0 +113 100 1.0 +117 100 1.0 +119 100 1.0 +121 100 1.0 +131 100 1.0 +102 101 1.0 +103 101 1.0 +107 101 1.0 +113 101 1.0 +117 101 1.0 +119 101 1.0 +121 101 1.0 +130 101 1.0 +103 102 1.0 +107 102 1.0 +108 102 1.0 +109 102 1.0 +113 102 1.0 +127 102 1.0 +128 102 1.0 +132 102 1.0 +107 103 1.0 +112 103 1.0 +113 103 1.0 +106 104 1.0 +110 104 1.0 +106 105 1.0 +110 105 1.0 +116 105 1.0 +116 106 1.0 +130 108 1.0 +132 108 1.0 +133 108 1.0 +133 109 1.0 +133 113 1.0 +132 114 1.0 +133 114 1.0 +132 115 1.0 +133 115 1.0 +132 118 1.0 +133 118 1.0 +133 119 1.0 +132 120 1.0 +133 120 1.0 +132 122 1.0 +133 122 1.0 +125 123 1.0 +127 123 1.0 +129 123 1.0 +132 123 1.0 +133 123 1.0 +125 124 1.0 +127 124 1.0 +131 124 1.0 +131 125 1.0 +129 126 1.0 +133 126 1.0 +133 127 1.0 +131 128 1.0 +133 128 1.0 +132 129 1.0 +133 129 1.0 +132 130 1.0 +133 130 1.0 +132 131 1.0 +133 131 1.0 +133 132 1.0 +100 101 1.0 +100 102 1.0 +100 103 1.0 +100 104 1.0 +100 105 1.0 +100 106 1.0 +100 107 1.0 +100 108 1.0 +100 110 1.0 +100 111 1.0 +100 112 1.0 +100 113 1.0 +100 117 1.0 +100 119 1.0 +100 121 1.0 +100 131 1.0 +101 102 1.0 +101 103 1.0 +101 107 1.0 +101 113 1.0 +101 117 1.0 +101 119 1.0 +101 121 1.0 +101 130 1.0 +102 103 1.0 +102 107 1.0 +102 108 1.0 +102 109 1.0 +102 113 1.0 +102 127 1.0 +102 128 1.0 +102 132 1.0 +103 107 1.0 +103 112 1.0 +103 113 1.0 +104 106 1.0 +104 110 1.0 +105 106 1.0 +105 110 1.0 +105 116 1.0 +106 116 1.0 +108 130 1.0 +108 132 1.0 +108 133 1.0 +109 133 1.0 +113 133 1.0 +114 132 1.0 +114 133 1.0 +115 132 1.0 +115 133 1.0 +118 132 1.0 +118 133 1.0 +119 133 1.0 +120 132 1.0 +120 133 1.0 +122 132 1.0 +122 133 1.0 +123 125 1.0 +123 127 1.0 +123 129 1.0 +123 132 1.0 +123 133 1.0 +124 125 1.0 +124 127 1.0 +124 131 1.0 +125 131 1.0 +126 129 1.0 +126 133 1.0 +127 133 1.0 +128 131 1.0 +128 133 1.0 +129 132 1.0 +129 133 1.0 +130 132 1.0 +130 133 1.0 +131 132 1.0 +131 133 1.0 +132 133 1.0 diff --git a/docs/source/api.rst b/docs/source/api.rst index 84a01604d33..b194aa0e03c 100644 --- a/docs/source/api.rst +++ b/docs/source/api.rst @@ -14,12 +14,6 @@ Graph :members: :undoc-members: -Renumbering ------------ - -.. automodule:: cugraph.structure.renumber - :members: - :undoc-members: Symmetrize ---------- @@ -57,34 +51,34 @@ Katz Centrality Community ========= -Leiden +Ensemble clustering for graphs (ECG) +------------------------------------ + +.. automodule:: cugraph.community.ecg + :members: + :undoc-members: + +K-Truss ------- -.. automodule:: cugraph.community.leiden +.. automodule:: cugraph.community.ktruss_subgraph :members: :undoc-members: -Louvain +Leiden ------- -.. automodule:: cugraph.community.louvain +.. automodule:: cugraph.community.leiden :members: :undoc-members: -K-Truss +Louvain ------- -.. automodule:: cugraph.community.ktruss_subgraph +.. automodule:: cugraph.community.louvain :members: :undoc-members: -ECG ---- - -.. automodule:: cugraph.community.ecg - :members: - :undoc-members: - Spectral Clustering ------------------- @@ -117,9 +111,17 @@ Connected Components :members: :undoc-members: + Cores ===== +Core Number +----------- + +.. automodule:: cugraph.cores.core_number + :members: + :undoc-members: + K-Core ------ @@ -127,12 +129,6 @@ K-Core :members: :undoc-members: -Core Number ------------ - -.. automodule:: cugraph.cores.core_number - :members: - :undoc-members: Layout ====== @@ -144,9 +140,17 @@ Force Atlas 2 :members: :undoc-members: + Link Analysis ============= +HITS +--------- + +.. automodule:: cugraph.link_analysis.hits + :members: + :undoc-members: + Pagerank --------- @@ -154,6 +158,7 @@ Pagerank :members: :undoc-members: + Link Prediction =============== @@ -179,6 +184,7 @@ Overlap Coefficient :members: :undoc-members: + Traversal ========= @@ -196,13 +202,9 @@ Single-source-shortest-path :members: :undoc-members: + Utilities ========= -R-mat Graph Generation ----------------------- -.. automodule:: cugraph.utilities.grmat - :members: - :undoc-members: diff --git a/docs/source/conf.py b/docs/source/conf.py index 6b144e93c94..0c8a0316278 100644 --- a/docs/source/conf.py +++ b/docs/source/conf.py @@ -188,3 +188,11 @@ def setup(app): app.add_stylesheet('params.css') + +from recommonmark.parser import CommonMarkParser + +source_parsers = { + '.md': CommonMarkParser, +} + +source_suffix = ['.rst', '.md'] \ No newline at end of file diff --git a/docs/source/cugraph_blogs.rst b/docs/source/cugraph_blogs.rst new file mode 100644 index 00000000000..a9954aee5cb --- /dev/null +++ b/docs/source/cugraph_blogs.rst @@ -0,0 +1,48 @@ + +cuGraph BLOGS and Presentations +************************************************ + +The RAPIDS team blogs at https://medium.com/rapids-ai, and many of +these blog posts provide deeper dives into features from cuGraph. +Here, we've selected just a few that are of particular interest to cuGraph users: + + +BLOGS +============== + +2019 +------- + * `RAPIDS cuGraph `_ + * `RAPIDS cuGraph — The vision and journey to version 1.0 and beyond `_ + * `RAPIDS cuGraph : multi-GPU PageRank `_ + * `Similarity in graphs: Jaccard versus the Overlap Coefficient `_ + + +2020 +------ + * `Status of RAPIDS cuGraph — Refactoring Code And Rethinking Graphs `_ + * `Tackling Large Graphs with RAPIDS cuGraph and CUDA Unified Memory on GPUs `_ + + +Media +=============== + * `Nvidia Rapids cuGraph: Making graph analysis ubiquitous `_ + * `RAPIDS cuGraph – Accelerating all your Graph needs `_ + +Academic Papers +=============== +* S. Kang, A. Fender, J. Eaton, B. Rees: Computing PageRank Scores of Web Crawl Data Using DGX A100 Clusters. In IEEE HPEC, Sep. 2020 + + +Other BLOGS +======================== +* `4 graph algorithms on steroids for data scientists with cugraph + `_ +* `Where should I walk `_ +* `Where really are the parking spots? `_ +* `Accelerating Single Cell Genomic Analysis using RAPIDS `_ + + + + +Copyright (c) 2020, NVIDIA CORPORATION. diff --git a/docs/source/cugraph_intro.rst b/docs/source/cugraph_intro.rst new file mode 100644 index 00000000000..cd2d750e35f --- /dev/null +++ b/docs/source/cugraph_intro.rst @@ -0,0 +1,13 @@ + +cuGraph Intro +------------------------------ + + + +Graph Type + + +Algorithms + + +Using diff --git a/docs/source/cugraph_ref.rst b/docs/source/cugraph_ref.rst new file mode 100644 index 00000000000..0742d5f5aa7 --- /dev/null +++ b/docs/source/cugraph_ref.rst @@ -0,0 +1,51 @@ +References +=============== + +Algorithms +************** + +Betweenness Centrality +------------------------- +- Brandes, U. (2001). A faster algorithm for betweenness centrality. Journal of mathematical sociology, 25(2), 163-177. +- Brandes, U. (2008). On variants of shortest-path betweenness centrality and their generic computation. Social Networks, 30(2), 136-145. +- McLaughlin, A., & Bader, D. A. (2018). Accelerating GPU betweenness centrality. Communications of the ACM, 61(8), 85-92. + + +Katz +------------------------- + +- J. Cohen, *Trusses: Cohesive subgraphs for social network analysis* National security agency technical report, 2008 +- O. Green, J. Fox, E. Kim, F. Busato, et al. *Quickly Finding a Truss in a Haystack* IEEE High Performance Extreme Computing Conference (HPEC), 2017 https://doi.org/10.1109/HPEC.2017.8091038 +- O. Green, P. Yalamanchili, L.M. Munguia, “*ast Triangle Counting on GPU* Irregular Applications: Architectures and Algorithms (IA3), 2014 + + + + + + +Data Sets +************** + +karate + - W. W. Zachary, *An information flow model for conflict and fission in small groups*, Journal of Anthropological Research 33, 452-473 (1977). +dolphins + - D. Lusseau, K. Schneider, O. J. Boisseau, P. Haase, E. Slooten, and S. M. Dawson, + *The bottlenose dolphin community of Doubtful Sound features a large proportion of long-lasting associations*, + Behavioral Ecology and Sociobiology 54, 396-405 (2003). +netscience + - M. E. J. Newman, + *Finding community structure in networks using the eigenvectors of matrices*, + Preprint physics/0605087 (2006). +email-Eu-core + - Hao Yin, Austin R. Benson, Jure Leskovec, and David F. Gleich. + *Local Higher-order Graph Clustering.* + In Proceedings of the 23rd ACM SIGKDD International Conference on Knowledge Discovery and Data Mining. 2017. + - J. Leskovec, J. Kleinberg and C. Faloutsos. + *Graph Evolution: Densification and Shrinking Diameters*. + ACM Transactions on Knowledge Discovery from Data (ACM TKDD), 1(1), 2007. http://www.cs.cmu.edu/~jure/pubs/powergrowth-tkdd.pdf +polbooks + - V. Krebs, unpublished, http://www.orgnet.com/. + + + + diff --git a/docs/source/index.rst b/docs/source/index.rst index 47360b88f91..2cd95e7f129 100644 --- a/docs/source/index.rst +++ b/docs/source/index.rst @@ -7,6 +7,10 @@ Welcome to cugraph's documentation! api.rst dask-cugraph.rst + cugraph_intro.rst + cugraph_blogs.rst + cugraph_ref.rst + Indices and tables ================== diff --git a/python/cugraph/community/ecg.py b/python/cugraph/community/ecg.py index 5030a2475b7..85d97b50a8e 100644 --- a/python/cugraph/community/ecg.py +++ b/python/cugraph/community/ecg.py @@ -47,20 +47,20 @@ def ecg(input_graph, min_weight=0.05, ensemble_size=16): GPU data frame of size V containing two columns, the vertex id and the partition id it is assigned to. - df['vertex'] : cudf.Series + df[vertex] : cudf.Series Contains the vertex identifiers - df['partition'] : cudf.Series + df[partition] : cudf.Series Contains the partition assigned to the vertices Examples -------- - >>> M = cudf.read_csv('datasets/karate.csv', - delimiter = ' ', + >>> M = cudf.read_csv('datasets/karate.csv', delimiter = ' ', dtype=['int32', 'int32', 'float32'], header=None) >>> G = cugraph.Graph() >>> G.from_cudf_edgelist(M, source='0', destination='1', edge_attr='2') >>> parts = cugraph.ecg(G) + """ parts = ecg_wrapper.ecg(input_graph, min_weight, ensemble_size) diff --git a/python/cugraph/layout/force_atlas2.py b/python/cugraph/layout/force_atlas2.py index 4a61e2a345b..047431b713a 100644 --- a/python/cugraph/layout/force_atlas2.py +++ b/python/cugraph/layout/force_atlas2.py @@ -83,6 +83,7 @@ def force_atlas2( callback: GraphBasedDimRedCallback An instance of GraphBasedDimRedCallback class to intercept the internal state of positions while they are being trained. + Example of callback usage: from cugraph.internals import GraphBasedDimRedCallback class CustomCallback(GraphBasedDimRedCallback): diff --git a/python/cugraph/structure/graph.py b/python/cugraph/structure/graph.py index f579a32cab6..c918cd44ae2 100644 --- a/python/cugraph/structure/graph.py +++ b/python/cugraph/structure/graph.py @@ -767,20 +767,12 @@ def get_two_hop_neighbors(self): Returns ------- df : cudf.DataFrame - df['first'] : cudf.Series + df[first] : cudf.Series the first vertex id of a pair, if an external vertex id is defined by only one column - df['second'] : cudf.Series + df[second] : cudf.Series the second vertex id of a pair, if an external vertex id is defined by only one column - df['*_first'] : cudf.Series - the first vertex id of a pair, column 0 of the external - vertex id will be represented as '0_first', column 1 as - '1_first', etc. - df['*_second'] : cudf.Series - the second vertex id of a pair, column 0 of the external - vertex id will be represented as '0_first', column 1 as - '1_first', etc. """ if self.distributed: raise Exception("Not supported for distributed graph") @@ -874,10 +866,11 @@ def in_degree(self, vertex_subset=None): vertices (vertex_subset) containing the in_degree. The ordering is relative to the adjacency list, or that given by the specified vertex_subset. - df['vertex'] : cudf.Series + + df[vertex] : cudf.Series The vertex IDs (will be identical to vertex_subset if specified). - df['degree'] : cudf.Series + df[degree] : cudf.Series The computed in-degree of the corresponding vertex. Examples @@ -912,10 +905,11 @@ def out_degree(self, vertex_subset=None): vertices (vertex_subset) containing the out_degree. The ordering is relative to the adjacency list, or that given by the specified vertex_subset. - df['vertex'] : cudf.Series + + df[vertex] : cudf.Series The vertex IDs (will be identical to vertex_subset if specified). - df['degree'] : cudf.Series + df[degree] : cudf.Series The computed out-degree of the corresponding vertex. Examples @@ -933,15 +927,16 @@ def out_degree(self, vertex_subset=None): def degree(self, vertex_subset=None): """ - Compute vertex degree. By default, this method computes vertex + Compute vertex degree, which is the total number of edges incident + to a vertex (both in and out edges). By default, this method computes degrees for the entire set of vertices. If vertex_subset is provided, - this method optionally filters out all but those listed in + then this method optionally filters out all but those listed in vertex_subset. Parameters ---------- vertex_subset : cudf.Series or iterable container, optional - A container of vertices for displaying corresponding degree. If not + a container of vertices for displaying corresponding degree. If not set, degrees are computed for the entire set of vertices. Returns @@ -951,6 +946,7 @@ def degree(self, vertex_subset=None): vertices (vertex_subset) containing the degree. The ordering is relative to the adjacency list, or that given by the specified vertex_subset. + df['vertex'] : cudf.Series The vertex IDs (will be identical to vertex_subset if specified). @@ -963,7 +959,8 @@ def degree(self, vertex_subset=None): >>> dtype=['int32', 'int32', 'float32'], header=None) >>> G = cugraph.Graph() >>> G.from_cudf_edgelist(M, '0', '1') - >>> df = G.degree([0,9,12]) + >>> all_df = G.degree() + >>> subset_df = G.degree([0,9,12]) """ if self.distributed: diff --git a/python/cugraph/tests/dask/mg_context.py b/python/cugraph/tests/dask/mg_context.py index e062cfd842b..9a7ea2ace67 100644 --- a/python/cugraph/tests/dask/mg_context.py +++ b/python/cugraph/tests/dask/mg_context.py @@ -1,3 +1,16 @@ +# Copyright (c) 2019-2020, 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. + import time import os diff --git a/python/cugraph/tests/dask/test_mg_batch_betweenness_centrality.py b/python/cugraph/tests/dask/test_mg_batch_betweenness_centrality.py index 92782e7a0b3..ccb0c94b020 100644 --- a/python/cugraph/tests/dask/test_mg_batch_betweenness_centrality.py +++ b/python/cugraph/tests/dask/test_mg_batch_betweenness_centrality.py @@ -1,5 +1,17 @@ +# Copyright (c) 2019-2020, 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. + import pytest -import cugraph.tests.utils as utils import numpy as np from cugraph.tests.dask.mg_context import (MGContext, @@ -24,7 +36,7 @@ # ============================================================================= # Parameters # ============================================================================= -DATASETS = utils.DATASETS_1 +DATASETS = ['../datasets/karate.csv'] MG_DEVICE_COUNT_OPTIONS = [1, 2, 3, 4] RESULT_DTYPE_OPTIONS = [np.float64] diff --git a/python/cugraph/tests/dask/test_mg_batch_edge_betweenness_centrality.py b/python/cugraph/tests/dask/test_mg_batch_edge_betweenness_centrality.py index d4906ca04ef..01023839d06 100644 --- a/python/cugraph/tests/dask/test_mg_batch_edge_betweenness_centrality.py +++ b/python/cugraph/tests/dask/test_mg_batch_edge_betweenness_centrality.py @@ -1,5 +1,17 @@ +# Copyright (c) 2019-2020, 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. + import pytest -import cugraph.tests.utils as utils import numpy as np @@ -24,8 +36,8 @@ # ============================================================================= # Parameters # ============================================================================= -DATASETS = utils.DATASETS_1 -MG_DEVICE_COUNT_OPTIONS = [1, 2, 3, 4] +DATASETS = ['../datasets/karate.csv'] +MG_DEVICE_COUNT_OPTIONS = [1, 2, 4] RESULT_DTYPE_OPTIONS = [np.float64] diff --git a/python/cugraph/tests/dask/test_mg_renumber.py b/python/cugraph/tests/dask/test_mg_renumber.py index cf6fada2def..ceeeeb77a5a 100644 --- a/python/cugraph/tests/dask/test_mg_renumber.py +++ b/python/cugraph/tests/dask/test_mg_renumber.py @@ -45,7 +45,7 @@ def client_connection(): # Test all combinations of default/managed and pooled/non-pooled allocation -@pytest.mark.parametrize("graph_file", utils.DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_UNRENUMBERED) def test_mg_renumber(graph_file, client_connection): gc.collect() @@ -85,7 +85,7 @@ def test_mg_renumber(graph_file, client_connection): # Test all combinations of default/managed and pooled/non-pooled allocation -@pytest.mark.parametrize("graph_file", utils.DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_UNRENUMBERED) def test_mg_renumber2(graph_file, client_connection): gc.collect() @@ -120,7 +120,7 @@ def test_mg_renumber2(graph_file, client_connection): # Test all combinations of default/managed and pooled/non-pooled allocation -@pytest.mark.parametrize("graph_file", utils.DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_UNRENUMBERED) def test_mg_renumber3(graph_file, client_connection): gc.collect() diff --git a/python/cugraph/tests/dask/test_mg_replication.py b/python/cugraph/tests/dask/test_mg_replication.py index b74e6e9a15f..061bcf83f20 100644 --- a/python/cugraph/tests/dask/test_mg_replication.py +++ b/python/cugraph/tests/dask/test_mg_replication.py @@ -18,8 +18,9 @@ import cugraph.dask.structure.replication as replication import cugraph.tests.utils as utils import pytest +import gc -DATASETS_OPTIONS = utils.DATASETS_1 +DATASETS_OPTIONS = utils.DATASETS_SMALL DIRECTED_GRAPH_OPTIONS = [False, True] MG_DEVICE_COUNT_OPTIONS = [1, 2, 3, 4] @@ -28,6 +29,7 @@ @pytest.mark.parametrize("mg_device_count", MG_DEVICE_COUNT_OPTIONS) def test_replicate_cudf_dataframe_with_weights(input_data_path, mg_device_count): + gc.collect() skip_if_not_enough_devices(mg_device_count) df = cudf.read_csv(input_data_path, delimiter=' ', @@ -45,6 +47,7 @@ def test_replicate_cudf_dataframe_with_weights(input_data_path, @pytest.mark.parametrize("mg_device_count", MG_DEVICE_COUNT_OPTIONS) def test_replicate_cudf_dataframe_no_weights(input_data_path, mg_device_count): + gc.collect() skip_if_not_enough_devices(mg_device_count) df = cudf.read_csv(input_data_path, delimiter=' ', @@ -62,6 +65,7 @@ def test_replicate_cudf_dataframe_no_weights(input_data_path, @pytest.mark.parametrize("mg_device_count", MG_DEVICE_COUNT_OPTIONS) def test_replicate_cudf_series(input_data_path, mg_device_count): + gc.collect() skip_if_not_enough_devices(mg_device_count) df = cudf.read_csv(input_data_path, delimiter=' ', @@ -85,6 +89,7 @@ def test_replicate_cudf_series(input_data_path, @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) @pytest.mark.parametrize("mg_device_count", MG_DEVICE_COUNT_OPTIONS) def test_enable_batch_no_context(graph_file, directed, mg_device_count): + gc.collect() skip_if_not_enough_devices(mg_device_count) G = utils.generate_cugraph_graph_from_file(graph_file, directed) assert G.batch_enabled is False, "Internal property should be False" @@ -97,6 +102,7 @@ def test_enable_batch_no_context(graph_file, directed, mg_device_count): @pytest.mark.parametrize("mg_device_count", MG_DEVICE_COUNT_OPTIONS) def test_enable_batch_no_context_view_adj(graph_file, directed, mg_device_count): + gc.collect() skip_if_not_enough_devices(mg_device_count) G = utils.generate_cugraph_graph_from_file(graph_file, directed) assert G.batch_enabled is False, "Internal property should be False" @@ -108,6 +114,7 @@ def test_enable_batch_no_context_view_adj(graph_file, directed, @pytest.mark.parametrize("mg_device_count", MG_DEVICE_COUNT_OPTIONS) def test_enable_batch_context_then_views(graph_file, directed, mg_device_count): + gc.collect() skip_if_not_enough_devices(mg_device_count) G = utils.generate_cugraph_graph_from_file(graph_file, directed) with MGContext(mg_device_count): @@ -131,6 +138,7 @@ def test_enable_batch_context_then_views(graph_file, directed, @pytest.mark.parametrize("mg_device_count", MG_DEVICE_COUNT_OPTIONS) def test_enable_batch_view_then_context(graph_file, directed, mg_device_count): + gc.collect() skip_if_not_enough_devices(mg_device_count) G = utils.generate_cugraph_graph_from_file(graph_file, directed) @@ -158,6 +166,7 @@ def test_enable_batch_view_then_context(graph_file, directed, @pytest.mark.parametrize("mg_device_count", MG_DEVICE_COUNT_OPTIONS) def test_enable_batch_context_no_context_views(graph_file, directed, mg_device_count): + gc.collect() skip_if_not_enough_devices(mg_device_count) G = utils.generate_cugraph_graph_from_file(graph_file, directed) with MGContext(mg_device_count): @@ -177,6 +186,7 @@ def test_enable_batch_context_no_context_views(graph_file, directed, @pytest.mark.parametrize("mg_device_count", MG_DEVICE_COUNT_OPTIONS) def test_enable_batch_edgelist_replication(graph_file, directed, mg_device_count): + gc.collect() skip_if_not_enough_devices(mg_device_count) G = utils.generate_cugraph_graph_from_file(graph_file, directed) with MGContext(mg_device_count): @@ -192,6 +202,7 @@ def test_enable_batch_edgelist_replication(graph_file, directed, @pytest.mark.parametrize("mg_device_count", MG_DEVICE_COUNT_OPTIONS) def test_enable_batch_adjlist_replication_weights(graph_file, directed, mg_device_count): + gc.collect() skip_if_not_enough_devices(mg_device_count) df = cudf.read_csv(graph_file, delimiter=' ', @@ -224,6 +235,7 @@ def test_enable_batch_adjlist_replication_weights(graph_file, directed, @pytest.mark.parametrize("mg_device_count", MG_DEVICE_COUNT_OPTIONS) def test_enable_batch_adjlist_replication_no_weights(graph_file, directed, mg_device_count): + gc.collect() skip_if_not_enough_devices(mg_device_count) df = cudf.read_csv(graph_file, delimiter=' ', diff --git a/python/cugraph/tests/test_betweenness_centrality.py b/python/cugraph/tests/test_betweenness_centrality.py index 165ef56e549..1ef1601edd5 100644 --- a/python/cugraph/tests/test_betweenness_centrality.py +++ b/python/cugraph/tests/test_betweenness_centrality.py @@ -41,10 +41,6 @@ NORMALIZED_OPTIONS = [False, True] DEFAULT_EPSILON = 0.0001 -DATASETS = ["../datasets/karate.csv", "../datasets/netscience.csv"] - -UNRENUMBERED_DATASETS = ["../datasets/karate.csv"] - SUBSET_SIZE_OPTIONS = [4, None] SUBSET_SEED_OPTIONS = [42] @@ -298,7 +294,7 @@ def prepare_test(): # ============================================================================= # Tests # ============================================================================= -@pytest.mark.parametrize("graph_file", DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_SMALL) @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) @pytest.mark.parametrize("subset_size", SUBSET_SIZE_OPTIONS) @pytest.mark.parametrize("normalized", NORMALIZED_OPTIONS) @@ -330,7 +326,7 @@ def test_betweenness_centrality( compare_scores(sorted_df, first_key="cu_bc", second_key="ref_bc") -@pytest.mark.parametrize("graph_file", DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_SMALL) @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) @pytest.mark.parametrize("subset_size", [None]) @pytest.mark.parametrize("normalized", NORMALIZED_OPTIONS) @@ -371,7 +367,7 @@ def test_betweenness_centrality_k_full( # the function operating the comparison inside is first proceeding # to a random sampling over the number of vertices (thus direct offsets) # in the graph structure instead of actual vertices identifiers -@pytest.mark.parametrize("graph_file", UNRENUMBERED_DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_UNRENUMBERED) @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) @pytest.mark.parametrize("subset_size", SUBSET_SIZE_OPTIONS) @pytest.mark.parametrize("normalized", NORMALIZED_OPTIONS) @@ -407,7 +403,7 @@ def test_betweenness_centrality_fixed_sample( compare_scores(sorted_df, first_key="cu_bc", second_key="ref_bc") -@pytest.mark.parametrize("graph_file", DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_SMALL) @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) @pytest.mark.parametrize("subset_size", SUBSET_SIZE_OPTIONS) @pytest.mark.parametrize("normalized", NORMALIZED_OPTIONS) @@ -445,7 +441,7 @@ def test_betweenness_centrality_weight_except( compare_scores(sorted_df, first_key="cu_bc", second_key="ref_bc") -@pytest.mark.parametrize("graph_file", DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_SMALL) @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) @pytest.mark.parametrize("normalized", NORMALIZED_OPTIONS) @pytest.mark.parametrize("subset_size", SUBSET_SIZE_OPTIONS) diff --git a/python/cugraph/tests/test_bfs.py b/python/cugraph/tests/test_bfs.py index 8f76e25031e..8eb175ad66d 100644 --- a/python/cugraph/tests/test_bfs.py +++ b/python/cugraph/tests/test_bfs.py @@ -235,7 +235,7 @@ def _compare_bfs_spc(G, Gnx, source): # ============================================================================= # Tests # ============================================================================= -@pytest.mark.parametrize("graph_file", utils.DATASETS_5) +@pytest.mark.parametrize("graph_file", utils.DATASETS) @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) @pytest.mark.parametrize("seed", SUBSET_SEED_OPTIONS) def test_bfs(graph_file, directed, seed): @@ -257,7 +257,7 @@ def test_bfs_spc(graph_file, directed, seed): ) -@pytest.mark.parametrize("graph_file", utils.TINY_DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_SMALL) @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) def test_bfs_spc_full(graph_file, directed): """Test BFS traversal on every vertex with shortest path counting""" diff --git a/python/cugraph/tests/test_connectivity.py b/python/cugraph/tests/test_connectivity.py index b33b6f8e9a3..508be9bb58d 100644 --- a/python/cugraph/tests/test_connectivity.py +++ b/python/cugraph/tests/test_connectivity.py @@ -108,6 +108,15 @@ def cugraph_strong_call(cu_M): return label_vertex_dict +def which_cluster_idx(_cluster, _find_vertex): + idx = -1 + for i in range(len(_cluster)): + if _find_vertex in _cluster[i]: + idx = i + break + return idx + + # Test all combinations of default/managed and pooled/non-pooled allocation @pytest.mark.parametrize("graph_file", utils.DATASETS) def test_weak_cc(graph_file): @@ -142,7 +151,13 @@ def test_weak_cc(graph_file): # Compare vertices of largest component nx_vertices = sorted(lst_nx_components[0]) - cg_vertices = sorted(lst_cg_components[0]) + first_vert = nx_vertices[0] + + idx = which_cluster_idx(lst_cg_components, first_vert) + assert idx != -1, "Check for Nx vertex in cuGraph results failed" + + cg_vertices = sorted(lst_cg_components[idx]) + assert nx_vertices == cg_vertices @@ -167,7 +182,7 @@ def test_strong_cc(graph_file): nx_n_components = len(netx_labels) cg_n_components = len(cugraph_labels) - # Comapre number of components + # Comapre number of components found assert nx_n_components == cg_n_components lst_nx_components = sorted(netx_labels, key=len, reverse=True) @@ -181,6 +196,12 @@ def test_strong_cc(graph_file): assert lst_nx_components_lens == lst_cg_components_lens # Compare vertices of largest component + # note that there might be more than one largest component nx_vertices = sorted(lst_nx_components[0]) - cg_vertices = sorted(lst_cg_components[0]) + first_vert = nx_vertices[0] + + idx = which_cluster_idx(lst_cg_components, first_vert) + assert idx != -1, "Check for Nx vertex in cuGraph results failed" + + cg_vertices = sorted(lst_cg_components[idx]) assert nx_vertices == cg_vertices diff --git a/python/cugraph/tests/test_edge_betweenness_centrality.py b/python/cugraph/tests/test_edge_betweenness_centrality.py index c32f58f1156..e23fdc210ff 100644 --- a/python/cugraph/tests/test_edge_betweenness_centrality.py +++ b/python/cugraph/tests/test_edge_betweenness_centrality.py @@ -42,10 +42,6 @@ NORMALIZED_OPTIONS = [False, True] DEFAULT_EPSILON = 0.0001 -DATASETS = ["../datasets/karate.csv", "../datasets/netscience.csv"] - -UNRENUMBERED_DATASETS = ["../datasets/karate.csv"] - SUBSET_SIZE_OPTIONS = [4, None] SUBSET_SEED_OPTIONS = [42] @@ -297,7 +293,7 @@ def prepare_test(): gc.collect() -@pytest.mark.parametrize("graph_file", DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_SMALL) @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) @pytest.mark.parametrize("subset_size", SUBSET_SIZE_OPTIONS) @pytest.mark.parametrize("normalized", NORMALIZED_OPTIONS) @@ -326,7 +322,7 @@ def test_edge_betweenness_centrality( compare_scores(sorted_df, first_key="cu_bc", second_key="ref_bc") -@pytest.mark.parametrize("graph_file", DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_SMALL) @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) @pytest.mark.parametrize("subset_size", [None]) @pytest.mark.parametrize("normalized", NORMALIZED_OPTIONS) @@ -364,7 +360,7 @@ def test_edge_betweenness_centrality_k_full( # the function operating the comparison inside is first proceeding # to a random sampling over the number of vertices (thus direct offsets) # in the graph structure instead of actual vertices identifiers -@pytest.mark.parametrize("graph_file", UNRENUMBERED_DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_UNRENUMBERED) @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) @pytest.mark.parametrize("subset_size", SUBSET_SIZE_OPTIONS) @pytest.mark.parametrize("normalized", NORMALIZED_OPTIONS) @@ -397,7 +393,7 @@ def test_edge_betweenness_centrality_fixed_sample( compare_scores(sorted_df, first_key="cu_bc", second_key="ref_bc") -@pytest.mark.parametrize("graph_file", DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_SMALL) @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) @pytest.mark.parametrize("subset_size", SUBSET_SIZE_OPTIONS) @pytest.mark.parametrize("normalized", NORMALIZED_OPTIONS) @@ -432,7 +428,7 @@ def test_edge_betweenness_centrality_weight_except( compare_scores(sorted_df, first_key="cu_bc", second_key="ref_bc") -@pytest.mark.parametrize("graph_file", DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_SMALL) @pytest.mark.parametrize("directed", DIRECTED_GRAPH_OPTIONS) @pytest.mark.parametrize("normalized", NORMALIZED_OPTIONS) @pytest.mark.parametrize("subset_size", SUBSET_SIZE_OPTIONS) diff --git a/python/cugraph/tests/test_graph.py b/python/cugraph/tests/test_graph.py index 15aa96ec287..44c856cf3dc 100644 --- a/python/cugraph/tests/test_graph.py +++ b/python/cugraph/tests/test_graph.py @@ -479,7 +479,7 @@ def test_consolidation(graph_file): # Test -@pytest.mark.parametrize('graph_file', utils.DATASETS_2) +@pytest.mark.parametrize('graph_file', utils.DATASETS_SMALL) def test_two_hop_neighbors(graph_file): gc.collect() @@ -593,7 +593,7 @@ def test_number_of_vertices(graph_file): # Test -@pytest.mark.parametrize("graph_file", utils.DATASETS_2) +@pytest.mark.parametrize("graph_file", utils.DATASETS_SMALL) def test_to_directed(graph_file): gc.collect() @@ -622,12 +622,14 @@ def test_to_directed(graph_file): # Test -@pytest.mark.parametrize("graph_file", utils.DATASETS_2) +@pytest.mark.parametrize("graph_file", utils.DATASETS_SMALL) def test_to_undirected(graph_file): gc.collect() + # Read data and then convert to directed by dropped some edges cu_M = utils.read_csv_file(graph_file) cu_M = cu_M[cu_M["0"] <= cu_M["1"]].reset_index(drop=True) + M = utils.read_csv_for_nx(graph_file) M = M[M["0"] <= M["1"]] assert len(cu_M) == len(M) @@ -635,6 +637,7 @@ def test_to_undirected(graph_file): # cugraph add_edge_list DiG = cugraph.DiGraph() DiG.from_cudf_edgelist(cu_M, source="0", destination="1") + DiGnx = nx.from_pandas_edgelist( M, source="0", target="1", create_using=nx.DiGraph() ) @@ -655,7 +658,7 @@ def test_to_undirected(graph_file): # Test -@pytest.mark.parametrize("graph_file", utils.DATASETS_2) +@pytest.mark.parametrize("graph_file", utils.DATASETS) def test_has_edge(graph_file): gc.collect() @@ -672,7 +675,7 @@ def test_has_edge(graph_file): # Test -@pytest.mark.parametrize("graph_file", utils.DATASETS_2) +@pytest.mark.parametrize("graph_file", utils.DATASETS) def test_has_node(graph_file): gc.collect() @@ -688,7 +691,7 @@ def test_has_node(graph_file): # Test all combinations of default/managed and pooled/non-pooled allocation -@pytest.mark.parametrize('graph_file', utils.DATASETS_2) +@pytest.mark.parametrize('graph_file', utils.DATASETS) def test_bipartite_api(graph_file): # This test only tests the functionality of adding set of nodes and # retrieving them. The datasets currently used are not truly bipartite. diff --git a/python/cugraph/tests/test_hits.py b/python/cugraph/tests/test_hits.py index 77471f57601..c8a9274e078 100644 --- a/python/cugraph/tests/test_hits.py +++ b/python/cugraph/tests/test_hits.py @@ -86,8 +86,6 @@ def networkx_call(M, max_iter, tol): return pr -DATASETS = ["../datasets/dolphins.csv", "../datasets/karate.csv"] - MAX_ITERATIONS = [50] TOLERANCE = [1.0e-06] @@ -95,7 +93,7 @@ def networkx_call(M, max_iter, tol): # Test all combinations of default/managed and pooled/non-pooled allocation -@pytest.mark.parametrize("graph_file", DATASETS) +@pytest.mark.parametrize("graph_file", utils.DATASETS_UNDIRECTED) @pytest.mark.parametrize("max_iter", MAX_ITERATIONS) @pytest.mark.parametrize("tol", TOLERANCE) def test_hits(graph_file, max_iter, tol): diff --git a/python/cugraph/tests/test_leiden.py b/python/cugraph/tests/test_leiden.py index a2b7424484d..7f7b4b577fe 100644 --- a/python/cugraph/tests/test_leiden.py +++ b/python/cugraph/tests/test_leiden.py @@ -62,7 +62,7 @@ def cugraph_louvain(cu_M, edgevals=False): return parts, mod -@pytest.mark.parametrize("graph_file", utils.DATASETS_3) +@pytest.mark.parametrize("graph_file", utils.DATASETS) def test_leiden(graph_file): gc.collect() diff --git a/python/cugraph/tests/test_sssp.py b/python/cugraph/tests/test_sssp.py index 5c3d6293dcd..3c3b575fdb5 100644 --- a/python/cugraph/tests/test_sssp.py +++ b/python/cugraph/tests/test_sssp.py @@ -94,7 +94,7 @@ def networkx_call(M, source, edgevals=False): # Test -@pytest.mark.parametrize("graph_file", utils.DATASETS_4) +@pytest.mark.parametrize("graph_file", utils.DATASETS) @pytest.mark.parametrize("source", SOURCES) def test_sssp(graph_file, source): print("DOING test_sssp : " + graph_file + "\n\n\n") @@ -127,7 +127,7 @@ def test_sssp(graph_file, source): # Test -@pytest.mark.parametrize("graph_file", utils.DATASETS_1) +@pytest.mark.parametrize("graph_file", utils.DATASETS) @pytest.mark.parametrize("source", SOURCES) def test_sssp_edgevals(graph_file, source): gc.collect() @@ -159,7 +159,7 @@ def test_sssp_edgevals(graph_file, source): assert err == 0 -@pytest.mark.parametrize("graph_file", utils.DATASETS_1) +@pytest.mark.parametrize("graph_file", utils.DATASETS) @pytest.mark.parametrize("source", SOURCES) def test_sssp_data_type_conversion(graph_file, source): gc.collect() diff --git a/python/cugraph/tests/test_subgraph_extraction.py b/python/cugraph/tests/test_subgraph_extraction.py index 9d9631c0f6d..9192495c6b2 100644 --- a/python/cugraph/tests/test_subgraph_extraction.py +++ b/python/cugraph/tests/test_subgraph_extraction.py @@ -71,7 +71,7 @@ def nx_call(M, verts, directed=True): # Test all combinations of default/managed and pooled/non-pooled allocation -@pytest.mark.parametrize("graph_file", utils.DATASETS_4) +@pytest.mark.parametrize("graph_file", utils.DATASETS) def test_subgraph_extraction_DiGraph(graph_file): gc.collect() @@ -88,7 +88,7 @@ def test_subgraph_extraction_DiGraph(graph_file): # Test all combinations of default/managed and pooled/non-pooled allocation -@pytest.mark.parametrize("graph_file", utils.DATASETS_4) +@pytest.mark.parametrize("graph_file", utils.DATASETS) def test_subgraph_extraction_Graph(graph_file): gc.collect() diff --git a/python/cugraph/tests/utils.py b/python/cugraph/tests/utils.py index fa0631f9c30..e68f934c619 100644 --- a/python/cugraph/tests/utils.py +++ b/python/cugraph/tests/utils.py @@ -20,29 +20,15 @@ from cugraph.dask.common.mg_utils import (get_client) # -# Datasets are numbered based on the number of elements in the array +# Datasets # -DATASETS_1 = ['../datasets/netscience.csv'] +DATASETS_UNDIRECTED = ['../datasets/karate.csv', '../datasets/dolphins.csv'] +DATASETS_UNRENUMBERED = ['../datasets/karate-disjoint.csv'] -DATASETS_2 = ['../datasets/karate.csv', - '../datasets/dolphins.csv'] - -DATASETS_3 = ['../datasets/karate.csv', - '../datasets/dolphins.csv', - '../datasets/email-Eu-core.csv'] - -# FIXME: netscience.csv causes NetworkX pagerank to throw an exception. -# (networkx/algorithms/link_analysis/pagerank_alg.py:152: KeyError: 1532) -DATASETS_4 = ['../datasets/karate.csv', - '../datasets/dolphins.csv', - '../datasets/netscience.csv', - '../datasets/email-Eu-core.csv'] - -DATASETS_5 = ['../datasets/karate.csv', - '../datasets/dolphins.csv', - '../datasets/polbooks.csv', - '../datasets/netscience.csv', - '../datasets/email-Eu-core.csv'] +DATASETS = ['../datasets/karate-disjoint.csv', + '../datasets/dolphins.csv', + '../datasets/netscience.csv'] +# '../datasets/email-Eu-core.csv'] STRONGDATASETS = ['../datasets/dolphins.csv', '../datasets/netscience.csv', @@ -53,19 +39,9 @@ ('../datasets/netscience.csv', '../datasets/ref/ktruss/netscience.csv')] -TINY_DATASETS = ['../datasets/karate.csv', - '../datasets/dolphins.csv', - '../datasets/polbooks.csv'] - -SMALL_DATASETS = ['../datasets/netscience.csv', - '../datasets/email-Eu-core.csv'] - -UNRENUMBERED_DATASETS = ['../datasets/karate.csv'] - - -# define the base for tests to use -DATASETS = DATASETS_3 -DATASETS_UNDIRECTED = DATASETS_2 +DATASETS_SMALL = ['../datasets/karate.csv', + '../datasets/dolphins.csv', + '../datasets/polbooks.csv'] def read_csv_for_nx(csv_file, read_weights_in_sp=True):