From 3f1c7b5ef8cd296aa5ad012cf855abdd2dfeb84b Mon Sep 17 00:00:00 2001 From: Chuck Hastings <45364586+ChuckHastings@users.noreply.github.com> Date: Mon, 20 Nov 2023 21:17:24 -0500 Subject: [PATCH] Update C API graph creation function signatures (#3982) Updating the C API graph creation functions to support the following: * Add support for isolated vertices * Add MG optimization to support multiple device arrays per rank as input and concatenate them internally * Add MG optimization to internally compute the number of edges via allreduce rather than requiring it as an input parameter (this can be expensive to compute in python) This PR implements these features. Some simple tests exist to check for isolate vertices (by running pagerank which generates a different result if the graph has isolated vertices). A simple test for multiple input arrays exists for the MG case. Closes #3947 Closes #3974 Authors: - Chuck Hastings (https://github.com/ChuckHastings) - Naim (https://github.com/naimnv) Approvers: - Naim (https://github.com/naimnv) - Joseph Nke (https://github.com/jnke2016) - Seunghwa Kang (https://github.com/seunghwak) URL: https://github.com/rapidsai/cugraph/pull/3982 --- cpp/CMakeLists.txt | 2 + cpp/include/cugraph/graph_functions.hpp | 67 +++ cpp/include/cugraph_c/graph.h | 193 ++++++- cpp/include/cugraph_c/resource_handle.h | 12 + cpp/src/c_api/graph_mg.cpp | 511 +++++++++++------- cpp/src/c_api/graph_sg.cpp | 170 +++++- cpp/src/c_api/resource_handle.cpp | 9 +- cpp/src/structure/detail/structure_utils.cuh | 61 ++- cpp/src/structure/remove_multi_edges.cu | 92 ++++ cpp/src/structure/remove_multi_edges_impl.cuh | 310 +++++++++++ cpp/src/structure/remove_self_loops.cu | 92 ++++ cpp/src/structure/remove_self_loops_impl.cuh | 94 ++++ cpp/tests/c_api/create_graph_test.c | 498 ++++++++++++++++- cpp/tests/c_api/mg_create_graph_test.c | 400 +++++++++++++- 14 files changed, 2265 insertions(+), 246 deletions(-) create mode 100644 cpp/src/structure/remove_multi_edges.cu create mode 100644 cpp/src/structure/remove_multi_edges_impl.cuh create mode 100644 cpp/src/structure/remove_self_loops.cu create mode 100644 cpp/src/structure/remove_self_loops_impl.cuh diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 626d62cffa5..836d5569ef7 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -202,6 +202,8 @@ set(CUGRAPH_SOURCES src/community/detail/mis_mg.cu src/detail/utility_wrappers.cu src/structure/graph_view_mg.cu + src/structure/remove_self_loops.cu + src/structure/remove_multi_edges.cu src/utilities/path_retrieval.cu src/structure/legacy/graph.cu src/linear_assignment/legacy/hungarian.cu diff --git a/cpp/include/cugraph/graph_functions.hpp b/cpp/include/cugraph/graph_functions.hpp index 5c1e9d5311f..6a75a420bf8 100644 --- a/cpp/include/cugraph/graph_functions.hpp +++ b/cpp/include/cugraph/graph_functions.hpp @@ -973,4 +973,71 @@ renumber_sampled_edgelist( label_offsets, bool do_expensive_check = false); +/** + * @brief Remove self loops from an edge list + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam edge_t Type of edge identifiers. Needs to be an integral type. + * @tparam weight_t Type of edge weight. Currently float and double are supported. + * @tparam edge_type_t Type of edge type. Needs to be an integral type. + * + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param edgelist_srcs List of source vertex ids + * @param edgelist_dsts List of destination vertex ids + * @param edgelist_weights Optional list of edge weights + * @param edgelist_edge_ids Optional list of edge ids + * @param edgelist_edge_types Optional list of edge types + * @return Tuple of vectors storing edge sources, destinations, optional weights, + * optional edge ids, optional edge types. + */ +template +std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_self_loops(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +/** + * @brief Remove all but one edge when a multi-edge exists. Note that this function does not use + * stable methods. When a multi-edge exists, one of the edges will remain, there is no + * guarantee on which one will remain. + * + * In an MG context it is assumed that edges have been shuffled to the proper GPU, + * in which case any multi-edges will be on the same GPU. + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam edge_t Type of edge identifiers. Needs to be an integral type. + * @tparam weight_t Type of edge weight. Currently float and double are supported. + * @tparam edge_type_t Type of edge type. Needs to be an integral type. + * + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param edgelist_srcs List of source vertex ids + * @param edgelist_dsts List of destination vertex ids + * @param edgelist_weights Optional list of edge weights + * @param edgelist_edge_ids Optional list of edge ids + * @param edgelist_edge_types Optional list of edge types + * @return Tuple of vectors storing edge sources, destinations, optional weights, + * optional edge ids, optional edge types. + */ +template +std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_multi_edges(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + } // namespace cugraph diff --git a/cpp/include/cugraph_c/graph.h b/cpp/include/cugraph_c/graph.h index e910d8b1244..88176a9c1b6 100644 --- a/cpp/include/cugraph_c/graph.h +++ b/cpp/include/cugraph_c/graph.h @@ -35,10 +35,11 @@ typedef struct { bool_t is_multigraph; } cugraph_graph_properties_t; -// FIXME: Add support for specifying isolated vertices /** * @brief Construct an SG graph * + * @deprecated This API will be deleted, use cugraph_graph_create_sg instead + * * @param [in] handle Handle for accessing resources * @param [in] properties Properties of the constructed graph * @param [in] src Device array containing the source vertex ids. @@ -51,11 +52,11 @@ typedef struct { argument that can be NULL if edge types are not used. * @param [in] store_transposed If true create the graph initially in transposed format * @param [in] renumber If true, renumber vertices to make an efficient data structure. - * If false, do not renumber. Renumbering is required if the vertices are not sequential - * integer values from 0 to num_vertices. + * If false, do not renumber. Renumbering enables some significant optimizations within + * the graph primitives library, so it is strongly encouraged. Renumbering is required if + * the vertices are not sequential integer values from 0 to num_vertices. * @param [in] do_expensive_check If true, do expensive checks to validate the input data * is consistent with software assumptions. If false bypass these checks. - * @param [in] properties Properties of the graph * @param [out] graph A pointer to the graph object * @param [out] error Pointer to an error object storing details of any error. Will * be populated if error code is not CUGRAPH_SUCCESS @@ -76,9 +77,63 @@ cugraph_error_code_t cugraph_sg_graph_create( cugraph_graph_t** graph, cugraph_error_t** error); +/** + * @brief Construct an SG graph + * + * @param [in] handle Handle for accessing resources + * @param [in] properties Properties of the constructed graph + * @param [in] vertices Optional device array containing a list of vertex ids + * (specify NULL if we should create vertex ids from the + * unique contents of @p src and @p dst) + * @param [in] src Device array containing the source vertex ids. + * @param [in] dst Device array containing the destination vertex ids + * @param [in] weights Device array containing the edge weights. Note that an unweighted + * graph can be created by passing weights == NULL. + * @param [in] edge_ids Device array containing the edge ids for each edge. Optional + argument that can be NULL if edge ids are not used. + * @param [in] edge_type_ids Device array containing the edge types for each edge. Optional + argument that can be NULL if edge types are not used. + * @param [in] store_transposed If true create the graph initially in transposed format + * @param [in] renumber If true, renumber vertices to make an efficient data structure. + * If false, do not renumber. Renumbering enables some significant optimizations within + * the graph primitives library, so it is strongly encouraged. Renumbering is required if + * the vertices are not sequential integer values from 0 to num_vertices. + * @param [in] drop_self_loops If true, drop any self loops that exist in the provided edge list. + * @param [in] drop_multi_edges If true, drop any multi edges that exist in the provided edge list. + * Note that setting this flag will arbitrarily select one instance of a multi edge to be the + * edge that survives. If the edges have properties that should be honored (e.g. sum the + weights, + * or take the maximum weight), the caller should do that on not rely on this flag. + * @param [in] do_expensive_check If true, do expensive checks to validate the input data + * is consistent with software assumptions. If false bypass these checks. + * @param [out] graph A pointer to the graph object + * @param [out] error Pointer to an error object storing details of any error. Will + * be populated if error code is not CUGRAPH_SUCCESS + * + * @return error code + */ +cugraph_error_code_t cugraph_graph_create_sg( + const cugraph_resource_handle_t* handle, + const cugraph_graph_properties_t* properties, + const cugraph_type_erased_device_array_view_t* vertices, + const cugraph_type_erased_device_array_view_t* src, + const cugraph_type_erased_device_array_view_t* dst, + const cugraph_type_erased_device_array_view_t* weights, + const cugraph_type_erased_device_array_view_t* edge_ids, + const cugraph_type_erased_device_array_view_t* edge_type_ids, + bool_t store_transposed, + bool_t renumber, + bool_t drop_self_loops, + bool_t drop_multi_edges, + bool_t do_expensive_check, + cugraph_graph_t** graph, + cugraph_error_t** error); + /** * @brief Construct an SG graph from a CSR input * + * @deprecated This API will be deleted, use cugraph_graph_create_sg_from_csr instead + * * @param [in] handle Handle for accessing resources * @param [in] properties Properties of the constructed graph * @param [in] offsets Device array containing the CSR offsets array @@ -91,11 +146,11 @@ cugraph_error_code_t cugraph_sg_graph_create( argument that can be NULL if edge types are not used. * @param [in] store_transposed If true create the graph initially in transposed format * @param [in] renumber If true, renumber vertices to make an efficient data structure. - * If false, do not renumber. Renumbering is required if the vertices are not sequential - * integer values from 0 to num_vertices. + * If false, do not renumber. Renumbering enables some significant optimizations within + * the graph primitives library, so it is strongly encouraged. Renumbering is required if + * the vertices are not sequential integer values from 0 to num_vertices. * @param [in] do_expensive_check If true, do expensive checks to validate the input data * is consistent with software assumptions. If false bypass these checks. - * @param [in] properties Properties of the graph * @param [out] graph A pointer to the graph object * @param [out] error Pointer to an error object storing details of any error. Will * be populated if error code is not CUGRAPH_SUCCESS @@ -117,18 +172,50 @@ cugraph_error_code_t cugraph_sg_graph_create_from_csr( cugraph_error_t** error); /** - * @brief Destroy an SG graph + * @brief Construct an SG graph from a CSR input * - * @param [in] graph A pointer to the graph object to destroy + * @param [in] handle Handle for accessing resources + * @param [in] properties Properties of the constructed graph + * @param [in] offsets Device array containing the CSR offsets array + * @param [in] indices Device array containing the destination vertex ids + * @param [in] weights Device array containing the edge weights. Note that an unweighted + * graph can be created by passing weights == NULL. + * @param [in] edge_ids Device array containing the edge ids for each edge. Optional + argument that can be NULL if edge ids are not used. + * @param [in] edge_type_ids Device array containing the edge types for each edge. Optional + argument that can be NULL if edge types are not used. + * @param [in] store_transposed If true create the graph initially in transposed format + * @param [in] renumber If true, renumber vertices to make an efficient data structure. + * If false, do not renumber. Renumbering enables some significant optimizations within + * the graph primitives library, so it is strongly encouraged. Renumbering is required if + * the vertices are not sequential integer values from 0 to num_vertices. + * @param [in] do_expensive_check If true, do expensive checks to validate the input data + * is consistent with software assumptions. If false bypass these checks. + * @param [out] graph A pointer to the graph object + * @param [out] error Pointer to an error object storing details of any error. Will + * be populated if error code is not CUGRAPH_SUCCESS + * + * @return error code */ -// FIXME: This should probably just be cugraph_graph_free -// but didn't want to confuse with original cugraph_free_graph -void cugraph_sg_graph_free(cugraph_graph_t* graph); +cugraph_error_code_t cugraph_graph_create_sg_from_csr( + const cugraph_resource_handle_t* handle, + const cugraph_graph_properties_t* properties, + const cugraph_type_erased_device_array_view_t* offsets, + const cugraph_type_erased_device_array_view_t* indices, + const cugraph_type_erased_device_array_view_t* weights, + const cugraph_type_erased_device_array_view_t* edge_ids, + const cugraph_type_erased_device_array_view_t* edge_type_ids, + bool_t store_transposed, + bool_t renumber, + bool_t do_expensive_check, + cugraph_graph_t** graph, + cugraph_error_t** error); -// FIXME: Add support for specifying isolated vertices /** * @brief Construct an MG graph * + * @deprecated This API will be deleted, use cugraph_graph_create_mg instead + * * @param [in] handle Handle for accessing resources * @param [in] properties Properties of the constructed graph * @param [in] src Device array containing the source vertex ids @@ -165,13 +252,89 @@ cugraph_error_code_t cugraph_mg_graph_create( cugraph_graph_t** graph, cugraph_error_t** error); +/** + * @brief Construct an MG graph + * + * @param [in] handle Handle for accessing resources + * @param [in] properties Properties of the constructed graph + * @param [in] vertices List of device arrays containing the unique vertex ids. + * If NULL we will construct this internally using the unique + * entries specified in src and dst + * All entries in this list will be concatenated on this GPU + * into a single array. + * @param [in] src List of device array containing the source vertex ids + * All entries in this list will be concatenated on this GPU + * into a single array. + * @param [in] dst List of device array containing the destination vertex ids + * All entries in this list will be concatenated on this GPU + * into a single array. + * @param [in] weights List of device array containing the edge weights. Note that an + * unweighted graph can be created by passing weights == NULL. If a weighted graph is to be + * created, the weights device array should be created on each rank, but the pointer can be NULL and + * the size 0 if there are no inputs provided by this rank All entries in this list will be + * concatenated on this GPU into a single array. + * @param [in] edge_ids List of device array containing the edge ids for each edge. Optional + * argument that can be NULL if edge ids are not used. + * All entries in this list will be concatenated on this GPU + * into a single array. + * @param [in] edge_type_ids List of device array containing the edge types for each edge. + * Optional argument that can be NULL if edge types are not used. All entries in this list will be + * concatenated on this GPU into a single array. + * @param [in] store_transposed If true create the graph initially in transposed format + * @param [in] num_arrays The number of arrays specified in @p vertices, @p src, @p dst, @p + * weights, @p edge_ids and @p edge_type_ids + * @param [in] drop_self_loops If true, drop any self loops that exist in the provided edge list. + * @param [in] drop_multi_edges If true, drop any multi edges that exist in the provided edge list. + * Note that setting this flag will arbitrarily select one instance of a multi edge to be the + * edge that survives. If the edges have properties that should be honored (e.g. sum the + * weights, or take the maximum weight), the caller should do that on not rely on this flag. + * @param [in] do_expensive_check If true, do expensive checks to validate the input data + * is consistent with software assumptions. If false bypass these checks. + * @param [out] graph A pointer to the graph object + * @param [out] error Pointer to an error object storing details of any error. Will + * be populated if error code is not CUGRAPH_SUCCESS + * @return error code + */ +cugraph_error_code_t cugraph_graph_create_mg( + cugraph_resource_handle_t const* handle, + cugraph_graph_properties_t const* properties, + cugraph_type_erased_device_array_view_t const* const* vertices, + cugraph_type_erased_device_array_view_t const* const* src, + cugraph_type_erased_device_array_view_t const* const* dst, + cugraph_type_erased_device_array_view_t const* const* weights, + cugraph_type_erased_device_array_view_t const* const* edge_ids, + cugraph_type_erased_device_array_view_t const* const* edge_type_ids, + bool_t store_transposed, + size_t num_arrays, + bool_t drop_self_loops, + bool_t drop_multi_edges, + bool_t do_expensive_check, + cugraph_graph_t** graph, + cugraph_error_t** error); + +/** + * @brief Destroy an graph + * + * @param [in] graph A pointer to the graph object to destroy + */ +void cugraph_graph_free(cugraph_graph_t* graph); + +/** + * @brief Destroy an SG graph + * + * @deprecated This API will be deleted, use cugraph_graph_free instead + * + * @param [in] graph A pointer to the graph object to destroy + */ +void cugraph_sg_graph_free(cugraph_graph_t* graph); + /** * @brief Destroy an MG graph * + * @deprecated This API will be deleted, use cugraph_graph_free instead + * * @param [in] graph A pointer to the graph object to destroy */ -// FIXME: This should probably just be cugraph_graph_free -// but didn't want to confuse with original cugraph_free_graph void cugraph_mg_graph_free(cugraph_graph_t* graph); /** diff --git a/cpp/include/cugraph_c/resource_handle.h b/cpp/include/cugraph_c/resource_handle.h index a239c24afe9..0e45102aae2 100644 --- a/cpp/include/cugraph_c/resource_handle.h +++ b/cpp/include/cugraph_c/resource_handle.h @@ -57,6 +57,18 @@ typedef struct cugraph_resource_handle_ { */ cugraph_resource_handle_t* cugraph_create_resource_handle(void* raft_handle); +/** + * @brief get comm_size from resource handle + * + * If the resource handle has been configured for multi-gpu, this will return + * the comm_size for this cluster. If the resource handle has not been configured for + * multi-gpu this will always return 1. + * + * @param [in] handle Handle for accessing resources + * @return comm_size + */ +int cugraph_resource_handle_get_comm_size(const cugraph_resource_handle_t* handle); + /** * @brief get rank from resource handle * diff --git a/cpp/src/c_api/graph_mg.cpp b/cpp/src/c_api/graph_mg.cpp index f50c7c08fb6..5413949e3a3 100644 --- a/cpp/src/c_api/graph_mg.cpp +++ b/cpp/src/c_api/graph_mg.cpp @@ -31,40 +31,85 @@ namespace { +template +rmm::device_uvector concatenate( + raft::handle_t const& handle, + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* values, + size_t num_arrays) +{ + size_t num_values = std::transform_reduce( + values, values + num_arrays, size_t{0}, std::plus{}, [](auto p) { return p->size_; }); + + rmm::device_uvector results(num_values, handle.get_stream()); + size_t concat_pos{0}; + + for (size_t i = 0; i < num_arrays; ++i) { + raft::copy(results.data() + concat_pos, + values[i]->as_type(), + values[i]->size_, + handle.get_stream()); + concat_pos += values[i]->size_; + } + + return results; +} + struct create_graph_functor : public cugraph::c_api::abstract_functor { raft::handle_t const& handle_; cugraph_graph_properties_t const* properties_; - cugraph::c_api::cugraph_type_erased_device_array_view_t const* src_; - cugraph::c_api::cugraph_type_erased_device_array_view_t const* dst_; - cugraph::c_api::cugraph_type_erased_device_array_view_t const* weights_; - cugraph::c_api::cugraph_type_erased_device_array_view_t const* edge_ids_; - cugraph::c_api::cugraph_type_erased_device_array_view_t const* edge_type_ids_; - bool_t renumber_; - bool_t check_; + cugraph_data_type_id_t vertex_type_; cugraph_data_type_id_t edge_type_; + cugraph_data_type_id_t weight_type_; + cugraph_data_type_id_t edge_type_id_type_; + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* vertices_; + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* src_; + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* dst_; + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* weights_; + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* edge_ids_; + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* edge_type_ids_; + size_t num_arrays_; + bool_t renumber_; + bool_t drop_self_loops_; + bool_t drop_multi_edges_; + bool_t do_expensive_check_; cugraph::c_api::cugraph_graph_t* result_{}; - create_graph_functor(raft::handle_t const& handle, - cugraph_graph_properties_t const* properties, - cugraph::c_api::cugraph_type_erased_device_array_view_t const* src, - cugraph::c_api::cugraph_type_erased_device_array_view_t const* dst, - cugraph::c_api::cugraph_type_erased_device_array_view_t const* weights, - cugraph::c_api::cugraph_type_erased_device_array_view_t const* edge_ids, - cugraph::c_api::cugraph_type_erased_device_array_view_t const* edge_type_ids, - bool_t renumber, - bool_t check, - cugraph_data_type_id_t edge_type) + create_graph_functor( + raft::handle_t const& handle, + cugraph_graph_properties_t const* properties, + cugraph_data_type_id_t vertex_type, + cugraph_data_type_id_t edge_type, + cugraph_data_type_id_t weight_type, + cugraph_data_type_id_t edge_type_id_type, + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* vertices, + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* src, + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* dst, + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* weights, + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* edge_ids, + cugraph::c_api::cugraph_type_erased_device_array_view_t const* const* edge_type_ids, + size_t num_arrays, + bool_t renumber, + bool_t drop_self_loops, + bool_t drop_multi_edges, + bool_t do_expensive_check) : abstract_functor(), properties_(properties), + vertex_type_(vertex_type), + edge_type_(edge_type), + weight_type_(weight_type), + edge_type_id_type_(edge_type_id_type), handle_(handle), + vertices_(vertices), src_(src), dst_(dst), weights_(weights), edge_ids_(edge_ids), edge_type_ids_(edge_type_ids), + num_arrays_(num_arrays), renumber_(renumber), - check_(check), - edge_type_(edge_type) + drop_self_loops_(drop_self_loops), + drop_multi_edges_(drop_multi_edges), + do_expensive_check_(do_expensive_check) { } @@ -96,49 +141,27 @@ struct create_graph_functor : public cugraph::c_api::abstract_functor { edge_type_id_t>> new_edge_types{std::nullopt}; - rmm::device_uvector edgelist_srcs(src_->size_, handle_.get_stream()); - rmm::device_uvector edgelist_dsts(dst_->size_, handle_.get_stream()); + std::optional> vertex_list = + vertices_ ? std::make_optional(concatenate(handle_, vertices_, num_arrays_)) + : std::nullopt; - raft::copy( - edgelist_srcs.data(), src_->as_type(), src_->size_, handle_.get_stream()); - raft::copy( - edgelist_dsts.data(), dst_->as_type(), dst_->size_, handle_.get_stream()); + rmm::device_uvector edgelist_srcs = + concatenate(handle_, src_, num_arrays_); + rmm::device_uvector edgelist_dsts = + concatenate(handle_, dst_, num_arrays_); std::optional> edgelist_weights = - weights_ - ? std::make_optional(rmm::device_uvector(weights_->size_, handle_.get_stream())) - : std::nullopt; - - if (edgelist_weights) { - raft::copy(edgelist_weights->data(), - weights_->as_type(), - weights_->size_, - handle_.get_stream()); - } + weights_ ? std::make_optional(concatenate(handle_, weights_, num_arrays_)) + : std::nullopt; std::optional> edgelist_edge_ids = - edge_ids_ - ? std::make_optional(rmm::device_uvector(edge_ids_->size_, handle_.get_stream())) - : std::nullopt; - - if (edgelist_edge_ids) { - raft::copy(edgelist_edge_ids->data(), - edge_ids_->as_type(), - edge_ids_->size_, - handle_.get_stream()); - } + edge_ids_ ? std::make_optional(concatenate(handle_, edge_ids_, num_arrays_)) + : std::nullopt; std::optional> edgelist_edge_types = - edge_type_ids_ ? std::make_optional(rmm::device_uvector( - edge_type_ids_->size_, handle_.get_stream())) - : std::nullopt; - - if (edgelist_edge_types) { - raft::copy(edgelist_edge_types->data(), - edge_type_ids_->as_type(), - edge_type_ids_->size_, - handle_.get_stream()); - } + edge_type_ids_ + ? std::make_optional(concatenate(handle_, edge_type_ids_, num_arrays_)) + : std::nullopt; std::tie(store_transposed ? edgelist_dsts : edgelist_srcs, store_transposed ? edgelist_srcs : edgelist_dsts, @@ -153,6 +176,11 @@ struct create_graph_functor : public cugraph::c_api::abstract_functor { std::move(edgelist_edge_ids), std::move(edgelist_edge_types)); + if (vertex_list) { + vertex_list = cugraph::detail::shuffle_ext_vertices_to_local_gpu_by_vertex_partitioning( + handle_, std::move(*vertex_list)); + } + auto graph = new cugraph::graph_t(handle_); rmm::device_uvector* number_map = @@ -170,6 +198,28 @@ struct create_graph_functor : public cugraph::c_api::abstract_functor { cugraph::graph_view_t, edge_type_id_t>(handle_); + if (drop_self_loops_) { + std::tie( + edgelist_srcs, edgelist_dsts, edgelist_weights, edgelist_edge_ids, edgelist_edge_types) = + cugraph::remove_self_loops(handle_, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::move(edgelist_weights), + std::move(edgelist_edge_ids), + std::move(edgelist_edge_types)); + } + + if (drop_multi_edges_) { + std::tie( + edgelist_srcs, edgelist_dsts, edgelist_weights, edgelist_edge_ids, edgelist_edge_types) = + cugraph::remove_multi_edges(handle_, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::move(edgelist_weights), + std::move(edgelist_edge_ids), + std::move(edgelist_edge_types)); + } + std::tie(*graph, new_edge_weights, new_edge_ids, new_edge_types, new_number_map) = cugraph::create_graph_from_edgelist( handle_, - std::nullopt, + std::move(vertex_list), std::move(edgelist_srcs), std::move(edgelist_dsts), std::move(edgelist_weights), @@ -187,7 +237,7 @@ struct create_graph_functor : public cugraph::c_api::abstract_functor { std::move(edgelist_edge_types), cugraph::graph_properties_t{properties_->is_symmetric, properties_->is_multigraph}, renumber_, - check_); + do_expensive_check_); if (renumber_) { *number_map = std::move(new_number_map.value()); @@ -204,90 +254,39 @@ struct create_graph_functor : public cugraph::c_api::abstract_functor { if (new_edge_types) { *edge_types = std::move(new_edge_types.value()); } // Set up return - auto result = new cugraph::c_api::cugraph_graph_t{ - src_->type_, - edge_type_, - weights_ ? weights_->type_ : cugraph_data_type_id_t::FLOAT32, - edge_type_ids_ ? edge_type_ids_->type_ : cugraph_data_type_id_t::INT32, - store_transposed, - multi_gpu, - graph, - number_map, - new_edge_weights ? edge_weights : nullptr, - new_edge_ids ? edge_ids : nullptr, - new_edge_types ? edge_types : nullptr}; + auto result = new cugraph::c_api::cugraph_graph_t{vertex_type_, + edge_type_, + weight_type_, + edge_type_id_type_, + store_transposed, + multi_gpu, + graph, + number_map, + new_edge_weights ? edge_weights : nullptr, + new_edge_ids ? edge_ids : nullptr, + new_edge_types ? edge_types : nullptr}; result_ = reinterpret_cast(result); } } }; -struct destroy_graph_functor : public cugraph::c_api::abstract_functor { - void* graph_; - void* number_map_; - void* edge_weights_; - void* edge_ids_; - void* edge_types_; - - destroy_graph_functor( - void* graph, void* number_map, void* edge_weights, void* edge_ids, void* edge_types) - : abstract_functor(), - graph_(graph), - number_map_(number_map), - edge_weights_(edge_weights), - edge_ids_(edge_ids), - edge_types_(edge_types) - { - } - - template - void operator()() - { - auto internal_graph_pointer = - reinterpret_cast*>(graph_); - - delete internal_graph_pointer; - - auto internal_number_map_pointer = - reinterpret_cast*>(number_map_); - - delete internal_number_map_pointer; - - auto internal_edge_weight_pointer = reinterpret_cast< - cugraph::edge_property_t, - weight_t>*>(edge_weights_); - if (internal_edge_weight_pointer) { delete internal_edge_weight_pointer; } - - auto internal_edge_id_pointer = reinterpret_cast< - cugraph::edge_property_t, - edge_t>*>(edge_ids_); - if (internal_edge_id_pointer) { delete internal_edge_id_pointer; } - - auto internal_edge_type_pointer = reinterpret_cast< - cugraph::edge_property_t, - edge_type_id_t>*>(edge_types_); - if (internal_edge_type_pointer) { delete internal_edge_type_pointer; } - } -}; - } // namespace -extern "C" cugraph_error_code_t cugraph_mg_graph_create( - const cugraph_resource_handle_t* handle, - const cugraph_graph_properties_t* properties, - const cugraph_type_erased_device_array_view_t* src, - const cugraph_type_erased_device_array_view_t* dst, - const cugraph_type_erased_device_array_view_t* weights, - const cugraph_type_erased_device_array_view_t* edge_ids, - const cugraph_type_erased_device_array_view_t* edge_type_ids, +extern "C" cugraph_error_code_t cugraph_graph_create_mg( + cugraph_resource_handle_t const* handle, + cugraph_graph_properties_t const* properties, + cugraph_type_erased_device_array_view_t const* const* vertices, + cugraph_type_erased_device_array_view_t const* const* src, + cugraph_type_erased_device_array_view_t const* const* dst, + cugraph_type_erased_device_array_view_t const* const* weights, + cugraph_type_erased_device_array_view_t const* const* edge_ids, + cugraph_type_erased_device_array_view_t const* const* edge_type_ids, bool_t store_transposed, - size_t num_edges, - bool_t check, + size_t num_arrays, + bool_t drop_self_loops, + bool_t drop_multi_edges, + bool_t do_expensive_check, cugraph_graph_t** graph, cugraph_error_t** error) { @@ -298,87 +297,198 @@ extern "C" cugraph_error_code_t cugraph_mg_graph_create( *error = nullptr; auto p_handle = reinterpret_cast(handle); + auto p_vertices = + reinterpret_cast( + vertices); auto p_src = - reinterpret_cast(src); + reinterpret_cast(src); auto p_dst = - reinterpret_cast(dst); + reinterpret_cast(dst); auto p_weights = - reinterpret_cast(weights); + reinterpret_cast( + weights); auto p_edge_ids = - reinterpret_cast(edge_ids); + reinterpret_cast( + edge_ids); auto p_edge_type_ids = - reinterpret_cast(edge_type_ids); + reinterpret_cast( + edge_type_ids); + + size_t local_num_edges{0}; + + // + // Determine the type of vertex, weight, edge_type_id across + // multiple input arrays and acros multiple GPUs. Also compute + // the number of edges so we can determine what type to use for + // edge_t + // + cugraph_data_type_id_t vertex_type{cugraph_data_type_id_t::NTYPES}; + cugraph_data_type_id_t weight_type{cugraph_data_type_id_t::NTYPES}; + + for (size_t i = 0; i < num_arrays; ++i) { + CAPI_EXPECTS(p_src[i]->size_ == p_dst[i]->size_, + CUGRAPH_INVALID_INPUT, + "Invalid input arguments: src size != dst size.", + *error); + + CAPI_EXPECTS(p_src[i]->type_ == p_dst[i]->type_, + CUGRAPH_INVALID_INPUT, + "Invalid input arguments: src type != dst type.", + *error); + + CAPI_EXPECTS((p_vertices == nullptr) || (p_src[i]->type_ == p_vertices[i]->type_), + CUGRAPH_INVALID_INPUT, + "Invalid input arguments: src type != vertices type.", + *error); + + CAPI_EXPECTS((weights == nullptr) || (p_weights[i]->size_ == p_src[i]->size_), + CUGRAPH_INVALID_INPUT, + "Invalid input arguments: src size != weights size.", + *error); + + local_num_edges += p_src[i]->size_; + + if (vertex_type == cugraph_data_type_id_t::NTYPES) vertex_type = p_src[i]->type_; + + if (weights != nullptr) { + if (weight_type == cugraph_data_type_id_t::NTYPES) weight_type = p_weights[i]->type_; + } - CAPI_EXPECTS(p_src->size_ == p_dst->size_, - CUGRAPH_INVALID_INPUT, - "Invalid input arguments: src size != dst size.", - *error); - CAPI_EXPECTS(p_src->type_ == p_dst->type_, + CAPI_EXPECTS(p_src[i]->type_ == vertex_type, + CUGRAPH_INVALID_INPUT, + "Invalid input arguments: all vertex types must match", + *error); + + CAPI_EXPECTS((weights == nullptr) || (p_weights[i]->type_ == weight_type), + CUGRAPH_INVALID_INPUT, + "Invalid input arguments: all weight types must match", + *error); + } + + size_t num_edges = cugraph::host_scalar_allreduce(p_handle->handle_->get_comms(), + local_num_edges, + raft::comms::op_t::SUM, + p_handle->handle_->get_stream()); + + auto vertex_types = cugraph::host_scalar_allgather( + p_handle->handle_->get_comms(), static_cast(vertex_type), p_handle->handle_->get_stream()); + + auto weight_types = cugraph::host_scalar_allgather( + p_handle->handle_->get_comms(), static_cast(weight_type), p_handle->handle_->get_stream()); + + if (vertex_type == cugraph_data_type_id_t::NTYPES) { + // Only true if this GPU had no vertex arrays + vertex_type = static_cast( + *std::min_element(vertex_types.begin(), vertex_types.end())); + } + + if (weight_type == cugraph_data_type_id_t::NTYPES) { + // Only true if this GPU had no weight arrays + weight_type = static_cast( + *std::min_element(weight_types.begin(), weight_types.end())); + } + + CAPI_EXPECTS(std::all_of(vertex_types.begin(), + vertex_types.end(), + [vertex_type](auto t) { return vertex_type == static_cast(t); }), CUGRAPH_INVALID_INPUT, - "Invalid input arguments: src type != dst type.", + "different vertex type used on different GPUs", *error); - CAPI_EXPECTS((weights == nullptr) || (p_weights->size_ == p_src->size_), + CAPI_EXPECTS(std::all_of(weight_types.begin(), + weight_types.end(), + [weight_type](auto t) { return weight_type == static_cast(t); }), CUGRAPH_INVALID_INPUT, - "Invalid input arguments: src size != weights size.", + "different weight type used on different GPUs", *error); cugraph_data_type_id_t edge_type; - cugraph_data_type_id_t weight_type; if (num_edges < int32_threshold) { - edge_type = p_src->type_; + edge_type = static_cast(vertex_types[0]); } else { edge_type = cugraph_data_type_id_t::INT64; } - if (weights != nullptr) { - weight_type = p_weights->type_; - } else { + if (weight_type == cugraph_data_type_id_t::NTYPES) { weight_type = cugraph_data_type_id_t::FLOAT32; } - CAPI_EXPECTS((edge_ids == nullptr) || (p_edge_ids->type_ == edge_type), - CUGRAPH_INVALID_INPUT, - "Invalid input arguments: Edge id type must match edge type", - *error); + cugraph_data_type_id_t edge_type_id_type{cugraph_data_type_id_t::NTYPES}; - CAPI_EXPECTS((edge_ids == nullptr) || (p_edge_ids->size_ == p_src->size_), - CUGRAPH_INVALID_INPUT, - "Invalid input arguments: src size != edge id prop size", - *error); + for (size_t i = 0; i < num_arrays; ++i) { + CAPI_EXPECTS((edge_ids == nullptr) || (p_edge_ids[i]->type_ == edge_type), + CUGRAPH_INVALID_INPUT, + "Invalid input arguments: Edge id type must match edge type", + *error); - CAPI_EXPECTS((edge_type_ids == nullptr) || (p_edge_type_ids->size_ == p_src->size_), - CUGRAPH_INVALID_INPUT, - "Invalid input arguments: src size != edge type prop size", - *error); + CAPI_EXPECTS((edge_ids == nullptr) || (p_edge_ids[i]->size_ == p_src[i]->size_), + CUGRAPH_INVALID_INPUT, + "Invalid input arguments: src size != edge id prop size", + *error); + + if (edge_type_ids != nullptr) { + CAPI_EXPECTS(p_edge_type_ids[i]->size_ == p_src[i]->size_, + CUGRAPH_INVALID_INPUT, + "Invalid input arguments: src size != edge type prop size", + *error); + + if (edge_type_id_type == cugraph_data_type_id_t::NTYPES) + edge_type_id_type = p_edge_type_ids[i]->type_; + + CAPI_EXPECTS(p_edge_type_ids[i]->type_ == edge_type_id_type, + CUGRAPH_INVALID_INPUT, + "Invalid input arguments: src size != edge type prop size", + *error); + } + } + + auto edge_type_id_types = cugraph::host_scalar_allgather(p_handle->handle_->get_comms(), + static_cast(edge_type_id_type), + p_handle->handle_->get_stream()); + + if (edge_type_id_type == cugraph_data_type_id_t::NTYPES) { + // Only true if this GPU had no edge_type_id arrays + edge_type_id_type = static_cast( + *std::min_element(edge_type_id_types.begin(), edge_type_id_types.end())); + } + + CAPI_EXPECTS( + std::all_of(edge_type_id_types.begin(), + edge_type_id_types.end(), + [edge_type_id_type](auto t) { return edge_type_id_type == static_cast(t); }), + CUGRAPH_INVALID_INPUT, + "different edge_type_id type used on different GPUs", + *error); - cugraph_data_type_id_t edge_type_id_type; - if (edge_type_ids == nullptr) { + if (edge_type_id_type == cugraph_data_type_id_t::NTYPES) { edge_type_id_type = cugraph_data_type_id_t::INT32; - } else { - edge_type_id_type = p_edge_type_ids->type_; } + // + // Now we know enough to create the graph + // create_graph_functor functor(*p_handle->handle_, properties, + vertex_type, + edge_type, + weight_type, + edge_type_id_type, + p_vertices, p_src, p_dst, p_weights, p_edge_ids, p_edge_type_ids, + num_arrays, bool_t::TRUE, - check, - edge_type); + drop_self_loops, + drop_multi_edges, + do_expensive_check); try { - cugraph::c_api::vertex_dispatcher(p_src->type_, - edge_type, - weight_type, - edge_type_id_type, - store_transposed, - multi_gpu, - functor); + cugraph::c_api::vertex_dispatcher( + vertex_type, edge_type, weight_type, edge_type_id_type, store_transposed, multi_gpu, functor); if (functor.error_code_ != CUGRAPH_SUCCESS) { *error = reinterpret_cast(functor.error_.release()); @@ -394,25 +504,38 @@ extern "C" cugraph_error_code_t cugraph_mg_graph_create( return CUGRAPH_SUCCESS; } +extern "C" cugraph_error_code_t cugraph_mg_graph_create( + cugraph_resource_handle_t const* handle, + cugraph_graph_properties_t const* properties, + cugraph_type_erased_device_array_view_t const* src, + cugraph_type_erased_device_array_view_t const* dst, + cugraph_type_erased_device_array_view_t const* weights, + cugraph_type_erased_device_array_view_t const* edge_ids, + cugraph_type_erased_device_array_view_t const* edge_type_ids, + bool_t store_transposed, + size_t num_edges, + bool_t do_expensive_check, + cugraph_graph_t** graph, + cugraph_error_t** error) +{ + return cugraph_graph_create_mg(handle, + properties, + NULL, + &src, + &dst, + &weights, + &edge_ids, + &edge_type_ids, + store_transposed, + 1, + FALSE, + FALSE, + do_expensive_check, + graph, + error); +} + extern "C" void cugraph_mg_graph_free(cugraph_graph_t* ptr_graph) { - if (ptr_graph != NULL) { - auto internal_pointer = reinterpret_cast(ptr_graph); - - destroy_graph_functor functor(internal_pointer->graph_, - internal_pointer->number_map_, - internal_pointer->edge_weights_, - internal_pointer->edge_ids_, - internal_pointer->edge_types_); - - cugraph::c_api::vertex_dispatcher(internal_pointer->vertex_type_, - internal_pointer->edge_type_, - internal_pointer->weight_type_, - internal_pointer->edge_type_id_type_, - internal_pointer->store_transposed_, - internal_pointer->multi_gpu_, - functor); - - delete internal_pointer; - } + if (ptr_graph != NULL) { cugraph_graph_free(ptr_graph); } } diff --git a/cpp/src/c_api/graph_sg.cpp b/cpp/src/c_api/graph_sg.cpp index 9536869f123..7793458b53a 100644 --- a/cpp/src/c_api/graph_sg.cpp +++ b/cpp/src/c_api/graph_sg.cpp @@ -33,35 +33,44 @@ namespace { struct create_graph_functor : public cugraph::c_api::abstract_functor { raft::handle_t const& handle_; cugraph_graph_properties_t const* properties_; + cugraph::c_api::cugraph_type_erased_device_array_view_t const* vertices_; cugraph::c_api::cugraph_type_erased_device_array_view_t const* src_; cugraph::c_api::cugraph_type_erased_device_array_view_t const* dst_; cugraph::c_api::cugraph_type_erased_device_array_view_t const* weights_; cugraph::c_api::cugraph_type_erased_device_array_view_t const* edge_ids_; cugraph::c_api::cugraph_type_erased_device_array_view_t const* edge_type_ids_; bool_t renumber_; + bool_t drop_self_loops_; + bool_t drop_multi_edges_; bool_t do_expensive_check_; cugraph_data_type_id_t edge_type_; cugraph::c_api::cugraph_graph_t* result_{}; create_graph_functor(raft::handle_t const& handle, cugraph_graph_properties_t const* properties, + cugraph::c_api::cugraph_type_erased_device_array_view_t const* vertices, cugraph::c_api::cugraph_type_erased_device_array_view_t const* src, cugraph::c_api::cugraph_type_erased_device_array_view_t const* dst, cugraph::c_api::cugraph_type_erased_device_array_view_t const* weights, cugraph::c_api::cugraph_type_erased_device_array_view_t const* edge_ids, cugraph::c_api::cugraph_type_erased_device_array_view_t const* edge_type_ids, bool_t renumber, + bool_t drop_self_loops, + bool_t drop_multi_edges, bool_t do_expensive_check, cugraph_data_type_id_t edge_type) : abstract_functor(), properties_(properties), handle_(handle), + vertices_(vertices), src_(src), dst_(dst), weights_(weights), edge_ids_(edge_ids), edge_type_ids_(edge_type_ids), renumber_(renumber), + drop_self_loops_(drop_self_loops), + drop_multi_edges_(drop_multi_edges), do_expensive_check_(do_expensive_check), edge_type_(edge_type) { @@ -99,6 +108,18 @@ struct create_graph_functor : public cugraph::c_api::abstract_functor { edge_type_id_t>> new_edge_types{std::nullopt}; + std::optional> vertex_list = + vertices_ ? std::make_optional( + rmm::device_uvector(vertices_->size_, handle_.get_stream())) + : std::nullopt; + + if (vertex_list) { + raft::copy(vertex_list->data(), + vertices_->as_type(), + vertices_->size_, + handle_.get_stream()); + } + rmm::device_uvector edgelist_srcs(src_->size_, handle_.get_stream()); rmm::device_uvector edgelist_dsts(dst_->size_, handle_.get_stream()); @@ -160,6 +181,28 @@ struct create_graph_functor : public cugraph::c_api::abstract_functor { cugraph::graph_view_t, edge_type_id_t>(handle_); + if (drop_self_loops_) { + std::tie( + edgelist_srcs, edgelist_dsts, edgelist_weights, edgelist_edge_ids, edgelist_edge_types) = + cugraph::remove_self_loops(handle_, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::move(edgelist_weights), + std::move(edgelist_edge_ids), + std::move(edgelist_edge_types)); + } + + if (drop_multi_edges_) { + std::tie( + edgelist_srcs, edgelist_dsts, edgelist_weights, edgelist_edge_ids, edgelist_edge_types) = + cugraph::remove_multi_edges(handle_, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::move(edgelist_weights), + std::move(edgelist_edge_ids), + std::move(edgelist_edge_types)); + } + std::tie(*graph, new_edge_weights, new_edge_ids, new_edge_types, new_number_map) = cugraph::create_graph_from_edgelist( handle_, - std::nullopt, + std::move(vertex_list), std::move(edgelist_srcs), std::move(edgelist_dsts), std::move(edgelist_weights), @@ -279,6 +322,12 @@ struct create_graph_csr_functor : public cugraph::c_api::abstract_functor { edge_type_id_t>> new_edge_types{std::nullopt}; + std::optional> vertex_list = std::make_optional( + rmm::device_uvector(offsets_->size_ - 1, handle_.get_stream())); + + cugraph::detail::sequence_fill( + handle_.get_stream(), vertex_list->data(), vertex_list->size(), vertex_t{0}); + rmm::device_uvector edgelist_srcs(0, handle_.get_stream()); rmm::device_uvector edgelist_dsts(indices_->size_, handle_.get_stream()); @@ -354,7 +403,7 @@ struct create_graph_csr_functor : public cugraph::c_api::abstract_functor { store_transposed, multi_gpu>( handle_, - std::nullopt, + std::move(vertex_list), std::move(edgelist_srcs), std::move(edgelist_dsts), std::move(edgelist_weights), @@ -452,9 +501,10 @@ struct destroy_graph_functor : public cugraph::c_api::abstract_functor { } // namespace -extern "C" cugraph_error_code_t cugraph_sg_graph_create( +extern "C" cugraph_error_code_t cugraph_graph_create_sg( const cugraph_resource_handle_t* handle, const cugraph_graph_properties_t* properties, + const cugraph_type_erased_device_array_view_t* vertices, const cugraph_type_erased_device_array_view_t* src, const cugraph_type_erased_device_array_view_t* dst, const cugraph_type_erased_device_array_view_t* weights, @@ -462,6 +512,8 @@ extern "C" cugraph_error_code_t cugraph_sg_graph_create( const cugraph_type_erased_device_array_view_t* edge_type_ids, bool_t store_transposed, bool_t renumber, + bool_t drop_self_loops, + bool_t drop_multi_edges, bool_t do_expensive_check, cugraph_graph_t** graph, cugraph_error_t** error) @@ -473,6 +525,8 @@ extern "C" cugraph_error_code_t cugraph_sg_graph_create( *error = nullptr; auto p_handle = reinterpret_cast(handle); + auto p_vertices = + reinterpret_cast(vertices); auto p_src = reinterpret_cast(src); auto p_dst = @@ -488,6 +542,12 @@ extern "C" cugraph_error_code_t cugraph_sg_graph_create( CUGRAPH_INVALID_INPUT, "Invalid input arguments: src size != dst size.", *error); + + CAPI_EXPECTS((p_vertices == nullptr) || (p_src->type_ == p_vertices->type_), + CUGRAPH_INVALID_INPUT, + "Invalid input arguments: src type != vertices type.", + *error); + CAPI_EXPECTS(p_src->type_ == p_dst->type_, CUGRAPH_INVALID_INPUT, "Invalid input arguments: src type != dst type.", @@ -533,12 +593,15 @@ extern "C" cugraph_error_code_t cugraph_sg_graph_create( ::create_graph_functor functor(*p_handle->handle_, properties, + p_vertices, p_src, p_dst, p_weights, p_edge_ids, p_edge_type_ids, renumber, + drop_self_loops, + drop_multi_edges, do_expensive_check, edge_type); @@ -565,7 +628,38 @@ extern "C" cugraph_error_code_t cugraph_sg_graph_create( return CUGRAPH_SUCCESS; } -cugraph_error_code_t cugraph_sg_graph_create_from_csr( +extern "C" cugraph_error_code_t cugraph_sg_graph_create( + const cugraph_resource_handle_t* handle, + const cugraph_graph_properties_t* properties, + const cugraph_type_erased_device_array_view_t* src, + const cugraph_type_erased_device_array_view_t* dst, + const cugraph_type_erased_device_array_view_t* weights, + const cugraph_type_erased_device_array_view_t* edge_ids, + const cugraph_type_erased_device_array_view_t* edge_type_ids, + bool_t store_transposed, + bool_t renumber, + bool_t do_expensive_check, + cugraph_graph_t** graph, + cugraph_error_t** error) +{ + return cugraph_graph_create_sg(handle, + properties, + NULL, + src, + dst, + weights, + edge_ids, + edge_type_ids, + store_transposed, + renumber, + FALSE, + FALSE, + do_expensive_check, + graph, + error); +} + +cugraph_error_code_t cugraph_graph_create_sg_from_csr( const cugraph_resource_handle_t* handle, const cugraph_graph_properties_t* properties, const cugraph_type_erased_device_array_view_t* offsets, @@ -662,23 +756,55 @@ cugraph_error_code_t cugraph_sg_graph_create_from_csr( return CUGRAPH_SUCCESS; } -extern "C" void cugraph_sg_graph_free(cugraph_graph_t* ptr_graph) +cugraph_error_code_t cugraph_sg_graph_create_from_csr( + const cugraph_resource_handle_t* handle, + const cugraph_graph_properties_t* properties, + const cugraph_type_erased_device_array_view_t* offsets, + const cugraph_type_erased_device_array_view_t* indices, + const cugraph_type_erased_device_array_view_t* weights, + const cugraph_type_erased_device_array_view_t* edge_ids, + const cugraph_type_erased_device_array_view_t* edge_type_ids, + bool_t store_transposed, + bool_t renumber, + bool_t do_expensive_check, + cugraph_graph_t** graph, + cugraph_error_t** error) { - auto internal_pointer = reinterpret_cast(ptr_graph); - - destroy_graph_functor functor(internal_pointer->graph_, - internal_pointer->number_map_, - internal_pointer->edge_weights_, - internal_pointer->edge_ids_, - internal_pointer->edge_types_); - - cugraph::c_api::vertex_dispatcher(internal_pointer->vertex_type_, - internal_pointer->edge_type_, - internal_pointer->weight_type_, - internal_pointer->edge_type_id_type_, - internal_pointer->store_transposed_, - internal_pointer->multi_gpu_, - functor); - - delete internal_pointer; + return cugraph_graph_create_sg_from_csr(handle, + properties, + offsets, + indices, + weights, + edge_ids, + edge_type_ids, + store_transposed, + renumber, + do_expensive_check, + graph, + error); } + +extern "C" void cugraph_graph_free(cugraph_graph_t* ptr_graph) +{ + if (ptr_graph != NULL) { + auto internal_pointer = reinterpret_cast(ptr_graph); + + destroy_graph_functor functor(internal_pointer->graph_, + internal_pointer->number_map_, + internal_pointer->edge_weights_, + internal_pointer->edge_ids_, + internal_pointer->edge_types_); + + cugraph::c_api::vertex_dispatcher(internal_pointer->vertex_type_, + internal_pointer->edge_type_, + internal_pointer->weight_type_, + internal_pointer->edge_type_id_type_, + internal_pointer->store_transposed_, + internal_pointer->multi_gpu_, + functor); + + delete internal_pointer; + } +} + +extern "C" void cugraph_sg_graph_free(cugraph_graph_t* ptr_graph) { cugraph_graph_free(ptr_graph); } diff --git a/cpp/src/c_api/resource_handle.cpp b/cpp/src/c_api/resource_handle.cpp index 767a6f0add6..75b9537ef49 100644 --- a/cpp/src/c_api/resource_handle.cpp +++ b/cpp/src/c_api/resource_handle.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -41,3 +41,10 @@ extern "C" int cugraph_resource_handle_get_rank(const cugraph_resource_handle_t* auto& comm = internal->handle_->get_comms(); return static_cast(comm.get_rank()); } + +extern "C" int cugraph_resource_handle_get_comm_size(const cugraph_resource_handle_t* handle) +{ + auto internal = reinterpret_cast(handle); + auto& comm = internal->handle_->get_comms(); + return static_cast(comm.get_size()); +} diff --git a/cpp/src/structure/detail/structure_utils.cuh b/cpp/src/structure/detail/structure_utils.cuh index 01fbccaa53e..c49b62e4543 100644 --- a/cpp/src/structure/detail/structure_utils.cuh +++ b/cpp/src/structure/detail/structure_utils.cuh @@ -21,6 +21,7 @@ #include #include #include +#include #include #include @@ -33,6 +34,7 @@ #include #include #include +#include #include #include #include @@ -496,6 +498,63 @@ void sort_adjacency_list(raft::handle_t const& handle, } } -} // namespace detail +template +std::tuple> mark_entries(raft::handle_t const& handle, + size_t num_entries, + comparison_t comparison) +{ + rmm::device_uvector marked_entries(cugraph::packed_bool_size(num_entries), + handle.get_stream()); + + thrust::tabulate(handle.get_thrust_policy(), + marked_entries.begin(), + marked_entries.end(), + [comparison, num_entries] __device__(size_t idx) { + auto word = cugraph::packed_bool_empty_mask(); + size_t start_index = idx * cugraph::packed_bools_per_word(); + size_t bits_in_this_word = + (start_index + cugraph::packed_bools_per_word() < num_entries) + ? cugraph::packed_bools_per_word() + : (num_entries - start_index); + + for (size_t bit = 0; bit < bits_in_this_word; ++bit) { + if (comparison(start_index + bit)) word |= cugraph::packed_bool_mask(bit); + } + + return word; + }); + + size_t bit_count = thrust::transform_reduce( + handle.get_thrust_policy(), + marked_entries.begin(), + marked_entries.end(), + [] __device__(auto word) { return __popc(word); }, + size_t{0}, + thrust::plus()); + + return std::make_tuple(bit_count, std::move(marked_entries)); +} +template +rmm::device_uvector remove_flagged_elements(raft::handle_t const& handle, + rmm::device_uvector&& vector, + raft::device_span remove_flags, + size_t remove_count) +{ + rmm::device_uvector result(vector.size() - remove_count, handle.get_stream()); + + thrust::copy_if( + handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(vector.size()), + thrust::make_transform_output_iterator(result.begin(), + indirection_t{vector.data()}), + [remove_flags] __device__(size_t i) { + return !(remove_flags[cugraph::packed_bool_offset(i)] & cugraph::packed_bool_mask(i)); + }); + + return result; +} + +} // namespace detail } // namespace cugraph diff --git a/cpp/src/structure/remove_multi_edges.cu b/cpp/src/structure/remove_multi_edges.cu new file mode 100644 index 00000000000..ba07d068c0e --- /dev/null +++ b/cpp/src/structure/remove_multi_edges.cu @@ -0,0 +1,92 @@ +/* + * Copyright (c) 2023, 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 + +namespace cugraph { + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_multi_edges(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_multi_edges(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_multi_edges(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_multi_edges(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_multi_edges(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_multi_edges(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +} // namespace cugraph diff --git a/cpp/src/structure/remove_multi_edges_impl.cuh b/cpp/src/structure/remove_multi_edges_impl.cuh new file mode 100644 index 00000000000..ab6b1fba8eb --- /dev/null +++ b/cpp/src/structure/remove_multi_edges_impl.cuh @@ -0,0 +1,310 @@ +/* + * Copyright (c) 2023, 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 +// FIXME: mem_frugal_partition should probably not be in shuffle_comm.hpp +// It's used here without any notion of shuffling +#include + +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +#include +#include + +namespace cugraph { + +namespace detail { + +template +struct hash_src_dst_pair { + int32_t num_groups; + + int32_t __device__ operator()(thrust::tuple t) const + { + vertex_t pair[2]; + pair[0] = thrust::get<0>(t); + pair[1] = thrust::get<1>(t); + cuco::detail::MurmurHash3_32 hash_func{}; + return hash_func.compute_hash(reinterpret_cast(pair), 2 * sizeof(vertex_t)) % + num_groups; + } +}; + +template +std::tuple, rmm::device_uvector> group_multi_edges( + raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + size_t mem_frugal_threshold) +{ + auto pair_first = thrust::make_zip_iterator(edgelist_srcs.begin(), edgelist_dsts.begin()); + + if (edgelist_srcs.size() > mem_frugal_threshold) { + // FIXME: Tuning parameter to address high frequency multi-edges + // Defaulting to 2 which makes the code easier. If + // num_groups > 2 we can evaluate whether to find a good + // midpoint to do 2 sorts, or if we should do more than 2 sorts. + const size_t num_groups{2}; + + auto group_counts = groupby_and_count(pair_first, + pair_first + edgelist_srcs.size(), + hash_src_dst_pair{}, + num_groups, + mem_frugal_threshold, + handle.get_stream()); + + std::vector h_group_counts(group_counts.size()); + raft::update_host( + h_group_counts.data(), group_counts.data(), group_counts.size(), handle.get_stream()); + + thrust::sort(handle.get_thrust_policy(), pair_first, pair_first + h_group_counts[0]); + thrust::sort(handle.get_thrust_policy(), + pair_first + h_group_counts[0], + pair_first + edgelist_srcs.size()); + } else { + thrust::sort(handle.get_thrust_policy(), pair_first, pair_first + edgelist_srcs.size()); + } + + return std::make_tuple(std::move(edgelist_srcs), std::move(edgelist_dsts)); +} + +template +std::tuple, + rmm::device_uvector, + decltype(allocate_dataframe_buffer(size_t{0}, rmm::cuda_stream_view{}))> +group_multi_edges( + raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + decltype(allocate_dataframe_buffer(0, rmm::cuda_stream_view{}))&& edgelist_values, + size_t mem_frugal_threshold) +{ + auto pair_first = thrust::make_zip_iterator(edgelist_srcs.begin(), edgelist_dsts.begin()); + auto value_first = get_dataframe_buffer_begin(edgelist_values); + + if (edgelist_srcs.size() > mem_frugal_threshold) { + // FIXME: Tuning parameter to address high frequency multi-edges + // Defaulting to 2 which makes the code easier. If + // num_groups > 2 we can evaluate whether to find a good + // midpoint to do 2 sorts, or if we should do more than 2 sorts. + const size_t num_groups{2}; + + auto group_counts = groupby_and_count(pair_first, + pair_first + edgelist_srcs.size(), + value_first, + hash_src_dst_pair{}, + num_groups, + mem_frugal_threshold, + handle.get_stream()); + + std::vector h_group_counts(group_counts.size()); + raft::update_host( + h_group_counts.data(), group_counts.data(), group_counts.size(), handle.get_stream()); + + thrust::sort_by_key(handle.get_thrust_policy(), + pair_first, + pair_first + h_group_counts[0], + get_dataframe_buffer_begin(edgelist_values)); + thrust::sort_by_key(handle.get_thrust_policy(), + pair_first + h_group_counts[0], + pair_first + edgelist_srcs.size(), + get_dataframe_buffer_begin(edgelist_values) + h_group_counts[0]); + } else { + thrust::sort_by_key(handle.get_thrust_policy(), + pair_first, + pair_first + edgelist_srcs.size(), + get_dataframe_buffer_begin(edgelist_values)); + } + + return std::make_tuple( + std::move(edgelist_srcs), std::move(edgelist_dsts), std::move(edgelist_values)); +} + +} // namespace detail + +template +std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_multi_edges(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types) +{ + auto total_global_mem = handle.get_device_properties().totalGlobalMem; + size_t element_size = sizeof(vertex_t) * 2; + if (edgelist_weights) { element_size += sizeof(weight_t); } + if (edgelist_edge_ids) { element_size += sizeof(edge_t); } + if (edgelist_edge_types) { element_size += sizeof(edge_type_t); } + + auto constexpr mem_frugal_ratio = + 0.25; // if the expected temporary buffer size exceeds the mem_frugal_ratio of the + // total_global_mem, switch to the memory frugal approach + auto mem_frugal_threshold = + static_cast(static_cast(total_global_mem / element_size) * mem_frugal_ratio); + + if (edgelist_weights) { + if (edgelist_edge_ids) { + if (edgelist_edge_types) { + std::forward_as_tuple(edgelist_srcs, + edgelist_dsts, + std::tie(edgelist_weights, edgelist_edge_ids, edgelist_edge_types)) = + detail::group_multi_edges>( + handle, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::make_tuple(std::move(*edgelist_weights), + std::move(*edgelist_edge_ids), + std::move(*edgelist_edge_types)), + mem_frugal_threshold); + } else { + std::forward_as_tuple( + edgelist_srcs, edgelist_dsts, std::tie(edgelist_weights, edgelist_edge_ids)) = + detail::group_multi_edges>( + handle, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::make_tuple(std::move(*edgelist_weights), std::move(*edgelist_edge_ids)), + mem_frugal_threshold); + } + } else { + if (edgelist_edge_types) { + std::forward_as_tuple( + edgelist_srcs, edgelist_dsts, std::tie(edgelist_weights, edgelist_edge_types)) = + detail::group_multi_edges>( + handle, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::make_tuple(std::move(*edgelist_weights), std::move(*edgelist_edge_types)), + mem_frugal_threshold); + } else { + std::forward_as_tuple(edgelist_srcs, edgelist_dsts, std::tie(edgelist_weights)) = + detail::group_multi_edges>( + handle, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::make_tuple(std::move(*edgelist_weights)), + mem_frugal_threshold); + } + } + } else { + if (edgelist_edge_ids) { + if (edgelist_edge_types) { + std::forward_as_tuple( + edgelist_srcs, edgelist_dsts, std::tie(edgelist_edge_ids, edgelist_edge_types)) = + detail::group_multi_edges>( + handle, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::make_tuple(std::move(*edgelist_edge_ids), std::move(*edgelist_edge_types)), + mem_frugal_threshold); + } else { + std::forward_as_tuple(edgelist_srcs, edgelist_dsts, std::tie(edgelist_edge_ids)) = + detail::group_multi_edges>( + handle, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::make_tuple(std::move(*edgelist_edge_ids)), + mem_frugal_threshold); + } + } else { + if (edgelist_edge_types) { + std::forward_as_tuple(edgelist_srcs, edgelist_dsts, std::tie(edgelist_edge_types)) = + detail::group_multi_edges>( + handle, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::make_tuple(std::move(*edgelist_edge_types)), + mem_frugal_threshold); + } else { + std::tie(edgelist_srcs, edgelist_dsts) = detail::group_multi_edges( + handle, std::move(edgelist_srcs), std::move(edgelist_dsts), mem_frugal_threshold); + } + } + } + + auto [multi_edge_count, multi_edges_to_delete] = + detail::mark_entries(handle, + edgelist_srcs.size(), + [d_edgelist_srcs = edgelist_srcs.data(), + d_edgelist_dsts = edgelist_dsts.data()] __device__(auto idx) { + return (idx > 0) && (d_edgelist_srcs[idx - 1] == d_edgelist_srcs[idx]) && + (d_edgelist_dsts[idx - 1] == d_edgelist_dsts[idx]); + }); + + if (multi_edge_count > 0) { + edgelist_srcs = detail::remove_flagged_elements( + handle, + std::move(edgelist_srcs), + raft::device_span{multi_edges_to_delete.data(), multi_edges_to_delete.size()}, + multi_edge_count); + edgelist_dsts = detail::remove_flagged_elements( + handle, + std::move(edgelist_dsts), + raft::device_span{multi_edges_to_delete.data(), multi_edges_to_delete.size()}, + multi_edge_count); + + if (edgelist_weights) + edgelist_weights = detail::remove_flagged_elements( + handle, + std::move(*edgelist_weights), + raft::device_span{multi_edges_to_delete.data(), + multi_edges_to_delete.size()}, + multi_edge_count); + + if (edgelist_edge_ids) + edgelist_edge_ids = detail::remove_flagged_elements( + handle, + std::move(*edgelist_edge_ids), + raft::device_span{multi_edges_to_delete.data(), + multi_edges_to_delete.size()}, + multi_edge_count); + + if (edgelist_edge_types) + edgelist_edge_types = detail::remove_flagged_elements( + handle, + std::move(*edgelist_edge_types), + raft::device_span{multi_edges_to_delete.data(), + multi_edges_to_delete.size()}, + multi_edge_count); + } + + return std::make_tuple(std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::move(edgelist_weights), + std::move(edgelist_edge_ids), + std::move(edgelist_edge_types)); +} + +} // namespace cugraph diff --git a/cpp/src/structure/remove_self_loops.cu b/cpp/src/structure/remove_self_loops.cu new file mode 100644 index 00000000000..8a66c1e05e3 --- /dev/null +++ b/cpp/src/structure/remove_self_loops.cu @@ -0,0 +1,92 @@ +/* + * Copyright (c) 2023, 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 + +namespace cugraph { + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_self_loops(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_self_loops(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_self_loops(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_self_loops(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_self_loops(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_self_loops(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types); + +} // namespace cugraph diff --git a/cpp/src/structure/remove_self_loops_impl.cuh b/cpp/src/structure/remove_self_loops_impl.cuh new file mode 100644 index 00000000000..161ffeae28e --- /dev/null +++ b/cpp/src/structure/remove_self_loops_impl.cuh @@ -0,0 +1,94 @@ +/* + * Copyright (c) 2023, 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 +#include + +namespace cugraph { + +template +std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +remove_self_loops(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::optional>&& edgelist_weights, + std::optional>&& edgelist_edge_ids, + std::optional>&& edgelist_edge_types) +{ + auto [self_loop_count, self_loops_to_delete] = + detail::mark_entries(handle, + edgelist_srcs.size(), + [d_srcs = edgelist_srcs.data(), d_dsts = edgelist_dsts.data()] __device__( + size_t i) { return d_srcs[i] == d_dsts[i]; }); + + if (self_loop_count > 0) { + edgelist_srcs = detail::remove_flagged_elements( + handle, + std::move(edgelist_srcs), + raft::device_span{self_loops_to_delete.data(), self_loops_to_delete.size()}, + self_loop_count); + edgelist_dsts = detail::remove_flagged_elements( + handle, + std::move(edgelist_dsts), + raft::device_span{self_loops_to_delete.data(), self_loops_to_delete.size()}, + self_loop_count); + + if (edgelist_weights) + edgelist_weights = detail::remove_flagged_elements( + handle, + std::move(*edgelist_weights), + raft::device_span{self_loops_to_delete.data(), self_loops_to_delete.size()}, + self_loop_count); + + if (edgelist_edge_ids) + edgelist_edge_ids = detail::remove_flagged_elements( + handle, + std::move(*edgelist_edge_ids), + raft::device_span{self_loops_to_delete.data(), self_loops_to_delete.size()}, + self_loop_count); + + if (edgelist_edge_types) + edgelist_edge_types = detail::remove_flagged_elements( + handle, + std::move(*edgelist_edge_types), + raft::device_span{self_loops_to_delete.data(), self_loops_to_delete.size()}, + self_loop_count); + } + + return std::make_tuple(std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::move(edgelist_weights), + std::move(edgelist_edge_ids), + std::move(edgelist_edge_types)); +} + +} // namespace cugraph diff --git a/cpp/tests/c_api/create_graph_test.c b/cpp/tests/c_api/create_graph_test.c index 736db761ebd..11da2eb8589 100644 --- a/cpp/tests/c_api/create_graph_test.c +++ b/cpp/tests/c_api/create_graph_test.c @@ -91,8 +91,9 @@ int test_create_sg_graph_simple() handle, wgt_view, (byte_t*)h_wgt, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt copy_from_host failed."); - ret_code = cugraph_sg_graph_create(handle, + ret_code = cugraph_graph_create_sg(handle, &properties, + NULL, src_view, dst_view, wgt_view, @@ -101,11 +102,13 @@ int test_create_sg_graph_simple() FALSE, FALSE, FALSE, + FALSE, + FALSE, &graph, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); - cugraph_sg_graph_free(graph); + cugraph_graph_free(graph); cugraph_type_erased_device_array_view_free(wgt_view); cugraph_type_erased_device_array_view_free(dst_view); @@ -300,7 +303,7 @@ int test_create_sg_graph_csr() } cugraph_sample_result_free(result); - cugraph_sg_graph_free(graph); + cugraph_graph_free(graph); cugraph_type_erased_device_array_view_free(wgt_view); cugraph_type_erased_device_array_view_free(indices_view); cugraph_type_erased_device_array_view_free(offsets_view); @@ -382,8 +385,9 @@ int test_create_sg_graph_symmetric_error() handle, wgt_view, (byte_t*)h_wgt, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt copy_from_host failed."); - ret_code = cugraph_sg_graph_create(handle, + ret_code = cugraph_graph_create_sg(handle, &properties, + NULL, src_view, dst_view, wgt_view, @@ -391,19 +395,500 @@ int test_create_sg_graph_symmetric_error() NULL, FALSE, FALSE, + FALSE, + FALSE, TRUE, &graph, &ret_error); TEST_ASSERT(test_ret_value, ret_code != CUGRAPH_SUCCESS, "graph creation succeeded but should have failed."); - if (ret_code == CUGRAPH_SUCCESS) cugraph_sg_graph_free(graph); + if (ret_code == CUGRAPH_SUCCESS) cugraph_graph_free(graph); + + cugraph_type_erased_device_array_view_free(wgt_view); + cugraph_type_erased_device_array_view_free(dst_view); + cugraph_type_erased_device_array_view_free(src_view); + cugraph_type_erased_device_array_free(wgt); + cugraph_type_erased_device_array_free(dst); + cugraph_type_erased_device_array_free(src); + + cugraph_free_resource_handle(handle); + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int test_create_sg_graph_with_isolated_vertices() +{ + int test_ret_value = 0; + + typedef int32_t vertex_t; + typedef int32_t edge_t; + typedef float weight_t; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error; + size_t num_edges = 8; + size_t num_vertices = 7; + double alpha = 0.95; + double epsilon = 0.0001; + size_t max_iterations = 20; + + vertex_t h_vertices[] = { 0, 1, 2, 3, 4, 5, 6 }; + vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; + weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f}; + weight_t h_result[] = { 0.0859168, 0.158029, 0.0616337, 0.179675, 0.113239, 0.339873, 0.0616337 }; + + cugraph_resource_handle_t* handle = NULL; + cugraph_graph_t* graph = NULL; + cugraph_graph_properties_t properties; + + properties.is_symmetric = FALSE; + properties.is_multigraph = FALSE; + + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + + handle = cugraph_create_resource_handle(NULL); + TEST_ASSERT(test_ret_value, handle != NULL, "resource handle creation failed."); + + cugraph_type_erased_device_array_t* vertices; + cugraph_type_erased_device_array_t* src; + cugraph_type_erased_device_array_t* dst; + cugraph_type_erased_device_array_t* wgt; + cugraph_type_erased_device_array_view_t* vertices_view; + cugraph_type_erased_device_array_view_t* src_view; + cugraph_type_erased_device_array_view_t* dst_view; + cugraph_type_erased_device_array_view_t* wgt_view; + + ret_code = + cugraph_type_erased_device_array_create(handle, num_vertices, vertex_tid, &vertices, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "vertices create failed."); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, vertex_tid, &src, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src create failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, vertex_tid, &dst, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst create failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, weight_tid, &wgt, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt create failed."); + + vertices_view = cugraph_type_erased_device_array_view(vertices); + src_view = cugraph_type_erased_device_array_view(src); + dst_view = cugraph_type_erased_device_array_view(dst); + wgt_view = cugraph_type_erased_device_array_view(wgt); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, vertices_view, (byte_t*)h_vertices, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "vertices copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, src_view, (byte_t*)h_src, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, dst_view, (byte_t*)h_dst, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, wgt_view, (byte_t*)h_wgt, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt copy_from_host failed."); + + ret_code = cugraph_graph_create_sg(handle, + &properties, + vertices_view, + src_view, + dst_view, + wgt_view, + NULL, + NULL, + FALSE, + FALSE, + FALSE, + FALSE, + FALSE, + &graph, + &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + + cugraph_centrality_result_t* result = NULL; + + // To verify we will call pagerank + ret_code = cugraph_pagerank(handle, + graph, + NULL, + NULL, + NULL, + NULL, + alpha, + epsilon, + max_iterations, + FALSE, + &result, + &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "cugraph_pagerank failed."); + TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + cugraph_type_erased_device_array_view_t* result_vertices; + cugraph_type_erased_device_array_view_t* pageranks; + + result_vertices = cugraph_centrality_result_get_vertices(result); + pageranks = cugraph_centrality_result_get_values(result); + + vertex_t h_result_vertices[num_vertices]; + weight_t h_pageranks[num_vertices]; + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_vertices, result_vertices, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_pageranks, pageranks, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + for (int i = 0; (i < num_vertices) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + nearlyEqual(h_result[h_result_vertices[i]], h_pageranks[i], 0.001), + "pagerank results don't match"); + } + + cugraph_centrality_result_free(result); + cugraph_graph_free(graph); + + cugraph_type_erased_device_array_view_free(wgt_view); + cugraph_type_erased_device_array_view_free(dst_view); + cugraph_type_erased_device_array_view_free(src_view); + cugraph_type_erased_device_array_view_free(vertices_view); + cugraph_type_erased_device_array_free(wgt); + cugraph_type_erased_device_array_free(dst); + cugraph_type_erased_device_array_free(src); + cugraph_type_erased_device_array_free(vertices); + + cugraph_free_resource_handle(handle); + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int test_create_sg_graph_csr_with_isolated() +{ + int test_ret_value = 0; + + typedef int32_t vertex_t; + typedef int32_t edge_t; + typedef float weight_t; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error; + size_t num_edges = 8; + size_t num_vertices = 7; + double alpha = 0.95; + double epsilon = 0.0001; + size_t max_iterations = 20; + + /* + vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; + */ + edge_t h_offsets[] = {0, 1, 3, 6, 7, 8, 8, 8}; + vertex_t h_indices[] = {1, 3, 4, 0, 1, 3, 5, 5}; + vertex_t h_start[] = {0, 1, 2, 3, 4, 5}; + weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f}; + weight_t h_result[] = { 0.0859168, 0.158029, 0.0616337, 0.179675, 0.113239, 0.339873, 0.0616337 }; + + cugraph_resource_handle_t* handle = NULL; + cugraph_graph_t* graph = NULL; + cugraph_graph_properties_t properties; + + properties.is_symmetric = FALSE; + properties.is_multigraph = FALSE; + + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + + handle = cugraph_create_resource_handle(NULL); + TEST_ASSERT(test_ret_value, handle != NULL, "resource handle creation failed."); + + cugraph_type_erased_device_array_t* offsets; + cugraph_type_erased_device_array_t* indices; + cugraph_type_erased_device_array_t* wgt; + cugraph_type_erased_device_array_view_t* offsets_view; + cugraph_type_erased_device_array_view_t* indices_view; + cugraph_type_erased_device_array_view_t* wgt_view; + + ret_code = cugraph_type_erased_device_array_create( + handle, num_vertices + 1, vertex_tid, &offsets, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "offsets create failed."); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, vertex_tid, &indices, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "indices create failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, weight_tid, &wgt, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt create failed."); + + offsets_view = cugraph_type_erased_device_array_view(offsets); + indices_view = cugraph_type_erased_device_array_view(indices); + wgt_view = cugraph_type_erased_device_array_view(wgt); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, offsets_view, (byte_t*)h_offsets, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "offsets copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, indices_view, (byte_t*)h_indices, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "indices copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, wgt_view, (byte_t*)h_wgt, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt copy_from_host failed."); + + ret_code = cugraph_sg_graph_create_from_csr(handle, + &properties, + offsets_view, + indices_view, + wgt_view, + NULL, + NULL, + FALSE, + FALSE, + FALSE, + &graph, + &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + + cugraph_centrality_result_t* result = NULL; + + // To verify we will call pagerank + ret_code = cugraph_pagerank(handle, + graph, + NULL, + NULL, + NULL, + NULL, + alpha, + epsilon, + max_iterations, + FALSE, + &result, + &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "cugraph_pagerank failed."); + TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + cugraph_type_erased_device_array_view_t* result_vertices; + cugraph_type_erased_device_array_view_t* pageranks; + + result_vertices = cugraph_centrality_result_get_vertices(result); + pageranks = cugraph_centrality_result_get_values(result); + + vertex_t h_result_vertices[num_vertices]; + weight_t h_pageranks[num_vertices]; + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_vertices, result_vertices, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_pageranks, pageranks, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + for (int i = 0; (i < num_vertices) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + nearlyEqual(h_result[h_result_vertices[i]], h_pageranks[i], 0.001), + "pagerank results don't match"); + } + + cugraph_centrality_result_free(result); + cugraph_graph_free(graph); + cugraph_type_erased_device_array_view_free(wgt_view); + cugraph_type_erased_device_array_view_free(indices_view); + cugraph_type_erased_device_array_view_free(offsets_view); + cugraph_type_erased_device_array_free(wgt); + cugraph_type_erased_device_array_free(indices); + cugraph_type_erased_device_array_free(offsets); + + cugraph_free_resource_handle(handle); + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int test_create_sg_graph_with_isolated_vertices_multi_input() +{ + int test_ret_value = 0; + + typedef int32_t vertex_t; + typedef int32_t edge_t; + typedef float weight_t; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error; + size_t num_edges = 66; + size_t num_vertices = 7; + double alpha = 0.95; + double epsilon = 0.0001; + size_t max_iterations = 20; + + vertex_t h_vertices[] = { 0, 1, 2, 3, 4, 5, 6 }; + vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 4, 4, 4, 5, + 0, 1, 1, 2, 2, 2, 3, 4, 4, 4, 5, + 0, 1, 1, 2, 2, 2, 3, 4, 4, 4, 5, + 0, 1, 1, 2, 2, 2, 3, 4, 4, 4, 5, + 0, 1, 1, 2, 2, 2, 3, 4, 4, 4, 5, + 0, 1, 1, 2, 2, 2, 3, 4, 4, 4, 5}; + vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5, 5, 5, 5, + 1, 3, 4, 0, 1, 3, 5, 5, 5, 5, 5, + 1, 3, 4, 0, 1, 3, 5, 5, 5, 5, 5, + 1, 3, 4, 0, 1, 3, 5, 5, 5, 5, 5, + 1, 3, 4, 0, 1, 3, 5, 5, 5, 5, 5, + 1, 3, 4, 0, 1, 3, 5, 5, 5, 5, 5}; + weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 3.2f, 3.2f, 1.7f, + 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 3.2f, 3.2f, 1.7f, + 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 3.2f, 3.2f, 1.7f, + 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 3.2f, 3.2f, 1.7f, + 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 3.2f, 3.2f, 1.7f, + 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 3.2f, 3.2f, 1.7f}; + weight_t h_result[] = { 0.0859168, 0.158029, 0.0616337, 0.179675, 0.113239, 0.339873, 0.0616337 }; + + cugraph_resource_handle_t* handle = NULL; + cugraph_graph_t* graph = NULL; + cugraph_graph_properties_t properties; + + properties.is_symmetric = FALSE; + properties.is_multigraph = FALSE; + + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + + handle = cugraph_create_resource_handle(NULL); + TEST_ASSERT(test_ret_value, handle != NULL, "resource handle creation failed."); + + cugraph_type_erased_device_array_t* vertices; + cugraph_type_erased_device_array_t* src; + cugraph_type_erased_device_array_t* dst; + cugraph_type_erased_device_array_t* wgt; + cugraph_type_erased_device_array_view_t* vertices_view; + cugraph_type_erased_device_array_view_t* src_view; + cugraph_type_erased_device_array_view_t* dst_view; + cugraph_type_erased_device_array_view_t* wgt_view; + + ret_code = + cugraph_type_erased_device_array_create(handle, num_vertices, vertex_tid, &vertices, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "vertices create failed."); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, vertex_tid, &src, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src create failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, vertex_tid, &dst, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst create failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, weight_tid, &wgt, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt create failed."); + + vertices_view = cugraph_type_erased_device_array_view(vertices); + src_view = cugraph_type_erased_device_array_view(src); + dst_view = cugraph_type_erased_device_array_view(dst); + wgt_view = cugraph_type_erased_device_array_view(wgt); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, vertices_view, (byte_t*)h_vertices, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "vertices copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, src_view, (byte_t*)h_src, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, dst_view, (byte_t*)h_dst, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, wgt_view, (byte_t*)h_wgt, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt copy_from_host failed."); + + ret_code = cugraph_graph_create_sg(handle, + &properties, + vertices_view, + src_view, + dst_view, + wgt_view, + NULL, + NULL, + FALSE, + FALSE, + TRUE, + TRUE, + FALSE, + &graph, + &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + + cugraph_centrality_result_t* result = NULL; + + // To verify we will call pagerank + ret_code = cugraph_pagerank(handle, + graph, + NULL, + NULL, + NULL, + NULL, + alpha, + epsilon, + max_iterations, + FALSE, + &result, + &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "cugraph_pagerank failed."); + TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + cugraph_type_erased_device_array_view_t* result_vertices; + cugraph_type_erased_device_array_view_t* pageranks; + + result_vertices = cugraph_centrality_result_get_vertices(result); + pageranks = cugraph_centrality_result_get_values(result); + + vertex_t h_result_vertices[num_vertices]; + weight_t h_pageranks[num_vertices]; + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_vertices, result_vertices, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_pageranks, pageranks, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + for (int i = 0; (i < num_vertices) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + nearlyEqual(h_result[h_result_vertices[i]], h_pageranks[i], 0.001), + "pagerank results don't match"); + } + + cugraph_centrality_result_free(result); + cugraph_graph_free(graph); cugraph_type_erased_device_array_view_free(wgt_view); cugraph_type_erased_device_array_view_free(dst_view); cugraph_type_erased_device_array_view_free(src_view); + cugraph_type_erased_device_array_view_free(vertices_view); cugraph_type_erased_device_array_free(wgt); cugraph_type_erased_device_array_free(dst); cugraph_type_erased_device_array_free(src); + cugraph_type_erased_device_array_free(vertices); cugraph_free_resource_handle(handle); cugraph_error_free(ret_error); @@ -419,5 +904,8 @@ int main(int argc, char** argv) result |= RUN_TEST(test_create_sg_graph_simple); result |= RUN_TEST(test_create_sg_graph_csr); result |= RUN_TEST(test_create_sg_graph_symmetric_error); + result |= RUN_TEST(test_create_sg_graph_with_isolated_vertices); + result |= RUN_TEST(test_create_sg_graph_csr_with_isolated); + result |= RUN_TEST(test_create_sg_graph_with_isolated_vertices_multi_input); return result; } diff --git a/cpp/tests/c_api/mg_create_graph_test.c b/cpp/tests/c_api/mg_create_graph_test.c index 4c8f2f22982..fec319d1881 100644 --- a/cpp/tests/c_api/mg_create_graph_test.c +++ b/cpp/tests/c_api/mg_create_graph_test.c @@ -17,6 +17,8 @@ #include "c_test_utils.h" /* RUN_TEST */ #include "mg_test_utils.h" /* RUN_TEST */ +#include + #include #include #include @@ -41,7 +43,7 @@ int test_create_mg_graph_simple(const cugraph_resource_handle_t* handle) vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f}; - cugraph_graph_t* p_graph = NULL; + cugraph_graph_t* graph = NULL; cugraph_graph_properties_t properties; properties.is_symmetric = FALSE; @@ -94,21 +96,25 @@ int test_create_mg_graph_simple(const cugraph_resource_handle_t* handle) handle, wgt_view, (byte_t*)h_wgt, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt copy_from_host failed."); - ret_code = cugraph_mg_graph_create(handle, + ret_code = cugraph_graph_create_mg(handle, &properties, - src_view, - dst_view, - wgt_view, + NULL, + (cugraph_type_erased_device_array_view_t const* const*) &src_view, + (cugraph_type_erased_device_array_view_t const* const*) &dst_view, + (cugraph_type_erased_device_array_view_t const* const*) &wgt_view, NULL, NULL, FALSE, - num_edges, + 1, + FALSE, + FALSE, TRUE, - &p_graph, + &graph, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); - cugraph_mg_graph_free(p_graph); + cugraph_graph_free(graph); cugraph_type_erased_device_array_view_free(wgt_view); cugraph_type_erased_device_array_view_free(dst_view); @@ -122,6 +128,382 @@ int test_create_mg_graph_simple(const cugraph_resource_handle_t* handle) return test_ret_value; } +int test_create_mg_graph_multiple_edge_lists(const cugraph_resource_handle_t* handle) +{ + int test_ret_value = 0; + + typedef int32_t vertex_t; + typedef int32_t edge_t; + typedef float weight_t; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error; + size_t num_edges = 8; + size_t num_vertices = 7; + + double alpha = 0.95; + double epsilon = 0.0001; + size_t max_iterations = 20; + + vertex_t h_vertices[] = { 0, 1, 2, 3, 4, 5, 6 }; + vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; + weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f}; + weight_t h_result[] = { 0.0859168, 0.158029, 0.0616337, 0.179675, 0.113239, 0.339873, 0.0616337 }; + + cugraph_graph_t* graph = NULL; + cugraph_graph_properties_t properties; + + properties.is_symmetric = FALSE; + properties.is_multigraph = FALSE; + + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + + const size_t num_local_arrays = 2; + + cugraph_type_erased_device_array_t* vertices[num_local_arrays]; + cugraph_type_erased_device_array_t* src[num_local_arrays]; + cugraph_type_erased_device_array_t* dst[num_local_arrays]; + cugraph_type_erased_device_array_t* wgt[num_local_arrays]; + cugraph_type_erased_device_array_view_t* vertices_view[num_local_arrays]; + cugraph_type_erased_device_array_view_t* src_view[num_local_arrays]; + cugraph_type_erased_device_array_view_t* dst_view[num_local_arrays]; + cugraph_type_erased_device_array_view_t* wgt_view[num_local_arrays]; + + int my_rank = cugraph_resource_handle_get_rank(handle); + int comm_size = cugraph_resource_handle_get_comm_size(handle); + + size_t local_num_vertices = (num_vertices + comm_size - 1) / comm_size; + size_t local_start_vertex = my_rank * local_num_vertices; + size_t local_num_edges = (num_edges + comm_size - 1) / comm_size; + size_t local_start_edge = my_rank * local_num_edges; + + local_num_edges = (local_num_edges < (num_edges - local_start_edge)) ? local_num_edges : (num_edges - local_start_edge); + local_num_vertices = (local_num_vertices < (num_vertices - local_start_vertex)) ? local_num_vertices : (num_vertices - local_start_vertex); + + for (size_t i = 0 ; i < num_local_arrays ; ++i) { + size_t vertex_count = (local_num_vertices + num_local_arrays - 1) / num_local_arrays; + size_t vertex_start = i * vertex_count; + vertex_count = (vertex_count < (local_num_vertices - vertex_start)) ? vertex_count : (local_num_vertices - vertex_start); + + ret_code = + cugraph_type_erased_device_array_create(handle, vertex_count, vertex_tid, vertices + i, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "vertices create failed."); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + size_t edge_count = (local_num_edges + num_local_arrays - 1) / num_local_arrays; + size_t edge_start = i * edge_count; + edge_count = (edge_count < (local_num_edges - edge_start)) ? edge_count : (local_num_edges - edge_start); + + ret_code = + cugraph_type_erased_device_array_create(handle, edge_count, vertex_tid, src + i, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src create failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, edge_count, vertex_tid, dst + i, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst create failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, edge_count, weight_tid, wgt + i, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt create failed."); + + vertices_view[i] = cugraph_type_erased_device_array_view(vertices[i]); + src_view[i] = cugraph_type_erased_device_array_view(src[i]); + dst_view[i] = cugraph_type_erased_device_array_view(dst[i]); + wgt_view[i] = cugraph_type_erased_device_array_view(wgt[i]); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, vertices_view[i], (byte_t*)(h_vertices + local_start_vertex + vertex_start), &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, src_view[i], (byte_t*)(h_src + local_start_edge + edge_start), &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, dst_view[i], (byte_t*)(h_dst + local_start_edge + edge_start), &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, wgt_view[i], (byte_t*)(h_wgt + local_start_edge + edge_start), &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt copy_from_host failed."); + } + + ret_code = cugraph_graph_create_mg(handle, + &properties, + (cugraph_type_erased_device_array_view_t const* const*) vertices_view, + (cugraph_type_erased_device_array_view_t const* const*) src_view, + (cugraph_type_erased_device_array_view_t const* const*) dst_view, + (cugraph_type_erased_device_array_view_t const* const*) wgt_view, + NULL, + NULL, + FALSE, + num_local_arrays, + FALSE, + FALSE, + TRUE, + &graph, + &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + // + // Now call pagerank and check results... + // + cugraph_centrality_result_t* result = NULL; + + ret_code = cugraph_pagerank(handle, + graph, + NULL, + NULL, + NULL, + NULL, + alpha, + epsilon, + max_iterations, + FALSE, + &result, + &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "cugraph_pagerank failed."); + + // NOTE: Because we get back vertex ids and pageranks, we can simply compare + // the returned values with the expected results for the entire + // graph. Each GPU will have a subset of the total vertices, so + // they will do a subset of the comparisons. + cugraph_type_erased_device_array_view_t* result_vertices; + cugraph_type_erased_device_array_view_t* pageranks; + + result_vertices = cugraph_centrality_result_get_vertices(result); + pageranks = cugraph_centrality_result_get_values(result); + + size_t num_local_vertices = cugraph_type_erased_device_array_view_size(result_vertices); + + vertex_t h_result_vertices[num_local_vertices]; + weight_t h_pageranks[num_local_vertices]; + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_vertices, result_vertices, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_pageranks, pageranks, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + for (int i = 0; (i < num_local_vertices) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + nearlyEqual(h_result[h_result_vertices[i]], h_pageranks[i], 0.001), + "pagerank results don't match"); + } + + cugraph_centrality_result_free(result); + cugraph_graph_free(graph); + + for (size_t i = 0 ; i < num_local_arrays ; ++i) { + cugraph_type_erased_device_array_view_free(wgt_view[i]); + cugraph_type_erased_device_array_view_free(dst_view[i]); + cugraph_type_erased_device_array_view_free(src_view[i]); + cugraph_type_erased_device_array_view_free(vertices_view[i]); + cugraph_type_erased_device_array_free(wgt[i]); + cugraph_type_erased_device_array_free(dst[i]); + cugraph_type_erased_device_array_free(src[i]); + cugraph_type_erased_device_array_free(vertices[i]); + } + + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int test_create_mg_graph_multiple_edge_lists_multi_edge(const cugraph_resource_handle_t* handle) +{ + int test_ret_value = 0; + + typedef int32_t vertex_t; + typedef int32_t edge_t; + typedef float weight_t; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error; + size_t num_edges = 11; + size_t num_vertices = 7; + + double alpha = 0.95; + double epsilon = 0.0001; + size_t max_iterations = 20; + + vertex_t h_vertices[] = { 0, 1, 2, 3, 4, 5, 6 }; + vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 4, 4, 4, 5}; + vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5, 5, 5, 5}; + weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 3.2f, 3.2f, 1.1f}; + weight_t h_result[] = { 0.0859168, 0.158029, 0.0616337, 0.179675, 0.113239, 0.339873, 0.0616337 }; + + cugraph_graph_t* graph = NULL; + cugraph_graph_properties_t properties; + + properties.is_symmetric = FALSE; + properties.is_multigraph = FALSE; + + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + + const size_t num_local_arrays = 2; + + cugraph_type_erased_device_array_t* vertices[num_local_arrays]; + cugraph_type_erased_device_array_t* src[num_local_arrays]; + cugraph_type_erased_device_array_t* dst[num_local_arrays]; + cugraph_type_erased_device_array_t* wgt[num_local_arrays]; + cugraph_type_erased_device_array_view_t* vertices_view[num_local_arrays]; + cugraph_type_erased_device_array_view_t* src_view[num_local_arrays]; + cugraph_type_erased_device_array_view_t* dst_view[num_local_arrays]; + cugraph_type_erased_device_array_view_t* wgt_view[num_local_arrays]; + + int my_rank = cugraph_resource_handle_get_rank(handle); + int comm_size = cugraph_resource_handle_get_comm_size(handle); + + size_t local_num_vertices = (num_vertices + comm_size - 1) / comm_size; + size_t local_start_vertex = my_rank * local_num_vertices; + size_t local_num_edges = (num_edges + comm_size - 1) / comm_size; + size_t local_start_edge = my_rank * local_num_edges; + + local_num_edges = (local_num_edges < (num_edges - local_start_edge)) ? local_num_edges : (num_edges - local_start_edge); + local_num_vertices = (local_num_vertices < (num_vertices - local_start_vertex)) ? local_num_vertices : (num_vertices - local_start_vertex); + + for (size_t i = 0 ; i < num_local_arrays ; ++i) { + size_t vertex_count = (local_num_vertices + num_local_arrays - 1) / num_local_arrays; + size_t vertex_start = i * vertex_count; + vertex_count = (vertex_count < (local_num_vertices - vertex_start)) ? vertex_count : (local_num_vertices - vertex_start); + + ret_code = + cugraph_type_erased_device_array_create(handle, vertex_count, vertex_tid, vertices + i, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "vertices create failed."); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + size_t edge_count = (local_num_edges + num_local_arrays - 1) / num_local_arrays; + size_t edge_start = i * edge_count; + edge_count = (edge_count < (local_num_edges - edge_start)) ? edge_count : (local_num_edges - edge_start); + + ret_code = + cugraph_type_erased_device_array_create(handle, edge_count, vertex_tid, src + i, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src create failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, edge_count, vertex_tid, dst + i, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst create failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, edge_count, weight_tid, wgt + i, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt create failed."); + + vertices_view[i] = cugraph_type_erased_device_array_view(vertices[i]); + src_view[i] = cugraph_type_erased_device_array_view(src[i]); + dst_view[i] = cugraph_type_erased_device_array_view(dst[i]); + wgt_view[i] = cugraph_type_erased_device_array_view(wgt[i]); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, vertices_view[i], (byte_t*)(h_vertices + local_start_vertex + vertex_start), &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, src_view[i], (byte_t*)(h_src + local_start_edge + edge_start), &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, dst_view[i], (byte_t*)(h_dst + local_start_edge + edge_start), &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, wgt_view[i], (byte_t*)(h_wgt + local_start_edge + edge_start), &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt copy_from_host failed."); + } + + ret_code = cugraph_graph_create_mg(handle, + &properties, + (cugraph_type_erased_device_array_view_t const* const*) vertices_view, + (cugraph_type_erased_device_array_view_t const* const*) src_view, + (cugraph_type_erased_device_array_view_t const* const*) dst_view, + (cugraph_type_erased_device_array_view_t const* const*) wgt_view, + NULL, + NULL, + FALSE, + num_local_arrays, + TRUE, + TRUE, + TRUE, + &graph, + &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + // + // Now call pagerank and check results... + // + cugraph_centrality_result_t* result = NULL; + + ret_code = cugraph_pagerank(handle, + graph, + NULL, + NULL, + NULL, + NULL, + alpha, + epsilon, + max_iterations, + FALSE, + &result, + &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "cugraph_pagerank failed."); + + // NOTE: Because we get back vertex ids and pageranks, we can simply compare + // the returned values with the expected results for the entire + // graph. Each GPU will have a subset of the total vertices, so + // they will do a subset of the comparisons. + cugraph_type_erased_device_array_view_t* result_vertices; + cugraph_type_erased_device_array_view_t* pageranks; + + result_vertices = cugraph_centrality_result_get_vertices(result); + pageranks = cugraph_centrality_result_get_values(result); + + size_t num_local_vertices = cugraph_type_erased_device_array_view_size(result_vertices); + + vertex_t h_result_vertices[num_local_vertices]; + weight_t h_pageranks[num_local_vertices]; + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_vertices, result_vertices, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_pageranks, pageranks, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + for (int i = 0; (i < num_local_vertices) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + nearlyEqual(h_result[h_result_vertices[i]], h_pageranks[i], 0.001), + "pagerank results don't match"); + } + + cugraph_centrality_result_free(result); + cugraph_graph_free(graph); + + for (size_t i = 0 ; i < num_local_arrays ; ++i) { + cugraph_type_erased_device_array_view_free(wgt_view[i]); + cugraph_type_erased_device_array_view_free(dst_view[i]); + cugraph_type_erased_device_array_view_free(src_view[i]); + cugraph_type_erased_device_array_view_free(vertices_view[i]); + cugraph_type_erased_device_array_free(wgt[i]); + cugraph_type_erased_device_array_free(dst[i]); + cugraph_type_erased_device_array_free(src[i]); + cugraph_type_erased_device_array_free(vertices[i]); + } + + cugraph_error_free(ret_error); + + return test_ret_value; +} + /******************************************************************************/ int main(int argc, char** argv) @@ -131,6 +513,8 @@ int main(int argc, char** argv) int result = 0; result |= RUN_MG_TEST(test_create_mg_graph_simple, handle); + result |= RUN_MG_TEST(test_create_mg_graph_multiple_edge_lists, handle); + result |= RUN_MG_TEST(test_create_mg_graph_multiple_edge_lists_multi_edge, handle); cugraph_free_resource_handle(handle); free_mg_raft_handle(raft_handle);