From 5bf07fb4bd76481f9077e21bfaccac73b749b325 Mon Sep 17 00:00:00 2001 From: Chuck Hastings <45364586+ChuckHastings@users.noreply.github.com> Date: Fri, 22 Jul 2022 09:45:22 -0400 Subject: [PATCH] Define API for MG random walk (#2407) This PR defines the API for MG random walk in the C API and the C++ API. C and C++ tests are defined, although some of the code is ifdef'ed out since there is not a working implementation here. Authors: - Chuck Hastings (https://github.com/ChuckHastings) Approvers: - Seunghwa Kang (https://github.com/seunghwak) - Joseph Nke (https://github.com/jnke2016) - Alex Barghi (https://github.com/alexbarghi-nv) URL: https://github.com/rapidsai/cugraph/pull/2407 --- cpp/CMakeLists.txt | 2 + cpp/include/cugraph/algorithms.hpp | 130 ++++++ cpp/include/cugraph_c/sampling_algorithms.h | 71 ++++ cpp/src/c_api/random_walks.cpp | 380 +++++++++++++++++- cpp/src/sampling/random_walks_mg.cu | 195 +++++++++ cpp/src/sampling/random_walks_sg.cu | 195 +++++++++ cpp/tests/CMakeLists.txt | 13 +- cpp/tests/c_api/mg_random_walks_test.c | 423 ++++++++++++++++++++ cpp/tests/c_api/sg_random_walks_test.c | 410 +++++++++++++++++++ cpp/tests/sampling/mg_random_walks_test.cu | 225 +++++++++++ cpp/tests/sampling/sg_random_walks_test.cu | 237 +++++++++++ 11 files changed, 2259 insertions(+), 22 deletions(-) create mode 100644 cpp/src/sampling/random_walks_mg.cu create mode 100644 cpp/src/sampling/random_walks_sg.cu create mode 100644 cpp/tests/c_api/mg_random_walks_test.c create mode 100644 cpp/tests/c_api/sg_random_walks_test.c create mode 100644 cpp/tests/sampling/mg_random_walks_test.cu create mode 100644 cpp/tests/sampling/sg_random_walks_test.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 913aab7bd5a..5398b526a9e 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -195,6 +195,8 @@ set(CUGRAPH_SOURCES src/community/legacy/extract_subgraph_by_vertex.cu src/community/legacy/egonet.cu src/sampling/random_walks.cu + src/sampling/random_walks_sg.cu + src/sampling/random_walks_mg.cu src/sampling/detail/sampling_utils_mg.cu src/sampling/detail/sampling_utils_sg.cu src/sampling/uniform_neighbor_sampling_mg.cpp diff --git a/cpp/include/cugraph/algorithms.hpp b/cpp/include/cugraph/algorithms.hpp index e4eb5f7030f..1b8c112bb2c 100644 --- a/cpp/include/cugraph/algorithms.hpp +++ b/cpp/include/cugraph/algorithms.hpp @@ -1350,6 +1350,136 @@ random_walks(raft::handle_t const& handle, bool use_padding = false, std::unique_ptr sampling_strategy = nullptr); +/** + * @brief returns uniform random walks from starting sources, where each path is of given + * maximum length. + * + * @p start_vertices can contain duplicates, in which case different random walks will + * be generated for each instance. + * + * If the graph is weighted, the return contains edge weights. If the graph is unweighted then + * the returned value will be std::nullopt. + * + * @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 weights. Needs to be a floating point type. + * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param graph_view graph view to operate on + * @param start_vertices Device span defining the starting vertices + * @param max_length maximum length of random walk + * @param seed (optional, defaults to system time), seed for random number generation + * @return tuple containing device vectors of vertices and the edge weights (if + * the graph is weighted)
+ * For each input selector there will be (max_length+1) elements in the + * vertex vector with the starting vertex followed by the subsequent + * vertices in the random walk. If a path terminates before max_length, + * the vertices will be populated with invalid_vertex_id + * (-1 for signed vertex_t, std::numeric_limits::max() for an + * unsigned vertex_t type)
+ * For each input selector there will be max_length elements in the weights + * vector with the edge weight for the edge in the path. If a path + * terminates before max_length the subsequent edge weights will be + * set to weight_t{0}. + */ +// FIXME: Do I care about transposed or not? I want to be able to operate in either +// direction. +template +std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed = std::numeric_limits::max()); + +/** + * @brief returns biased random walks from starting sources, where each path is of given + * maximum length. + * + * The next vertex is biased based on the edge weights. The probability of traversing a + * departing edge will be the edge weight divided by the sum of the departing edge weights. + * + * @p start_vertices can contain duplicates, in which case different random walks will + * be generated for each instance. + * + * @throws cugraph::logic_error if the graph is unweighted + * + * @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 weights. Needs to be a floating point type. + * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param graph_view graph view to operate on + * @param start_vertices Device span defining the starting vertices + * @param max_length maximum length of random walk + * @param seed (optional, defaults to system time), seed for random number generation + * @return tuple containing device vectors of vertices and the edge weights
+ * For each input selector there will be (max_length+1) elements in the + * vertex vector with the starting vertex followed by the subsequent + * vertices in the random walk. If a path terminates before max_length, + * the vertices will be populated with invalid_vertex_id + * (-1 for signed vertex_t, std::numeric_limits::max() for an + * unsigned vertex_t type)
+ * For each input selector there will be max_length elements in the weights + * vector with the edge weight for the edge in the path. If a path + * terminates before max_length the subsequent edge weights will be + * set to weight_t{0}. + */ +template +std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed = std::numeric_limits::max()); + +/** + * @brief returns biased random walks with node2vec biases from starting sources, + * where each path is of given maximum length. + * + * @p start_vertices can contain duplicates, in which case different random walks will + * be generated for each instance. + * + * If the graph is weighted, the return contains edge weights and the node2vec computation + * will utilize the edge weights. If the graph is unweighted then the return will not contain + * edge weights and the node2vec computation will assume an edge weight of 1 for all edges. + * + * @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 weights. Needs to be a floating point type. + * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param graph_view graph view to operate on + * @param start_vertices Device span defining the starting vertices + * @param max_length maximum length of random walk + * @param p node2vec return parameter + * @param q node2vec in-out parameter + * @param seed (optional, defaults to system time), seed for random number generation + * @return tuple containing device vectors of vertices and the edge weights
+ * For each input selector there will be (max_length+1) elements in the + * vertex vector with the starting vertex followed by the subsequent + * vertices in the random walk. If a path terminates before max_length, + * the vertices will be populated with invalid_vertex_id + * (-1 for signed vertex_t, std::numeric_limits::max() for an + * unsigned vertex_t type)
+ * For each input selector there will be max_length elements in the weights + * vector with the edge weight for the edge in the path. If a path + * terminates before max_length the subsequent edge weights will be + * set to weight_t{0}. + */ +template +std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + weight_t p, + weight_t q, + uint64_t seed = std::numeric_limits::max()); + #ifndef NO_CUGRAPH_OPS /** * @brief generate sub-sampled graph as an adjacency list (CSR format) given input graph, diff --git a/cpp/include/cugraph_c/sampling_algorithms.h b/cpp/include/cugraph_c/sampling_algorithms.h index dbefac81742..e923b91a2a9 100644 --- a/cpp/include/cugraph_c/sampling_algorithms.h +++ b/cpp/include/cugraph_c/sampling_algorithms.h @@ -36,8 +36,78 @@ typedef struct { int32_t align_; } cugraph_random_walk_result_t; +/** + * @brief Compute uniform random walks + * + * @param [in] handle Handle for accessing resources + * @param [in] graph Pointer to graph. NOTE: Graph might be modified if the storage + * needs to be transposed + * @param [in] start_vertices Array of source vertices + * @param [in] max_length Maximum length of the generated path + * @param [in] result Output from the node2vec call + * @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_uniform_random_walks( + const cugraph_resource_handle_t* handle, + cugraph_graph_t* graph, + const cugraph_type_erased_device_array_view_t* start_vertices, + size_t max_length, + cugraph_random_walk_result_t** result, + cugraph_error_t** error); + +/** + * @brief Compute biased random walks + * + * @param [in] handle Handle for accessing resources + * @param [in] graph Pointer to graph. NOTE: Graph might be modified if the storage + * needs to be transposed + * @param [in] start_vertices Array of source vertices + * @param [in] max_length Maximum length of the generated path + * @param [in] result Output from the node2vec call + * @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_biased_random_walks( + const cugraph_resource_handle_t* handle, + cugraph_graph_t* graph, + const cugraph_type_erased_device_array_view_t* start_vertices, + size_t max_length, + cugraph_random_walk_result_t** result, + cugraph_error_t** error); + +/** + * @brief Compute random walks using the node2vec framework. + * + * @param [in] handle Handle for accessing resources + * @param [in] graph Pointer to graph. NOTE: Graph might be modified if the storage + * needs to be transposed + * @param [in] start_vertices Array of source vertices + * @param [in] max_length Maximum length of the generated path + * @param [in] compress_result If true, return the paths as a compressed sparse row matrix, + * otherwise return as a dense matrix + * @param [in] p The return parameter + * @param [in] q The in/out parameter + * @param [in] result Output from the node2vec call + * @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_node2vec_random_walks( + const cugraph_resource_handle_t* handle, + cugraph_graph_t* graph, + const cugraph_type_erased_device_array_view_t* start_vertices, + size_t max_length, + double p, + double q, + cugraph_random_walk_result_t** result, + cugraph_error_t** error); + /** * @brief Compute random walks using the node2vec framework. + * @deprecated This call should be replaced with cugraph_node2vec_random_walks * * @param [in] handle Handle for accessing resources * @param [in] graph Pointer to graph. NOTE: Graph might be modified if the storage @@ -94,6 +164,7 @@ cugraph_type_erased_device_array_view_t* cugraph_random_walk_result_get_weights( /** * @brief If the random walk result is compressed, get the path sizes + * @deprecated This call will no longer be relevant once the new node2vec are called * * @param [in] result The result from a random walk algorithm * @return type erased array pointing to the path sizes in device memory diff --git a/cpp/src/c_api/random_walks.cpp b/cpp/src/c_api/random_walks.cpp index 65c94ed6734..bc13c6ce3f3 100644 --- a/cpp/src/c_api/random_walks.cpp +++ b/cpp/src/c_api/random_walks.cpp @@ -40,8 +40,8 @@ struct cugraph_random_walk_result_t { struct node2vec_functor : public abstract_functor { raft::handle_t const& handle_; cugraph_graph_t* graph_{nullptr}; - cugraph_type_erased_device_array_view_t const* sources_{nullptr}; - size_t max_depth_{0}; + cugraph_type_erased_device_array_view_t const* start_vertices_{nullptr}; + size_t max_length_{0}; bool compress_result_{false}; double p_{0}; double q_{0}; @@ -49,17 +49,18 @@ struct node2vec_functor : public abstract_functor { node2vec_functor(::cugraph_resource_handle_t const* handle, ::cugraph_graph_t* graph, - ::cugraph_type_erased_device_array_view_t const* sources, - size_t max_depth, + ::cugraph_type_erased_device_array_view_t const* start_vertices, + size_t max_length, bool compress_result, double p, double q) : abstract_functor(), handle_(*reinterpret_cast(handle)->handle_), graph_(reinterpret_cast(graph)), - sources_( - reinterpret_cast(sources)), - max_depth_(max_depth), + start_vertices_( + reinterpret_cast( + start_vertices)), + max_length_(max_length), compress_result_(compress_result), p_(p), q_(q) @@ -95,16 +96,18 @@ struct node2vec_functor : public abstract_functor { auto number_map = reinterpret_cast*>(graph_->number_map_); - rmm::device_uvector sources(sources_->size_, handle_.get_stream()); - raft::copy( - sources.data(), sources_->as_type(), sources.size(), handle_.get_stream()); + rmm::device_uvector start_vertices(start_vertices_->size_, handle_.get_stream()); + raft::copy(start_vertices.data(), + start_vertices_->as_type(), + start_vertices.size(), + handle_.get_stream()); // - // Need to renumber sources + // Need to renumber start_vertices // renumber_ext_vertices(handle_, - sources.data(), - sources.size(), + start_vertices.data(), + start_vertices.size(), number_map->data(), graph_view.local_vertex_partition_range_first(), graph_view.local_vertex_partition_range_last(), @@ -112,13 +115,13 @@ struct node2vec_functor : public abstract_functor { // FIXME: Forcing this to edge_t for now. What should it really be? // Seems like it should be the smallest size that can accommodate - // max_depth_ * sources_->size_ + // max_length_ * start_vertices_->size_ auto [paths, weights, sizes] = cugraph::random_walks( handle_, graph_view, - sources.data(), - static_cast(sources.size()), - static_cast(max_depth_), + start_vertices.data(), + static_cast(start_vertices.size()), + static_cast(max_length_), !compress_result_, // std::make_unique(2, p_, q_, false)); std::make_unique(cugraph::sampling_strategy_t::NODE2VEC, p_, q_)); @@ -131,7 +134,7 @@ struct node2vec_functor : public abstract_functor { result_ = new cugraph_random_walk_result_t{ compress_result_, - max_depth_, + max_length_, new cugraph_type_erased_device_array_t(paths, graph_->vertex_type_), new cugraph_type_erased_device_array_t(weights, graph_->weight_type_), new cugraph_type_erased_device_array_t(sizes, graph_->vertex_type_)}; @@ -142,10 +145,304 @@ struct node2vec_functor : public abstract_functor { } // namespace c_api } // namespace cugraph +namespace { + +struct uniform_random_walks_functor : public cugraph::c_api::abstract_functor { + raft::handle_t const& handle_; + cugraph::c_api::cugraph_graph_t* graph_{nullptr}; + cugraph::c_api::cugraph_type_erased_device_array_view_t const* start_vertices_{nullptr}; + size_t max_length_{0}; + size_t seed_{0}; + cugraph::c_api::cugraph_random_walk_result_t* result_{nullptr}; + + uniform_random_walks_functor(cugraph_resource_handle_t const* handle, + cugraph_graph_t* graph, + cugraph_type_erased_device_array_view_t const* start_vertices, + size_t max_length) + : abstract_functor(), + handle_(*reinterpret_cast(handle)->handle_), + graph_(reinterpret_cast(graph)), + start_vertices_( + reinterpret_cast( + start_vertices)), + max_length_(max_length) + { + } + + template + void operator()() + { + // FIXME: Think about how to handle SG vice MG + if constexpr (!cugraph::is_candidate::value) { + unsupported(); + } else if constexpr (multi_gpu) { + unsupported(); + } else { + // random walks expects store_transposed == false + if constexpr (store_transposed) { + error_code_ = cugraph::c_api:: + transpose_storage( + handle_, graph_, error_.get()); + if (error_code_ != CUGRAPH_SUCCESS) return; + } + + auto graph = + reinterpret_cast*>( + graph_->graph_); + + auto graph_view = graph->view(); + + auto number_map = reinterpret_cast*>(graph_->number_map_); + + rmm::device_uvector start_vertices(start_vertices_->size_, handle_.get_stream()); + raft::copy(start_vertices.data(), + start_vertices_->as_type(), + start_vertices.size(), + handle_.get_stream()); + + // + // Need to renumber start_vertices + // + cugraph::renumber_ext_vertices( + handle_, + start_vertices.data(), + start_vertices.size(), + number_map->data(), + graph_view.local_vertex_partition_range_first(), + graph_view.local_vertex_partition_range_last(), + false); + + auto [paths, weights] = cugraph::uniform_random_walks( + handle_, + graph_view, + raft::device_span{start_vertices.data(), start_vertices.size()}, + max_length_, + seed_); + + // + // Need to unrenumber the vertices in the resulting paths + // + cugraph::unrenumber_local_int_vertices( + handle_, paths.data(), paths.size(), number_map->data(), 0, paths.size() - 1, false); + + result_ = new cugraph::c_api::cugraph_random_walk_result_t{ + false, + max_length_, + new cugraph::c_api::cugraph_type_erased_device_array_t(paths, graph_->vertex_type_), + weights + ? new cugraph::c_api::cugraph_type_erased_device_array_t(*weights, graph_->weight_type_) + : nullptr, + nullptr}; + } + } +}; + +struct biased_random_walks_functor : public cugraph::c_api::abstract_functor { + raft::handle_t const& handle_; + cugraph::c_api::cugraph_graph_t* graph_{nullptr}; + cugraph::c_api::cugraph_type_erased_device_array_view_t const* start_vertices_{nullptr}; + size_t max_length_{0}; + cugraph::c_api::cugraph_random_walk_result_t* result_{nullptr}; + uint64_t seed_{0}; + + biased_random_walks_functor(cugraph_resource_handle_t const* handle, + cugraph_graph_t* graph, + cugraph_type_erased_device_array_view_t const* start_vertices, + size_t max_length) + : abstract_functor(), + handle_(*reinterpret_cast(handle)->handle_), + graph_(reinterpret_cast(graph)), + start_vertices_( + reinterpret_cast( + start_vertices)), + max_length_(max_length) + { + } + + template + void operator()() + { + // FIXME: Think about how to handle SG vice MG + if constexpr (!cugraph::is_candidate::value) { + unsupported(); + } else if constexpr (multi_gpu) { + unsupported(); + } else { + // random walks expects store_transposed == false + if constexpr (store_transposed) { + error_code_ = cugraph::c_api:: + transpose_storage( + handle_, graph_, error_.get()); + if (error_code_ != CUGRAPH_SUCCESS) return; + } + + auto graph = + reinterpret_cast*>( + graph_->graph_); + + auto graph_view = graph->view(); + + auto number_map = reinterpret_cast*>(graph_->number_map_); + + rmm::device_uvector start_vertices(start_vertices_->size_, handle_.get_stream()); + raft::copy(start_vertices.data(), + start_vertices_->as_type(), + start_vertices.size(), + handle_.get_stream()); + + // + // Need to renumber start_vertices + // + cugraph::renumber_ext_vertices( + handle_, + start_vertices.data(), + start_vertices.size(), + number_map->data(), + graph_view.local_vertex_partition_range_first(), + graph_view.local_vertex_partition_range_last(), + false); + + auto [paths, weights] = cugraph::biased_random_walks( + handle_, + graph_view, + raft::device_span{start_vertices.data(), start_vertices.size()}, + max_length_, + seed_); + + // + // Need to unrenumber the vertices in the resulting paths + // + cugraph::unrenumber_local_int_vertices( + handle_, paths.data(), paths.size(), number_map->data(), 0, paths.size() - 1, false); + + result_ = new cugraph::c_api::cugraph_random_walk_result_t{ + false, + max_length_, + new cugraph::c_api::cugraph_type_erased_device_array_t(paths, graph_->vertex_type_), + weights + ? new cugraph::c_api::cugraph_type_erased_device_array_t(*weights, graph_->weight_type_) + : nullptr, + nullptr}; + } + } +}; + +struct node2vec_random_walks_functor : public cugraph::c_api::abstract_functor { + raft::handle_t const& handle_; + cugraph::c_api::cugraph_graph_t* graph_{nullptr}; + cugraph::c_api::cugraph_type_erased_device_array_view_t const* start_vertices_{nullptr}; + size_t max_length_{0}; + double p_{0}; + double q_{0}; + uint64_t seed_{0}; + cugraph::c_api::cugraph_random_walk_result_t* result_{nullptr}; + + node2vec_random_walks_functor(cugraph_resource_handle_t const* handle, + cugraph_graph_t* graph, + cugraph_type_erased_device_array_view_t const* start_vertices, + size_t max_length, + double p, + double q) + : abstract_functor(), + handle_(*reinterpret_cast(handle)->handle_), + graph_(reinterpret_cast(graph)), + start_vertices_( + reinterpret_cast( + start_vertices)), + max_length_(max_length), + p_(p), + q_(q) + { + } + + template + void operator()() + { + // FIXME: Think about how to handle SG vice MG + if constexpr (!cugraph::is_candidate::value) { + unsupported(); + } else if constexpr (multi_gpu) { + unsupported(); + } else { + // random walks expects store_transposed == false + if constexpr (store_transposed) { + error_code_ = cugraph::c_api:: + transpose_storage( + handle_, graph_, error_.get()); + if (error_code_ != CUGRAPH_SUCCESS) return; + } + + auto graph = + reinterpret_cast*>( + graph_->graph_); + + auto graph_view = graph->view(); + + auto number_map = reinterpret_cast*>(graph_->number_map_); + + rmm::device_uvector start_vertices(start_vertices_->size_, handle_.get_stream()); + raft::copy(start_vertices.data(), + start_vertices_->as_type(), + start_vertices.size(), + handle_.get_stream()); + + // + // Need to renumber start_vertices + // + cugraph::renumber_ext_vertices( + handle_, + start_vertices.data(), + start_vertices.size(), + number_map->data(), + graph_view.local_vertex_partition_range_first(), + graph_view.local_vertex_partition_range_last(), + false); + + auto [paths, weights] = cugraph::node2vec_random_walks( + handle_, + graph_view, + raft::device_span{start_vertices.data(), start_vertices.size()}, + max_length_, + static_cast(p_), + static_cast(q_), + seed_); + + // + // Need to unrenumber the vertices in the resulting paths + // + cugraph::unrenumber_local_int_vertices( + handle_, paths.data(), paths.size(), number_map->data(), 0, paths.size() - 1, false); + + result_ = new cugraph::c_api::cugraph_random_walk_result_t{ + false, + max_length_, + new cugraph::c_api::cugraph_type_erased_device_array_t(paths, graph_->vertex_type_), + weights + ? new cugraph::c_api::cugraph_type_erased_device_array_t(*weights, graph_->weight_type_) + : nullptr, + nullptr}; + } + } +}; + +} // anonymous namespace + cugraph_error_code_t cugraph_node2vec(const cugraph_resource_handle_t* handle, cugraph_graph_t* graph, - const cugraph_type_erased_device_array_view_t* sources, - size_t max_depth, + const cugraph_type_erased_device_array_view_t* start_vertices, + size_t max_length, bool_t compress_results, double p, double q, @@ -153,7 +450,7 @@ cugraph_error_code_t cugraph_node2vec(const cugraph_resource_handle_t* handle, cugraph_error_t** error) { cugraph::c_api::node2vec_functor functor( - handle, graph, sources, max_depth, compress_results, p, q); + handle, graph, start_vertices, max_length, compress_results, p, q); return cugraph::c_api::run_algorithm(graph, functor, result, error); } @@ -196,3 +493,44 @@ void cugraph_random_walk_result_free(cugraph_random_walk_result_t* result) delete internal_pointer->weights_; delete internal_pointer; } + +cugraph_error_code_t cugraph_uniform_random_walks( + const cugraph_resource_handle_t* handle, + cugraph_graph_t* graph, + const cugraph_type_erased_device_array_view_t* start_vertices, + size_t max_length, + cugraph_random_walk_result_t** result, + cugraph_error_t** error) +{ + uniform_random_walks_functor functor(handle, graph, start_vertices, max_length); + + return cugraph::c_api::run_algorithm(graph, functor, result, error); +} + +cugraph_error_code_t cugraph_biased_random_walks( + const cugraph_resource_handle_t* handle, + cugraph_graph_t* graph, + const cugraph_type_erased_device_array_view_t* start_vertices, + size_t max_length, + cugraph_random_walk_result_t** result, + cugraph_error_t** error) +{ + biased_random_walks_functor functor(handle, graph, start_vertices, max_length); + + return cugraph::c_api::run_algorithm(graph, functor, result, error); +} + +cugraph_error_code_t cugraph_node2vec_random_walks( + const cugraph_resource_handle_t* handle, + cugraph_graph_t* graph, + const cugraph_type_erased_device_array_view_t* start_vertices, + size_t max_length, + double p, + double q, + cugraph_random_walk_result_t** result, + cugraph_error_t** error) +{ + node2vec_random_walks_functor functor(handle, graph, start_vertices, max_length, p, q); + + return cugraph::c_api::run_algorithm(graph, functor, result, error); +} diff --git a/cpp/src/sampling/random_walks_mg.cu b/cpp/src/sampling/random_walks_mg.cu new file mode 100644 index 00000000000..ff4a07cb93c --- /dev/null +++ b/cpp/src/sampling/random_walks_mg.cu @@ -0,0 +1,195 @@ +/* + * Copyright (c) 2022, 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 { + +// FIXME: Temporarily here until random_walks_impl.cuh is ready with the real implementation +template +std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed) +{ + CUGRAPH_FAIL("Not Implemented"); +} + +template +std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed) +{ + CUGRAPH_FAIL("Not Implemented"); +} + +template +std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + weight_t p, + weight_t q, + uint64_t seed) +{ + CUGRAPH_FAIL("Not Implemented"); +} + +template std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + float p, + float q, + uint64_t seed); + +template std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + float p, + float q, + uint64_t seed); + +template std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + float p, + float q, + uint64_t seed); + +template std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + double p, + double q, + uint64_t seed); + +template std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + double p, + double q, + uint64_t seed); + +template std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + double p, + double q, + uint64_t seed); + +} // namespace cugraph diff --git a/cpp/src/sampling/random_walks_sg.cu b/cpp/src/sampling/random_walks_sg.cu new file mode 100644 index 00000000000..e7634795eb6 --- /dev/null +++ b/cpp/src/sampling/random_walks_sg.cu @@ -0,0 +1,195 @@ +/* + * Copyright (c) 2022, 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 { + +// FIXME: Temporarily here until random_walks_impl.cuh is ready with the real implementation +template +std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed) +{ + CUGRAPH_FAIL("Not Implemented"); +} + +template +std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed) +{ + CUGRAPH_FAIL("Not Implemented"); +} + +template +std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + weight_t p, + weight_t q, + uint64_t seed) +{ + CUGRAPH_FAIL("Not Implemented"); +} + +template std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +uniform_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +biased_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + uint64_t seed); + +template std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + float p, + float q, + uint64_t seed); + +template std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + float p, + float q, + uint64_t seed); + +template std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + float p, + float q, + uint64_t seed); + +template std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + double p, + double q, + uint64_t seed); + +template std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + double p, + double q, + uint64_t seed); + +template std::tuple, std::optional>> +node2vec_random_walks(raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_length, + double p, + double q, + uint64_t seed); + +} // namespace cugraph diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 80fb2580c03..dead8173993 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -452,9 +452,14 @@ ConfigureTest(EIGENVECTOR_CENTRALITY_TEST centrality/eigenvector_centrality_test # - WEAKLY CONNECTED COMPONENTS tests ------------------------------------------------------------- ConfigureTest(WEAKLY_CONNECTED_COMPONENTS_TEST components/weakly_connected_components_test.cpp) +################################################################################################### +# - Legacy RANDOM_WALKS tests ---------------------------------------------------------------------------- +ConfigureTest(LEGACY_RANDOM_WALKS_TEST sampling/random_walks_test.cu) + ################################################################################################### # - RANDOM_WALKS tests ---------------------------------------------------------------------------- -ConfigureTest(RANDOM_WALKS_TEST sampling/random_walks_test.cu) +# FIXME: Rename to random_walks_test.cu once the legacy implementation is deleted +ConfigureTest(RANDOM_WALKS_TEST sampling/sg_random_walks_test.cu) ################################################################################################### ConfigureTest(RANDOM_WALKS_LOW_LEVEL_TEST sampling/rw_low_level_test.cu) @@ -618,6 +623,10 @@ if(BUILD_CUGRAPH_MG_TESTS) ########################################################################################### # - MG NBR SAMPLING tests ----------------------------------------------------------------- ConfigureTestMG(MG_UNIFORM_NEIGHBOR_SAMPLING_TEST sampling/mg_uniform_neighbor_sampling.cu) + + ########################################################################################### + # - RANDOM_WALKS tests -------------------------------------------------------------------- + ConfigureTestMG(MG_RANDOM_WALKS_TEST sampling/mg_random_walks_test.cu) ########################################################################################### # - MG C API tests ------------------------------------------------------------------------ @@ -631,6 +640,7 @@ if(BUILD_CUGRAPH_MG_TESTS) ConfigureCTestMG(MG_CAPI_EIGENVECTOR_CENTRALITY c_api/mg_eigenvector_centrality_test.c c_api/mg_test_utils.cpp) ConfigureCTestMG(MG_CAPI_HITS c_api/mg_hits_test.c c_api/mg_test_utils.cpp) ConfigureCTestMG(MG_CAPI_UNIFORM_NEIGHBOR_SAMPLE c_api/mg_uniform_neighbor_sample_test.c c_api/mg_test_utils.cpp) + ConfigureCTestMG(MG_CAPI_RANDOM_WALKS c_api/mg_random_walks_test.c c_api/mg_test_utils.cpp) ConfigureCTestMG(MG_CAPI_TRIANGLE_COUNT c_api/mg_triangle_count_test.c c_api/mg_test_utils.cpp) ConfigureCTestMG(MG_CAPI_LOUVAIN c_api/mg_louvain_test.c c_api/mg_test_utils.cpp) ConfigureCTestMG(MG_CAPI_CORE_NUMBER c_api/mg_core_number_test.c c_api/mg_test_utils.cpp) @@ -684,6 +694,7 @@ ConfigureCTest(CAPI_NODE2VEC_TEST c_api/node2vec_test.c) ConfigureCTest(CAPI_WEAKLY_CONNECTED_COMPONENTS c_api/weakly_connected_components_test.c) ConfigureCTest(CAPI_STRONGLY_CONNECTED_COMPONENTS c_api/strongly_connected_components_test.c) ConfigureCTest(CAPI_UNIFORM_NEIGHBOR_SAMPLE c_api/uniform_neighbor_sample_test.c) +ConfigureCTest(CAPI_RANDOM_WALKS c_api/sg_random_walks_test.c) ConfigureCTest(CAPI_TRIANGLE_COUNT c_api/triangle_count_test.c) ConfigureCTest(CAPI_LOUVAIN c_api/louvain_test.c) ConfigureCTest(CAPI_CORE_NUMBER c_api/core_number_test.c) diff --git a/cpp/tests/c_api/mg_random_walks_test.c b/cpp/tests/c_api/mg_random_walks_test.c new file mode 100644 index 00000000000..fe512bda30d --- /dev/null +++ b/cpp/tests/c_api/mg_random_walks_test.c @@ -0,0 +1,423 @@ +/* + * Copyright (c) 2022, 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 "mg_test_utils.h" /* RUN_TEST */ + +#include +#include + +#include + +typedef int32_t vertex_t; +typedef int32_t edge_t; +typedef float weight_t; + +int generic_uniform_random_walks_test(const cugraph_resource_handle_t* handle, + vertex_t* h_src, + vertex_t* h_dst, + weight_t* h_wgt, + size_t num_vertices, + size_t num_edges, + vertex_t* h_start, + size_t num_starts, + size_t max_depth, + bool_t store_transposed) +{ + int test_ret_value = 0; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + cugraph_graph_t* graph = NULL; + cugraph_random_walk_result_t* result = NULL; + + cugraph_type_erased_device_array_t* d_start = NULL; + cugraph_type_erased_device_array_view_t* d_start_view = NULL; + + ret_code = create_mg_test_graph( + handle, h_src, h_dst, h_wgt, num_edges, store_transposed, FALSE, &graph, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_starts, INT32, &d_start, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "d_start create failed."); + + d_start_view = cugraph_type_erased_device_array_view(d_start); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, d_start_view, (byte_t*)h_start, &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "start copy_from_host failed."); + + ret_code = cugraph_uniform_random_walks(handle, graph, d_start_view, FALSE, &result, &ret_error); + +#if 1 + TEST_ASSERT( + test_ret_value, ret_code != CUGRAPH_SUCCESS, "uniform_random_walks should have failed") +#else + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "uniform_random_walks failed."); + + cugraph_type_erased_device_array_view_t* verts; + cugraph_type_erased_device_array_view_t* wgts; + + verts = cugraph_random_walk_result_get_paths(result); + wgts = cugraph_random_walk_result_get_weights(result); + + size_t verts_size = cugraph_type_erased_device_array_view_size(verts); + size_t wgts_size = cugraph_type_erased_device_array_view_size(wgts); + + vertex_t h_result_verts[verts_size]; + vertex_t h_result_wgts[wgts_size]; + + ret_code = + cugraph_type_erased_device_array_view_copy_to_host(handle, (byte_t*)h_verts, verts, &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_result_wgts, wgts, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + // NOTE: The C++ tester does a more thorough validation. For our purposes + // here we will do a simpler validation, merely checking that all edges + // are actually part of the graph + weight_t M[num_vertices][num_vertices]; + + for (int i = 0; i < num_vertices; ++i) + for (int j = 0; j < num_vertices; ++j) + M[i][j] = -1; + + for (int i = 0; i < num_edges; ++i) + M[h_src[i]][h_dst[i]] = h_wgt[i]; + + TEST_ASSERT(test_ret_value, + cugraph_random_walk_result_get_max_path_length() == max_depth, + "path length does not match"); + + for (int i = 0; (i < num_starts) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + M[h_start[i]][h_result_verts[i * (max_depth + 1)]] == h_result_wgts[i * max_depth], + "uniform_random_walks got edge that doesn't exist"); + for (size_t j = 1; j < cugraph_random_walk_result_get_max_path_length(); ++j) + TEST_ASSERT( + test_ret_value, + M[h_start[i * (max_depth + 1) + j - 1]][h_result_verts[i * (max_depth + 1) + j]] == + h_result_wgts[i * max_depth + j - 1], + "uniform_random_walks got edge that doesn't exist"); + } + + cugraph_random_walk_result_free(result); +#endif + + cugraph_mg_graph_free(graph); + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int generic_biased_random_walks_test(const cugraph_resource_handle_t* handle, + vertex_t* h_src, + vertex_t* h_dst, + weight_t* h_wgt, + size_t num_vertices, + size_t num_edges, + vertex_t* h_start, + size_t num_starts, + size_t max_depth, + bool_t store_transposed) +{ + int test_ret_value = 0; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + cugraph_graph_t* graph = NULL; + cugraph_random_walk_result_t* result = NULL; + + cugraph_type_erased_device_array_t* d_start = NULL; + cugraph_type_erased_device_array_view_t* d_start_view = NULL; + + ret_code = create_mg_test_graph( + handle, h_src, h_dst, h_wgt, num_edges, store_transposed, FALSE, &graph, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_starts, INT32, &d_start, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "d_start create failed."); + + d_start_view = cugraph_type_erased_device_array_view(d_start); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, d_start_view, (byte_t*)h_start, &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "start copy_from_host failed."); + + ret_code = cugraph_biased_random_walks(handle, graph, d_start_view, FALSE, &result, &ret_error); + +#if 1 + TEST_ASSERT(test_ret_value, ret_code != CUGRAPH_SUCCESS, "biased_random_walks should have failed") +#else + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "biased_random_walks failed."); + + cugraph_type_erased_device_array_view_t* verts; + cugraph_type_erased_device_array_view_t* wgts; + + verts = cugraph_random_walk_result_get_paths(result); + wgts = cugraph_random_walk_result_get_weights(result); + + size_t verts_size = cugraph_type_erased_device_array_view_size(verts); + size_t wgts_size = cugraph_type_erased_device_array_view_size(wgts); + + vertex_t h_result_verts[verts_size]; + vertex_t h_result_wgts[wgts_size]; + + ret_code = + cugraph_type_erased_device_array_view_copy_to_host(handle, (byte_t*)h_verts, verts, &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_result_wgts, wgts, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + // NOTE: The C++ tester does a more thorough validation. For our purposes + // here we will do a simpler validation, merely checking that all edges + // are actually part of the graph + weight_t M[num_vertices][num_vertices]; + + for (int i = 0; i < num_vertices; ++i) + for (int j = 0; j < num_vertices; ++j) + M[i][j] = -1; + + for (int i = 0; i < num_edges; ++i) + M[h_src[i]][h_dst[i]] = h_wgt[i]; + + TEST_ASSERT(test_ret_value, + cugraph_random_walk_result_get_max_path_length() == max_depth, + "path length does not match"); + + for (int i = 0; (i < num_starts) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + M[h_start[i]][h_result_verts[i * (max_depth + 1)]] == h_result_wgts[i * max_depth], + "biased_random_walks got edge that doesn't exist"); + for (size_t j = 1; j < cugraph_random_walk_result_get_max_path_length(); ++j) + TEST_ASSERT( + test_ret_value, + M[h_start[i * (max_depth + 1) + j - 1]][h_result_verts[i * (max_depth + 1) + j]] == + h_result_wgts[i * max_depth + j - 1], + "biased_random_walks got edge that doesn't exist"); + } + + cugraph_random_walk_result_free(result); +#endif + + cugraph_mg_graph_free(graph); + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int generic_node2vec_random_walks_test(const cugraph_resource_handle_t* handle, + vertex_t* h_src, + vertex_t* h_dst, + weight_t* h_wgt, + size_t num_vertices, + size_t num_edges, + vertex_t* h_start, + size_t num_starts, + size_t max_depth, + float p, + float q, + bool_t store_transposed) +{ + int test_ret_value = 0; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + cugraph_graph_t* graph = NULL; + cugraph_random_walk_result_t* result = NULL; + + cugraph_type_erased_device_array_t* d_start = NULL; + cugraph_type_erased_device_array_view_t* d_start_view = NULL; + + ret_code = create_mg_test_graph( + handle, h_src, h_dst, h_wgt, num_edges, store_transposed, FALSE, &graph, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_starts, INT32, &d_start, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "d_start create failed."); + + d_start_view = cugraph_type_erased_device_array_view(d_start); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, d_start_view, (byte_t*)h_start, &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "start copy_from_host failed."); + + ret_code = + cugraph_node2vec_random_walks(handle, graph, d_start_view, FALSE, p, q, &result, &ret_error); + +#if 1 + TEST_ASSERT( + test_ret_value, ret_code != CUGRAPH_SUCCESS, "node2vec_random_walks should have failed") +#else + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "node2vec_random_walks failed."); + + cugraph_type_erased_device_array_view_t* verts; + cugraph_type_erased_device_array_view_t* wgts; + + verts = cugraph_random_walk_result_get_paths(result); + wgts = cugraph_random_walk_result_get_weights(result); + + size_t verts_size = cugraph_type_erased_device_array_view_size(verts); + size_t wgts_size = cugraph_type_erased_device_array_view_size(wgts); + + vertex_t h_result_verts[verts_size]; + vertex_t h_result_wgts[wgts_size]; + + ret_code = + cugraph_type_erased_device_array_view_copy_to_host(handle, (byte_t*)h_verts, verts, &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_result_wgts, wgts, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + // NOTE: The C++ tester does a more thorough validation. For our purposes + // here we will do a simpler validation, merely checking that all edges + // are actually part of the graph + weight_t M[num_vertices][num_vertices]; + + for (int i = 0; i < num_vertices; ++i) + for (int j = 0; j < num_vertices; ++j) + M[i][j] = -1; + + for (int i = 0; i < num_edges; ++i) + M[h_src[i]][h_dst[i]] = h_wgt[i]; + + TEST_ASSERT(test_ret_value, + cugraph_random_walk_result_get_max_path_length() == max_depth, + "path length does not match"); + + for (int i = 0; (i < num_starts) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + M[h_start[i]][h_result_verts[i * (max_depth + 1)]] == h_result_wgts[i * max_depth], + "node2vec_random_walks got edge that doesn't exist"); + for (size_t j = 1; j < cugraph_random_walk_result_get_max_path_length(); ++j) + TEST_ASSERT( + test_ret_value, + M[h_start[i * (max_depth + 1) + j - 1]][h_result_verts[i * (max_depth + 1) + j]] == + h_result_wgts[i * max_depth + j - 1], + "node2vec_random_walks got edge that doesn't exist"); + } + + cugraph_random_walk_result_free(result); +#endif + + cugraph_mg_graph_free(graph); + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int test_uniform_random_walks(const cugraph_resource_handle_t* handle) +{ + size_t num_edges = 8; + size_t num_vertices = 6; + size_t num_starts = 2; + + vertex_t src[] = {0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; + weight_t wgt[] = {1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0}; + vertex_t start[] = {2, 2}; + + return generic_uniform_random_walks_test( + handle, src, dst, wgt, num_vertices, num_edges, start, num_starts, FALSE, FALSE); +} + +int test_biased_random_walks(const cugraph_resource_handle_t* handle) +{ + size_t num_edges = 8; + size_t num_vertices = 6; + size_t num_starts = 2; + + vertex_t src[] = {0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; + weight_t wgt[] = {1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0}; + vertex_t start[] = {2, 2}; + + return generic_biased_random_walks_test( + handle, src, dst, wgt, num_vertices, num_edges, start, num_starts, FALSE, FALSE); +} + +int test_node2vec_random_walks(const cugraph_resource_handle_t* handle) +{ + size_t num_edges = 8; + size_t num_vertices = 6; + size_t num_starts = 2; + + vertex_t src[] = {0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; + weight_t wgt[] = {1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0}; + vertex_t start[] = {2, 2}; + + weight_t p = 5; + weight_t q = 8; + + return generic_node2vec_random_walks_test( + handle, src, dst, wgt, num_vertices, num_edges, start, num_starts, p, q, FALSE, FALSE); +} + +int main(int argc, char** argv) +{ + // Set up MPI: + int comm_rank; + int comm_size; + int num_gpus_per_node; + cudaError_t status; + int mpi_status; + int result = 0; + cugraph_resource_handle_t* handle = NULL; + cugraph_error_t* ret_error; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + int prows = 1; + + C_MPI_TRY(MPI_Init(&argc, &argv)); + C_MPI_TRY(MPI_Comm_rank(MPI_COMM_WORLD, &comm_rank)); + C_MPI_TRY(MPI_Comm_size(MPI_COMM_WORLD, &comm_size)); + C_CUDA_TRY(cudaGetDeviceCount(&num_gpus_per_node)); + C_CUDA_TRY(cudaSetDevice(comm_rank % num_gpus_per_node)); + + void* raft_handle = create_raft_handle(prows); + handle = cugraph_create_resource_handle(raft_handle); + + if (result == 0) { + result |= RUN_MG_TEST(test_uniform_random_walks, handle); + result |= RUN_MG_TEST(test_biased_random_walks, handle); + result |= RUN_MG_TEST(test_node2vec_random_walks, handle); + + cugraph_free_resource_handle(handle); + } + + free_raft_handle(raft_handle); + + C_MPI_TRY(MPI_Finalize()); + + return result; +} diff --git a/cpp/tests/c_api/sg_random_walks_test.c b/cpp/tests/c_api/sg_random_walks_test.c new file mode 100644 index 00000000000..cd771b5f457 --- /dev/null +++ b/cpp/tests/c_api/sg_random_walks_test.c @@ -0,0 +1,410 @@ +/* + * Copyright (c) 2022, 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 "c_test_utils.h" /* RUN_TEST */ + +#include +#include + +#include + +typedef int32_t vertex_t; +typedef int32_t edge_t; +typedef float weight_t; + +int generic_uniform_random_walks_test(vertex_t* h_src, + vertex_t* h_dst, + weight_t* h_wgt, + size_t num_vertices, + size_t num_edges, + vertex_t* h_start, + size_t num_starts, + size_t max_depth, + bool_t renumber, + bool_t store_transposed) +{ + int test_ret_value = 0; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + cugraph_resource_handle_t* handle = NULL; + cugraph_graph_t* graph = NULL; + cugraph_random_walk_result_t* result = NULL; + + cugraph_type_erased_device_array_t* d_start = NULL; + cugraph_type_erased_device_array_view_t* d_start_view = NULL; + + handle = cugraph_create_resource_handle(NULL); + TEST_ASSERT(test_ret_value, handle != NULL, "resource handle creation failed."); + + ret_code = create_test_graph( + handle, h_src, h_dst, h_wgt, num_edges, store_transposed, renumber, FALSE, &graph, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_starts, INT32, &d_start, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "d_start create failed."); + + d_start_view = cugraph_type_erased_device_array_view(d_start); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, d_start_view, (byte_t*)h_start, &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "start copy_from_host failed."); + + ret_code = + cugraph_uniform_random_walks(handle, graph, d_start_view, max_depth, &result, &ret_error); + +#if 1 + TEST_ASSERT( + test_ret_value, ret_code != CUGRAPH_SUCCESS, "uniform_random_walks should have failed") +#else + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "uniform_random_walks failed."); + + cugraph_type_erased_device_array_view_t* verts; + cugraph_type_erased_device_array_view_t* wgts; + + verts = cugraph_random_walk_result_get_paths(result); + wgts = cugraph_random_walk_result_get_weights(result); + + size_t verts_size = cugraph_type_erased_device_array_view_size(verts); + size_t wgts_size = cugraph_type_erased_device_array_view_size(wgts); + + vertex_t h_result_verts[verts_size]; + vertex_t h_result_wgts[wgts_size]; + + ret_code = + cugraph_type_erased_device_array_view_copy_to_host(handle, (byte_t*)h_verts, verts, &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_result_wgts, wgts, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + // NOTE: The C++ tester does a more thorough validation. For our purposes + // here we will do a simpler validation, merely checking that all edges + // are actually part of the graph + weight_t M[num_vertices][num_vertices]; + + for (int i = 0; i < num_vertices; ++i) + for (int j = 0; j < num_vertices; ++j) + M[i][j] = -1; + + for (int i = 0; i < num_edges; ++i) + M[h_src[i]][h_dst[i]] = h_wgt[i]; + + TEST_ASSERT(test_ret_value, + cugraph_random_walk_result_get_max_path_length() == max_depth, + "path length does not match"); + + for (int i = 0; (i < num_starts) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + M[h_start[i]][h_result_verts[i * (max_depth + 1)]] == h_result_wgts[i * max_depth], + "uniform_random_walks got edge that doesn't exist"); + for (size_t j = 1; j < cugraph_random_walk_result_get_max_path_length(); ++j) + TEST_ASSERT( + test_ret_value, + M[h_start[i * (max_depth + 1) + j - 1]][h_result_verts[i * (max_depth + 1) + j]] == + h_result_wgts[i * max_depth + j - 1], + "uniform_random_walks got edge that doesn't exist"); + } + + cugraph_random_walk_result_free(result); +#endif + + cugraph_sg_graph_free(graph); + cugraph_free_resource_handle(handle); + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int generic_biased_random_walks_test(vertex_t* h_src, + vertex_t* h_dst, + weight_t* h_wgt, + size_t num_vertices, + size_t num_edges, + vertex_t* h_start, + size_t num_starts, + size_t max_depth, + bool_t renumber, + bool_t store_transposed) +{ + int test_ret_value = 0; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + cugraph_resource_handle_t* handle = NULL; + cugraph_graph_t* graph = NULL; + cugraph_random_walk_result_t* result = NULL; + + cugraph_type_erased_device_array_t* d_start = NULL; + cugraph_type_erased_device_array_view_t* d_start_view = NULL; + + handle = cugraph_create_resource_handle(NULL); + TEST_ASSERT(test_ret_value, handle != NULL, "resource handle creation failed."); + + ret_code = create_test_graph( + handle, h_src, h_dst, h_wgt, num_edges, store_transposed, renumber, FALSE, &graph, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_starts, INT32, &d_start, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "d_start create failed."); + + d_start_view = cugraph_type_erased_device_array_view(d_start); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, d_start_view, (byte_t*)h_start, &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "start copy_from_host failed."); + + ret_code = + cugraph_biased_random_walks(handle, graph, d_start_view, max_depth, &result, &ret_error); + +#if 1 + TEST_ASSERT(test_ret_value, ret_code != CUGRAPH_SUCCESS, "biased_random_walks should have failed") +#else + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "biased_random_walks failed."); + + cugraph_type_erased_device_array_view_t* verts; + cugraph_type_erased_device_array_view_t* wgts; + + verts = cugraph_random_walk_result_get_paths(result); + wgts = cugraph_random_walk_result_get_weights(result); + + size_t verts_size = cugraph_type_erased_device_array_view_size(verts); + size_t wgts_size = cugraph_type_erased_device_array_view_size(wgts); + + vertex_t h_result_verts[verts_size]; + vertex_t h_result_wgts[wgts_size]; + + ret_code = + cugraph_type_erased_device_array_view_copy_to_host(handle, (byte_t*)h_verts, verts, &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_result_wgts, wgts, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + // NOTE: The C++ tester does a more thorough validation. For our purposes + // here we will do a simpler validation, merely checking that all edges + // are actually part of the graph + weight_t M[num_vertices][num_vertices]; + + for (int i = 0; i < num_vertices; ++i) + for (int j = 0; j < num_vertices; ++j) + M[i][j] = -1; + + for (int i = 0; i < num_edges; ++i) + M[h_src[i]][h_dst[i]] = h_wgt[i]; + + TEST_ASSERT(test_ret_value, + cugraph_random_walk_result_get_max_path_length() == max_depth, + "path length does not match"); + + for (int i = 0; (i < num_starts) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + M[h_start[i]][h_result_verts[i * (max_depth + 1)]] == h_result_wgts[i * max_depth], + "biased_random_walks got edge that doesn't exist"); + for (size_t j = 1; j < cugraph_random_walk_result_get_max_path_length(); ++j) + TEST_ASSERT( + test_ret_value, + M[h_start[i * (max_depth + 1) + j - 1]][h_result_verts[i * (max_depth + 1) + j]] == + h_result_wgts[i * max_depth + j - 1], + "biased_random_walks got edge that doesn't exist"); + } + + cugraph_random_walk_result_free(result); +#endif + + cugraph_sg_graph_free(graph); + cugraph_free_resource_handle(handle); + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int generic_node2vec_random_walks_test(vertex_t* h_src, + vertex_t* h_dst, + weight_t* h_wgt, + size_t num_vertices, + size_t num_edges, + vertex_t* h_start, + size_t num_starts, + size_t max_depth, + weight_t p, + weight_t q, + bool_t renumber, + bool_t store_transposed) +{ + int test_ret_value = 0; + + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + cugraph_resource_handle_t* handle = NULL; + cugraph_graph_t* graph = NULL; + cugraph_random_walk_result_t* result = NULL; + + cugraph_type_erased_device_array_t* d_start = NULL; + cugraph_type_erased_device_array_view_t* d_start_view = NULL; + + handle = cugraph_create_resource_handle(NULL); + TEST_ASSERT(test_ret_value, handle != NULL, "resource handle creation failed."); + + ret_code = create_test_graph( + handle, h_src, h_dst, h_wgt, num_edges, store_transposed, renumber, FALSE, &graph, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_starts, INT32, &d_start, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "d_start create failed."); + + d_start_view = cugraph_type_erased_device_array_view(d_start); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, d_start_view, (byte_t*)h_start, &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "start copy_from_host failed."); + + ret_code = cugraph_node2vec_random_walks( + handle, graph, d_start_view, max_depth, p, q, &result, &ret_error); + +#if 1 + TEST_ASSERT( + test_ret_value, ret_code != CUGRAPH_SUCCESS, "node2vec_random_walks should have failed") +#else + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "node2vec_random_walks failed."); + + cugraph_type_erased_device_array_view_t* verts; + cugraph_type_erased_device_array_view_t* wgts; + + verts = cugraph_random_walk_result_get_paths(result); + wgts = cugraph_random_walk_result_get_weights(result); + + size_t verts_size = cugraph_type_erased_device_array_view_size(verts); + size_t wgts_size = cugraph_type_erased_device_array_view_size(wgts); + + vertex_t h_result_verts[verts_size]; + vertex_t h_result_wgts[wgts_size]; + + ret_code = + cugraph_type_erased_device_array_view_copy_to_host(handle, (byte_t*)h_verts, verts, &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_result_wgts, wgts, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + // NOTE: The C++ tester does a more thorough validation. For our purposes + // here we will do a simpler validation, merely checking that all edges + // are actually part of the graph + weight_t M[num_vertices][num_vertices]; + + for (int i = 0; i < num_vertices; ++i) + for (int j = 0; j < num_vertices; ++j) + M[i][j] = -1; + + for (int i = 0; i < num_edges; ++i) + M[h_src[i]][h_dst[i]] = h_wgt[i]; + + TEST_ASSERT(test_ret_value, + cugraph_random_walk_result_get_max_path_length() == max_depth, + "path length does not match"); + + for (int i = 0; (i < num_starts) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + M[h_start[i]][h_result_verts[i * (max_depth + 1)]] == h_result_wgts[i * max_depth], + "node2vec_random_walks got edge that doesn't exist"); + for (size_t j = 1; j < cugraph_random_walk_result_get_max_path_length(); ++j) + TEST_ASSERT( + test_ret_value, + M[h_start[i * (max_depth + 1) + j - 1]][h_result_verts[i * (max_depth + 1) + j]] == + h_result_wgts[i * max_depth + j - 1], + "node2vec_random_walks got edge that doesn't exist"); + } + + cugraph_random_walk_result_free(result); +#endif + + cugraph_sg_graph_free(graph); + cugraph_free_resource_handle(handle); + cugraph_error_free(ret_error); + + return test_ret_value; +} + +int test_uniform_random_walks() +{ + size_t num_edges = 8; + size_t num_vertices = 6; + size_t num_starts = 2; + + vertex_t src[] = {0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; + weight_t wgt[] = {0, 1, 2, 3, 4, 5, 6, 7}; + vertex_t start[] = {2, 2}; + + return generic_uniform_random_walks_test( + src, dst, wgt, num_vertices, num_edges, start, num_starts, TRUE, FALSE, FALSE); +} + +int test_biased_random_walks() +{ + size_t num_edges = 8; + size_t num_vertices = 6; + size_t num_starts = 2; + + vertex_t src[] = {0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; + weight_t wgt[] = {0, 1, 2, 3, 4, 5, 6, 7}; + vertex_t start[] = {2, 2}; + + return generic_biased_random_walks_test( + src, dst, wgt, num_vertices, num_edges, start, num_starts, TRUE, FALSE, FALSE); +} + +int test_node2vec_random_walks() +{ + size_t num_edges = 8; + size_t num_vertices = 6; + size_t num_starts = 2; + + vertex_t src[] = {0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; + weight_t wgt[] = {0, 1, 2, 3, 4, 5, 6, 7}; + vertex_t start[] = {2, 2}; + weight_t p = 5; + weight_t q = 9; + + return generic_node2vec_random_walks_test( + src, dst, wgt, num_vertices, num_edges, start, num_starts, p, q, TRUE, FALSE, FALSE); +} + +int main(int argc, char** argv) +{ + int result = 0; + result |= RUN_TEST(test_uniform_random_walks); + result |= RUN_TEST(test_biased_random_walks); + result |= RUN_TEST(test_node2vec_random_walks); + return result; +} diff --git a/cpp/tests/sampling/mg_random_walks_test.cu b/cpp/tests/sampling/mg_random_walks_test.cu new file mode 100644 index 00000000000..b2a225e7b74 --- /dev/null +++ b/cpp/tests/sampling/mg_random_walks_test.cu @@ -0,0 +1,225 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +#include +#include + +#include + +struct UniformRandomWalks_Usecase { + bool test_weighted{false}; + uint64_t seed{0}; + bool check_correctness{false}; + + template + std::tuple, std::optional>> + operator()(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_depth) + { + return cugraph::uniform_random_walks(handle, graph_view, start_vertices, max_depth, seed); + } +}; + +struct BiasedRandomWalks_Usecase { + bool test_weighted{true}; + uint64_t seed{0}; + bool check_correctness{false}; + + template + std::tuple, std::optional>> + operator()(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_depth) + { + return cugraph::biased_random_walks(handle, graph_view, start_vertices, max_depth, seed); + } +}; + +struct Node2VecRandomWalks_Usecase { + double p{1}; + double q{1}; + bool test_weighted{false}; + uint64_t seed{0}; + bool check_correctness{false}; + + template + std::tuple, std::optional>> + operator()(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_depth) + { + return cugraph::node2vec_random_walks( + handle, graph_view, start_vertices, max_depth, p, q, seed); + } +}; + +template +class Tests_RandomWalks : public ::testing::TestWithParam { + public: + Tests_RandomWalks() {} + static void SetupTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(tuple_t const& param) + { + raft::handle_t handle{}; + HighResClock hr_clock{}; + + raft::comms::initialize_mpi_comms(&handle, MPI_COMM_WORLD); + auto& comm = handle.get_comms(); + auto const comm_size = comm.get_size(); + + auto row_comm_size = static_cast(sqrt(static_cast(comm_size))); + while (comm_size % row_comm_size != 0) { + --row_comm_size; + } + + cugraph::partition_2d::subcomm_factory_t + subcomm_factory(handle, row_comm_size); + + auto [randomwalks_usecase, input_usecase] = param; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_clock.start(); + } + + bool renumber{true}; + auto [graph, d_renumber_map_labels] = + cugraph::test::construct_graph( + handle, input_usecase, randomwalks_usecase.test_weighted, renumber); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + double elapsed_time{0.0}; + hr_clock.stop(&elapsed_time); + std::cout << "construct_graph took " << elapsed_time * 1e-6 << " s.\n"; + } + + auto graph_view = graph.view(); + + edge_t num_paths = 10; + rmm::device_uvector d_start(num_paths, handle.get_stream()); + + thrust::tabulate(handle.get_thrust_policy(), + d_start.begin(), + d_start.end(), + [num_vertices = graph_view.number_of_vertices()] __device__(auto idx) { + return (idx % num_vertices); + }); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_clock.start(); + } + +#if 0 + auto [vertices, weights] = randomwalks_usecase( + handle, graph_view, raft::device_span{d_start.data(), d_start.size()}, size_t{10}); +#else + EXPECT_THROW( + randomwalks_usecase(handle, + graph_view, + raft::device_span{d_start.data(), d_start.size()}, + size_t{10}), + std::exception); +#endif + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + double elapsed_time{0.0}; + hr_clock.stop(&elapsed_time); + std::cout << "PageRank took " << elapsed_time * 1e-6 << " s.\n"; + } + + if (randomwalks_usecase.check_correctness) { +#if 0 + // FIXME: Need an MG test +#endif + } + } +}; + +using Tests_UniformRandomWalks_File = + Tests_RandomWalks>; +using Tests_UniformRandomWalks_Rmat = + Tests_RandomWalks>; +using Tests_BiasedRandomWalks_File = + Tests_RandomWalks>; +using Tests_BiasedRandomWalks_Rmat = + Tests_RandomWalks>; +using Tests_Node2VecRandomWalks_File = + Tests_RandomWalks>; +using Tests_Node2VecRandomWalks_Rmat = + Tests_RandomWalks>; + +TEST_P(Tests_UniformRandomWalks_File, Initialize_i32_i32_f) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +INSTANTIATE_TEST_SUITE_P( + simple_test, + Tests_UniformRandomWalks_File, + ::testing::Combine( + ::testing::Values(UniformRandomWalks_Usecase{}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), + cugraph::test::File_Usecase("test/datasets/web-Google.mtx"), + cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), + cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + simple_test, + Tests_BiasedRandomWalks_File, + ::testing::Combine( + ::testing::Values(BiasedRandomWalks_Usecase{}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), + cugraph::test::File_Usecase("test/datasets/web-Google.mtx"), + cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), + cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + simple_test, + Tests_Node2VecRandomWalks_File, + ::testing::Combine( + ::testing::Values(Node2VecRandomWalks_Usecase{4, 8}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), + cugraph::test::File_Usecase("test/datasets/web-Google.mtx"), + cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), + cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); + +CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/sampling/sg_random_walks_test.cu b/cpp/tests/sampling/sg_random_walks_test.cu new file mode 100644 index 00000000000..51ced4ee700 --- /dev/null +++ b/cpp/tests/sampling/sg_random_walks_test.cu @@ -0,0 +1,237 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include +#include +#include + +#include +#include +#include +#include + +#include + +#if 0 +#include "cuda_profiler_api.h" + +#include +#include +#include +#include + +#include + +#include + +#include "random_walks_utils.cuh" + +#include +#include +#include +#include +#include +#include +#endif + +struct UniformRandomWalks_Usecase { + bool test_weighted{false}; + uint64_t seed{0}; + bool check_correctness{false}; + + template + std::tuple, std::optional>> + operator()(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_depth) + { + return cugraph::uniform_random_walks(handle, graph_view, start_vertices, max_depth, seed); + } +}; + +struct BiasedRandomWalks_Usecase { + bool test_weighted{true}; + uint64_t seed{0}; + bool check_correctness{false}; + + template + std::tuple, std::optional>> + operator()(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_depth) + { + return cugraph::biased_random_walks(handle, graph_view, start_vertices, max_depth, seed); + } +}; + +struct Node2VecRandomWalks_Usecase { + double p{1}; + double q{1}; + bool test_weighted{false}; + uint64_t seed{0}; + bool check_correctness{false}; + + template + std::tuple, std::optional>> + operator()(raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::device_span start_vertices, + size_t max_depth) + { + return cugraph::node2vec_random_walks( + handle, graph_view, start_vertices, max_depth, p, q, seed); + } +}; + +template +class Tests_RandomWalks : public ::testing::TestWithParam { + public: + Tests_RandomWalks() {} + static void SetupTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(tuple_t const& param) + { + raft::handle_t handle{}; + HighResClock hr_clock{}; + + auto [randomwalks_usecase, input_usecase] = param; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_clock.start(); + } + + // TODO: Do I need to turn renumber off? It's off in the old test + bool renumber{true}; + auto [graph, d_renumber_map_labels] = + cugraph::test::construct_graph( + handle, input_usecase, randomwalks_usecase.test_weighted, renumber); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + double elapsed_time{0.0}; + hr_clock.stop(&elapsed_time); + std::cout << "construct_graph took " << elapsed_time * 1e-6 << " s.\n"; + } + + auto graph_view = graph.view(); + + edge_t num_paths = 10; + rmm::device_uvector d_start(num_paths, handle.get_stream()); + + thrust::tabulate(handle.get_thrust_policy(), + d_start.begin(), + d_start.end(), + [num_vertices = graph_view.number_of_vertices()] __device__(auto idx) { + return (idx % num_vertices); + }); + + edge_t max_depth{10}; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_clock.start(); + } + +#if 0 + auto [vertices, weights] = randomwalks_usecase( + handle, graph_view, raft::device_span{d_start.data(), d_start.size()}, size_t{10}); +#else + EXPECT_THROW( + randomwalks_usecase(handle, + graph_view, + raft::device_span{d_start.data(), d_start.size()}, + size_t{10}), + std::exception); +#endif + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + double elapsed_time{0.0}; + hr_clock.stop(&elapsed_time); + std::cout << "PageRank took " << elapsed_time * 1e-6 << " s.\n"; + } + + if (randomwalks_usecase.check_correctness) { +#if 0 + bool test_all_paths = + cugraph::test::host_check_rw_paths(handle, graph_view, vertices, weights); + + ASSERT_TRUE(test_all_paths); +#endif + } + } +}; + +using Tests_UniformRandomWalks_File = + Tests_RandomWalks>; +using Tests_UniformRandomWalks_Rmat = + Tests_RandomWalks>; +using Tests_BiasedRandomWalks_File = + Tests_RandomWalks>; +using Tests_BiasedRandomWalks_Rmat = + Tests_RandomWalks>; +using Tests_Node2VecRandomWalks_File = + Tests_RandomWalks>; +using Tests_Node2VecRandomWalks_Rmat = + Tests_RandomWalks>; + +TEST_P(Tests_UniformRandomWalks_File, Initialize_i32_i32_f) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +INSTANTIATE_TEST_SUITE_P( + simple_test, + Tests_UniformRandomWalks_File, + ::testing::Combine( + ::testing::Values(UniformRandomWalks_Usecase{}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), + cugraph::test::File_Usecase("test/datasets/web-Google.mtx"), + cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), + cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + simple_test, + Tests_BiasedRandomWalks_File, + ::testing::Combine( + ::testing::Values(BiasedRandomWalks_Usecase{}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), + cugraph::test::File_Usecase("test/datasets/web-Google.mtx"), + cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), + cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + simple_test, + Tests_Node2VecRandomWalks_File, + ::testing::Combine( + ::testing::Values(Node2VecRandomWalks_Usecase{4, 8}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), + cugraph::test::File_Usecase("test/datasets/web-Google.mtx"), + cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), + cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); + +CUGRAPH_TEST_PROGRAM_MAIN()