From f0dbb4cc0b59c9c8f7a4be01439b1c57821a948b Mon Sep 17 00:00:00 2001 From: Charles Hastings Date: Wed, 8 Feb 2023 08:53:22 -0800 Subject: [PATCH 1/2] remove legacy sampling implementation, no longer used --- cpp/include/cugraph/algorithms.hpp | 48 ---- cpp/include/cugraph_c/sampling_algorithms.h | 39 ---- cpp/src/c_api/uniform_neighbor_sampling.cpp | 155 +------------ cpp/src/sampling/detail/graph_functions.hpp | 79 ------- .../sampling/detail/sampling_utils_impl.cuh | 211 ------------------ cpp/src/sampling/detail/sampling_utils_mg.cu | 123 ---------- cpp/src/sampling/detail/sampling_utils_sg.cu | 177 --------------- .../uniform_neighbor_sampling_impl.hpp | 121 +--------- .../sampling/uniform_neighbor_sampling_mg.cpp | 72 ------ .../sampling/uniform_neighbor_sampling_sg.cpp | 72 ------ cpp/tests/c_api/create_graph_test.c | 77 ++++--- .../c_api/uniform_neighbor_sample_test.c | 165 -------------- .../pylibcugraph/_cugraph_c/algorithms.pxd | 13 -- .../pylibcugraph/uniform_neighbor_sample.pyx | 1 - 14 files changed, 43 insertions(+), 1310 deletions(-) diff --git a/cpp/include/cugraph/algorithms.hpp b/cpp/include/cugraph/algorithms.hpp index 4bb39922fd3..daa892950e9 100644 --- a/cpp/include/cugraph/algorithms.hpp +++ b/cpp/include/cugraph/algorithms.hpp @@ -1735,54 +1735,6 @@ k_core(raft::handle_t const& handle, std::optional> core_numbers, bool do_expensive_check = false); -/** - * @brief Uniform Neighborhood Sampling. - * - * @deprecated This function should be replaced with uniform_neighbor_sample. Input of the - * new function adds an optional parameter, output has a number of extra fields. - * - * This function traverses from a set of starting vertices, traversing outgoing edges and - * randomly selects from these outgoing neighbors to extract a subgraph. - * - * Output from this function a set of tuples (src, dst, weight, count), identifying the randomly - * selected edges. src is the source vertex, dst is the destination vertex, weight is the weight - * of the edge and count identifies the number of times this edge was encountered during the - * sampling of this graph (so it is >= 1). - * - * @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 object to generate NBR Sampling on. - * @param edge_weight_view Optional view object holding edge weights for @p graph_view. - * @param starting_vertices Device span of starting vertex IDs for the NBR Sampling. - * @param fan_out Host span defining branching out (fan-out) degree per source vertex for each - * level - * @param with_replacement boolean flag specifying if random sampling is done with replacement - * (true); or, without replacement (false); default = true; - * @param seed A seed to initialize the random number generator - * @return tuple device vectors (vertex_t source_vertex, vertex_t destination_vertex, weight_t - * weight, edge_t count) - */ -template -std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement = true, - uint64_t seed = 0); - /** * @brief Uniform Neighborhood Sampling. * diff --git a/cpp/include/cugraph_c/sampling_algorithms.h b/cpp/include/cugraph_c/sampling_algorithms.h index 7a9bd93b079..edb5f0a8445 100644 --- a/cpp/include/cugraph_c/sampling_algorithms.h +++ b/cpp/include/cugraph_c/sampling_algorithms.h @@ -187,35 +187,6 @@ typedef struct { int32_t align_; } cugraph_sample_result_t; -/** - * @brief Uniform Neighborhood Sampling - * @deprecated This call should be replaced with cugraph_uniform_neighborhood_sampling - * - * @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 Device array of start vertices for the sampling - * @param [in] fanout Host array defining the fan out at each step in the sampling algorithm - * @param [in] with_replacement - * Boolean value. If true selection of edges is done with - * replacement. If false selection is done without replacement. - * @param [in] do_expensive_check - * A flag to run expensive checks for input arguments (if set to true) - * @param [in] result Output from the uniform_neighbor_sample 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_neighbor_sample( - const cugraph_resource_handle_t* handle, - cugraph_graph_t* graph, - const cugraph_type_erased_device_array_view_t* start, - const cugraph_type_erased_host_array_view_t* fan_out, - bool_t with_replacement, - bool_t do_expensive_check, - cugraph_sample_result_t** result, - cugraph_error_t** error); - /** * @brief Uniform Neighborhood Sampling * @@ -331,16 +302,6 @@ cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_hop( cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_index( const cugraph_sample_result_t* result); -/** - * @brief Get the transaction counts from the sampling algorithm result - * - * @param [in] result The result from a sampling algorithm - * @return type erased host array pointing to the counts - */ -// FIXME: This will be obsolete when the older mechanism is removed -cugraph_type_erased_host_array_view_t* cugraph_sample_result_get_counts( - const cugraph_sample_result_t* result); - /** * @brief Free a sampling result * diff --git a/cpp/src/c_api/uniform_neighbor_sampling.cpp b/cpp/src/c_api/uniform_neighbor_sampling.cpp index 0383823317c..1160f2326be 100644 --- a/cpp/src/c_api/uniform_neighbor_sampling.cpp +++ b/cpp/src/c_api/uniform_neighbor_sampling.cpp @@ -40,8 +40,6 @@ struct cugraph_sample_result_t { cugraph_type_erased_device_array_t* wgt_{nullptr}; cugraph_type_erased_device_array_t* hop_{nullptr}; cugraph_type_erased_device_array_t* label_{nullptr}; - // FIXME: Will be deleted once experimental replaces current - cugraph_type_erased_host_array_t* count_{nullptr}; }; } // namespace c_api @@ -49,117 +47,6 @@ struct cugraph_sample_result_t { namespace { -struct uniform_neighbor_sampling_functor_deprecate : 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_{nullptr}; - cugraph::c_api::cugraph_type_erased_host_array_view_t const* fan_out_{nullptr}; - bool with_replacement_{false}; - bool do_expensive_check_{false}; - cugraph::c_api::cugraph_sample_result_t* result_{nullptr}; - - uniform_neighbor_sampling_functor_deprecate(cugraph_resource_handle_t const* handle, - cugraph_graph_t* graph, - cugraph_type_erased_device_array_view_t const* start, - cugraph_type_erased_host_array_view_t const* fan_out, - bool with_replacement, - bool do_expensive_check) - : abstract_functor(), - handle_(*reinterpret_cast(handle)->handle_), - graph_(reinterpret_cast(graph)), - start_( - reinterpret_cast(start)), - fan_out_( - reinterpret_cast(fan_out)), - with_replacement_(with_replacement), - do_expensive_check_(do_expensive_check) - { - } - - template - void operator()() - { - // FIXME: Think about how to handle SG vice MG - if constexpr (!cugraph::is_candidate::value) { - unsupported(); - } else { - // uniform_nbr_sample 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 edge_weights = reinterpret_cast< - cugraph::edge_property_t, - weight_t>*>(graph_->edge_weights_); - - auto number_map = reinterpret_cast*>(graph_->number_map_); - - rmm::device_uvector start(start_->size_, handle_.get_stream()); - raft::copy(start.data(), start_->as_type(), start.size(), handle_.get_stream()); - - // - // Need to renumber sources - // - cugraph::renumber_ext_vertices( - handle_, - start.data(), - start.size(), - number_map->data(), - graph_view.local_vertex_partition_range_first(), - graph_view.local_vertex_partition_range_last(), - false); - - auto&& [srcs, dsts, weights, counts] = cugraph::uniform_nbr_sample( - handle_, - graph_view, - (edge_weights != nullptr) ? std::make_optional(edge_weights->view()) : std::nullopt, - raft::device_span(start.data(), start.size()), - raft::host_span(fan_out_->as_type(), fan_out_->size_), - with_replacement_); - - std::vector vertex_partition_lasts = graph_view.vertex_partition_range_lasts(); - - cugraph::unrenumber_int_vertices(handle_, - srcs.data(), - srcs.size(), - number_map->data(), - vertex_partition_lasts, - do_expensive_check_); - - cugraph::unrenumber_int_vertices(handle_, - dsts.data(), - dsts.size(), - number_map->data(), - vertex_partition_lasts, - do_expensive_check_); - - result_ = new cugraph::c_api::cugraph_sample_result_t{ - new cugraph::c_api::cugraph_type_erased_device_array_t(srcs, graph_->vertex_type_), - new cugraph::c_api::cugraph_type_erased_device_array_t(dsts, graph_->vertex_type_), - new cugraph::c_api::cugraph_type_erased_device_array_t( - weights, graph_->weight_type_), // needs to be edge id... - nullptr, - nullptr, - nullptr, - nullptr, - nullptr}; - } - } -}; - struct uniform_neighbor_sampling_functor : public cugraph::c_api::abstract_functor { raft::handle_t const& handle_; cugraph::c_api::cugraph_graph_t* graph_{nullptr}; @@ -302,36 +189,6 @@ struct uniform_neighbor_sampling_functor : public cugraph::c_api::abstract_funct } // namespace -extern "C" cugraph_error_code_t cugraph_uniform_neighbor_sample( - const cugraph_resource_handle_t* handle, - cugraph_graph_t* graph, - const cugraph_type_erased_device_array_view_t* start, - const cugraph_type_erased_host_array_view_t* fan_out, - bool_t with_replacement, - bool_t do_expensive_check, - cugraph_sample_result_t** result, - cugraph_error_t** error) -{ - CAPI_EXPECTS( - reinterpret_cast(graph)->vertex_type_ == - reinterpret_cast(start) - ->type_, - CUGRAPH_INVALID_INPUT, - "vertex type of graph and start must match", - *error); - - CAPI_EXPECTS( - reinterpret_cast(fan_out) - ->type_ == INT32, - CUGRAPH_INVALID_INPUT, - "fan_out should be of type int", - *error); - - uniform_neighbor_sampling_functor_deprecate functor{ - handle, graph, start, fan_out, with_replacement, do_expensive_check}; - return cugraph::c_api::run_algorithm(graph, functor, result, error); -} - extern "C" cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_sources( const cugraph_sample_result_t* result) { @@ -404,13 +261,6 @@ extern "C" cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_in internal_pointer->edge_id_->view()); } -extern "C" cugraph_type_erased_host_array_view_t* cugraph_sample_result_get_counts( - const cugraph_sample_result_t* result) -{ - auto internal_pointer = reinterpret_cast(result); - return reinterpret_cast(internal_pointer->count_->view()); -} - extern "C" cugraph_error_code_t cugraph_test_uniform_neighborhood_sample_result_create( const cugraph_resource_handle_t* handle, const cugraph_type_erased_device_array_view_t* srcs, @@ -639,8 +489,8 @@ extern "C" cugraph_error_code_t cugraph_test_sample_result_create( reinterpret_cast(new_device_wgt.release()), reinterpret_cast( new_device_label.release()), - reinterpret_cast(new_device_hop.release()), - nullptr}); + reinterpret_cast( + new_device_hop.release())}); return CUGRAPH_SUCCESS; } @@ -655,7 +505,6 @@ extern "C" void cugraph_sample_result_free(cugraph_sample_result_t* result) delete internal_pointer->wgt_; delete internal_pointer->hop_; delete internal_pointer->label_; - delete internal_pointer->count_; delete internal_pointer; } diff --git a/cpp/src/sampling/detail/graph_functions.hpp b/cpp/src/sampling/detail/graph_functions.hpp index 59e97b1bfde..5dae06a5e89 100644 --- a/cpp/src/sampling/detail/graph_functions.hpp +++ b/cpp/src/sampling/detail/graph_functions.hpp @@ -25,58 +25,6 @@ namespace detail { // in implementation, naming and documentation. We should review these and // consider updating things to support an arbitrary value for store_transposed -/** - * @brief Gather active majors across gpus in a column communicator - * - * Collect all the vertex ids and client gpu ids to be processed by every gpu in - * the column communicator and sort the list. - * - * @tparam vertex_t Type of vertex indices. - * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and - * handles to various CUDA libraries) to run graph algorithms. - * @param d_in Device vector containing vertices local to this GPU - * @return Device vector containing all the vertices that are to be processed by every gpu - * in the column communicator - */ -template -rmm::device_uvector allgather_active_majors(raft::handle_t const& handle, - rmm::device_uvector&& d_in); - -// FIXME: Need docs if this function survives -template -std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -count_and_remove_duplicates(raft::handle_t const& handle, - rmm::device_uvector&& src, - rmm::device_uvector&& dst, - rmm::device_uvector&& wgt); - -/** - * @brief Gather edge list for specified vertices - * - * Collect all the edges that are present in the adjacency lists on the current gpu - * - * @tparam GraphViewType Type of the passed non-owning graph object. - * @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 Non-owning graph object. - * @param active_majors Device vector containing all the vertex id that are processed by - * gpus in the column communicator - * @return A tuple of device vector containing the majors, minors and weights gathered locally - */ -template -std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - const rmm::device_uvector& active_majors, - bool do_expensive_check = false); - /** * @brief Gather edge list for specified vertices * @@ -114,33 +62,6 @@ gather_one_hop_edgelist( std::optional> const& active_major_labels, bool do_expensive_check = false); -/** - * @brief Randomly sample edges from the adjacency list of specified vertices - * - * @tparam GraphViewType Type of the passed non-owning graph object. - * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and - * handles to various CUDA libraries) to run graph algorithms. - * @param rng_state Random number generator state - * @param graph_view Non-owning graph object. - * @param active_majors Device vector containing all the vertex id that are processed by - * gpus in the column communicator - * @param fanout How many edges to sample for each vertex - * @param with_replacement If true sample with replacement, otherwise sample without replacement - * @param invalid_vertex_id Value to use for an invalid vertex - * @return A tuple of device vector containing the majors, minors and weights gathered locally - */ -template -std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - /** * @brief Randomly sample edges from the adjacency list of specified vertices * diff --git a/cpp/src/sampling/detail/sampling_utils_impl.cuh b/cpp/src/sampling/detail/sampling_utils_impl.cuh index 31c4513ff77..37f7281bbe2 100644 --- a/cpp/src/sampling/detail/sampling_utils_impl.cuh +++ b/cpp/src/sampling/detail/sampling_utils_impl.cuh @@ -37,49 +37,6 @@ namespace cugraph { namespace detail { -// FIXME: Need to move this to a publicly available function -// and refactor now true edge ids are available -template -std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -count_and_remove_duplicates(raft::handle_t const& handle, - rmm::device_uvector&& src, - rmm::device_uvector&& dst, - rmm::device_uvector&& wgt) -{ - auto tuple_iter_begin = - thrust::make_zip_iterator(thrust::make_tuple(src.begin(), dst.begin(), wgt.begin())); - - thrust::sort(handle.get_thrust_policy(), tuple_iter_begin, tuple_iter_begin + src.size()); - - auto num_uniques = - thrust::count_if(handle.get_thrust_policy(), - thrust::make_counting_iterator(size_t{0}), - thrust::make_counting_iterator(src.size()), - detail::is_first_in_run_t{tuple_iter_begin}); - - rmm::device_uvector result_src(num_uniques, handle.get_stream()); - rmm::device_uvector result_dst(num_uniques, handle.get_stream()); - rmm::device_uvector result_wgt(num_uniques, handle.get_stream()); - rmm::device_uvector result_count(num_uniques, handle.get_stream()); - - rmm::device_uvector count(src.size(), handle.get_stream()); - thrust::fill(handle.get_thrust_policy(), count.begin(), count.end(), edge_t{1}); - - thrust::reduce_by_key(handle.get_thrust_policy(), - tuple_iter_begin, - tuple_iter_begin + src.size(), - count.begin(), - thrust::make_zip_iterator(thrust::make_tuple( - result_src.begin(), result_dst.begin(), result_wgt.begin())), - result_count.begin()); - - return std::make_tuple( - std::move(result_src), std::move(result_dst), std::move(result_wgt), std::move(result_count)); -} - struct return_edges_with_properties_e_op { template auto __host__ __device__ operator()(key_t optionally_tagged_src, @@ -163,174 +120,6 @@ struct sample_edges_op_t { } }; -// FIXME: Old Implementation, phase out -template -std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - const rmm::device_uvector& active_majors, - bool do_expensive_check) -{ - // FIXME: add as a template parameter - using tag_t = void; - - cugraph::vertex_frontier_t vertex_frontier(handle, 1); - - vertex_frontier.bucket(0).insert(active_majors.begin(), active_majors.end()); - - rmm::device_uvector majors(0, handle.get_stream()); - rmm::device_uvector minors(0, handle.get_stream()); - std::optional> weights{std::nullopt}; - if (edge_weight_view) { - std::tie(majors, minors, weights) = - cugraph::extract_transform_v_frontier_outgoing_e(handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - *edge_weight_view, - return_edges_with_properties_e_op{}, - do_expensive_check); - } else { - std::tie(majors, minors) = - cugraph::extract_transform_v_frontier_outgoing_e(handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - edge_dummy_property_t{}.view(), - return_edges_with_properties_e_op{}, - do_expensive_check); - } - return std::make_tuple(std::move(majors), std::move(minors), std::move(weights)); -} - -// FIXME: Old Implementation, phase out -template -std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement) -{ - // FIXME: add as a template parameter - using tag_t = void; - - cugraph::vertex_frontier_t vertex_frontier(handle, 1); - - vertex_frontier.bucket(0).insert(active_majors.begin(), active_majors.end()); - - rmm::device_uvector majors(0, handle.get_stream()); - 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), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - *edge_weight_view, - sample_edges_op_t{}, - rng_state, - fanout, - with_replacement, - std::make_optional>( - cugraph::invalid_vertex_id::value, - cugraph::invalid_vertex_id::value, - weight_t{}), - true); - majors = std::move(std::get<0>(sample_e_op_results)); - minors = std::move(std::get<1>(sample_e_op_results)); - weights = std::move(std::get<2>(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), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - edge_dummy_property_t{}.view(), - sample_edges_op_t{}, - rng_state, - fanout, - with_replacement, - std::make_optional>( - cugraph::invalid_vertex_id::value, cugraph::invalid_vertex_id::value), - true); - majors = std::move(std::get<0>(sample_e_op_results)); - minors = std::move(std::get<1>(sample_e_op_results)); - } - - // - // FIXME: Debugging status, EOD 9/18/22 - // 2) Need to consider the case of a directed graph where a vertex is a sink but is selected - // as a seed for sampling. Output degree is 0, so there can be no departing vertices. ALso - // consider case where output degree is 1 and want 2 edges without replacement. - // 3) Finally... can I switch to using cugraph::invalid_vertex_id instead of - // number_of_vertices()? 4) I'm close, I should do the code cleanup. - // - if (weights) { - auto edge_first = thrust::make_zip_iterator( - thrust::make_tuple(majors.begin(), minors.begin(), (*weights).begin())); - auto end_iter = - thrust::remove_if(handle.get_thrust_policy(), - edge_first, - edge_first + majors.size(), - [] __device__(auto tuple) { - auto v1 = thrust::get<0>(tuple); - auto v2 = thrust::get<1>(tuple); - - return ((v1 == cugraph::invalid_vertex_id::value) || - (v1 == cugraph::invalid_vertex_id::value)); - }); - - size_t new_size = thrust::distance(edge_first, end_iter); - - if (new_size != majors.size()) { - majors.resize(new_size, handle.get_stream()); - majors.shrink_to_fit(handle.get_stream()); - minors.resize(new_size, handle.get_stream()); - minors.shrink_to_fit(handle.get_stream()); - (*weights).resize(new_size, handle.get_stream()); - (*weights).shrink_to_fit(handle.get_stream()); - } - } else { - auto edge_first = thrust::make_zip_iterator(thrust::make_tuple(majors.begin(), minors.begin())); - auto end_iter = - thrust::remove_if(handle.get_thrust_policy(), - edge_first, - edge_first + majors.size(), - [] __device__(auto tuple) { - auto v1 = thrust::get<0>(tuple); - auto v2 = thrust::get<1>(tuple); - - return ((v1 == cugraph::invalid_vertex_id::value) || - (v2 == cugraph::invalid_vertex_id::value)); - }); - - size_t new_size = thrust::distance(edge_first, end_iter); - - if (new_size != majors.size()) { - majors.resize(new_size, handle.get_stream()); - majors.shrink_to_fit(handle.get_stream()); - minors.resize(new_size, handle.get_stream()); - minors.shrink_to_fit(handle.get_stream()); - } - } - - return std::make_tuple(std::move(majors), std::move(minors), std::move(weights)); -} - template , - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector const& active_majors, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector const& active_majors, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector const& active_majors, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector const& active_majors, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector const& active_majors, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector const& active_majors, - bool do_expensive_check); - template std::tuple, rmm::device_uvector, std::optional>, @@ -184,72 +127,6 @@ gather_one_hop_edgelist( std::optional> const& active_major_labels, bool do_expensive_check); -template std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - template std::tuple, rmm::device_uvector, std::optional>, diff --git a/cpp/src/sampling/detail/sampling_utils_sg.cu b/cpp/src/sampling/detail/sampling_utils_sg.cu index 5ce9f509030..2dcfe5e070b 100644 --- a/cpp/src/sampling/detail/sampling_utils_sg.cu +++ b/cpp/src/sampling/detail/sampling_utils_sg.cu @@ -19,117 +19,6 @@ namespace cugraph { namespace detail { -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -count_and_remove_duplicates(raft::handle_t const& handle, - rmm::device_uvector&& src, - rmm::device_uvector&& dst, - rmm::device_uvector&& wgt); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -count_and_remove_duplicates(raft::handle_t const& handle, - rmm::device_uvector&& src, - rmm::device_uvector&& dst, - rmm::device_uvector&& wgt); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -count_and_remove_duplicates(raft::handle_t const& handle, - rmm::device_uvector&& src, - rmm::device_uvector&& dst, - rmm::device_uvector&& wgt); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -count_and_remove_duplicates(raft::handle_t const& handle, - rmm::device_uvector&& src, - rmm::device_uvector&& dst, - rmm::device_uvector&& wgt); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -count_and_remove_duplicates(raft::handle_t const& handle, - rmm::device_uvector&& src, - rmm::device_uvector&& dst, - rmm::device_uvector&& wgt); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -count_and_remove_duplicates(raft::handle_t const& handle, - rmm::device_uvector&& src, - rmm::device_uvector&& dst, - rmm::device_uvector&& wgt); - -template std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector const& active_majors, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector const& active_majors, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector const& active_majors, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector const& active_majors, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector const& active_majors, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector const& active_majors, - bool do_expensive_check); - template std::tuple, rmm::device_uvector, std::optional>, @@ -238,72 +127,6 @@ gather_one_hop_edgelist( std::optional> const& active_major_labels, bool do_expensive_check); -template std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::random::RngState& rng_state, - rmm::device_uvector const& active_majors, - size_t fanout, - bool with_replacement); - template std::tuple, rmm::device_uvector, std::optional>, diff --git a/cpp/src/sampling/uniform_neighbor_sampling_impl.hpp b/cpp/src/sampling/uniform_neighbor_sampling_impl.hpp index 096898257d8..7a1aca736b5 100644 --- a/cpp/src/sampling/uniform_neighbor_sampling_impl.hpp +++ b/cpp/src/sampling/uniform_neighbor_sampling_impl.hpp @@ -29,100 +29,6 @@ namespace cugraph { namespace detail { -template -std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample_impl( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - rmm::device_uvector& d_in, - raft::host_span h_fan_out, - bool with_replacement, - uint64_t seed) -{ -#ifdef NO_CUGRAPH_OPS - CUGRAPH_FAIL( - "uniform_nbr_sampl_impl not supported in this configuration, built with NO_CUGRAPH_OPS"); -#else - - CUGRAPH_EXPECTS(h_fan_out.size() > 0, - "Invalid input argument: number of levels must be non-zero."); - - rmm::device_uvector d_result_src(0, handle.get_stream()); - rmm::device_uvector d_result_dst(0, handle.get_stream()); - auto d_result_weights = - thrust::make_optional(rmm::device_uvector(0, handle.get_stream())); - - int32_t hop{0}; - size_t comm_size{1}; - - if constexpr (multi_gpu) { - seed += handle.get_comms().get_rank(); - comm_size = handle.get_comms().get_size(); - } - - for (auto&& k_level : h_fan_out) { - // prep step for extracting out-degs(sources): - if constexpr (multi_gpu) { - d_in = shuffle_int_vertices_to_local_gpu_by_vertex_partitioning( - handle, std::move(d_in), graph_view.vertex_partition_range_lasts()); - } - - rmm::device_uvector d_out_src(0, handle.get_stream()); - rmm::device_uvector d_out_dst(0, handle.get_stream()); - auto d_out_weights = std::make_optional(rmm::device_uvector(0, handle.get_stream())); - - if (k_level > 0) { - raft::random::RngState rng_state(seed); - seed += d_in.size() * k_level * comm_size; - - std::tie(d_out_src, d_out_dst, d_out_weights) = sample_edges(handle, - graph_view, - edge_weight_view, - rng_state, - d_in, - static_cast(k_level), - with_replacement); - } else { - std::tie(d_out_src, d_out_dst, d_out_weights) = - gather_one_hop_edgelist(handle, graph_view, edge_weight_view, d_in); - } - - // resize accumulators: - auto old_sz = d_result_dst.size(); - auto add_sz = d_out_dst.size(); - auto new_sz = old_sz + add_sz; - - d_result_src.resize(new_sz, handle.get_stream()); - d_result_dst.resize(new_sz, handle.get_stream()); - d_result_weights->resize(new_sz, handle.get_stream()); - - raft::copy( - d_result_src.begin() + old_sz, d_out_src.begin(), d_out_src.size(), handle.get_stream()); - raft::copy( - d_result_dst.begin() + old_sz, d_out_dst.begin(), d_out_dst.size(), handle.get_stream()); - raft::copy(d_result_weights->begin() + old_sz, - d_out_weights->begin(), - d_out_weights->size(), - handle.get_stream()); - - d_in = std::move(d_out_dst); - - ++hop; - } - - return count_and_remove_duplicates( - handle, std::move(d_result_src), std::move(d_result_dst), std::move(*d_result_weights)); -#endif -} - template 0, @@ -284,31 +190,6 @@ uniform_neighbor_sample_impl( } // namespace detail -template -std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed) -{ - rmm::device_uvector d_start_vs(starting_vertices.size(), handle.get_stream()); - raft::copy( - d_start_vs.data(), starting_vertices.data(), starting_vertices.size(), handle.get_stream()); - - return detail::uniform_nbr_sample_impl( - handle, graph_view, edge_weight_view, d_start_vs, fan_out, with_replacement, seed); -} - template , - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed); - template std::tuple, rmm::device_uvector, std::optional>, diff --git a/cpp/src/sampling/uniform_neighbor_sampling_sg.cpp b/cpp/src/sampling/uniform_neighbor_sampling_sg.cpp index 1504fbe17c1..04958d9dfee 100644 --- a/cpp/src/sampling/uniform_neighbor_sampling_sg.cpp +++ b/cpp/src/sampling/uniform_neighbor_sampling_sg.cpp @@ -20,78 +20,6 @@ namespace cugraph { -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed); - -template std::tuple, - rmm::device_uvector, - rmm::device_uvector, - rmm::device_uvector> -uniform_nbr_sample(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - raft::device_span starting_vertices, - raft::host_span fan_out, - bool with_replacement, - uint64_t seed); - template std::tuple, rmm::device_uvector, std::optional>, diff --git a/cpp/tests/c_api/create_graph_test.c b/cpp/tests/c_api/create_graph_test.c index e4560f25fc5..586df59c1b2 100644 --- a/cpp/tests/c_api/create_graph_test.c +++ b/cpp/tests/c_api/create_graph_test.c @@ -41,8 +41,8 @@ int test_create_sg_graph_simple() vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5}; weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f}; - cugraph_resource_handle_t* p_handle = NULL; - cugraph_graph_t* p_graph = NULL; + cugraph_resource_handle_t* handle = NULL; + cugraph_graph_t* graph = NULL; cugraph_graph_properties_t properties; properties.is_symmetric = FALSE; @@ -52,8 +52,8 @@ int test_create_sg_graph_simple() data_type_id_t edge_tid = INT32; data_type_id_t weight_tid = FLOAT32; - p_handle = cugraph_create_resource_handle(NULL); - TEST_ASSERT(test_ret_value, p_handle != NULL, "resource handle creation failed."); + handle = cugraph_create_resource_handle(NULL); + TEST_ASSERT(test_ret_value, handle != NULL, "resource handle creation failed."); cugraph_type_erased_device_array_t* src; cugraph_type_erased_device_array_t* dst; @@ -63,16 +63,16 @@ int test_create_sg_graph_simple() cugraph_type_erased_device_array_view_t* wgt_view; ret_code = - cugraph_type_erased_device_array_create(p_handle, num_edges, vertex_tid, &src, &ret_error); + cugraph_type_erased_device_array_create(handle, num_edges, vertex_tid, &src, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src create failed."); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); ret_code = - cugraph_type_erased_device_array_create(p_handle, num_edges, vertex_tid, &dst, &ret_error); + cugraph_type_erased_device_array_create(handle, num_edges, vertex_tid, &dst, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst create failed."); ret_code = - cugraph_type_erased_device_array_create(p_handle, num_edges, weight_tid, &wgt, &ret_error); + cugraph_type_erased_device_array_create(handle, num_edges, weight_tid, &wgt, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt create failed."); src_view = cugraph_type_erased_device_array_view(src); @@ -80,18 +80,18 @@ int test_create_sg_graph_simple() wgt_view = cugraph_type_erased_device_array_view(wgt); ret_code = cugraph_type_erased_device_array_view_copy_from_host( - p_handle, src_view, (byte_t*)h_src, &ret_error); + handle, src_view, (byte_t*)h_src, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src copy_from_host failed."); ret_code = cugraph_type_erased_device_array_view_copy_from_host( - p_handle, dst_view, (byte_t*)h_dst, &ret_error); + handle, dst_view, (byte_t*)h_dst, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst copy_from_host failed."); ret_code = cugraph_type_erased_device_array_view_copy_from_host( - p_handle, wgt_view, (byte_t*)h_wgt, &ret_error); + handle, wgt_view, (byte_t*)h_wgt, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt copy_from_host failed."); - ret_code = cugraph_sg_graph_create(p_handle, + ret_code = cugraph_sg_graph_create(handle, &properties, src_view, dst_view, @@ -101,11 +101,11 @@ int test_create_sg_graph_simple() FALSE, FALSE, FALSE, - &p_graph, + &graph, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); - cugraph_sg_graph_free(p_graph); + cugraph_sg_graph_free(graph); cugraph_type_erased_device_array_view_free(wgt_view); cugraph_type_erased_device_array_view_free(dst_view); @@ -114,7 +114,7 @@ int test_create_sg_graph_simple() cugraph_type_erased_device_array_free(dst); cugraph_type_erased_device_array_free(src); - cugraph_free_resource_handle(p_handle); + cugraph_free_resource_handle(handle); cugraph_error_free(ret_error); return test_ret_value; @@ -142,8 +142,8 @@ int test_create_sg_graph_csr() vertex_t h_start[] = {0, 1, 2, 3, 4, 5}; weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f}; - cugraph_resource_handle_t* p_handle = NULL; - cugraph_graph_t* p_graph = NULL; + cugraph_resource_handle_t* handle = NULL; + cugraph_graph_t* graph = NULL; cugraph_graph_properties_t properties; properties.is_symmetric = FALSE; @@ -153,8 +153,8 @@ int test_create_sg_graph_csr() data_type_id_t edge_tid = INT32; data_type_id_t weight_tid = FLOAT32; - p_handle = cugraph_create_resource_handle(NULL); - TEST_ASSERT(test_ret_value, p_handle != NULL, "resource handle creation failed."); + handle = cugraph_create_resource_handle(NULL); + TEST_ASSERT(test_ret_value, handle != NULL, "resource handle creation failed."); cugraph_type_erased_device_array_t* offsets; cugraph_type_erased_device_array_t* indices; @@ -164,16 +164,16 @@ int test_create_sg_graph_csr() cugraph_type_erased_device_array_view_t* wgt_view; ret_code = cugraph_type_erased_device_array_create( - p_handle, num_vertices + 1, vertex_tid, &offsets, &ret_error); + handle, num_vertices + 1, vertex_tid, &offsets, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "offsets create failed."); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); ret_code = - cugraph_type_erased_device_array_create(p_handle, num_edges, vertex_tid, &indices, &ret_error); + cugraph_type_erased_device_array_create(handle, num_edges, vertex_tid, &indices, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "indices create failed."); ret_code = - cugraph_type_erased_device_array_create(p_handle, num_edges, weight_tid, &wgt, &ret_error); + cugraph_type_erased_device_array_create(handle, num_edges, weight_tid, &wgt, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt create failed."); offsets_view = cugraph_type_erased_device_array_view(offsets); @@ -181,18 +181,18 @@ int test_create_sg_graph_csr() wgt_view = cugraph_type_erased_device_array_view(wgt); ret_code = cugraph_type_erased_device_array_view_copy_from_host( - p_handle, offsets_view, (byte_t*)h_offsets, &ret_error); + handle, offsets_view, (byte_t*)h_offsets, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "offsets copy_from_host failed."); ret_code = cugraph_type_erased_device_array_view_copy_from_host( - p_handle, indices_view, (byte_t*)h_indices, &ret_error); + handle, indices_view, (byte_t*)h_indices, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "indices copy_from_host failed."); ret_code = cugraph_type_erased_device_array_view_copy_from_host( - p_handle, wgt_view, (byte_t*)h_wgt, &ret_error); + handle, wgt_view, (byte_t*)h_wgt, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt copy_from_host failed."); - ret_code = cugraph_sg_graph_create_from_csr(p_handle, + ret_code = cugraph_sg_graph_create_from_csr(handle, &properties, offsets_view, indices_view, @@ -202,7 +202,7 @@ int test_create_sg_graph_csr() FALSE, FALSE, FALSE, - &p_graph, + &graph, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); @@ -227,15 +227,19 @@ int test_create_sg_graph_csr() h_fan_out_view = cugraph_type_erased_host_array_view_create(fan_out, 1, INT32); ret_code = - cugraph_type_erased_device_array_create(p_handle, num_vertices, INT32, &d_start, &ret_error); + cugraph_type_erased_device_array_create(handle, num_vertices, 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( - p_handle, d_start_view, (byte_t*)h_start, &ret_error); + handle, d_start_view, (byte_t*)h_start, &ret_error); - ret_code = cugraph_uniform_neighbor_sample( - p_handle, p_graph, d_start_view, h_fan_out_view, FALSE, FALSE, &result, &ret_error); + cugraph_rng_state_t *rng_state; + ret_code = cugraph_rng_state_create(handle, 0, &rng_state, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "rng_state create failed."); + + ret_code = cugraph_uniform_neighbor_sample_with_edge_properties( + handle, graph, d_start_view, NULL, h_fan_out_view, rng_state, FALSE, FALSE, &result, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, "uniform_neighbor_sample failed."); @@ -246,8 +250,7 @@ int test_create_sg_graph_csr() srcs = cugraph_sample_result_get_sources(result); dsts = cugraph_sample_result_get_destinations(result); - // FIXME: This is a hack, once we get support for weight and index this should be cleaned up - wgts = cugraph_sample_result_get_index(result); + wgts = cugraph_sample_result_get_edge_weight(result); size_t result_size = cugraph_type_erased_device_array_view_size(srcs); @@ -256,15 +259,15 @@ int test_create_sg_graph_csr() weight_t h_result_wgts[result_size]; ret_code = cugraph_type_erased_device_array_view_copy_to_host( - p_handle, (byte_t*)h_result_srcs, srcs, &ret_error); + handle, (byte_t*)h_result_srcs, srcs, &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( - p_handle, (byte_t*)h_result_dsts, dsts, &ret_error); + handle, (byte_t*)h_result_dsts, dsts, &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( - p_handle, (byte_t*)h_result_wgts, wgts, &ret_error); + handle, (byte_t*)h_result_wgts, wgts, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); TEST_ASSERT(test_ret_value, result_size == num_edges, "number of edges does not match"); @@ -276,7 +279,7 @@ int test_create_sg_graph_csr() } cugraph_sample_result_free(result); - cugraph_sg_graph_free(p_graph); + cugraph_sg_graph_free(graph); cugraph_type_erased_device_array_view_free(wgt_view); cugraph_type_erased_device_array_view_free(indices_view); cugraph_type_erased_device_array_view_free(offsets_view); @@ -284,7 +287,7 @@ int test_create_sg_graph_csr() cugraph_type_erased_device_array_free(indices); cugraph_type_erased_device_array_free(offsets); - cugraph_free_resource_handle(p_handle); + cugraph_free_resource_handle(handle); cugraph_error_free(ret_error); return test_ret_value; diff --git a/cpp/tests/c_api/uniform_neighbor_sample_test.c b/cpp/tests/c_api/uniform_neighbor_sample_test.c index e8fe2b04da9..7864f8c966c 100644 --- a/cpp/tests/c_api/uniform_neighbor_sample_test.c +++ b/cpp/tests/c_api/uniform_neighbor_sample_test.c @@ -113,169 +113,6 @@ int create_test_graph_with_edge_ids(const cugraph_resource_handle_t* p_handle, return test_ret_value; } -int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle, - vertex_t* h_src, - vertex_t* h_dst, - edge_t* h_ids, - size_t num_vertices, - size_t num_edges, - vertex_t* h_start, - size_t num_starts, - int* fan_out, - size_t max_depth, - bool_t with_replacement, - 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_graph_t* graph = NULL; - cugraph_sample_result_t* result = NULL; - - cugraph_type_erased_device_array_t* d_start = NULL; - cugraph_type_erased_device_array_view_t* d_start_view = NULL; - cugraph_type_erased_host_array_view_t* h_fan_out_view = NULL; - - ret_code = create_test_graph_with_edge_ids( - handle, h_src, h_dst, h_ids, 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."); - - h_fan_out_view = cugraph_type_erased_host_array_view_create(fan_out, max_depth, INT32); - - ret_code = cugraph_uniform_neighbor_sample( - handle, graph, d_start_view, h_fan_out_view, with_replacement, FALSE, &result, &ret_error); - -#ifdef NO_CUGRAPH_OPS - TEST_ASSERT( - test_ret_value, ret_code != CUGRAPH_SUCCESS, "uniform_neighbor_sample 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_neighbor_sample failed."); - - cugraph_type_erased_device_array_view_t* srcs; - cugraph_type_erased_device_array_view_t* dsts; - cugraph_type_erased_device_array_view_t* index; - - srcs = cugraph_sample_result_get_sources(result); - dsts = cugraph_sample_result_get_destinations(result); - index = cugraph_sample_result_get_index(result); - - size_t result_size = cugraph_type_erased_device_array_view_size(srcs); - - vertex_t h_srcs[result_size]; - vertex_t h_dsts[result_size]; - edge_t h_index[result_size]; - - ret_code = - cugraph_type_erased_device_array_view_copy_to_host(handle, (byte_t*)h_srcs, srcs, &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_dsts, dsts, &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_index, index, &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 - edge_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_ids[i]; - - for (int i = 0; (i < result_size) && (test_ret_value == 0); ++i) { - TEST_ASSERT(test_ret_value, - M[h_srcs[i]][h_dsts[i]] == h_index[i], - "uniform_neighbor_sample got edge that doesn't exist"); - } - - cugraph_sample_result_free(result); -#endif - - cugraph_type_erased_host_array_view_free(h_fan_out_view); - cugraph_sg_graph_free(graph); - cugraph_error_free(ret_error); - - return test_ret_value; -} - -int test_uniform_neighbor_sample(const cugraph_resource_handle_t* handle) -{ - size_t num_edges = 8; - size_t num_vertices = 6; - size_t fan_out_size = 2; - 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}; - edge_t edge_ids[] = {1, 2, 3, 4, 5, 6, 7, 8}; - vertex_t start[] = {2, 2}; - int fan_out[] = {1, 2}; - - return generic_uniform_neighbor_sample_test(handle, - src, - dst, - edge_ids, - num_vertices, - num_edges, - start, - num_starts, - fan_out, - fan_out_size, - TRUE, - FALSE, - FALSE); -} - -int test_uniform_neighbor_sample_all_neighbors(const cugraph_resource_handle_t* handle) -{ - size_t num_edges = 8; - size_t num_vertices = 6; - size_t fan_out_size = 1; - 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}; - edge_t edge_ids[] = {0, 1, 2, 3, 4, 5, 6, 7}; - vertex_t start[] = {2}; - int fan_out[] = {-1}; - - return generic_uniform_neighbor_sample_test(handle, - src, - dst, - edge_ids, - num_vertices, - num_edges, - start, - num_starts, - fan_out, - fan_out_size, - TRUE, - FALSE, - FALSE); -} - int test_uniform_neighbor_sample_with_properties(const cugraph_resource_handle_t* handle) { data_type_id_t vertex_tid = INT32; @@ -457,8 +294,6 @@ int main(int argc, char** argv) handle = cugraph_create_resource_handle(NULL); int result = 0; - result |= RUN_TEST_NEW(test_uniform_neighbor_sample, handle); - result |= RUN_TEST_NEW(test_uniform_neighbor_sample_all_neighbors, handle); result |= RUN_TEST_NEW(test_uniform_neighbor_sample_with_properties, handle); cugraph_free_resource_handle(handle); diff --git a/python/pylibcugraph/pylibcugraph/_cugraph_c/algorithms.pxd b/python/pylibcugraph/pylibcugraph/_cugraph_c/algorithms.pxd index e20b68693fc..b3a5d07dd9d 100644 --- a/python/pylibcugraph/pylibcugraph/_cugraph_c/algorithms.pxd +++ b/python/pylibcugraph/pylibcugraph/_cugraph_c/algorithms.pxd @@ -227,19 +227,6 @@ cdef extern from "cugraph_c/algorithms.h": cugraph_error_t** error ) - # uniform neighborhood sampling - cdef cugraph_error_code_t \ - cugraph_uniform_neighbor_sample( - const cugraph_resource_handle_t* handle, - cugraph_graph_t* graph, - const cugraph_type_erased_device_array_view_t* start, - const cugraph_type_erased_host_array_view_t* fan_out, - bool_t without_replacement, - bool_t do_expensive_check, - cugraph_sample_result_t** result, - cugraph_error_t** error - ) - # uniform random walks cdef cugraph_error_code_t \ cugraph_uniform_random_walks( diff --git a/python/pylibcugraph/pylibcugraph/uniform_neighbor_sample.pyx b/python/pylibcugraph/pylibcugraph/uniform_neighbor_sample.pyx index 0b353a1b4c4..a0b4473eb3b 100644 --- a/python/pylibcugraph/pylibcugraph/uniform_neighbor_sample.pyx +++ b/python/pylibcugraph/pylibcugraph/uniform_neighbor_sample.pyx @@ -37,7 +37,6 @@ from pylibcugraph._cugraph_c.graph cimport ( cugraph_graph_t, ) from pylibcugraph._cugraph_c.algorithms cimport ( - cugraph_uniform_neighbor_sample, cugraph_sample_result_t, cugraph_sample_result_get_sources, cugraph_sample_result_get_destinations, From 6b5d7507d82fc0f6db60f58eb29ed8bf803f02f4 Mon Sep 17 00:00:00 2001 From: Charles Hastings Date: Thu, 9 Feb 2023 14:15:40 -0800 Subject: [PATCH 2/2] fix MG tester --- .../c_api/mg_uniform_neighbor_sample_test.c | 26 ++++++++++++------- 1 file changed, 17 insertions(+), 9 deletions(-) diff --git a/cpp/tests/c_api/mg_uniform_neighbor_sample_test.c b/cpp/tests/c_api/mg_uniform_neighbor_sample_test.c index fd808ee6317..824c9c7ee90 100644 --- a/cpp/tests/c_api/mg_uniform_neighbor_sample_test.c +++ b/cpp/tests/c_api/mg_uniform_neighbor_sample_test.c @@ -51,6 +51,13 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle cugraph_type_erased_device_array_view_t* d_start_view = NULL; cugraph_type_erased_host_array_view_t* h_fan_out_view = NULL; + int rank = cugraph_resource_handle_get_rank(handle); + + cugraph_rng_state_t* rng_state; + ret_code = cugraph_rng_state_create(handle, rank, &rng_state, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "rng_state create failed."); + TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + ret_code = create_mg_test_graph_with_edge_ids( handle, h_src, h_dst, h_idx, num_edges, store_transposed, FALSE, &graph, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); @@ -67,8 +74,16 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle h_fan_out_view = cugraph_type_erased_host_array_view_create(fan_out, max_depth, INT32); - ret_code = cugraph_uniform_neighbor_sample( - handle, graph, d_start_view, h_fan_out_view, with_replacement, FALSE, &result, &ret_error); + ret_code = cugraph_uniform_neighbor_sample_with_edge_properties(handle, + graph, + d_start_view, + NULL, + h_fan_out_view, + rng_state, + with_replacement, + FALSE, + &result, + &ret_error); #ifdef NO_CUGRAPH_OPS TEST_ASSERT( @@ -79,18 +94,15 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle cugraph_type_erased_device_array_view_t* srcs; cugraph_type_erased_device_array_view_t* dsts; - cugraph_type_erased_device_array_view_t* index; srcs = cugraph_sample_result_get_sources(result); dsts = cugraph_sample_result_get_destinations(result); - index = cugraph_sample_result_get_index(result); size_t result_size = cugraph_type_erased_device_array_view_size(srcs); vertex_t h_srcs[result_size]; vertex_t h_dsts[result_size]; int h_labels[result_size]; - edge_t h_index[result_size]; size_t* h_counts; ret_code = @@ -101,10 +113,6 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle cugraph_type_erased_device_array_view_copy_to_host(handle, (byte_t*)h_dsts, dsts, &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_index, index, &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