diff --git a/cpp/src/components/weakly_connected_components_impl.cuh b/cpp/src/components/weakly_connected_components_impl.cuh index 5d553edf7de..a46e1eba6fb 100644 --- a/cpp/src/components/weakly_connected_components_impl.cuh +++ b/cpp/src/components/weakly_connected_components_impl.cuh @@ -466,8 +466,8 @@ void weakly_connected_components_impl(raft::handle_t const& handle, // 2-3. initialize vertex frontier, edge_buffer, and edge_dst_components (if // multi-gpu) - vertex_frontier_t vertex_frontier(handle, - num_buckets); + vertex_frontier_t vertex_frontier( + handle, num_buckets); vertex_t next_candidate_offset{0}; edge_t edge_count{0}; @@ -533,7 +533,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, auto max_pushes = GraphViewType::is_multi_gpu ? compute_num_out_nbrs_from_frontier( - handle, level_graph_view, vertex_frontier, bucket_idx_cur) + handle, level_graph_view, vertex_frontier.bucket(bucket_idx_cur)) : edge_count; // FIXME: if we use cuco::static_map (no duplicates, ideally we need static_set), edge_buffer @@ -545,8 +545,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, auto new_frontier_tagged_vertex_buffer = transform_reduce_v_frontier_outgoing_e_by_dst( handle, level_graph_view, - vertex_frontier, - bucket_idx_cur, + vertex_frontier.bucket(bucket_idx_cur), edge_src_dummy_property_t{}.view(), edge_dst_dummy_property_t{}.view(), [col_components = diff --git a/cpp/src/cores/core_number_impl.cuh b/cpp/src/cores/core_number_impl.cuh index 808a0cb98a3..9a335a0a41c 100644 --- a/cpp/src/cores/core_number_impl.cuh +++ b/cpp/src/cores/core_number_impl.cuh @@ -161,7 +161,7 @@ void core_number(raft::handle_t const& handle, constexpr size_t bucket_idx_next = 1; constexpr size_t num_buckets = 2; - vertex_frontier_t vertex_frontier(handle, num_buckets); + vertex_frontier_t vertex_frontier(handle, num_buckets); edge_dst_property_t, edge_t> dst_core_numbers(handle, graph_view); @@ -213,8 +213,7 @@ void core_number(raft::handle_t const& handle, transform_reduce_v_frontier_outgoing_e_by_dst( handle, graph_view, - vertex_frontier, - bucket_idx_cur, + vertex_frontier.bucket(bucket_idx_cur), edge_src_dummy_property_t{}.view(), dst_core_numbers.view(), [k, delta] __device__(vertex_t src, vertex_t dst, auto, auto dst_val) { diff --git a/cpp/src/prims/transform_reduce_v_frontier_outgoing_e_by_dst.cuh b/cpp/src/prims/transform_reduce_v_frontier_outgoing_e_by_dst.cuh index d43de416693..bbe15b76370 100644 --- a/cpp/src/prims/transform_reduce_v_frontier_outgoing_e_by_dst.cuh +++ b/cpp/src/prims/transform_reduce_v_frontier_outgoing_e_by_dst.cuh @@ -792,12 +792,11 @@ auto sort_and_reduce_buffer_elements( } // namespace detail -template +template typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( raft::handle_t const& handle, GraphViewType const& graph_view, - VertexFrontierType const& frontier, - size_t cur_frontier_bucket_idx) + VertexFrontierBucketType const& frontier) { static_assert(!GraphViewType::is_storage_transposed, "GraphViewType should support the push model."); @@ -805,25 +804,23 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( using vertex_t = typename GraphViewType::vertex_type; using edge_t = typename GraphViewType::edge_type; using weight_t = typename GraphViewType::weight_type; - using key_t = typename VertexFrontierType::key_type; + using key_t = typename VertexFrontierBucketType::key_type; edge_t ret{0}; - auto const& cur_frontier_bucket = frontier.bucket(cur_frontier_bucket_idx); vertex_t const* local_frontier_vertex_first{nullptr}; if constexpr (std::is_same_v) { - local_frontier_vertex_first = cur_frontier_bucket.begin(); + local_frontier_vertex_first = frontier.begin(); } else { - local_frontier_vertex_first = thrust::get<0>(cur_frontier_bucket.begin().get_iterator_tuple()); + local_frontier_vertex_first = thrust::get<0>(frontier.begin().get_iterator_tuple()); } std::vector local_frontier_sizes{}; if constexpr (GraphViewType::is_multi_gpu) { - auto& col_comm = handle.get_subcomm(cugraph::partition_2d::key_naming_t().col_name()); - local_frontier_sizes = - host_scalar_allgather(col_comm, cur_frontier_bucket.size(), handle.get_stream()); + auto& col_comm = handle.get_subcomm(cugraph::partition_2d::key_naming_t().col_name()); + local_frontier_sizes = host_scalar_allgather(col_comm, frontier.size(), handle.get_stream()); } else { - local_frontier_sizes = std::vector{static_cast(cur_frontier_bucket.size())}; + local_frontier_sizes = std::vector{static_cast(frontier.size())}; } for (size_t i = 0; i < graph_view.number_of_local_edge_partitions(); ++i) { auto edge_partition = @@ -882,7 +879,7 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( ret += thrust::transform_reduce( handle.get_thrust_policy(), local_frontier_vertex_first, - local_frontier_vertex_first + cur_frontier_bucket.size(), + local_frontier_vertex_first + frontier.size(), [edge_partition] __device__(auto major) { auto major_offset = edge_partition.major_offset_from_major_nocheck(major); return edge_partition.local_degree(major_offset); @@ -900,12 +897,12 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( * outputs by (tagged-)destination ID. * * Edge functor outputs are thrust::optional objects and invalid if thrust::nullopt. Vertices are - * assumed to be tagged if VertexFrontierType::key_type is a tuple of a vertex type and a tag type - * (VertexFrontierType::key_type is identical to a vertex type otherwise). + * assumed to be tagged if VertexFrontierBucketType::key_type is a tuple of a vertex type and a tag + * type (VertexFrontierBucketType::key_type is identical to a vertex type otherwise). * * @tparam GraphViewType Type of the passed non-owning graph object. - * @tparam VertexFrontierType Type of the vertex frontier class which abstracts vertex frontier - * managements. + * @tparam VertexFrontierBucketType Type of the vertex frontier bucket class which abstracts the + * current (tagged-)vertex frontier. * @tparam EdgeSrcValueInputWrapper Type of the wrapper for edge source property values. * @tparam EdgeDstValueInputWrapper Type of the wrapper for edge destination property values. * @tparam EdgeOp Type of the quaternary (or quinary) edge operator. @@ -913,10 +910,7 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( * @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 frontier VertexFrontierType class object for vertex frontier managements. This object - * includes multiple bucket objects. - * @param cur_frontier_bucket_idx Index of the vertex frontier bucket holding vertices for the - * current iteration. + * @param frontier VertexFrontierBucketType class object for the current vertex frontier. * @param edge_src_value_input Wrapper used to access source input property values (for the edge * sources assigned to this process in multi-GPU). Use either cugraph::edge_src_property_t::view() * (if @p e_op needs to access source property values) or cugraph::edge_src_dummy_property_t::view() @@ -945,23 +939,22 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( * values (if ReduceOp::value_type is void). */ template std::conditional_t< !std::is_same_v, - std::tuple( + std::tuple( 0, rmm::cuda_stream_view{})), decltype(detail::allocate_optional_payload_buffer( 0, rmm::cuda_stream_view{}))>, - decltype( - allocate_dataframe_buffer(0, rmm::cuda_stream_view{}))> + decltype(allocate_dataframe_buffer( + 0, rmm::cuda_stream_view{}))> transform_reduce_v_frontier_outgoing_e_by_dst(raft::handle_t const& handle, GraphViewType const& graph_view, - VertexFrontierType const& frontier, - size_t cur_frontier_bucket_idx, + VertexFrontierBucketType const& frontier, EdgeSrcValueInputWrapper edge_src_value_input, EdgeDstValueInputWrapper edge_dst_value_input, EdgeOp e_op, @@ -974,7 +967,7 @@ transform_reduce_v_frontier_outgoing_e_by_dst(raft::handle_t const& handle, using vertex_t = typename GraphViewType::vertex_type; using edge_t = typename GraphViewType::edge_type; using weight_t = typename GraphViewType::weight_type; - using key_t = typename VertexFrontierType::key_type; + using key_t = typename VertexFrontierBucketType::key_type; using payload_t = typename ReduceOp::value_type; using edge_partition_src_input_device_view_t = std::conditional_t< @@ -998,15 +991,12 @@ transform_reduce_v_frontier_outgoing_e_by_dst(raft::handle_t const& handle, vertex_t, typename EdgeDstValueInputWrapper::value_iterator>>>; - CUGRAPH_EXPECTS(cur_frontier_bucket_idx < frontier.num_buckets(), - "Invalid input argument: invalid current bucket index."); - if (do_expensive_check) { // currently, nothing to do } - auto frontier_key_first = frontier.bucket(cur_frontier_bucket_idx).begin(); - auto frontier_key_last = frontier.bucket(cur_frontier_bucket_idx).end(); + auto frontier_key_first = frontier.begin(); + auto frontier_key_last = frontier.end(); // 1. fill the buffer @@ -1124,6 +1114,12 @@ transform_reduce_v_frontier_outgoing_e_by_dst(raft::handle_t const& handle, } if (segment_offsets) { + if constexpr (!VertexFrontierBucketType::is_sorted_unique) { + thrust::sort(handle.get_thrust_policy(), + edge_partition_frontier_src_first, + edge_partition_frontier_src_last); + } + static_assert(detail::num_sparse_segments_per_vertex_partition == 3); std::vector h_thresholds(detail::num_sparse_segments_per_vertex_partition + (graph_view.use_dcs() ? 1 : 0) - 1); diff --git a/cpp/src/prims/update_v_frontier.cuh b/cpp/src/prims/update_v_frontier.cuh index e28cce8d3ec..dd460e339b0 100644 --- a/cpp/src/prims/update_v_frontier.cuh +++ b/cpp/src/prims/update_v_frontier.cuh @@ -186,9 +186,9 @@ void update_v_frontier(raft::handle_t const& handle, VertexFrontierType& frontier, std::vector const& next_frontier_bucket_indices, VertexValueInputIterator vertex_value_input_first, - // FIXME: currently, it is undefined behavior if vertices in the frontier are - // tagged and the same vertex property is updated by multiple v_op - // invocations with the same vertex but with different tags. + // FIXME: currently, it is undefined behavior if there are more than one @p + // key_buffer elements with the same vertex ID and the same vertex property + // value is updated by multiple @p v_op invocations with the same vertex ID. VertexValueOutputIterator vertex_value_output_first, VertexOp v_op, bool do_expensive_check = false) @@ -313,9 +313,9 @@ void update_v_frontier(raft::handle_t const& handle, VertexFrontierType& frontier, std::vector const& next_frontier_bucket_indices, VertexValueInputIterator vertex_value_input_first, - // FIXME: currently, it is undefined behavior if vertices in the frontier are - // tagged and the same vertex property is updated by multiple v_op - // invocations with the same vertex but with different tags. + // FIXME: currently, it is undefined behavior if there are more than one @p + // key_buffer elements with the same vertex ID and the same vertex property + // value is updated by multiple @p v_op invocations with the same vertex ID. VertexValueOutputIterator vertex_value_output_first, VertexOp v_op, bool do_expensive_check = false) diff --git a/cpp/src/prims/vertex_frontier.cuh b/cpp/src/prims/vertex_frontier.cuh index 8b3b1ed3d91..72fec4518cf 100644 --- a/cpp/src/prims/vertex_frontier.cuh +++ b/cpp/src/prims/vertex_frontier.cuh @@ -46,24 +46,32 @@ namespace cugraph { -// stores unique key objects in the sorted (non-descending) order; key type is either vertex_t -// (tag_t == void) or thrust::tuple (tag_t != void) -template -class sorted_unique_key_bucket_t { +// key type is either vertex_t (tag_t == void) or thrust::tuple (tag_t != void) +// if sorted_unique is true, stores unique key objects in the sorted (non-descending) order. +// if false, there can be duplicates and the elements may not be sorted. +template +class key_bucket_t { + public: + using key_type = + std::conditional_t, vertex_t, thrust::tuple>; + static bool constexpr is_sorted_unique = sorted_unique; + static_assert(std::is_same_v || std::is_arithmetic_v); using optional_buffer_type = std:: conditional_t, std::byte /* dummy */, rmm::device_uvector>; - public: template >* = nullptr> - sorted_unique_key_bucket_t(raft::handle_t const& handle) + key_bucket_t(raft::handle_t const& handle) : handle_ptr_(&handle), vertices_(0, handle.get_stream()), tags_(std::byte{0}) { } template >* = nullptr> - sorted_unique_key_bucket_t(raft::handle_t const& handle) + key_bucket_t(raft::handle_t const& handle) : handle_ptr_(&handle), vertices_(0, handle.get_stream()), tags_(0, handle.get_stream()) { } @@ -126,21 +134,31 @@ class sorted_unique_key_bucket_t { std::is_same_v::value_type, vertex_t>); if (vertices_.size() > 0) { - rmm::device_uvector merged_vertices( - vertices_.size() + thrust::distance(vertex_first, vertex_last), handle_ptr_->get_stream()); - thrust::merge(handle_ptr_->get_thrust_policy(), - vertices_.begin(), - vertices_.end(), - vertex_first, - vertex_last, - merged_vertices.begin()); - merged_vertices.resize(thrust::distance(merged_vertices.begin(), - thrust::unique(handle_ptr_->get_thrust_policy(), - merged_vertices.begin(), - merged_vertices.end())), - handle_ptr_->get_stream()); - merged_vertices.shrink_to_fit(handle_ptr_->get_stream()); - vertices_ = std::move(merged_vertices); + if constexpr (sorted_unique) { + rmm::device_uvector merged_vertices( + vertices_.size() + thrust::distance(vertex_first, vertex_last), + handle_ptr_->get_stream()); + thrust::merge(handle_ptr_->get_thrust_policy(), + vertices_.begin(), + vertices_.end(), + vertex_first, + vertex_last, + merged_vertices.begin()); + merged_vertices.resize(thrust::distance(merged_vertices.begin(), + thrust::unique(handle_ptr_->get_thrust_policy(), + merged_vertices.begin(), + merged_vertices.end())), + handle_ptr_->get_stream()); + merged_vertices.shrink_to_fit(handle_ptr_->get_stream()); + vertices_ = std::move(merged_vertices); + } else { + auto cur_size = vertices_.size(); + vertices_.resize(cur_size + thrust::distance(vertex_first, vertex_last)); + thrust::copy(handle_ptr_->get_thrust_policy(), + vertex_first, + vertex_last, + vertices_.begin() + cur_size); + } } else { vertices_.resize(thrust::distance(vertex_first, vertex_last), handle_ptr_->get_stream()); thrust::copy(handle_ptr_->get_thrust_policy(), vertex_first, vertex_last, vertices_.begin()); @@ -164,30 +182,42 @@ class sorted_unique_key_bucket_t { thrust::tuple>); if (vertices_.size() > 0) { - rmm::device_uvector merged_vertices( - vertices_.size() + thrust::distance(key_first, key_last), handle_ptr_->get_stream()); - rmm::device_uvector merged_tags(merged_vertices.size(), handle_ptr_->get_stream()); - auto old_pair_first = - thrust::make_zip_iterator(thrust::make_tuple(vertices_.begin(), tags_.begin())); - auto merged_pair_first = - thrust::make_zip_iterator(thrust::make_tuple(merged_vertices.begin(), merged_tags.begin())); - thrust::merge(handle_ptr_->get_thrust_policy(), - old_pair_first, - old_pair_first + vertices_.size(), - key_first, - key_last, - merged_pair_first); - merged_vertices.resize( - thrust::distance(merged_pair_first, - thrust::unique(handle_ptr_->get_thrust_policy(), - merged_pair_first, - merged_pair_first + merged_vertices.size())), - handle_ptr_->get_stream()); - merged_tags.resize(merged_vertices.size(), handle_ptr_->get_stream()); - merged_vertices.shrink_to_fit(handle_ptr_->get_stream()); - merged_tags.shrink_to_fit(handle_ptr_->get_stream()); - vertices_ = std::move(merged_vertices); - tags_ = std::move(merged_tags); + if constexpr (sorted_unique) { + rmm::device_uvector merged_vertices( + vertices_.size() + thrust::distance(key_first, key_last), handle_ptr_->get_stream()); + rmm::device_uvector merged_tags(merged_vertices.size(), handle_ptr_->get_stream()); + auto old_pair_first = + thrust::make_zip_iterator(thrust::make_tuple(vertices_.begin(), tags_.begin())); + auto merged_pair_first = thrust::make_zip_iterator( + thrust::make_tuple(merged_vertices.begin(), merged_tags.begin())); + thrust::merge(handle_ptr_->get_thrust_policy(), + old_pair_first, + old_pair_first + vertices_.size(), + key_first, + key_last, + merged_pair_first); + merged_vertices.resize( + thrust::distance(merged_pair_first, + thrust::unique(handle_ptr_->get_thrust_policy(), + merged_pair_first, + merged_pair_first + merged_vertices.size())), + handle_ptr_->get_stream()); + merged_tags.resize(merged_vertices.size(), handle_ptr_->get_stream()); + merged_vertices.shrink_to_fit(handle_ptr_->get_stream()); + merged_tags.shrink_to_fit(handle_ptr_->get_stream()); + vertices_ = std::move(merged_vertices); + tags_ = std::move(merged_tags); + } else { + auto cur_size = vertices_.size(); + vertices_.resize(cur_size + thrust::distance(key_first, key_last)); + tags_.resize(vertices_.size()); + thrust::copy( + handle_ptr_->get_thrust_policy(), + key_first, + key_last, + thrust::make_zip_iterator(thrust::make_tuple(vertices_.begin(), tags_.begin())) + + cur_size); + } } else { vertices_.resize(thrust::distance(key_first, key_last), handle_ptr_->get_stream()); tags_.resize(thrust::distance(key_first, key_last), handle_ptr_->get_stream()); @@ -200,7 +230,7 @@ class sorted_unique_key_bucket_t { size_t size() const { return vertices_.size(); } - template + template std::enable_if_t aggregate_size() const { return host_scalar_allreduce(handle_ptr_->get_comms(), @@ -209,7 +239,7 @@ class sorted_unique_key_bucket_t { handle_ptr_->get_stream()); } - template + template std::enable_if_t aggregate_size() const { return vertices_.size(); @@ -229,34 +259,6 @@ class sorted_unique_key_bucket_t { if constexpr (!std::is_same_v) { tags_.shrink_to_fit(handle_ptr_->get_stream()); } } -// FIXME: to silence the spurious warning (missing return statement ...) due to the nvcc bug -// (https://stackoverflow.com/questions/64523302/cuda-missing-return-statement-at-end-of-non-void- -// function-in-constexpr-if-fun) -#if 1 - template >* = nullptr> - auto const begin() const - { - return vertices_.begin(); - } - - template >* = nullptr> - auto begin() - { - return vertices_.begin(); - } - - template >* = nullptr> - auto const begin() const - { - return thrust::make_zip_iterator(thrust::make_tuple(vertices_.begin(), tags_.begin())); - } - - template >* = nullptr> - auto begin() - { - return thrust::make_zip_iterator(thrust::make_tuple(vertices_.begin(), tags_.begin())); - } -#else auto const begin() const { if constexpr (std::is_same_v) { @@ -274,7 +276,6 @@ class sorted_unique_key_bucket_t { return thrust::make_zip_iterator(thrust::make_tuple(vertices_.begin(), tags_.begin())); } } -#endif auto const end() const { return begin() + vertices_.size(); } @@ -286,7 +287,10 @@ class sorted_unique_key_bucket_t { optional_buffer_type tags_; }; -template +template class vertex_frontier_t { static_assert(std::is_same_v || std::is_arithmetic_v); @@ -305,12 +309,13 @@ class vertex_frontier_t { size_t num_buckets() const { return buckets_.size(); } - sorted_unique_key_bucket_t& bucket(size_t bucket_idx) + key_bucket_t& bucket(size_t bucket_idx) { return buckets_[bucket_idx]; } - sorted_unique_key_bucket_t const& bucket(size_t bucket_idx) const + key_bucket_t const& bucket( + size_t bucket_idx) const { return buckets_[bucket_idx]; } @@ -462,7 +467,7 @@ class vertex_frontier_t { private: raft::handle_t const* handle_ptr_{nullptr}; - std::vector> buckets_{}; + std::vector> buckets_{}; }; } // namespace cugraph diff --git a/cpp/src/traversal/bfs_impl.cuh b/cpp/src/traversal/bfs_impl.cuh index d652a6d674b..90ce118a4d8 100644 --- a/cpp/src/traversal/bfs_impl.cuh +++ b/cpp/src/traversal/bfs_impl.cuh @@ -179,8 +179,8 @@ void bfs(raft::handle_t const& handle, constexpr size_t bucket_idx_next = 1; constexpr size_t num_buckets = 2; - vertex_frontier_t vertex_frontier(handle, - num_buckets); + vertex_frontier_t vertex_frontier(handle, + num_buckets); vertex_frontier.bucket(bucket_idx_cur).insert(sources, sources + n_sources); rmm::device_uvector visited_flags( @@ -234,8 +234,7 @@ void bfs(raft::handle_t const& handle, auto [new_frontier_vertex_buffer, predecessor_buffer] = transform_reduce_v_frontier_outgoing_e_by_dst(handle, push_graph_view, - vertex_frontier, - bucket_idx_cur, + vertex_frontier.bucket(bucket_idx_cur), edge_src_dummy_property_t{}.view(), edge_dst_dummy_property_t{}.view(), #if 1 diff --git a/cpp/src/traversal/sssp_impl.cuh b/cpp/src/traversal/sssp_impl.cuh index 9f7190d3632..e79b8aaa5ed 100644 --- a/cpp/src/traversal/sssp_impl.cuh +++ b/cpp/src/traversal/sssp_impl.cuh @@ -133,8 +133,8 @@ void sssp(raft::handle_t const& handle, constexpr size_t bucket_idx_far = 2; constexpr size_t num_buckets = 3; - vertex_frontier_t vertex_frontier(handle, - num_buckets); + vertex_frontier_t vertex_frontier(handle, + num_buckets); // 5. SSSP iteration @@ -169,8 +169,7 @@ void sssp(raft::handle_t const& handle, transform_reduce_v_frontier_outgoing_e_by_dst( handle, push_graph_view, - vertex_frontier, - bucket_idx_cur_near, + vertex_frontier.bucket(bucket_idx_cur_near), GraphViewType::is_multi_gpu ? edge_src_distances.view() : detail::edge_major_property_view_t(distances), diff --git a/cpp/tests/prims/mg_transform_reduce_v_frontier_outgoing_e_by_dst.cu b/cpp/tests/prims/mg_transform_reduce_v_frontier_outgoing_e_by_dst.cu index c41751cd365..45dce13c618 100644 --- a/cpp/tests/prims/mg_transform_reduce_v_frontier_outgoing_e_by_dst.cu +++ b/cpp/tests/prims/mg_transform_reduce_v_frontier_outgoing_e_by_dst.cu @@ -233,8 +233,8 @@ class Tests_MGTransformReduceVFrontierOutgoingEByDst constexpr size_t bucket_idx_cur = 0; constexpr size_t num_buckets = 1; - cugraph::vertex_frontier_t mg_vertex_frontier(*handle_, - num_buckets); + cugraph::vertex_frontier_t mg_vertex_frontier(*handle_, + num_buckets); mg_vertex_frontier.bucket(bucket_idx_cur) .insert(cugraph::get_dataframe_buffer_begin(mg_key_buffer), cugraph::get_dataframe_buffer_end(mg_key_buffer)); @@ -254,8 +254,7 @@ class Tests_MGTransformReduceVFrontierOutgoingEByDst mg_new_frontier_key_buffer = cugraph::transform_reduce_v_frontier_outgoing_e_by_dst( *handle_, mg_graph_view, - mg_vertex_frontier, - bucket_idx_cur, + mg_vertex_frontier.bucket(bucket_idx_cur), mg_src_properties.view(), mg_dst_properties.view(), e_op_t{}, @@ -265,8 +264,7 @@ class Tests_MGTransformReduceVFrontierOutgoingEByDst cugraph::transform_reduce_v_frontier_outgoing_e_by_dst( *handle_, mg_graph_view, - mg_vertex_frontier, - bucket_idx_cur, + mg_vertex_frontier.bucket(bucket_idx_cur), mg_src_properties.view(), mg_dst_properties.view(), e_op_t{}, @@ -395,8 +393,8 @@ class Tests_MGTransformReduceVFrontierOutgoingEByDst }); } - cugraph::vertex_frontier_t sg_vertex_frontier(*handle_, - num_buckets); + cugraph::vertex_frontier_t sg_vertex_frontier( + *handle_, num_buckets); sg_vertex_frontier.bucket(bucket_idx_cur) .insert(cugraph::get_dataframe_buffer_begin(sg_key_buffer), cugraph::get_dataframe_buffer_end(sg_key_buffer)); @@ -409,8 +407,7 @@ class Tests_MGTransformReduceVFrontierOutgoingEByDst sg_new_frontier_key_buffer = cugraph::transform_reduce_v_frontier_outgoing_e_by_dst( *handle_, sg_graph_view, - sg_vertex_frontier, - bucket_idx_cur, + sg_vertex_frontier.bucket(bucket_idx_cur), sg_src_properties.view(), sg_dst_properties.view(), e_op_t{}, @@ -420,8 +417,7 @@ class Tests_MGTransformReduceVFrontierOutgoingEByDst cugraph::transform_reduce_v_frontier_outgoing_e_by_dst( *handle_, sg_graph_view, - sg_vertex_frontier, - bucket_idx_cur, + sg_vertex_frontier.bucket(bucket_idx_cur), sg_src_properties.view(), sg_dst_properties.view(), e_op_t{},