diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 3e867643041..5ed83a5c467 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -201,6 +201,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);