diff --git a/cpp/include/cugraph/algorithms.hpp b/cpp/include/cugraph/algorithms.hpp index ed42460ed8e..7e5af4ac686 100644 --- a/cpp/include/cugraph/algorithms.hpp +++ b/cpp/include/cugraph/algorithms.hpp @@ -1579,11 +1579,11 @@ std:: template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed = std::numeric_limits::max()); + size_t max_length); /** * @brief returns biased random walks from starting sources, where each path is of given @@ -1623,11 +1623,11 @@ uniform_random_walks(raft::handle_t const& handle, template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed = std::numeric_limits::max()); + size_t max_length); /** * @brief returns biased random walks with node2vec biases from starting sources, @@ -1670,13 +1670,13 @@ biased_random_walks(raft::handle_t const& handle, template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, weight_t p, - weight_t q, - uint64_t seed = std::numeric_limits::max()); + weight_t q); #ifndef NO_CUGRAPH_OPS /** diff --git a/cpp/src/c_api/random_walks.cpp b/cpp/src/c_api/random_walks.cpp index b9a2c8e4f60..705d2108437 100644 --- a/cpp/src/c_api/random_walks.cpp +++ b/cpp/src/c_api/random_walks.cpp @@ -16,6 +16,7 @@ #include "c_api/abstract_functor.hpp" #include "c_api/graph.hpp" +#include "c_api/random.hpp" #include "c_api/resource_handle.hpp" #include "c_api/utils.hpp" @@ -153,10 +154,11 @@ namespace { struct uniform_random_walks_functor : public cugraph::c_api::abstract_functor { raft::handle_t const& handle_; + // FIXME: rng_state_ should be passed as a parameter + cugraph::c_api::cugraph_rng_state_t* rng_state_{nullptr}; 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, @@ -222,13 +224,17 @@ struct uniform_random_walks_functor : public cugraph::c_api::abstract_functor { graph_view.local_vertex_partition_range_last(), false); + // FIXME: remove once rng_state passed as parameter + rng_state_ = reinterpret_cast( + new cugraph::c_api::cugraph_rng_state_t{raft::random::RngState{0}}); + auto [paths, weights] = cugraph::uniform_random_walks( handle_, + rng_state_->rng_state_, graph_view, (edge_weights != nullptr) ? std::make_optional(edge_weights->view()) : std::nullopt, raft::device_span{start_vertices.data(), start_vertices.size()}, - max_length_, - seed_); + max_length_); // // Need to unrenumber the vertices in the resulting paths @@ -255,11 +261,12 @@ struct uniform_random_walks_functor : public cugraph::c_api::abstract_functor { struct biased_random_walks_functor : public cugraph::c_api::abstract_functor { raft::handle_t const& handle_; + // FIXME: rng_state_ should be passed as a parameter + cugraph::c_api::cugraph_rng_state_t* rng_state_{nullptr}; 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, @@ -326,13 +333,17 @@ struct biased_random_walks_functor : public cugraph::c_api::abstract_functor { graph_view.local_vertex_partition_range_last(), false); + // FIXME: remove once rng_state passed as parameter + rng_state_ = reinterpret_cast( + new cugraph::c_api::cugraph_rng_state_t{raft::random::RngState{0}}); + auto [paths, weights] = cugraph::biased_random_walks( handle_, + rng_state_->rng_state_, graph_view, edge_weights->view(), raft::device_span{start_vertices.data(), start_vertices.size()}, - max_length_, - seed_); + max_length_); // // Need to unrenumber the vertices in the resulting paths @@ -354,12 +365,13 @@ struct biased_random_walks_functor : public cugraph::c_api::abstract_functor { struct node2vec_random_walks_functor : public cugraph::c_api::abstract_functor { raft::handle_t const& handle_; + // FIXME: rng_state_ should be passed as a parameter + cugraph::c_api::cugraph_rng_state_t* rng_state_{nullptr}; 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, @@ -431,15 +443,19 @@ struct node2vec_random_walks_functor : public cugraph::c_api::abstract_functor { graph_view.local_vertex_partition_range_last(), false); + // FIXME: remove once rng_state passed as parameter + rng_state_ = reinterpret_cast( + new cugraph::c_api::cugraph_rng_state_t{raft::random::RngState{0}}); + auto [paths, weights] = cugraph::node2vec_random_walks( handle_, + rng_state_->rng_state_, graph_view, (edge_weights != nullptr) ? std::make_optional(edge_weights->view()) : std::nullopt, raft::device_span{start_vertices.data(), start_vertices.size()}, max_length_, static_cast(p_), - static_cast(q_), - seed_); + static_cast(q_)); // FIXME: Need to fix invalid_vtx issue here. We can't unrenumber max_vertex_id+1 // properly... diff --git a/cpp/src/prims/detail/sample_and_compute_local_nbr_indices.cuh b/cpp/src/prims/detail/sample_and_compute_local_nbr_indices.cuh index 7d4750c0554..a6a164d36c1 100644 --- a/cpp/src/prims/detail/sample_and_compute_local_nbr_indices.cuh +++ b/cpp/src/prims/detail/sample_and_compute_local_nbr_indices.cuh @@ -392,11 +392,11 @@ compute_unique_keys(raft::handle_t const& handle, cuda::proclaim_return_type( [unique_key_first = get_dataframe_buffer_begin(aggregate_local_frontier_unique_keys) + local_frontier_unique_key_displacements[i], - num_unique_keys = local_frontier_unique_key_sizes[i]] __device__(key_t key) { + unique_key_last = get_dataframe_buffer_begin(aggregate_local_frontier_unique_keys) + + local_frontier_unique_key_displacements[i] + + local_frontier_unique_key_sizes[i]] __device__(key_t key) { return static_cast(thrust::distance( - unique_key_first, - thrust::lower_bound( - thrust::seq, unique_key_first, unique_key_first + num_unique_keys, key))); + unique_key_first, thrust::find(thrust::seq, unique_key_first, unique_key_last, key))); })); } @@ -1759,8 +1759,7 @@ biased_sample_and_compute_local_nbr_indices( std::optional> key_indices{std::nullopt}; std::vector local_frontier_sample_offsets{}; if (with_replacement) { - // computet segmented inclusive sums (one segment per seed) - + // compute segmented inclusive sums (one segment per seed) auto unique_key_first = thrust::make_transform_iterator( thrust::make_counting_iterator(size_t{0}), cuda::proclaim_return_type( diff --git a/cpp/src/prims/per_v_random_select_transform_outgoing_e.cuh b/cpp/src/prims/per_v_random_select_transform_outgoing_e.cuh index 03514e52e6e..9d0f711d106 100644 --- a/cpp/src/prims/per_v_random_select_transform_outgoing_e.cuh +++ b/cpp/src/prims/per_v_random_select_transform_outgoing_e.cuh @@ -351,7 +351,7 @@ per_v_random_select_transform_e(raft::handle_t const& handle, uniform_sample_and_compute_local_nbr_indices( handle, graph_view, - (minor_comm_size > 1) ? get_dataframe_buffer_begin(*aggregate_local_frontier) + (minor_comm_size > 1) ? get_dataframe_buffer_cbegin(*aggregate_local_frontier) : frontier.begin(), local_frontier_displacements, local_frontier_sizes, @@ -363,7 +363,7 @@ per_v_random_select_transform_e(raft::handle_t const& handle, biased_sample_and_compute_local_nbr_indices( handle, graph_view, - (minor_comm_size > 1) ? get_dataframe_buffer_begin(*aggregate_local_frontier) + (minor_comm_size > 1) ? get_dataframe_buffer_cbegin(*aggregate_local_frontier) : frontier.begin(), edge_bias_src_value_input, edge_bias_dst_value_input, @@ -392,7 +392,7 @@ per_v_random_select_transform_e(raft::handle_t const& handle, graph_view.local_edge_partition_view(i)); auto edge_partition_frontier_key_first = - ((minor_comm_size > 1) ? get_dataframe_buffer_begin(*aggregate_local_frontier) + ((minor_comm_size > 1) ? get_dataframe_buffer_cbegin(*aggregate_local_frontier) : frontier.begin()) + local_frontier_displacements[i]; auto edge_partition_sample_local_nbr_index_first = diff --git a/cpp/src/sampling/random_walks_impl.cuh b/cpp/src/sampling/random_walks_impl.cuh index d582893d756..6c10fc473f3 100644 --- a/cpp/src/sampling/random_walks_impl.cuh +++ b/cpp/src/sampling/random_walks_impl.cuh @@ -17,7 +17,10 @@ #pragma once #include "detail/graph_partition_utils.cuh" +#include "prims/detail/nbr_intersection.cuh" #include "prims/per_v_random_select_transform_outgoing_e.cuh" +#include "prims/property_op_utils.cuh" +#include "prims/update_edge_src_dst_property.cuh" #include "prims/vertex_frontier.cuh" #include @@ -25,6 +28,7 @@ #include #include #include +#include #include #include #include @@ -46,13 +50,6 @@ namespace cugraph { namespace detail { -inline uint64_t get_current_time_nanoseconds() -{ - auto cur = std::chrono::steady_clock::now(); - return static_cast( - std::chrono::duration_cast(cur.time_since_epoch()).count()); -} - template struct sample_edges_op_t { template @@ -70,21 +67,129 @@ struct sample_edges_op_t { } }; +template +struct biased_random_walk_e_bias_op_t { + __device__ bias_t + operator()(vertex_t, vertex_t, bias_t src_out_weight_sum, thrust::nullopt_t, bias_t weight) const + { + return weight / src_out_weight_sum; + } +}; + +template +struct biased_sample_edges_op_t { + __device__ thrust::tuple operator()( + vertex_t, vertex_t dst, weight_t, thrust::nullopt_t, weight_t weight) const + { + return thrust::make_tuple(dst, weight); + } +}; + +template +struct node2vec_random_walk_e_bias_op_t { + bias_t p_{}; + bias_t q_{}; + raft::device_span intersection_offsets_{}; + raft::device_span intersection_indices_{}; + raft::device_span current_vertices_{}; + raft::device_span prev_vertices_{}; + + // Unweighted Bias Operator + template + __device__ std::enable_if_t, bias_t> operator()( + thrust::tuple tagged_src, + vertex_t dst, + thrust::nullopt_t, + thrust::nullopt_t, + thrust::nullopt_t) const + { + // Check tag (prev vert) for destination + if (dst == thrust::get<1>(tagged_src)) { return 1.0 / p_; } + // Search zipped vertices for tagged src + auto lower_itr = thrust::lower_bound( + thrust::seq, + thrust::make_zip_iterator(current_vertices_.begin(), prev_vertices_.begin()), + thrust::make_zip_iterator(current_vertices_.end(), prev_vertices_.end()), + tagged_src); + auto low_idx = thrust::distance( + thrust::make_zip_iterator(current_vertices_.begin(), prev_vertices_.begin()), lower_itr); + auto intersection_index_first = intersection_indices_.begin() + intersection_offsets_[low_idx]; + auto intersection_index_last = + intersection_indices_.begin() + intersection_offsets_[low_idx + 1]; + auto itr = + thrust::lower_bound(thrust::seq, intersection_index_first, intersection_index_last, dst); + return (itr != intersection_index_last && *itr == dst) ? 1.0 : 1.0 / q_; + } + + // Weighted Bias Operator + template + __device__ std::enable_if_t, bias_t> operator()( + thrust::tuple tagged_src, + vertex_t dst, + thrust::nullopt_t, + thrust::nullopt_t, + W) const + { + // Check tag (prev vert) for destination + if (dst == thrust::get<1>(tagged_src)) { return 1.0 / p_; } + // Search zipped vertices for tagged src + auto lower_itr = thrust::lower_bound( + thrust::seq, + thrust::make_zip_iterator(current_vertices_.begin(), prev_vertices_.begin()), + thrust::make_zip_iterator(current_vertices_.end(), prev_vertices_.end()), + tagged_src); + auto low_idx = thrust::distance( + thrust::make_zip_iterator(current_vertices_.begin(), prev_vertices_.begin()), lower_itr); + auto intersection_index_first = intersection_indices_.begin() + intersection_offsets_[low_idx]; + auto intersection_index_last = + intersection_indices_.begin() + intersection_offsets_[low_idx + 1]; + auto itr = + thrust::lower_bound(thrust::seq, intersection_index_first, intersection_index_last, dst); + return (itr != intersection_index_last && *itr == dst) ? 1.0 : 1.0 / q_; + } +}; + +template +struct node2vec_sample_edges_op_t { + template + __device__ std::enable_if_t, vertex_t> operator()( + thrust::tuple tagged_src, + vertex_t dst, + thrust::nullopt_t, + thrust::nullopt_t, + thrust::nullopt_t) const + { + return dst; + } + + template + __device__ std::enable_if_t, thrust::tuple> operator()( + thrust::tuple tagged_src, + vertex_t dst, + thrust::nullopt_t, + thrust::nullopt_t, + W w) const + { + return thrust::make_tuple(dst, w); + } +}; + template struct uniform_selector { - raft::random::RngState rng_state_; - - uniform_selector(uint64_t seed) : rng_state_(seed) {} + raft::random::RngState& rng_state_; + static constexpr bool is_second_order_ = false; template std::tuple, + std::optional>, std::optional>> follow_random_edge( raft::handle_t const& handle, GraphViewType const& graph_view, std::optional> edge_weight_view, - rmm::device_uvector const& current_vertices) + rmm::device_uvector&& current_vertices, + std::optional>&& previous_vertices) { using vertex_t = typename GraphViewType::vertex_type; @@ -133,30 +238,67 @@ struct uniform_selector { minors = std::move(sample_e_op_results); } - return std::make_tuple(std::move(minors), std::move(weights)); + return std::make_tuple(std::move(minors), std::move(previous_vertices), std::move(weights)); } }; template struct biased_selector { - uint64_t seed_{0}; + raft::random::RngState& rng_state_; + static constexpr bool is_second_order_ = false; template std::tuple, + std::optional>, std::optional>> follow_random_edge( raft::handle_t const& handle, GraphViewType const& graph_view, std::optional> edge_weight_view, - rmm::device_uvector const& current_vertices) + rmm::device_uvector&& current_vertices, + std::optional>&& previous_vertices) { - // To do biased sampling, I need out_weights instead of out_degrees. - // Then I generate a random float between [0, out_weights[v]). Then - // instead of making a decision based on the index I need to find - // upper_bound (or is it lower_bound) of the random number and - // the cumulative weight. - CUGRAPH_FAIL("biased sampling not implemented"); + // Create vertex frontier + using vertex_t = typename GraphViewType::vertex_type; + + using tag_t = void; + + cugraph::vertex_frontier_t vertex_frontier( + handle, 1); + + vertex_frontier.bucket(0).insert(current_vertices.begin(), current_vertices.end()); + + auto vertex_weight_sum = compute_out_weight_sums(handle, graph_view, *edge_weight_view); + edge_src_property_t edge_src_out_weight_sums(handle, graph_view); + update_edge_src_property(handle, + graph_view, + vertex_frontier.bucket(0).begin(), + vertex_frontier.bucket(0).end(), + vertex_weight_sum.data(), + edge_src_out_weight_sums.mutable_view()); + auto [sample_offsets, sample_e_op_results] = cugraph::per_v_random_select_transform_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_out_weight_sums.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + *edge_weight_view, + biased_random_walk_e_bias_op_t{}, + edge_src_out_weight_sums.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + *edge_weight_view, + biased_sample_edges_op_t{}, + rng_state_, + size_t{1}, + true, + std::make_optional( + thrust::make_tuple(vertex_t{cugraph::invalid_vertex_id::value}, weight_t{0.0}))); + + // Return results + return std::make_tuple(std::move(std::get<0>(sample_e_op_results)), + std::move(previous_vertices), + std::move(std::get<1>(sample_e_op_results))); } }; @@ -164,26 +306,232 @@ template struct node2vec_selector { weight_t p_; weight_t q_; - uint64_t seed_{0}; + raft::random::RngState& rng_state_; + static constexpr bool is_second_order_ = true; template std::tuple, + std::optional>, std::optional>> follow_random_edge( raft::handle_t const& handle, GraphViewType const& graph_view, std::optional> edge_weight_view, - rmm::device_uvector const& current_vertices) + rmm::device_uvector&& current_vertices, + std::optional>&& previous_vertices) { - // To do node2vec, I need the following: - // 1) transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v to compute the sum of the - // node2vec style weights - // 2) Generate a random number between [0, output_from_trdnioeebv[v]) - // 3) a sampling value that lets me pick the correct edge based on the same computation - // (essentially weighted sampling, but with a function that computes the weight rather - // than just using the edge weights) - CUGRAPH_FAIL("node2vec not implemented"); + // Create vertex frontier + using vertex_t = typename GraphViewType::vertex_type; + + using tag_t = vertex_t; + + // Zip previous and current vertices for nbr_intersection() + auto intersection_pairs = + thrust::make_zip_iterator(current_vertices.begin(), (*previous_vertices).begin()); + + auto [intersection_offsets, intersection_indices] = + detail::nbr_intersection(handle, + graph_view, + cugraph::edge_dummy_property_t{}.view(), + intersection_pairs, + intersection_pairs + current_vertices.size(), + std::array{true, true}, + false); + + rmm::device_uvector intersection_counts(size_t{0}, handle.get_stream()); + rmm::device_uvector aggregate_offsets(size_t{0}, handle.get_stream()); + rmm::device_uvector aggregate_currents(size_t{0}, handle.get_stream()); + rmm::device_uvector aggregate_previous(size_t{0}, handle.get_stream()); + rmm::device_uvector aggregate_indices(size_t{0}, handle.get_stream()); + + // Aggregate intersection data across minor comm + if constexpr (GraphViewType::is_multi_gpu) { + intersection_counts.resize(intersection_offsets.size(), handle.get_stream()); + thrust::adjacent_difference(handle.get_thrust_policy(), + intersection_offsets.begin(), + intersection_offsets.end(), + intersection_counts.begin()); + + auto recv_counts = cugraph::host_scalar_allgather( + handle.get_subcomm(cugraph::partition_manager::minor_comm_name()), + current_vertices.size(), + handle.get_stream()); + + std::vector displacements(recv_counts.size()); + std::exclusive_scan(recv_counts.begin(), recv_counts.end(), displacements.begin(), size_t{0}); + + aggregate_offsets.resize(displacements.back() + recv_counts.back() + 1, handle.get_stream()); + aggregate_offsets.set_element_to_zero_async(aggregate_offsets.size() - 1, + handle.get_stream()); + + cugraph::device_allgatherv(handle.get_subcomm(cugraph::partition_manager::minor_comm_name()), + intersection_counts.begin() + 1, + aggregate_offsets.begin(), + recv_counts, + displacements, + handle.get_stream()); + + thrust::exclusive_scan(handle.get_thrust_policy(), + aggregate_offsets.begin(), + aggregate_offsets.end(), + aggregate_offsets.begin()); + + aggregate_currents.resize(displacements.back() + recv_counts.back(), handle.get_stream()); + + cugraph::device_allgatherv(handle.get_subcomm(cugraph::partition_manager::minor_comm_name()), + current_vertices.begin(), + aggregate_currents.begin(), + recv_counts, + displacements, + handle.get_stream()); + + aggregate_previous.resize(displacements.back() + recv_counts.back(), handle.get_stream()); + + cugraph::device_allgatherv(handle.get_subcomm(cugraph::partition_manager::minor_comm_name()), + (*previous_vertices).begin(), + aggregate_previous.begin(), + recv_counts, + displacements, + handle.get_stream()); + + recv_counts = cugraph::host_scalar_allgather( + handle.get_subcomm(cugraph::partition_manager::minor_comm_name()), + intersection_offsets.back_element(handle.get_stream()), + handle.get_stream()); + + displacements.resize(recv_counts.size()); + std::exclusive_scan(recv_counts.begin(), recv_counts.end(), displacements.begin(), size_t{0}); + + aggregate_indices.resize(displacements.back() + recv_counts.back(), handle.get_stream()); + + cugraph::device_allgatherv(handle.get_subcomm(cugraph::partition_manager::minor_comm_name()), + intersection_indices.begin(), + aggregate_indices.begin(), + recv_counts, + displacements, + handle.get_stream()); + } + + cugraph::vertex_frontier_t vertex_frontier( + handle, 1); + vertex_frontier.bucket(0).insert( + thrust::make_zip_iterator(current_vertices.begin(), (*previous_vertices).begin()), + thrust::make_zip_iterator(current_vertices.end(), (*previous_vertices).end())); + + // Create data structs for results + rmm::device_uvector minors(0, handle.get_stream()); + std::optional> weights{std::nullopt}; + + if (edge_weight_view) { + auto [sample_offsets, sample_e_op_results] = + cugraph::per_v_random_select_transform_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + *edge_weight_view, + GraphViewType::is_multi_gpu + ? node2vec_random_walk_e_bias_op_t{p_, + q_, + raft::device_span( + aggregate_offsets.data(), + aggregate_offsets.size()), + raft::device_span( + aggregate_indices.data(), + aggregate_indices.size()), + raft::device_span( + aggregate_currents.data(), + aggregate_currents.size()), + raft::device_span( + aggregate_previous.data(), + aggregate_previous.size())} + : node2vec_random_walk_e_bias_op_t{p_, + q_, + raft::device_span( + intersection_offsets.data(), + intersection_offsets.size()), + raft::device_span( + intersection_indices.data(), + intersection_indices.size()), + raft::device_span< + vertex_t const>(current_vertices.data(), + current_vertices.size()), + raft::device_span( + (*previous_vertices).data(), + (*previous_vertices).size())}, + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + *edge_weight_view, + node2vec_sample_edges_op_t{}, + rng_state_, + size_t{1}, + true, + std::make_optional(thrust::make_tuple( + vertex_t{cugraph::invalid_vertex_id::value}, weight_t{0.0}))); + minors = std::move(std::get<0>(sample_e_op_results)); + weights = std::move(std::get<1>(sample_e_op_results)); + } else { + auto [sample_offsets, sample_e_op_results] = + cugraph::per_v_random_select_transform_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + cugraph::edge_dummy_property_t{}.view(), + GraphViewType::is_multi_gpu + ? node2vec_random_walk_e_bias_op_t{p_, + q_, + raft::device_span( + aggregate_offsets.data(), + aggregate_offsets.size()), + raft::device_span( + aggregate_indices.data(), + aggregate_indices.size()), + raft::device_span( + aggregate_currents.data(), + aggregate_currents.size()), + raft::device_span( + aggregate_previous.data(), + aggregate_previous.size())} + : node2vec_random_walk_e_bias_op_t{p_, + q_, + raft::device_span( + intersection_offsets.data(), + intersection_offsets.size()), + raft::device_span( + intersection_indices.data(), + intersection_indices.size()), + raft::device_span< + vertex_t const>(current_vertices.data(), + current_vertices.size()), + raft::device_span( + (*previous_vertices).data(), + (*previous_vertices).size())}, + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + cugraph::edge_dummy_property_t{}.view(), + node2vec_sample_edges_op_t{}, + rng_state_, + size_t{1}, + true, + std::make_optional(vertex_t{cugraph::invalid_vertex_id::value})); + minors = std::move(sample_e_op_results); + } + + *previous_vertices = std::move(current_vertices); + + return std::make_tuple(std::move(minors), std::move(previous_vertices), std::move(weights)); } }; @@ -221,6 +569,16 @@ random_walk_impl(raft::handle_t const& handle, ? std::make_optional>(0, handle.get_stream()) : std::nullopt; + auto previous_vertices = (random_selector.is_second_order_) + ? std::make_optional>( + current_vertices.size(), handle.get_stream()) + : std::nullopt; + if (previous_vertices) { + raft::copy((*previous_vertices).data(), + start_vertices.data(), + start_vertices.size(), + handle.get_stream()); + } raft::copy( current_vertices.data(), start_vertices.data(), start_vertices.size(), handle.get_stream()); detail::sequence_fill( @@ -255,25 +613,73 @@ random_walk_impl(raft::handle_t const& handle, auto& minor_comm = handle.get_subcomm(cugraph::partition_manager::minor_comm_name()); auto const minor_comm_size = minor_comm.get_size(); - // Shuffle vertices to correct GPU to compute random indices - std::forward_as_tuple(std::tie(current_vertices, current_gpu, current_position), - std::ignore) = - cugraph::groupby_gpu_id_and_shuffle_values( - handle.get_comms(), + if (previous_vertices) { + std::forward_as_tuple( + std::tie(current_vertices, current_gpu, current_position, previous_vertices), + std::ignore) = + cugraph::groupby_gpu_id_and_shuffle_values( + handle.get_comms(), + thrust::make_zip_iterator(current_vertices.begin(), + current_gpu.begin(), + current_position.begin(), + previous_vertices->begin()), + thrust::make_zip_iterator(current_vertices.end(), + current_gpu.end(), + current_position.end(), + previous_vertices->end()), + [key_func = + cugraph::detail::compute_gpu_id_from_int_vertex_t{ + {vertex_partition_range_lasts.begin(), vertex_partition_range_lasts.size()}, + major_comm_size, + minor_comm_size}] __device__(auto val) { return key_func(thrust::get<0>(val)); }, + handle.get_stream()); + } else { + // Shuffle vertices to correct GPU to compute random indices + std::forward_as_tuple(std::tie(current_vertices, current_gpu, current_position), + std::ignore) = + cugraph::groupby_gpu_id_and_shuffle_values( + handle.get_comms(), + thrust::make_zip_iterator( + current_vertices.begin(), current_gpu.begin(), current_position.begin()), + thrust::make_zip_iterator( + current_vertices.end(), current_gpu.end(), current_position.end()), + [key_func = + cugraph::detail::compute_gpu_id_from_int_vertex_t{ + {vertex_partition_range_lasts.begin(), vertex_partition_range_lasts.size()}, + major_comm_size, + minor_comm_size}] __device__(auto val) { return key_func(thrust::get<0>(val)); }, + handle.get_stream()); + } + } + + // Sort for nbr_intersection, must sort all together + if (previous_vertices) { + if constexpr (multi_gpu) { + thrust::sort(handle.get_thrust_policy(), + thrust::make_zip_iterator(current_vertices.begin(), + (*previous_vertices).begin(), + current_position.begin(), + current_gpu.begin()), + thrust::make_zip_iterator(current_vertices.end(), + (*previous_vertices).end(), + current_position.end(), + current_gpu.end())); + } else { + thrust::sort( + handle.get_thrust_policy(), thrust::make_zip_iterator( - current_vertices.begin(), current_gpu.begin(), current_position.begin()), + current_vertices.begin(), (*previous_vertices).begin(), current_position.begin()), thrust::make_zip_iterator( - current_vertices.end(), current_gpu.end(), current_position.end()), - [key_func = - cugraph::detail::compute_gpu_id_from_int_vertex_t{ - {vertex_partition_range_lasts.begin(), vertex_partition_range_lasts.size()}, - major_comm_size, - minor_comm_size}] __device__(auto val) { return key_func(thrust::get<0>(val)); }, - handle.get_stream()); + current_vertices.end(), (*previous_vertices).end(), current_position.end())); + } } - std::tie(current_vertices, new_weights) = - random_selector.follow_random_edge(handle, graph_view, edge_weight_view, current_vertices); + std::tie(current_vertices, previous_vertices, new_weights) = + random_selector.follow_random_edge(handle, + graph_view, + edge_weight_view, + std::move(current_vertices), + std::move(previous_vertices)); // FIXME: remove_if has a 32-bit overflow issue // (https://github.com/NVIDIA/thrust/issues/1302) Seems unlikely here (the goal of @@ -281,164 +687,244 @@ random_walk_impl(raft::handle_t const& handle, CUGRAPH_EXPECTS( current_vertices.size() < static_cast(std::numeric_limits::max()), "remove_if will fail, current_vertices.size() is too large"); - + size_t compacted_length{0}; if constexpr (multi_gpu) { if (result_weights) { - auto input_iter = thrust::make_zip_iterator(current_vertices.begin(), - new_weights->begin(), - current_gpu.begin(), - current_position.begin()); - - auto compacted_length = thrust::distance( - input_iter, - thrust::remove_if(handle.get_thrust_policy(), - input_iter, - input_iter + current_vertices.size(), - current_vertices.begin(), - [] __device__(auto dst) { - return (dst == cugraph::invalid_vertex_id::value); - })); - - current_vertices.resize(compacted_length, handle.get_stream()); - new_weights->resize(compacted_length, handle.get_stream()); - current_gpu.resize(compacted_length, handle.get_stream()); - current_position.resize(compacted_length, handle.get_stream()); - - // Shuffle back to original GPU - auto current_iter = thrust::make_zip_iterator(current_vertices.begin(), + if (previous_vertices) { + auto input_iter = thrust::make_zip_iterator(current_vertices.begin(), + new_weights->begin(), + current_gpu.begin(), + current_position.begin(), + previous_vertices->begin()); + + compacted_length = thrust::distance( + input_iter, + thrust::remove_if(handle.get_thrust_policy(), + input_iter, + input_iter + current_vertices.size(), + current_vertices.begin(), + [] __device__(auto dst) { + return (dst == cugraph::invalid_vertex_id::value); + })); + } else { + auto input_iter = thrust::make_zip_iterator(current_vertices.begin(), new_weights->begin(), current_gpu.begin(), current_position.begin()); - std::forward_as_tuple( - std::tie(current_vertices, *new_weights, current_gpu, current_position), std::ignore) = - cugraph::groupby_gpu_id_and_shuffle_values( - handle.get_comms(), - current_iter, - current_iter + current_vertices.size(), - [] __device__(auto val) { return thrust::get<2>(val); }, - handle.get_stream()); - - thrust::for_each( - handle.get_thrust_policy(), - thrust::make_zip_iterator( - current_vertices.begin(), new_weights->begin(), current_position.begin()), - thrust::make_zip_iterator( - current_vertices.end(), new_weights->end(), current_position.end()), - [result_verts = result_vertices.data(), - result_wgts = result_weights->data(), - level, - max_length] __device__(auto tuple) { - vertex_t v = thrust::get<0>(tuple); - weight_t w = thrust::get<1>(tuple); - size_t pos = thrust::get<2>(tuple); - result_verts[pos * (max_length + 1) + level + 1] = v; - result_wgts[pos * max_length + level] = w; - }); + compacted_length = thrust::distance( + input_iter, + thrust::remove_if(handle.get_thrust_policy(), + input_iter, + input_iter + current_vertices.size(), + current_vertices.begin(), + [] __device__(auto dst) { + return (dst == cugraph::invalid_vertex_id::value); + })); + } } else { - auto input_iter = thrust::make_zip_iterator( - current_vertices.begin(), current_gpu.begin(), current_position.begin()); - - auto compacted_length = thrust::distance( - input_iter, - thrust::remove_if(handle.get_thrust_policy(), - input_iter, - input_iter + current_vertices.size(), - current_vertices.begin(), - [] __device__(auto dst) { - return (dst == cugraph::invalid_vertex_id::value); - })); - - current_vertices.resize(compacted_length, handle.get_stream()); - current_gpu.resize(compacted_length, handle.get_stream()); - current_position.resize(compacted_length, handle.get_stream()); - - // Shuffle back to original GPU - auto current_iter = thrust::make_zip_iterator( - current_vertices.begin(), current_gpu.begin(), current_position.begin()); - - std::forward_as_tuple(std::tie(current_vertices, current_gpu, current_position), - std::ignore) = - cugraph::groupby_gpu_id_and_shuffle_values( - handle.get_comms(), - current_iter, - current_iter + current_vertices.size(), - [] __device__(auto val) { return thrust::get<1>(val); }, - handle.get_stream()); - - thrust::for_each( - handle.get_thrust_policy(), - thrust::make_zip_iterator(current_vertices.begin(), current_position.begin()), - thrust::make_zip_iterator(current_vertices.end(), current_position.end()), - [result_verts = result_vertices.data(), level, max_length] __device__(auto tuple) { - vertex_t v = thrust::get<0>(tuple); - size_t pos = thrust::get<1>(tuple); - result_verts[pos * (max_length + 1) + level + 1] = v; - }); + if (previous_vertices) { + auto input_iter = thrust::make_zip_iterator(current_vertices.begin(), + current_gpu.begin(), + current_position.begin(), + previous_vertices->begin()); + + compacted_length = thrust::distance( + input_iter, + thrust::remove_if(handle.get_thrust_policy(), + input_iter, + input_iter + current_vertices.size(), + current_vertices.begin(), + [] __device__(auto dst) { + return (dst == cugraph::invalid_vertex_id::value); + })); + } else { + auto input_iter = thrust::make_zip_iterator( + current_vertices.begin(), current_gpu.begin(), current_position.begin()); + + compacted_length = thrust::distance( + input_iter, + thrust::remove_if(handle.get_thrust_policy(), + input_iter, + input_iter + current_vertices.size(), + current_vertices.begin(), + [] __device__(auto dst) { + return (dst == cugraph::invalid_vertex_id::value); + })); + } } } else { if (result_weights) { - auto input_iter = thrust::make_zip_iterator( - current_vertices.begin(), new_weights->begin(), current_position.begin()); - - auto compacted_length = thrust::distance( - input_iter, - thrust::remove_if(handle.get_thrust_policy(), - input_iter, - input_iter + current_vertices.size(), - current_vertices.begin(), - [] __device__(auto dst) { - return (dst == cugraph::invalid_vertex_id::value); - })); - - current_vertices.resize(compacted_length, handle.get_stream()); - new_weights->resize(compacted_length, handle.get_stream()); - current_position.resize(compacted_length, handle.get_stream()); - - thrust::for_each( - handle.get_thrust_policy(), - thrust::make_zip_iterator( - current_vertices.begin(), new_weights->begin(), current_position.begin()), - thrust::make_zip_iterator( - current_vertices.end(), new_weights->end(), current_position.end()), - [result_verts = result_vertices.data(), - result_wgts = result_weights->data(), - level, - max_length] __device__(auto tuple) { - vertex_t v = thrust::get<0>(tuple); - weight_t w = thrust::get<1>(tuple); - size_t pos = thrust::get<2>(tuple); - result_verts[pos * (max_length + 1) + level + 1] = v; - result_wgts[pos * max_length + level] = w; - }); + if (previous_vertices) { + auto input_iter = thrust::make_zip_iterator(current_vertices.begin(), + new_weights->begin(), + current_position.begin(), + previous_vertices->begin()); + + compacted_length = thrust::distance( + input_iter, + thrust::remove_if(handle.get_thrust_policy(), + input_iter, + input_iter + current_vertices.size(), + current_vertices.begin(), + [] __device__(auto dst) { + return (dst == cugraph::invalid_vertex_id::value); + })); + } else { + auto input_iter = thrust::make_zip_iterator( + current_vertices.begin(), new_weights->begin(), current_position.begin()); + + compacted_length = thrust::distance( + input_iter, + thrust::remove_if(handle.get_thrust_policy(), + input_iter, + input_iter + current_vertices.size(), + current_vertices.begin(), + [] __device__(auto dst) { + return (dst == cugraph::invalid_vertex_id::value); + })); + } } else { - auto input_iter = - thrust::make_zip_iterator(current_vertices.begin(), current_position.begin()); - - auto compacted_length = thrust::distance( - input_iter, - thrust::remove_if(handle.get_thrust_policy(), - input_iter, - input_iter + current_vertices.size(), - current_vertices.begin(), - [] __device__(auto dst) { - return (dst == cugraph::invalid_vertex_id::value); - })); - - current_vertices.resize(compacted_length, handle.get_stream()); - current_position.resize(compacted_length, handle.get_stream()); - - thrust::for_each( - handle.get_thrust_policy(), - thrust::make_zip_iterator(current_vertices.begin(), current_position.begin()), - thrust::make_zip_iterator(current_vertices.end(), current_position.end()), - [result_verts = result_vertices.data(), level, max_length] __device__(auto tuple) { - vertex_t v = thrust::get<0>(tuple); - size_t pos = thrust::get<1>(tuple); - result_verts[pos * (max_length + 1) + level + 1] = v; - }); + if (previous_vertices) { + auto input_iter = thrust::make_zip_iterator( + current_vertices.begin(), current_position.begin(), previous_vertices->begin()); + + compacted_length = thrust::distance( + input_iter, + thrust::remove_if(handle.get_thrust_policy(), + input_iter, + input_iter + current_vertices.size(), + current_vertices.begin(), + [] __device__(auto dst) { + return (dst == cugraph::invalid_vertex_id::value); + })); + } else { + auto input_iter = + thrust::make_zip_iterator(current_vertices.begin(), current_position.begin()); + + compacted_length = thrust::distance( + input_iter, + thrust::remove_if(handle.get_thrust_policy(), + input_iter, + input_iter + current_vertices.size(), + current_vertices.begin(), + [] __device__(auto dst) { + return (dst == cugraph::invalid_vertex_id::value); + })); + } + } + } + + // Moved out of if statements to cut down on code duplication + current_vertices.resize(compacted_length, handle.get_stream()); + current_vertices.shrink_to_fit(handle.get_stream()); + current_position.resize(compacted_length, handle.get_stream()); + current_position.shrink_to_fit(handle.get_stream()); + if (result_weights) { + new_weights->resize(compacted_length, handle.get_stream()); + new_weights->shrink_to_fit(handle.get_stream()); + } + if (previous_vertices) { + previous_vertices->resize(compacted_length, handle.get_stream()); + previous_vertices->shrink_to_fit(handle.get_stream()); + } + if constexpr (multi_gpu) { + current_gpu.resize(compacted_length, handle.get_stream()); + current_gpu.shrink_to_fit(handle.get_stream()); + + // Shuffle back to original GPU + if (previous_vertices) { + if (result_weights) { + auto current_iter = thrust::make_zip_iterator(current_vertices.begin(), + new_weights->begin(), + current_gpu.begin(), + current_position.begin(), + previous_vertices->begin()); + + std::forward_as_tuple( + std::tie( + current_vertices, *new_weights, current_gpu, current_position, *previous_vertices), + std::ignore) = + cugraph::groupby_gpu_id_and_shuffle_values( + handle.get_comms(), + current_iter, + current_iter + current_vertices.size(), + [] __device__(auto val) { return thrust::get<2>(val); }, + handle.get_stream()); + } else { + auto current_iter = thrust::make_zip_iterator(current_vertices.begin(), + current_gpu.begin(), + current_position.begin(), + previous_vertices->begin()); + + std::forward_as_tuple( + std::tie(current_vertices, current_gpu, current_position, *previous_vertices), + std::ignore) = + cugraph::groupby_gpu_id_and_shuffle_values( + handle.get_comms(), + current_iter, + current_iter + current_vertices.size(), + [] __device__(auto val) { return thrust::get<1>(val); }, + handle.get_stream()); + } + } else { + if (result_weights) { + auto current_iter = thrust::make_zip_iterator(current_vertices.begin(), + new_weights->begin(), + current_gpu.begin(), + current_position.begin()); + + std::forward_as_tuple( + std::tie(current_vertices, *new_weights, current_gpu, current_position), std::ignore) = + cugraph::groupby_gpu_id_and_shuffle_values( + handle.get_comms(), + current_iter, + current_iter + current_vertices.size(), + [] __device__(auto val) { return thrust::get<2>(val); }, + handle.get_stream()); + } else { + auto current_iter = thrust::make_zip_iterator( + current_vertices.begin(), current_gpu.begin(), current_position.begin()); + + std::forward_as_tuple(std::tie(current_vertices, current_gpu, current_position), + std::ignore) = + cugraph::groupby_gpu_id_and_shuffle_values( + handle.get_comms(), + current_iter, + current_iter + current_vertices.size(), + [] __device__(auto val) { return thrust::get<1>(val); }, + handle.get_stream()); + } } } + + if (result_weights) { + thrust::for_each(handle.get_thrust_policy(), + thrust::make_zip_iterator( + current_vertices.begin(), new_weights->begin(), current_position.begin()), + thrust::make_zip_iterator( + current_vertices.end(), new_weights->end(), current_position.end()), + [result_verts = result_vertices.data(), + result_wgts = result_weights->data(), + level, + max_length] __device__(auto tuple) { + vertex_t v = thrust::get<0>(tuple); + weight_t w = thrust::get<1>(tuple); + size_t pos = thrust::get<2>(tuple); + result_verts[pos * (max_length + 1) + level + 1] = v; + result_wgts[pos * max_length + level] = w; + }); + } else { + thrust::for_each( + handle.get_thrust_policy(), + thrust::make_zip_iterator(current_vertices.begin(), current_position.begin()), + thrust::make_zip_iterator(current_vertices.end(), current_position.end()), + [result_verts = result_vertices.data(), level, max_length] __device__(auto tuple) { + vertex_t v = thrust::get<0>(tuple); + size_t pos = thrust::get<1>(tuple); + result_verts[pos * (max_length + 1) + level + 1] = v; + }); + } } return std::make_tuple(std::move(result_vertices), std::move(result_weights)); @@ -449,11 +935,11 @@ random_walk_impl(raft::handle_t const& handle, template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed) + size_t max_length) { CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); @@ -462,18 +948,17 @@ uniform_random_walks(raft::handle_t const& handle, edge_weight_view, start_vertices, max_length, - detail::uniform_selector( - (seed == 0 ? detail::get_current_time_nanoseconds() : seed))); + detail::uniform_selector{rng_state}); } template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed) + size_t max_length) { CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); @@ -483,30 +968,28 @@ biased_random_walks(raft::handle_t const& handle, std::optional>{edge_weight_view}, start_vertices, max_length, - detail::biased_selector{(seed == 0 ? detail::get_current_time_nanoseconds() : seed)}); + detail::biased_selector{rng_state}); } template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, weight_t p, - weight_t q, - uint64_t seed) + weight_t q) { CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); - return detail::random_walk_impl( - handle, - graph_view, - edge_weight_view, - start_vertices, - max_length, - detail::node2vec_selector{ - p, q, (seed == 0 ? detail::get_current_time_nanoseconds() : seed)}); + return detail::random_walk_impl(handle, + graph_view, + edge_weight_view, + start_vertices, + max_length, + detail::node2vec_selector{p, q, rng_state}); } } // namespace cugraph diff --git a/cpp/src/sampling/random_walks_mg_v32_e32.cu b/cpp/src/sampling/random_walks_mg_v32_e32.cu index 421d3e9c818..abe5386da1c 100644 --- a/cpp/src/sampling/random_walks_mg_v32_e32.cu +++ b/cpp/src/sampling/random_walks_mg_v32_e32.cu @@ -22,54 +22,54 @@ namespace cugraph { template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, float p, - float q, - uint64_t seed); + float q); template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, double p, - double q, - uint64_t seed); + double q); } // namespace cugraph diff --git a/cpp/src/sampling/random_walks_mg_v32_e64.cu b/cpp/src/sampling/random_walks_mg_v32_e64.cu index d38af65a505..b1bf1a19b77 100644 --- a/cpp/src/sampling/random_walks_mg_v32_e64.cu +++ b/cpp/src/sampling/random_walks_mg_v32_e64.cu @@ -22,54 +22,54 @@ namespace cugraph { template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, float p, - float q, - uint64_t seed); + float q); template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, double p, - double q, - uint64_t seed); + double q); } // namespace cugraph diff --git a/cpp/src/sampling/random_walks_mg_v64_e64.cu b/cpp/src/sampling/random_walks_mg_v64_e64.cu index 9dedc893242..13cc899e50d 100644 --- a/cpp/src/sampling/random_walks_mg_v64_e64.cu +++ b/cpp/src/sampling/random_walks_mg_v64_e64.cu @@ -22,54 +22,54 @@ namespace cugraph { template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, float p, - float q, - uint64_t seed); + float q); template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, double p, - double q, - uint64_t seed); + double q); } // namespace cugraph diff --git a/cpp/src/sampling/random_walks_sg_v32_e32.cu b/cpp/src/sampling/random_walks_sg_v32_e32.cu index 7b64d107250..383917c0248 100644 --- a/cpp/src/sampling/random_walks_sg_v32_e32.cu +++ b/cpp/src/sampling/random_walks_sg_v32_e32.cu @@ -22,54 +22,54 @@ namespace cugraph { template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, float p, - float q, - uint64_t seed); + float q); template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, double p, - double q, - uint64_t seed); + double q); } // namespace cugraph diff --git a/cpp/src/sampling/random_walks_sg_v32_e64.cu b/cpp/src/sampling/random_walks_sg_v32_e64.cu index d9ea09f36ef..98d2bb02d88 100644 --- a/cpp/src/sampling/random_walks_sg_v32_e64.cu +++ b/cpp/src/sampling/random_walks_sg_v32_e64.cu @@ -22,54 +22,54 @@ namespace cugraph { template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, float p, - float q, - uint64_t seed); + float q); template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, double p, - double q, - uint64_t seed); + double q); } // namespace cugraph diff --git a/cpp/src/sampling/random_walks_sg_v64_e64.cu b/cpp/src/sampling/random_walks_sg_v64_e64.cu index 0b9be107276..c139acec4b7 100644 --- a/cpp/src/sampling/random_walks_sg_v64_e64.cu +++ b/cpp/src/sampling/random_walks_sg_v64_e64.cu @@ -22,54 +22,54 @@ namespace cugraph { template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> uniform_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> biased_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, edge_property_view_t edge_weight_view, raft::device_span start_vertices, - size_t max_length, - uint64_t seed); + size_t max_length); template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, float p, - float q, - uint64_t seed); + float q); template std::tuple, std::optional>> node2vec_random_walks(raft::handle_t const& handle, + raft::random::RngState& rng_state, graph_view_t const& graph_view, std::optional> edge_weight_view, raft::device_span start_vertices, size_t max_length, double p, - double q, - uint64_t seed); + double q); } // namespace cugraph diff --git a/cpp/tests/c_api/sg_random_walks_test.c b/cpp/tests/c_api/sg_random_walks_test.c index 14108d91c04..a4a77b5775a 100644 --- a/cpp/tests/c_api/sg_random_walks_test.c +++ b/cpp/tests/c_api/sg_random_walks_test.c @@ -192,9 +192,6 @@ int generic_biased_random_walks_test(vertex_t* h_src, 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."); @@ -208,10 +205,10 @@ int generic_biased_random_walks_test(vertex_t* h_src, 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]; + weight_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); + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_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( @@ -231,23 +228,35 @@ int generic_biased_random_walks_test(vertex_t* h_src, 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, + cugraph_random_walk_result_get_max_path_length(result) == 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"); + TEST_ASSERT( + test_ret_value, h_start[i] == h_result_verts[i * (max_depth + 1)], "start of path not found"); + for (size_t j = 0; j < max_depth; ++j) { + int src_index = i * (max_depth + 1) + j; + int dst_index = src_index + 1; + if (h_result_verts[dst_index] < 0) { + if (h_result_verts[src_index] >= 0) { + int departing_count = 0; + for (int k = 0; k < num_vertices; ++k) { + if (M[h_result_verts[src_index]][k] >= 0) departing_count++; + } + TEST_ASSERT(test_ret_value, + departing_count == 0, + "biased_random_walks found no edge when an edge exists"); + } + } else { + TEST_ASSERT(test_ret_value, + M[h_result_verts[src_index]][h_result_verts[dst_index]] == + h_result_wgts[i * max_depth + j], + "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); @@ -302,10 +311,6 @@ int generic_node2vec_random_walks_test(vertex_t* h_src, 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."); @@ -319,10 +324,10 @@ int generic_node2vec_random_walks_test(vertex_t* h_src, 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]; + weight_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); + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_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( @@ -342,23 +347,35 @@ int generic_node2vec_random_walks_test(vertex_t* h_src, 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, + cugraph_random_walk_result_get_max_path_length(result) == 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 < max_depth; ++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"); + TEST_ASSERT( + test_ret_value, h_start[i] == h_result_verts[i * (max_depth + 1)], "start of path not found"); + for (size_t j = 0; j < max_depth; ++j) { + int src_index = i * (max_depth + 1) + j; + int dst_index = src_index + 1; + if (h_result_verts[dst_index] < 0) { + if (h_result_verts[src_index] >= 0) { + int departing_count = 0; + for (int k = 0; k < num_vertices; ++k) { + if (M[h_result_verts[src_index]][k] >= 0) departing_count++; + } + TEST_ASSERT(test_ret_value, + departing_count == 0, + "node2vec_random_walks found no edge when an edge exists"); + } + } else { + TEST_ASSERT(test_ret_value, + M[h_result_verts[src_index]][h_result_verts[dst_index]] == + h_result_wgts[i * max_depth + j], + "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); @@ -390,7 +407,7 @@ int test_biased_random_walks() 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}; + weight_t wgt[] = {1, 2, 3, 4, 5, 6, 7, 8}; vertex_t start[] = {2, 2}; return generic_biased_random_walks_test( diff --git a/cpp/tests/sampling/mg_random_walks_test.cpp b/cpp/tests/sampling/mg_random_walks_test.cpp index c2ad5c37e9e..e2415c08e60 100644 --- a/cpp/tests/sampling/mg_random_walks_test.cpp +++ b/cpp/tests/sampling/mg_random_walks_test.cpp @@ -44,8 +44,10 @@ struct UniformRandomWalks_Usecase { raft::device_span start_vertices, size_t max_depth) { + raft::random::RngState rng_state(static_cast(handle.get_comms().get_rank())); + return cugraph::uniform_random_walks( - handle, graph_view, edge_weight_view, start_vertices, max_depth, seed); + handle, rng_state, graph_view, edge_weight_view, start_vertices, max_depth); } bool expect_throw() { return false; } @@ -66,12 +68,13 @@ struct BiasedRandomWalks_Usecase { { CUGRAPH_EXPECTS(edge_weight_view.has_value(), "Biased random walk requires edge weights."); + raft::random::RngState rng_state(static_cast(handle.get_comms().get_rank())); + return cugraph::biased_random_walks( - handle, graph_view, *edge_weight_view, start_vertices, max_depth, seed); + handle, rng_state, graph_view, *edge_weight_view, start_vertices, max_depth); } - // FIXME: Not currently implemented - bool expect_throw() { return true; } + bool expect_throw() { return !test_weighted; } }; struct Node2VecRandomWalks_Usecase { @@ -89,18 +92,19 @@ struct Node2VecRandomWalks_Usecase { raft::device_span start_vertices, size_t max_depth) { + raft::random::RngState rng_state(static_cast(handle.get_comms().get_rank())); + return cugraph::node2vec_random_walks(handle, + rng_state, graph_view, edge_weight_view, start_vertices, max_depth, static_cast(p), - static_cast(q), - seed); + static_cast(q)); } - // FIXME: Not currently implemented - bool expect_throw() { return true; } + bool expect_throw() { return false; } }; template diff --git a/cpp/tests/sampling/random_walks_check.cuh b/cpp/tests/sampling/random_walks_check.cuh index 0fd73b5bba7..380b97a5b84 100644 --- a/cpp/tests/sampling/random_walks_check.cuh +++ b/cpp/tests/sampling/random_walks_check.cuh @@ -108,7 +108,7 @@ void random_walks_validate( (int)d, (float)w); } else { - printf("edge (%d,%d) NOT FOUND\n", (int)s, (int)d); + printf("edge (%d,%d), weight %g NOT FOUND\n", (int)s, (int)d, (float)w); } return 1; diff --git a/cpp/tests/sampling/sg_random_walks_test.cpp b/cpp/tests/sampling/sg_random_walks_test.cpp index 7409c2ab758..4bcfebc6d51 100644 --- a/cpp/tests/sampling/sg_random_walks_test.cpp +++ b/cpp/tests/sampling/sg_random_walks_test.cpp @@ -40,8 +40,10 @@ struct UniformRandomWalks_Usecase { raft::device_span start_vertices, size_t num_paths) { + raft::random::RngState rng_state(0); + return cugraph::uniform_random_walks( - handle, graph_view, edge_weight_view, start_vertices, num_paths, seed); + handle, rng_state, graph_view, edge_weight_view, start_vertices, num_paths); } bool expect_throw() { return false; } @@ -62,12 +64,13 @@ struct BiasedRandomWalks_Usecase { { CUGRAPH_EXPECTS(edge_weight_view.has_value(), "Biased random walk requires edge weights."); + raft::random::RngState rng_state(0); + return cugraph::biased_random_walks( - handle, graph_view, *edge_weight_view, start_vertices, num_paths, seed); + handle, rng_state, graph_view, *edge_weight_view, start_vertices, num_paths); } - // FIXME: Not currently implemented - bool expect_throw() { return true; } + bool expect_throw() { return !test_weighted; } }; struct Node2VecRandomWalks_Usecase { @@ -85,18 +88,19 @@ struct Node2VecRandomWalks_Usecase { raft::device_span start_vertices, size_t num_paths) { + raft::random::RngState rng_state(0); + return cugraph::node2vec_random_walks(handle, + rng_state, graph_view, edge_weight_view, start_vertices, num_paths, static_cast(p), - static_cast(q), - seed); + static_cast(q)); } - // FIXME: Not currently implemented - bool expect_throw() { return true; } + bool expect_throw() { return false; } }; template @@ -197,9 +201,6 @@ using Tests_Node2VecRandomWalks_File = using Tests_Node2VecRandomWalks_Rmat = Tests_RandomWalks>; -#if 0 -// FIXME: We should use these tests, gtest-1.11.0 makes it a runtime error -// to define and not instantiate these. TEST_P(Tests_UniformRandomWalks_File, Initialize_i32_i32_f) { run_current_test( @@ -211,7 +212,6 @@ TEST_P(Tests_UniformRandomWalks_Rmat, Initialize_i32_i32_f) run_current_test( override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); } -#endif TEST_P(Tests_BiasedRandomWalks_File, Initialize_i32_i32_f) { @@ -237,19 +237,12 @@ TEST_P(Tests_Node2VecRandomWalks_Rmat, Initialize_i32_i32_f) override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); } -#if 0 -// FIXME: Not sure why these are failing, but we're refactoring anyway. INSTANTIATE_TEST_SUITE_P( simple_test, Tests_UniformRandomWalks_File, - ::testing::Combine( - ::testing::Values(UniformRandomWalks_Usecase{false, 0, true}, - UniformRandomWalks_Usecase{true, 0, true}), - ::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")))); -#endif + ::testing::Combine(::testing::Values(UniformRandomWalks_Usecase{false, 0, true}, + UniformRandomWalks_Usecase{true, 0, true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); INSTANTIATE_TEST_SUITE_P( file_test, @@ -265,6 +258,16 @@ INSTANTIATE_TEST_SUITE_P( Node2VecRandomWalks_Usecase{4, 8, true, 0, true}), ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); +INSTANTIATE_TEST_SUITE_P( + file_large_test, + Tests_UniformRandomWalks_File, + ::testing::Combine( + ::testing::Values(UniformRandomWalks_Usecase{false, 0, true}, + UniformRandomWalks_Usecase{true, 0, true}), + ::testing::Values(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( file_large_test, Tests_BiasedRandomWalks_File, @@ -285,23 +288,20 @@ INSTANTIATE_TEST_SUITE_P( cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); -#if 0 -// FIXME: Not sure why these are failing, but we're refactoring anyway. INSTANTIATE_TEST_SUITE_P( rmat_small_test, Tests_UniformRandomWalks_Rmat, - ::testing::Combine(::testing::Values(UniformRandomWalks_Usecase{false, 0, true}, - UniformRandomWalks_Usecase{true, 0, true}), - ::testing::Values(cugraph::test::Rmat_Usecase( - 10, 16, 0.57, 0.19, 0.19, 0, false, false)))); + ::testing::Combine( + ::testing::Values(UniformRandomWalks_Usecase{false, 0, true}, + UniformRandomWalks_Usecase{true, 0, true}), + ::testing::Values(cugraph::test::Rmat_Usecase(10, 16, 0.57, 0.19, 0.19, 0, false, false)))); INSTANTIATE_TEST_SUITE_P( rmat_benchmark_test, Tests_UniformRandomWalks_Rmat, - ::testing::Combine(::testing::Values(UniformRandomWalks_Usecase{true, 0, false}), - ::testing::Values(cugraph::test::Rmat_Usecase( - 20, 32, 0.57, 0.19, 0.19, 0, false, false)))); -#endif + ::testing::Combine( + ::testing::Values(UniformRandomWalks_Usecase{true, 0, false}), + ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); INSTANTIATE_TEST_SUITE_P( rmat_small_test,