diff --git a/cpp/include/cugraph/graph_view.hpp b/cpp/include/cugraph/graph_view.hpp index 65c9e74a214..ee355cf9d55 100644 --- a/cpp/include/cugraph/graph_view.hpp +++ b/cpp/include/cugraph/graph_view.hpp @@ -59,10 +59,14 @@ class serializer_t; // forward... * * To be more specific, a GPU with (col_comm_rank, row_comm_rank) will be responsible for * col_comm_size rectangular partitions [a_i,b_i) by [c,d) where a_i = - * vertex_partition_offsets[row_comm_size * i + row_comm_rank] and b_i = - * vertex_partition_offsets[row_comm_size * i + row_comm_rank + 1]. c is - * vertex_partition_offsets[row_comm_size * col_comm_rank] and d = - * vertex_partition_offsests[row_comm_size * (col_comm_rank + 1)]. + * vertex_partition_range_offsets[row_comm_size * i + row_comm_rank] and b_i = + * vertex_partition_range_offsets[row_comm_size * i + row_comm_rank + 1]. c is + * vertex_partition_range_offsets[row_comm_size * col_comm_rank] and d = + * vertex_partition_range_offsets[row_comm_size * (col_comm_rank + 1)]. + * Here, vertex_partition_range_offsets (size = P + 1) stores the 1D partitioning of the vertex ID + * range [0, # vertices). The first P values store the beginning (inclusive) of each GPU partition. + * The last value marks the end (exclusive) of the last GPU partition (which coincides with # + * vertices). * * See E. G. Boman et. al., “Scalable matrix computations on large scale-free graphs using 2D graph * partitioning”, 2013 for additional detail. @@ -74,27 +78,29 @@ class partition_t { public: partition_t() = default; - partition_t(std::vector const& vertex_partition_offsets, + partition_t(std::vector const& vertex_partition_range_offsets, int row_comm_size, int col_comm_size, int row_comm_rank, int col_comm_rank) - : vertex_partition_offsets_(vertex_partition_offsets), + : vertex_partition_range_offsets_(vertex_partition_range_offsets), comm_rank_(col_comm_rank * row_comm_size + row_comm_rank), row_comm_size_(row_comm_size), col_comm_size_(col_comm_size), row_comm_rank_(row_comm_rank), col_comm_rank_(col_comm_rank) { - CUGRAPH_EXPECTS( - vertex_partition_offsets.size() == static_cast(row_comm_size * col_comm_size + 1), - "Invalid API parameter: erroneous vertex_partition_offsets.size()."); + CUGRAPH_EXPECTS(vertex_partition_range_offsets.size() == + static_cast(row_comm_size * col_comm_size + 1), + "Invalid API parameter: erroneous vertex_partition_range_offsets.size()."); + CUGRAPH_EXPECTS(std::is_sorted(vertex_partition_range_offsets_.begin(), + vertex_partition_range_offsets_.end()), + "Invalid API parameter: partition.vertex_partition_range_offsets values should " + "be non-descending."); CUGRAPH_EXPECTS( - std::is_sorted(vertex_partition_offsets_.begin(), vertex_partition_offsets_.end()), - "Invalid API parameter: partition.vertex_partition_offsets values should be non-descending."); - CUGRAPH_EXPECTS(vertex_partition_offsets_[0] == vertex_t{0}, - "Invalid API parameter: partition.vertex_partition_offsets[0] should be 0."); + vertex_partition_range_offsets_[0] == vertex_t{0}, + "Invalid API parameter: partition.vertex_partition_range_offsets[0] should be 0."); vertex_t start_offset{0}; edge_partition_major_value_start_offsets_.assign(number_of_local_edge_partitions(), 0); @@ -111,31 +117,31 @@ class partition_t { int col_comm_size() const { return col_comm_size_; } int comm_rank() const { return comm_rank_; } - std::vector const& vertex_partition_offsets() const + std::vector const& vertex_partition_range_offsets() const { - return vertex_partition_offsets_; + return vertex_partition_range_offsets_; } std::vector vertex_partition_range_lasts() const { - return std::vector(vertex_partition_offsets_.begin() + 1, - vertex_partition_offsets_.end()); + return std::vector(vertex_partition_range_offsets_.begin() + 1, + vertex_partition_range_offsets_.end()); } std::tuple local_vertex_partition_range() const { - return std::make_tuple(vertex_partition_offsets_[comm_rank_], - vertex_partition_offsets_[comm_rank_ + 1]); + return std::make_tuple(vertex_partition_range_offsets_[comm_rank_], + vertex_partition_range_offsets_[comm_rank_ + 1]); } vertex_t local_vertex_partition_range_first() const { - return vertex_partition_offsets_[comm_rank_]; + return vertex_partition_range_offsets_[comm_rank_]; } vertex_t local_vertex_partition_range_last() const { - return vertex_partition_offsets_[comm_rank_ + 1]; + return vertex_partition_range_offsets_[comm_rank_ + 1]; } vertex_t local_vertex_partition_range_size() const @@ -145,18 +151,18 @@ class partition_t { std::tuple vertex_partition_range(size_t partition_idx) const { - return std::make_tuple(vertex_partition_offsets_[partition_idx], - vertex_partition_offsets_[partition_idx + 1]); + return std::make_tuple(vertex_partition_range_offsets_[partition_idx], + vertex_partition_range_offsets_[partition_idx + 1]); } vertex_t vertex_partition_range_first(size_t partition_idx) const { - return vertex_partition_offsets_[partition_idx]; + return vertex_partition_range_offsets_[partition_idx]; } vertex_t vertex_partition_range_last(size_t partition_idx) const { - return vertex_partition_offsets_[partition_idx + 1]; + return vertex_partition_range_offsets_[partition_idx + 1]; } vertex_t vertex_partition_range_size(size_t partition_idx) const @@ -176,12 +182,12 @@ class partition_t { vertex_t local_edge_partition_major_range_first(size_t partition_idx) const { - return vertex_partition_offsets_[row_comm_size_ * partition_idx + row_comm_rank_]; + return vertex_partition_range_offsets_[row_comm_size_ * partition_idx + row_comm_rank_]; } vertex_t local_edge_partition_major_range_last(size_t partition_idx) const { - return vertex_partition_offsets_[row_comm_size_ * partition_idx + row_comm_rank_ + 1]; + return vertex_partition_range_offsets_[row_comm_size_ * partition_idx + row_comm_rank_ + 1]; } vertex_t local_edge_partition_major_range_size(size_t partition_idx) const @@ -205,12 +211,12 @@ class partition_t { vertex_t local_edge_partition_minor_range_first() const { - return vertex_partition_offsets_[col_comm_rank_ * row_comm_size_]; + return vertex_partition_range_offsets_[col_comm_rank_ * row_comm_size_]; } vertex_t local_edge_partition_minor_range_last() const { - return vertex_partition_offsets_[(col_comm_rank_ + 1) * row_comm_size_]; + return vertex_partition_range_offsets_[(col_comm_rank_ + 1) * row_comm_size_]; } vertex_t local_edge_partition_minor_range_size() const @@ -219,7 +225,7 @@ class partition_t { } private: - std::vector vertex_partition_offsets_{}; // size = P + 1 + std::vector vertex_partition_range_offsets_{}; // size = P + 1 int comm_rank_{0}; int row_comm_size_{0}; diff --git a/cpp/include/cugraph/prims/transform_reduce_by_src_dst_key_e.cuh b/cpp/include/cugraph/prims/per_src_dst_key_transform_reduce_e.cuh similarity index 93% rename from cpp/include/cugraph/prims/transform_reduce_by_src_dst_key_e.cuh rename to cpp/include/cugraph/prims/per_src_dst_key_transform_reduce_e.cuh index f126dbbc90b..7d6a2ad2e43 100644 --- a/cpp/include/cugraph/prims/transform_reduce_by_src_dst_key_e.cuh +++ b/cpp/include/cugraph/prims/per_src_dst_key_transform_reduce_e.cuh @@ -32,7 +32,7 @@ namespace cugraph { namespace detail { // FIXME: block size requires tuning -int32_t constexpr transform_reduce_by_src_dst_key_e_kernel_block_size = 128; +int32_t constexpr per_src_dst_key_transform_reduce_e_kernel_block_size = 128; template (major_range_first - edge_partition.major_range_first()); @@ -352,7 +352,7 @@ template std::tuple, decltype(allocate_dataframe_buffer(0, cudaStream_t{nullptr}))> -transform_reduce_by_src_dst_key_e( +per_src_dst_key_transform_reduce_e( raft::handle_t const& handle, GraphViewType const& graph_view, EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, @@ -414,7 +414,7 @@ transform_reduce_by_src_dst_key_e( if ((*segment_offsets)[1] > 0) { raft::grid_1d_block_t update_grid( (*segment_offsets)[1], - detail::transform_reduce_by_src_dst_key_e_kernel_block_size, + detail::per_src_dst_key_transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_by_src_dst_key_high_degree <<>>( @@ -431,7 +431,7 @@ transform_reduce_by_src_dst_key_e( if ((*segment_offsets)[2] - (*segment_offsets)[1] > 0) { raft::grid_1d_warp_t update_grid( (*segment_offsets)[2] - (*segment_offsets)[1], - detail::transform_reduce_by_src_dst_key_e_kernel_block_size, + detail::per_src_dst_key_transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_by_src_dst_key_mid_degree <<>>( @@ -448,7 +448,7 @@ transform_reduce_by_src_dst_key_e( if ((*segment_offsets)[3] - (*segment_offsets)[2] > 0) { raft::grid_1d_thread_t update_grid( (*segment_offsets)[3] - (*segment_offsets)[2], - detail::transform_reduce_by_src_dst_key_e_kernel_block_size, + detail::per_src_dst_key_transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_by_src_dst_key_low_degree <<>>( @@ -466,7 +466,7 @@ transform_reduce_by_src_dst_key_e( (*(edge_partition.dcs_nzd_vertex_count()) > 0)) { raft::grid_1d_thread_t update_grid( *(edge_partition.dcs_nzd_vertex_count()), - detail::transform_reduce_by_src_dst_key_e_kernel_block_size, + detail::per_src_dst_key_transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_by_src_dst_key_hypersparse <<>>( @@ -482,7 +482,7 @@ transform_reduce_by_src_dst_key_e( } else { raft::grid_1d_thread_t update_grid( edge_partition.major_range_size(), - detail::transform_reduce_by_src_dst_key_e_kernel_block_size, + detail::per_src_dst_key_transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_by_src_dst_key_low_degree @@ -602,25 +602,26 @@ template -auto transform_reduce_by_src_key_e(raft::handle_t const& handle, - GraphViewType const& graph_view, - EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, - EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, - EdgePartitionSrcKeyInputWrapper edge_partition_src_key_input, - EdgeOp e_op, - T init) +auto per_src_key_transform_reduce_e( + raft::handle_t const& handle, + GraphViewType const& graph_view, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgePartitionSrcKeyInputWrapper edge_partition_src_key_input, + EdgeOp e_op, + T init) { static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); static_assert(std::is_same::value); - return detail::transform_reduce_by_src_dst_key_e(handle, - graph_view, - edge_partition_src_value_input, - edge_partition_dst_value_input, - edge_partition_src_key_input, - e_op, - init); + return detail::per_src_dst_key_transform_reduce_e(handle, + graph_view, + edge_partition_src_value_input, + edge_partition_dst_value_input, + edge_partition_src_key_input, + e_op, + init); } // FIXME: EdgeOp & VertexOp in update_frontier_v_push_if_out_nbr concatenates push inidicator or @@ -673,25 +674,26 @@ template -auto transform_reduce_by_dst_key_e(raft::handle_t const& handle, - GraphViewType const& graph_view, - EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, - EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, - EdgePartitionDstKeyInputWrapper edge_partition_dst_key_input, - EdgeOp e_op, - T init) +auto per_dst_key_transform_reduce_e( + raft::handle_t const& handle, + GraphViewType const& graph_view, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgePartitionDstKeyInputWrapper edge_partition_dst_key_input, + EdgeOp e_op, + T init) { static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); static_assert(std::is_same::value); - return detail::transform_reduce_by_src_dst_key_e(handle, - graph_view, - edge_partition_src_value_input, - edge_partition_dst_value_input, - edge_partition_dst_key_input, - e_op, - init); + return detail::per_src_dst_key_transform_reduce_e(handle, + graph_view, + edge_partition_src_value_input, + edge_partition_dst_value_input, + edge_partition_dst_key_input, + e_op, + init); } } // namespace cugraph diff --git a/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh b/cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh similarity index 98% rename from cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh rename to cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh index ec92081f6b8..71094b8f898 100644 --- a/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh +++ b/cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh @@ -143,13 +143,12 @@ struct reduce_with_init_t { } // namespace detail /** - * @brief Iterate over every vertex's key-aggregated outgoing edges to update vertex properties. + * @brief Iterate over every vertex's destination key-aggregated outgoing edges to update vertex + * properties. * - * This function is inspired by thrust::transfrom_reduce() (iteration over the outgoing edges - * part) and thrust::copy() (update vertex properties part, take transform_reduce output as copy - * input). - * Unlike copy_v_transform_reduce_out_nbr, this function first aggregates outgoing edges by key to - * support two level reduction for every vertex. + * This function is inspired by thrust::transfrom_reduce(). + * Unlike per_v_transform_reduce_outgoing_e, this function first aggregates outgoing edges by + * destination keys to support two level reduction for every vertex. * * @tparam GraphViewType Type of the passed non-owning graph object. * @tparam EdgePartitionSrcValueInputWrapper Type of the wrapper for edge partition source property @@ -204,7 +203,7 @@ template -void copy_v_transform_reduce_key_aggregated_out_nbr( +void per_v_transform_reduce_dst_key_aggregated_outgoing_e( raft::handle_t const& handle, GraphViewType const& graph_view, EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, diff --git a/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh b/cpp/include/cugraph/prims/per_v_transform_reduce_incoming_outgoing_e.cuh similarity index 92% rename from cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh rename to cpp/include/cugraph/prims/per_v_transform_reduce_incoming_outgoing_e.cuh index 04fc797517a..bdbf79c4379 100644 --- a/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh +++ b/cpp/include/cugraph/prims/per_v_transform_reduce_incoming_outgoing_e.cuh @@ -44,7 +44,7 @@ namespace cugraph { namespace detail { -int32_t constexpr copy_v_transform_reduce_nbr_kernel_block_size = 512; +int32_t constexpr per_v_transform_reduce_e_kernel_block_size = 512; template -__global__ void copy_v_transform_reduce_nbr_hypersparse( +__global__ void per_v_transform_reduce_e_hypersparse( edge_partition_device_view_t -__global__ void copy_v_transform_reduce_nbr_low_degree( +__global__ void per_v_transform_reduce_e_low_degree( edge_partition_device_view_t -__global__ void copy_v_transform_reduce_nbr_mid_degree( +__global__ void per_v_transform_reduce_e_mid_degree( edge_partition_device_view_t(major_range_first - edge_partition.major_range_first()); @@ -308,7 +308,7 @@ __global__ void copy_v_transform_reduce_nbr_mid_degree( using WarpReduce = cub::WarpReduce; [[maybe_unused]] __shared__ typename WarpReduce::TempStorage - temp_storage[copy_v_transform_reduce_nbr_kernel_block_size / + temp_storage[per_v_transform_reduce_e_kernel_block_size / raft::warp_size()]; // relevant only if update_major == true [[maybe_unused]] property_op @@ -376,7 +376,7 @@ template -__global__ void copy_v_transform_reduce_nbr_high_degree( +__global__ void per_v_transform_reduce_e_high_degree( edge_partition_device_view_t(major_range_first - edge_partition.major_range_first()); auto idx = static_cast(blockIdx.x); - using BlockReduce = - cub::BlockReduce; + using BlockReduce = cub::BlockReduce; [[maybe_unused]] __shared__ typename BlockReduce::TempStorage temp_storage; // relevant only if update_major == true @@ -457,22 +456,23 @@ __global__ void copy_v_transform_reduce_nbr_high_degree( } } -template -void copy_v_transform_reduce_nbr(raft::handle_t const& handle, - GraphViewType const& graph_view, - EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, - EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, - EdgeOp e_op, - T init, - VertexValueOutputIterator vertex_value_output_first) +void per_v_transform_reduce_e(raft::handle_t const& handle, + GraphViewType const& graph_view, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgeOp e_op, + T init, + VertexValueOutputIterator vertex_value_output_first) { - constexpr auto update_major = (in == GraphViewType::is_storage_transposed); + constexpr auto update_major = (incoming == GraphViewType::is_storage_transposed); [[maybe_unused]] constexpr auto max_segments = detail::num_sparse_segments_per_vertex_partition + size_t{1}; using vertex_t = typename GraphViewType::vertex_type; @@ -650,7 +650,7 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, : handle.get_stream(); if constexpr (update_major) { // this is necessary as we don't visit every vertex in the // hypersparse segment in - // copy_v_transform_reduce_nbr_hypersparse + // per_v_transform_reduce_e_hypersparse thrust::fill(rmm::exec_policy(exec_stream), output_buffer + (*segment_offsets)[3], output_buffer + (*segment_offsets)[4], @@ -658,11 +658,11 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, } if (*(edge_partition.dcs_nzd_vertex_count()) > 0) { raft::grid_1d_thread_t update_grid(*(edge_partition.dcs_nzd_vertex_count()), - detail::copy_v_transform_reduce_nbr_kernel_block_size, + detail::per_v_transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); auto segment_output_buffer = output_buffer; if constexpr (update_major) { segment_output_buffer += (*segment_offsets)[3]; } - detail::copy_v_transform_reduce_nbr_hypersparse + detail::per_v_transform_reduce_e_hypersparse <<>>( edge_partition, edge_partition.major_range_first() + (*segment_offsets)[3], @@ -679,11 +679,11 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, (*stream_pool_indices).size()) : handle.get_stream(); raft::grid_1d_thread_t update_grid((*segment_offsets)[3] - (*segment_offsets)[2], - detail::copy_v_transform_reduce_nbr_kernel_block_size, + detail::per_v_transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); auto segment_output_buffer = output_buffer; if constexpr (update_major) { segment_output_buffer += (*segment_offsets)[2]; } - detail::copy_v_transform_reduce_nbr_low_degree + detail::per_v_transform_reduce_e_low_degree <<>>( edge_partition, edge_partition.major_range_first() + (*segment_offsets)[2], @@ -700,11 +700,11 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, (*stream_pool_indices).size()) : handle.get_stream(); raft::grid_1d_warp_t update_grid((*segment_offsets)[2] - (*segment_offsets)[1], - detail::copy_v_transform_reduce_nbr_kernel_block_size, + detail::per_v_transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); auto segment_output_buffer = output_buffer; if constexpr (update_major) { segment_output_buffer += (*segment_offsets)[1]; } - detail::copy_v_transform_reduce_nbr_mid_degree + detail::per_v_transform_reduce_e_mid_degree <<>>( edge_partition, edge_partition.major_range_first() + (*segment_offsets)[1], @@ -721,9 +721,9 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, (*stream_pool_indices).size()) : handle.get_stream(); raft::grid_1d_block_t update_grid((*segment_offsets)[1], - detail::copy_v_transform_reduce_nbr_kernel_block_size, + detail::per_v_transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); - detail::copy_v_transform_reduce_nbr_high_degree + detail::per_v_transform_reduce_e_high_degree <<>>( edge_partition, edge_partition.major_range_first(), @@ -737,9 +737,9 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, } else { if (edge_partition.major_range_size() > 0) { raft::grid_1d_thread_t update_grid(edge_partition.major_range_size(), - detail::copy_v_transform_reduce_nbr_kernel_block_size, + detail::per_v_transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); - detail::copy_v_transform_reduce_nbr_low_degree + detail::per_v_transform_reduce_e_low_degree <<>>( edge_partition, edge_partition.major_range_first(), @@ -883,10 +883,9 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, } // namespace detail /** - * @brief Iterate over the incoming edges to update vertex properties. + * @brief Iterate over every vertex's incoming edges to update vertex properties. * - * This function is inspired by thrust::transfrom_reduce() (iteration over the incoming edges part) - * and thrust::copy() (update vertex properties part, take transform_reduce output as copy input). + * This function is inspired by thrust::transfrom_reduce. * * @tparam GraphViewType Type of the passed non-owning graph object. * @tparam EdgePartitionSrcValueInputWrapper Type of the wrapper for edge partition source property @@ -924,7 +923,7 @@ template -void copy_v_transform_reduce_in_nbr( +void per_v_transform_reduce_incoming_e( raft::handle_t const& handle, GraphViewType const& graph_view, EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, @@ -933,21 +932,19 @@ void copy_v_transform_reduce_in_nbr( T init, VertexValueOutputIterator vertex_value_output_first) { - detail::copy_v_transform_reduce_nbr(handle, - graph_view, - edge_partition_src_value_input, - edge_partition_dst_value_input, - e_op, - init, - vertex_value_output_first); + detail::per_v_transform_reduce_e(handle, + graph_view, + edge_partition_src_value_input, + edge_partition_dst_value_input, + e_op, + init, + vertex_value_output_first); } /** - * @brief Iterate over the outgoing edges to update vertex properties. + * @brief Iterate over every vertex's outgoing edges to update vertex properties. * - * This function is inspired by thrust::transfrom_reduce() (iteration over the outgoing edges - * part) and thrust::copy() (update vertex properties part, take transform_reduce output as copy - * input). + * This function is inspired by thrust::transfrom_reduce(). * * @tparam GraphViewType Type of the passed non-owning graph object. * @tparam EdgePartitionSrcValueInputWrapper Type of the wrapper for edge partition source property @@ -985,7 +982,7 @@ template -void copy_v_transform_reduce_out_nbr( +void per_v_transform_reduce_outgoing_e( raft::handle_t const& handle, GraphViewType const& graph_view, EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, @@ -994,13 +991,13 @@ void copy_v_transform_reduce_out_nbr( T init, VertexValueOutputIterator vertex_value_output_first) { - detail::copy_v_transform_reduce_nbr(handle, - graph_view, - edge_partition_src_value_input, - edge_partition_dst_value_input, - e_op, - init, - vertex_value_output_first); + detail::per_v_transform_reduce_e(handle, + graph_view, + edge_partition_src_value_input, + edge_partition_dst_value_input, + e_op, + init, + vertex_value_output_first); } } // namespace cugraph diff --git a/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh b/cpp/include/cugraph/prims/update_v_frontier_from_outgoing_e.cuh similarity index 96% rename from cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh rename to cpp/include/cugraph/prims/update_v_frontier_from_outgoing_e.cuh index 027c2763e0e..8156dde81cc 100644 --- a/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh +++ b/cpp/include/cugraph/prims/update_v_frontier_from_outgoing_e.cuh @@ -55,7 +55,7 @@ namespace cugraph { namespace detail { -int32_t constexpr update_frontier_v_push_if_out_nbr_kernel_block_size = 512; +int32_t constexpr update_v_frontier_from_outgoing_e_kernel_block_size = 512; // we cannot use std::iterator_traits::value_type if Iterator is void* (reference to void // is not allowed) @@ -131,7 +131,7 @@ template -struct update_frontier_call_v_op_t { +struct update_v_frontier_call_v_op_t { VertexValueInputIterator vertex_value_input_first{}; VertexValueOutputIterator vertex_value_output_first{}; VertexOp v_op{}; @@ -216,7 +216,7 @@ template -__global__ void update_frontier_v_push_if_out_nbr_hypersparse( +__global__ void update_v_frontier_from_outgoing_e_hypersparse( edge_partition_device_view_t(tid); __shared__ edge_t - warp_local_degree_inclusive_sums[update_frontier_v_push_if_out_nbr_kernel_block_size]; + warp_local_degree_inclusive_sums[update_v_frontier_from_outgoing_e_kernel_block_size]; __shared__ edge_t - warp_key_local_edge_offsets[update_frontier_v_push_if_out_nbr_kernel_block_size]; + warp_key_local_edge_offsets[update_v_frontier_from_outgoing_e_kernel_block_size]; using WarpScan = cub::WarpScan; __shared__ typename WarpScan::TempStorage temp_storage; - __shared__ size_t buffer_warp_start_indices[update_frontier_v_push_if_out_nbr_kernel_block_size / + __shared__ size_t buffer_warp_start_indices[update_v_frontier_from_outgoing_e_kernel_block_size / raft::warp_size()]; auto indices = edge_partition.indices(); @@ -386,7 +386,7 @@ template -__global__ void update_frontier_v_push_if_out_nbr_low_degree( +__global__ void update_v_frontier_from_outgoing_e_low_degree( edge_partition_device_view_t(tid); __shared__ edge_t - warp_local_degree_inclusive_sums[update_frontier_v_push_if_out_nbr_kernel_block_size]; + warp_local_degree_inclusive_sums[update_v_frontier_from_outgoing_e_kernel_block_size]; __shared__ edge_t - warp_key_local_edge_offsets[update_frontier_v_push_if_out_nbr_kernel_block_size]; + warp_key_local_edge_offsets[update_v_frontier_from_outgoing_e_kernel_block_size]; using WarpScan = cub::WarpScan; __shared__ typename WarpScan::TempStorage temp_storage; - __shared__ size_t buffer_warp_start_indices[update_frontier_v_push_if_out_nbr_kernel_block_size / + __shared__ size_t buffer_warp_start_indices[update_v_frontier_from_outgoing_e_kernel_block_size / raft::warp_size()]; auto indices = edge_partition.indices(); @@ -546,7 +546,7 @@ template -__global__ void update_frontier_v_push_if_out_nbr_mid_degree( +__global__ void update_v_frontier_from_outgoing_e_mid_degree( edge_partition_device_view_t(tid / raft::warp_size()); - __shared__ size_t buffer_warp_start_indices[update_frontier_v_push_if_out_nbr_kernel_block_size / + __shared__ size_t buffer_warp_start_indices[update_v_frontier_from_outgoing_e_kernel_block_size / raft::warp_size()]; while (idx < static_cast(thrust::distance(key_first, key_last))) { auto key = *(key_first + idx); @@ -653,7 +653,7 @@ template -__global__ void update_frontier_v_push_if_out_nbr_high_degree( +__global__ void update_v_frontier_from_outgoing_e_high_degree( edge_partition_device_view_t(blockIdx.x); - using BlockScan = cub::BlockScan; + using BlockScan = cub::BlockScan; __shared__ typename BlockScan::TempStorage temp_storage; __shared__ size_t buffer_block_start_idx; @@ -705,9 +705,9 @@ __global__ void update_frontier_v_push_if_out_nbr_high_degree( thrust::tie(indices, weights, local_out_degree) = edge_partition.local_edges(src_offset); auto rounded_up_local_out_degree = ((static_cast(local_out_degree) + - (update_frontier_v_push_if_out_nbr_kernel_block_size - 1)) / - update_frontier_v_push_if_out_nbr_kernel_block_size) * - update_frontier_v_push_if_out_nbr_kernel_block_size; + (update_v_frontier_from_outgoing_e_kernel_block_size - 1)) / + update_v_frontier_from_outgoing_e_kernel_block_size) * + update_v_frontier_from_outgoing_e_kernel_block_size; for (size_t i = threadIdx.x; i < rounded_up_local_out_degree; i += blockDim.x) { e_op_result_t e_op_result{}; vertex_t dst{}; @@ -935,7 +935,7 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( // FIXME: this documentation needs to be updated due to (tagged-)vertex support /** * @brief Update (tagged-)vertex frontier and (tagged-)vertex property values iterating over the - * outgoing edges from the frontier. + * outgoing edges from the vertices in the current frontier. * * @tparam GraphViewType Type of the passed non-owning graph object. * @tparam VertexFrontierType Type of the vertex frontier class which abstracts vertex frontier @@ -994,7 +994,7 @@ template -void update_frontier_v_push_if_out_nbr( +void update_v_frontier_from_outgoing_e( raft::handle_t const& handle, GraphViewType const& graph_view, VertexFrontierType& frontier, @@ -1175,9 +1175,9 @@ void update_frontier_v_push_if_out_nbr( if (h_offsets[0] > 0) { raft::grid_1d_block_t update_grid( h_offsets[0], - detail::update_frontier_v_push_if_out_nbr_kernel_block_size, + detail::update_v_frontier_from_outgoing_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); - detail::update_frontier_v_push_if_out_nbr_high_degree + detail::update_v_frontier_from_outgoing_e_high_degree <<>>( edge_partition, get_dataframe_buffer_begin(edge_partition_frontier_key_buffer), @@ -1192,9 +1192,9 @@ void update_frontier_v_push_if_out_nbr( if (h_offsets[1] - h_offsets[0] > 0) { raft::grid_1d_warp_t update_grid( h_offsets[1] - h_offsets[0], - detail::update_frontier_v_push_if_out_nbr_kernel_block_size, + detail::update_v_frontier_from_outgoing_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); - detail::update_frontier_v_push_if_out_nbr_mid_degree + detail::update_v_frontier_from_outgoing_e_mid_degree <<>>( edge_partition, get_dataframe_buffer_begin(edge_partition_frontier_key_buffer) + h_offsets[0], @@ -1209,9 +1209,9 @@ void update_frontier_v_push_if_out_nbr( if (h_offsets[2] - h_offsets[1] > 0) { raft::grid_1d_thread_t update_grid( h_offsets[2] - h_offsets[1], - detail::update_frontier_v_push_if_out_nbr_kernel_block_size, + detail::update_v_frontier_from_outgoing_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); - detail::update_frontier_v_push_if_out_nbr_low_degree + detail::update_v_frontier_from_outgoing_e_low_degree <<>>( edge_partition, get_dataframe_buffer_begin(edge_partition_frontier_key_buffer) + h_offsets[1], @@ -1226,9 +1226,9 @@ void update_frontier_v_push_if_out_nbr( if (edge_partition.dcs_nzd_vertex_count() && (h_offsets[3] - h_offsets[2] > 0)) { raft::grid_1d_thread_t update_grid( h_offsets[3] - h_offsets[2], - detail::update_frontier_v_push_if_out_nbr_kernel_block_size, + detail::update_v_frontier_from_outgoing_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); - detail::update_frontier_v_push_if_out_nbr_hypersparse + detail::update_v_frontier_from_outgoing_e_hypersparse <<>>( edge_partition, edge_partition.major_range_first() + (*segment_offsets)[3], @@ -1245,10 +1245,10 @@ void update_frontier_v_push_if_out_nbr( if (edge_partition_frontier_size > 0) { raft::grid_1d_thread_t update_grid( edge_partition_frontier_size, - detail::update_frontier_v_push_if_out_nbr_kernel_block_size, + detail::update_v_frontier_from_outgoing_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); - detail::update_frontier_v_push_if_out_nbr_low_degree + detail::update_v_frontier_from_outgoing_e_low_degree <<>>( edge_partition, get_dataframe_buffer_begin(edge_partition_frontier_key_buffer), @@ -1380,12 +1380,12 @@ void update_frontier_v_push_if_out_nbr( get_dataframe_buffer_begin(key_buffer), get_dataframe_buffer_begin(key_buffer) + num_buffer_elements, bucket_indices.begin(), - detail::update_frontier_call_v_op_t{ + detail::update_v_frontier_call_v_op_t{ vertex_value_input_first, vertex_value_output_first, v_op, diff --git a/cpp/include/cugraph/utilities/cython.hpp b/cpp/include/cugraph/utilities/cython.hpp index 9d697dbf95a..8f6d860eba8 100644 --- a/cpp/include/cugraph/utilities/cython.hpp +++ b/cpp/include/cugraph/utilities/cython.hpp @@ -282,7 +282,7 @@ struct renum_tuple_t { // std::unique_ptr> get_partition_offsets_wrap(void) // const { - return std::make_unique>(part_.vertex_partition_offsets()); + return std::make_unique>(part_.vertex_partition_range_offsets()); } std::pair get_part_local_vertex_range() const diff --git a/cpp/include/cugraph_c/graph.h b/cpp/include/cugraph_c/graph.h index a4fde4dde22..7dd880a6919 100644 --- a/cpp/include/cugraph_c/graph.h +++ b/cpp/include/cugraph_c/graph.h @@ -83,11 +83,7 @@ void cugraph_sg_graph_free(cugraph_graph_t* graph); * @param [in] src Device array containing the source vertex ids * @param [in] dst Device array containing the destination vertex ids * @param [in] weights Device array containing the edge weights - * @param [in] vertex_partition_offsets Host array containing the offsets for each vertex - * partition - * @param [in] segment_offsets Host array containing the offsets for each segment * @param [in] store_transposed If true create the graph initially in transposed format - * @param [in] num_vertices Number of vertices * @param [in] num_edges Number of edges * @param [in] check If true, do expensive checks to validate the input data * is consistent with software assumptions. If false bypass these checks. diff --git a/cpp/src/centrality/katz_centrality_impl.cuh b/cpp/src/centrality/katz_centrality_impl.cuh index 1fa5c93c4b9..71e23edf65a 100644 --- a/cpp/src/centrality/katz_centrality_impl.cuh +++ b/cpp/src/centrality/katz_centrality_impl.cuh @@ -17,9 +17,9 @@ #include #include -#include #include #include +#include #include #include #include @@ -106,7 +106,7 @@ void katz_centrality(raft::handle_t const& handle, update_edge_partition_src_property( handle, pull_graph_view, old_katz_centralities, edge_partition_src_katz_centralities); - copy_v_transform_reduce_in_nbr( + per_v_transform_reduce_incoming_e( handle, pull_graph_view, edge_partition_src_katz_centralities.device_view(), diff --git a/cpp/src/community/legacy/louvain.cuh b/cpp/src/community/legacy/louvain.cuh index 19e26b31071..cb8383244c2 100644 --- a/cpp/src/community/legacy/louvain.cuh +++ b/cpp/src/community/legacy/louvain.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -230,7 +230,7 @@ class Louvain { weight_t* d_cluster_weights = cluster_weights_v_.data(); // - // MNMG: copy_v_transform_reduce_out_nbr, then copy + // MNMG: per_v_transform_reduce_outgoing_e, then copy // thrust::for_each( handle_.get_thrust_policy(), diff --git a/cpp/src/community/louvain.cuh b/cpp/src/community/louvain.cuh index a5551d76ae2..677496f5b3b 100644 --- a/cpp/src/community/louvain.cuh +++ b/cpp/src/community/louvain.cuh @@ -20,10 +20,10 @@ #include #include -#include -#include #include -#include +#include +#include +#include #include #include #include @@ -393,7 +393,7 @@ class Louvain { rmm::device_uvector cluster_subtract_v( current_graph_view_.local_vertex_partition_range_size(), handle_.get_stream()); - copy_v_transform_reduce_out_nbr( + per_v_transform_reduce_outgoing_e( handle_, current_graph_view_, graph_view_t::is_multi_gpu @@ -509,7 +509,7 @@ class Louvain { vertex_t, decltype(cluster_old_sum_subtract_pair_first)>(cluster_old_sum_subtract_pair_first)); - copy_v_transform_reduce_key_aggregated_out_nbr( + per_v_transform_reduce_dst_key_aggregated_outgoing_e( handle_, current_graph_view_, zipped_src_device_view, @@ -539,7 +539,7 @@ class Louvain { handle_, current_graph_view_, next_clusters_v_.begin(), dst_clusters_cache_); } - std::tie(cluster_keys_v_, cluster_weights_v_) = cugraph::transform_reduce_by_src_key_e( + std::tie(cluster_keys_v_, cluster_weights_v_) = cugraph::per_src_key_transform_reduce_e( handle_, current_graph_view_, dummy_property_t{}.device_view(), diff --git a/cpp/src/components/weakly_connected_components_impl.cuh b/cpp/src/components/weakly_connected_components_impl.cuh index fd2609eb242..861dc3652d9 100644 --- a/cpp/src/components/weakly_connected_components_impl.cuh +++ b/cpp/src/components/weakly_connected_components_impl.cuh @@ -21,7 +21,7 @@ #include #include #include -#include +#include #include #include #include @@ -523,7 +523,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, auto old_num_edge_inserts = num_edge_inserts.value(handle.get_stream()); resize_dataframe_buffer(edge_buffer, old_num_edge_inserts + max_pushes, handle.get_stream()); - update_frontier_v_push_if_out_nbr( + update_v_frontier_from_outgoing_e( handle, level_graph_view, vertex_frontier, diff --git a/cpp/src/cores/core_number_impl.cuh b/cpp/src/cores/core_number_impl.cuh index 6d077e7085c..7534db99ce0 100644 --- a/cpp/src/cores/core_number_impl.cuh +++ b/cpp/src/cores/core_number_impl.cuh @@ -20,7 +20,7 @@ #include #include #include -#include +#include #include #include @@ -195,7 +195,7 @@ void core_number(raft::handle_t const& handle, // mask-out/delete edges. if (graph_view.is_symmetric() || ((degree_type == k_core_degree_type_t::IN) || (degree_type == k_core_degree_type_t::INOUT))) { - update_frontier_v_push_if_out_nbr( + update_v_frontier_from_outgoing_e( handle, graph_view, vertex_frontier, diff --git a/cpp/src/link_analysis/hits_impl.cuh b/cpp/src/link_analysis/hits_impl.cuh index 5422c1d327f..11838af3d11 100644 --- a/cpp/src/link_analysis/hits_impl.cuh +++ b/cpp/src/link_analysis/hits_impl.cuh @@ -17,9 +17,9 @@ #include #include -#include #include #include +#include #include #include #include @@ -105,7 +105,7 @@ std::tuple hits(raft::handle_t const& handle, } for (size_t iter = 0; iter < max_iterations; ++iter) { // Update current destination authorities property - copy_v_transform_reduce_in_nbr( + per_v_transform_reduce_incoming_e( handle, graph_view, prev_src_hubs.device_view(), @@ -117,7 +117,7 @@ std::tuple hits(raft::handle_t const& handle, update_edge_partition_dst_property(handle, graph_view, authorities, curr_dst_auth); // Update current source hubs property - copy_v_transform_reduce_out_nbr( + per_v_transform_reduce_outgoing_e( handle, graph_view, dummy_property_t{}.device_view(), diff --git a/cpp/src/link_analysis/pagerank_impl.cuh b/cpp/src/link_analysis/pagerank_impl.cuh index 5d4c7bde1ed..b929bb07e30 100644 --- a/cpp/src/link_analysis/pagerank_impl.cuh +++ b/cpp/src/link_analysis/pagerank_impl.cuh @@ -17,10 +17,10 @@ #include #include -#include #include #include #include +#include #include #include #include @@ -243,7 +243,7 @@ void pagerank( static_cast(num_vertices) : result_t{0.0}; - copy_v_transform_reduce_in_nbr( + per_v_transform_reduce_incoming_e( handle, pull_graph_view, edge_partition_src_pageranks.device_view(), diff --git a/cpp/src/structure/graph_view_impl.cuh b/cpp/src/structure/graph_view_impl.cuh index f102d78694c..b7b4f61e5e6 100644 --- a/cpp/src/structure/graph_view_impl.cuh +++ b/cpp/src/structure/graph_view_impl.cuh @@ -21,8 +21,8 @@ #include #include #include -#include #include +#include #include #include #include @@ -190,7 +190,7 @@ rmm::device_uvector compute_minor_degrees( rmm::device_uvector minor_degrees(graph_view.local_vertex_partition_range_size(), handle.get_stream()); if (store_transposed) { - copy_v_transform_reduce_out_nbr( + per_v_transform_reduce_outgoing_e( handle, graph_view, dummy_property_t{}.device_view(), @@ -199,7 +199,7 @@ rmm::device_uvector compute_minor_degrees( edge_t{0}, minor_degrees.data()); } else { - copy_v_transform_reduce_in_nbr( + per_v_transform_reduce_incoming_e( handle, graph_view, dummy_property_t{}.device_view(), @@ -225,7 +225,7 @@ rmm::device_uvector compute_weight_sums( rmm::device_uvector weight_sums(graph_view.local_vertex_partition_range_size(), handle.get_stream()); if (major == store_transposed) { - copy_v_transform_reduce_in_nbr( + per_v_transform_reduce_incoming_e( handle, graph_view, dummy_property_t{}.device_view(), @@ -234,7 +234,7 @@ rmm::device_uvector compute_weight_sums( weight_t{0.0}, weight_sums.data()); } else { - copy_v_transform_reduce_out_nbr( + per_v_transform_reduce_outgoing_e( handle, graph_view, dummy_property_t{}.device_view(), diff --git a/cpp/src/structure/renumber_edgelist_impl.cuh b/cpp/src/structure/renumber_edgelist_impl.cuh index 59ee165d39b..2ec1559bbb4 100644 --- a/cpp/src/structure/renumber_edgelist_impl.cuh +++ b/cpp/src/structure/renumber_edgelist_impl.cuh @@ -696,14 +696,14 @@ renumber_edgelist( auto vertex_counts = host_scalar_allgather( comm, static_cast(renumber_map_labels.size()), handle.get_stream()); - std::vector vertex_partition_offsets(comm_size + 1, 0); + std::vector vertex_partition_range_offsets(comm_size + 1, 0); std::partial_sum( - vertex_counts.begin(), vertex_counts.end(), vertex_partition_offsets.begin() + 1); + vertex_counts.begin(), vertex_counts.end(), vertex_partition_range_offsets.begin() + 1); partition_t partition( - vertex_partition_offsets, row_comm_size, col_comm_size, row_comm_rank, col_comm_rank); + vertex_partition_range_offsets, row_comm_size, col_comm_size, row_comm_rank, col_comm_rank); - auto number_of_vertices = vertex_partition_offsets.back(); + auto number_of_vertices = vertex_partition_range_offsets.back(); auto number_of_edges = host_scalar_allreduce( comm, std::accumulate(edgelist_edge_counts.begin(), edgelist_edge_counts.end(), edge_t{0}), diff --git a/cpp/src/traversal/bfs_impl.cuh b/cpp/src/traversal/bfs_impl.cuh index 525328ab563..a16bb3bd49d 100644 --- a/cpp/src/traversal/bfs_impl.cuh +++ b/cpp/src/traversal/bfs_impl.cuh @@ -20,7 +20,7 @@ #include #include #include -#include +#include #include #include #include @@ -222,7 +222,7 @@ void bfs(raft::handle_t const& handle, e_op.prev_visited_flags = prev_visited_flags.data(); } - update_frontier_v_push_if_out_nbr( + update_v_frontier_from_outgoing_e( handle, push_graph_view, vertex_frontier, diff --git a/cpp/src/traversal/sssp_impl.cuh b/cpp/src/traversal/sssp_impl.cuh index c7f1e8b3748..bf45a08375e 100644 --- a/cpp/src/traversal/sssp_impl.cuh +++ b/cpp/src/traversal/sssp_impl.cuh @@ -22,7 +22,7 @@ #include #include #include -#include +#include #include #include #include @@ -161,7 +161,7 @@ void sssp(raft::handle_t const& handle, auto vertex_partition = vertex_partition_device_view_t( push_graph_view.local_vertex_partition_view()); - update_frontier_v_push_if_out_nbr( + update_v_frontier_from_outgoing_e( handle, push_graph_view, vertex_frontier, diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 2b3caf5ee49..e684de16e08 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -575,9 +575,9 @@ if(BUILD_CUGRAPH_MG_TESTS) ConfigureTestMG(MG_COUNT_IF_V_TEST prims/mg_count_if_v.cu) ########################################################################################### - # - MG PRIMS UPDATE_FRONTIER_V_PUSH_IF_OUT_NBR tests -------------------------------------- - ConfigureTestMG(MG_UPDATE_FRONTIER_V_PUSH_IF_OUT_NBR_TEST - prims/mg_update_frontier_v_push_if_out_nbr.cu) + # - MG PRIMS UPDATE_V_FRONTIER_FROM_OUTGOING_E tests -------------------------------------- + ConfigureTestMG(MG_UPDATE_V_FRONTIER_FROM_OUTGOING_E_TEST + prims/mg_update_v_frontier_from_outgoing_e.cu) ########################################################################################### # - MG PRIMS REDUCE_V tests --------------------------------------------------------------- @@ -596,9 +596,9 @@ if(BUILD_CUGRAPH_MG_TESTS) ConfigureTestMG(MG_COUNT_IF_E_TEST prims/mg_count_if_e.cu) ########################################################################################### - # - MG PRIMS COPY_V_TRANSFORM_REDUCE_OUT tests -------------------------------------------- - ConfigureTestMG(MG_COPY_V_TRANSFORM_REDUCE_INOUT_NBR_TEST - prims/mg_copy_v_transform_reduce_inout_nbr.cu) + # - MG PRIMS PER_V_TRANSFORM_REDUCE_INCOMING_OUTGOING_E tests ----------------------------- + ConfigureTestMG(MG_PER_V_TRANSFORM_REDUCE_INCOMING_OUTGOING_E_TEST + prims/mg_per_v_transform_reduce_incoming_outgoing_e.cu) ########################################################################################### # - MG PRIMS EXTRACT_IF_E tests ----------------------------------------------------------- diff --git a/cpp/tests/prims/mg_copy_v_transform_reduce_inout_nbr.cu b/cpp/tests/prims/mg_per_v_transform_reduce_incoming_outgoing_e.cu similarity index 90% rename from cpp/tests/prims/mg_copy_v_transform_reduce_inout_nbr.cu rename to cpp/tests/prims/mg_per_v_transform_reduce_incoming_outgoing_e.cu index 36692b8d9b1..b6a5100fb94 100644 --- a/cpp/tests/prims/mg_copy_v_transform_reduce_inout_nbr.cu +++ b/cpp/tests/prims/mg_per_v_transform_reduce_incoming_outgoing_e.cu @@ -28,8 +28,8 @@ #include #include #include -#include #include +#include #include #include @@ -220,17 +220,17 @@ struct Prims_Usecase { }; template -class Tests_MG_CopyVTransformReduceInOutNbr +class Tests_MG_PerVTransformReduceIncomingOutgoingE : public ::testing::TestWithParam> { public: - Tests_MG_CopyVTransformReduceInOutNbr() {} + Tests_MG_PerVTransformReduceIncomingOutgoingE() {} static void SetupTestCase() {} static void TearDownTestCase() {} virtual void SetUp() {} virtual void TearDown() {} - // Compare the results of copy_v_transform_reduce_out_nbr primitive + // Compare the results of per_v_transform_reduce_incoming|outgoing_e primitive template ( sg_graph_view.local_vertex_partition_range_size(), handle.get_stream()); - copy_v_transform_reduce_out_nbr( + per_v_transform_reduce_outgoing_e( handle, sg_graph_view, sg_row_prop.device_view(), @@ -390,7 +390,7 @@ class Tests_MG_CopyVTransformReduceInOutNbr auto global_in_result = cugraph::allocate_dataframe_buffer( sg_graph_view.local_vertex_partition_range_size(), handle.get_stream()); - copy_v_transform_reduce_in_nbr( + per_v_transform_reduce_incoming_e( handle, sg_graph_view, sg_row_prop.device_view(), @@ -419,19 +419,21 @@ class Tests_MG_CopyVTransformReduceInOutNbr } }; -using Tests_MG_CopyVTransformReduceInOutNbr_File = - Tests_MG_CopyVTransformReduceInOutNbr; -using Tests_MG_CopyVTransformReduceInOutNbr_Rmat = - Tests_MG_CopyVTransformReduceInOutNbr; +using Tests_MG_PerVTransformReduceIncomingOutgoingE_File = + Tests_MG_PerVTransformReduceIncomingOutgoingE; +using Tests_MG_PerVTransformReduceIncomingOutgoingE_Rmat = + Tests_MG_PerVTransformReduceIncomingOutgoingE; -TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_File, CheckInt32Int32FloatTupleIntFloatTransposeFalse) +TEST_P(Tests_MG_PerVTransformReduceIncomingOutgoingE_File, + CheckInt32Int32FloatTupleIntFloatTransposeFalse) { auto param = GetParam(); run_current_test, false>(std::get<0>(param), std::get<1>(param)); } -TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_Rmat, CheckInt32Int32FloatTupleIntFloatTransposeFalse) +TEST_P(Tests_MG_PerVTransformReduceIncomingOutgoingE_Rmat, + CheckInt32Int32FloatTupleIntFloatTransposeFalse) { auto param = GetParam(); run_current_test, false>( @@ -439,14 +441,16 @@ TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_Rmat, CheckInt32Int32FloatTupleIntF cugraph::test::override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); } -TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_File, CheckInt32Int32FloatTupleIntFloatTransposeTrue) +TEST_P(Tests_MG_PerVTransformReduceIncomingOutgoingE_File, + CheckInt32Int32FloatTupleIntFloatTransposeTrue) { auto param = GetParam(); run_current_test, true>(std::get<0>(param), std::get<1>(param)); } -TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_Rmat, CheckInt32Int32FloatTupleIntFloatTransposeTrue) +TEST_P(Tests_MG_PerVTransformReduceIncomingOutgoingE_Rmat, + CheckInt32Int32FloatTupleIntFloatTransposeTrue) { auto param = GetParam(); run_current_test, true>( @@ -454,13 +458,13 @@ TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_Rmat, CheckInt32Int32FloatTupleIntF cugraph::test::override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); } -TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_File, CheckInt32Int32FloatTransposeFalse) +TEST_P(Tests_MG_PerVTransformReduceIncomingOutgoingE_File, CheckInt32Int32FloatTransposeFalse) { auto param = GetParam(); run_current_test(std::get<0>(param), std::get<1>(param)); } -TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_Rmat, CheckInt32Int32FloatTransposeFalse) +TEST_P(Tests_MG_PerVTransformReduceIncomingOutgoingE_Rmat, CheckInt32Int32FloatTransposeFalse) { auto param = GetParam(); run_current_test( @@ -468,13 +472,13 @@ TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_Rmat, CheckInt32Int32FloatTranspose cugraph::test::override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); } -TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_File, CheckInt32Int32FloatTransposeTrue) +TEST_P(Tests_MG_PerVTransformReduceIncomingOutgoingE_File, CheckInt32Int32FloatTransposeTrue) { auto param = GetParam(); run_current_test(std::get<0>(param), std::get<1>(param)); } -TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_Rmat, CheckInt32Int32FloatTransposeTrue) +TEST_P(Tests_MG_PerVTransformReduceIncomingOutgoingE_Rmat, CheckInt32Int32FloatTransposeTrue) { auto param = GetParam(); run_current_test( @@ -484,7 +488,7 @@ TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_Rmat, CheckInt32Int32FloatTranspose INSTANTIATE_TEST_SUITE_P( file_test, - Tests_MG_CopyVTransformReduceInOutNbr_File, + Tests_MG_PerVTransformReduceIncomingOutgoingE_File, ::testing::Combine( ::testing::Values(Prims_Usecase{true}), ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), @@ -494,14 +498,14 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( rmat_small_test, - Tests_MG_CopyVTransformReduceInOutNbr_Rmat, + Tests_MG_PerVTransformReduceIncomingOutgoingE_Rmat, ::testing::Combine(::testing::Values(Prims_Usecase{true}), ::testing::Values(cugraph::test::Rmat_Usecase( 10, 16, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); INSTANTIATE_TEST_SUITE_P( rmat_large_test, - Tests_MG_CopyVTransformReduceInOutNbr_Rmat, + Tests_MG_PerVTransformReduceIncomingOutgoingE_Rmat, ::testing::Combine(::testing::Values(Prims_Usecase{false}), ::testing::Values(cugraph::test::Rmat_Usecase( 20, 32, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); diff --git a/cpp/tests/prims/mg_update_frontier_v_push_if_out_nbr.cu b/cpp/tests/prims/mg_update_v_frontier_from_outgoing_e.cu similarity index 91% rename from cpp/tests/prims/mg_update_frontier_v_push_if_out_nbr.cu rename to cpp/tests/prims/mg_update_v_frontier_from_outgoing_e.cu index 18e70bd693f..60680bba016 100644 --- a/cpp/tests/prims/mg_update_frontier_v_push_if_out_nbr.cu +++ b/cpp/tests/prims/mg_update_v_frontier_from_outgoing_e.cu @@ -30,7 +30,7 @@ #include #include #include -#include +#include #include #include @@ -86,17 +86,17 @@ struct Prims_Usecase { }; template -class Tests_MG_UpdateFrontierVPushIfOutNbr +class Tests_MG_UpdateVFrontierFromOutgoingE : public ::testing::TestWithParam> { public: - Tests_MG_UpdateFrontierVPushIfOutNbr() {} + Tests_MG_UpdateVFrontierFromOutgoingE() {} static void SetupTestCase() {} static void TearDownTestCase() {} virtual void SetUp() {} virtual void TearDown() {} - // Compare the results of update_frontier_v_push_if_out_nbr primitive + // Compare the results of update_v_frontier_from_outgoing_e primitive template void run_current_test(Prims_Usecase const& prims_usecase, input_usecase_t const& input_usecase) { @@ -122,7 +122,7 @@ class Tests_MG_UpdateFrontierVPushIfOutNbr constexpr bool is_multi_gpu = true; constexpr bool renumber = true; // needs to be true for multi gpu case constexpr bool store_transposed = - false; // needs to be false for using update_frontier_v_push_if_out_nbr + false; // needs to be false for using update_v_frontier_from_outgoing_e if (cugraph::test::g_perf) { RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement handle.get_comms().barrier(); @@ -190,7 +190,7 @@ class Tests_MG_UpdateFrontierVPushIfOutNbr } // prims call - update_frontier_v_push_if_out_nbr( + update_v_frontier_from_outgoing_e( handle, mg_graph_view, mg_vertex_frontier, @@ -216,7 +216,7 @@ class Tests_MG_UpdateFrontierVPushIfOutNbr handle.get_comms().barrier(); double elapsed_time{0.0}; hr_clock.stop(&elapsed_time); - std::cout << "MG update_frontier_v_push_if_out_nbr took " << elapsed_time * 1e-6 << " s.\n"; + std::cout << "MG update_v_frontier_from_outgoing_e took " << elapsed_time * 1e-6 << " s.\n"; } //// 4. compare SG & MG results @@ -272,7 +272,7 @@ class Tests_MG_UpdateFrontierVPushIfOutNbr num_buckets); sg_vertex_frontier.bucket(bucket_idx_cur).insert(sources.begin(), sources.end()); - update_frontier_v_push_if_out_nbr( + update_v_frontier_from_outgoing_e( handle, sg_graph_view, sg_vertex_frontier, @@ -306,19 +306,19 @@ class Tests_MG_UpdateFrontierVPushIfOutNbr } }; -using Tests_MG_update_frontier_v_push_if_out_nbr_File = - Tests_MG_UpdateFrontierVPushIfOutNbr; -using Tests_MG_update_frontier_v_push_if_out_nbr_Rmat = - Tests_MG_UpdateFrontierVPushIfOutNbr; +using Tests_MG_UpdateVFrontierFromOutgoingE_File = + Tests_MG_UpdateVFrontierFromOutgoingE; +using Tests_MG_UpdateVFrontierFromOutgoingE_Rmat = + Tests_MG_UpdateVFrontierFromOutgoingE; -TEST_P(Tests_MG_update_frontier_v_push_if_out_nbr_File, CheckInt32Int32FloatTupleIntFloat) +TEST_P(Tests_MG_UpdateVFrontierFromOutgoingE_File, CheckInt32Int32FloatTupleIntFloat) { auto param = GetParam(); run_current_test>(std::get<0>(param), std::get<1>(param)); } -TEST_P(Tests_MG_update_frontier_v_push_if_out_nbr_Rmat, CheckInt32Int32FloatTupleIntFloat) +TEST_P(Tests_MG_UpdateVFrontierFromOutgoingE_Rmat, CheckInt32Int32FloatTupleIntFloat) { auto param = GetParam(); run_current_test>( @@ -326,13 +326,13 @@ TEST_P(Tests_MG_update_frontier_v_push_if_out_nbr_Rmat, CheckInt32Int32FloatTupl cugraph::test::override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); } -TEST_P(Tests_MG_update_frontier_v_push_if_out_nbr_File, CheckInt32Int32Float) +TEST_P(Tests_MG_UpdateVFrontierFromOutgoingE_File, CheckInt32Int32Float) { auto param = GetParam(); run_current_test(std::get<0>(param), std::get<1>(param)); } -TEST_P(Tests_MG_update_frontier_v_push_if_out_nbr_Rmat, CheckInt32Int32Float) +TEST_P(Tests_MG_UpdateVFrontierFromOutgoingE_Rmat, CheckInt32Int32Float) { auto param = GetParam(); run_current_test( @@ -340,13 +340,13 @@ TEST_P(Tests_MG_update_frontier_v_push_if_out_nbr_Rmat, CheckInt32Int32Float) cugraph::test::override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); } -TEST_P(Tests_MG_update_frontier_v_push_if_out_nbr_File, CheckInt32Int64Float) +TEST_P(Tests_MG_UpdateVFrontierFromOutgoingE_File, CheckInt32Int64Float) { auto param = GetParam(); run_current_test(std::get<0>(param), std::get<1>(param)); } -TEST_P(Tests_MG_update_frontier_v_push_if_out_nbr_Rmat, CheckInt32Int64Float) +TEST_P(Tests_MG_UpdateVFrontierFromOutgoingE_Rmat, CheckInt32Int64Float) { auto param = GetParam(); run_current_test( @@ -354,13 +354,13 @@ TEST_P(Tests_MG_update_frontier_v_push_if_out_nbr_Rmat, CheckInt32Int64Float) cugraph::test::override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param))); } -TEST_P(Tests_MG_update_frontier_v_push_if_out_nbr_File, CheckInt64Int64Float) +TEST_P(Tests_MG_UpdateVFrontierFromOutgoingE_File, CheckInt64Int64Float) { auto param = GetParam(); run_current_test(std::get<0>(param), std::get<1>(param)); } -TEST_P(Tests_MG_update_frontier_v_push_if_out_nbr_Rmat, CheckInt64Int64Float) +TEST_P(Tests_MG_UpdateVFrontierFromOutgoingE_Rmat, CheckInt64Int64Float) { auto param = GetParam(); run_current_test( @@ -370,7 +370,7 @@ TEST_P(Tests_MG_update_frontier_v_push_if_out_nbr_Rmat, CheckInt64Int64Float) INSTANTIATE_TEST_SUITE_P( file_test, - Tests_MG_update_frontier_v_push_if_out_nbr_File, + Tests_MG_UpdateVFrontierFromOutgoingE_File, ::testing::Combine( ::testing::Values(Prims_Usecase{true}), ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), @@ -380,7 +380,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( rmat_small_test, - Tests_MG_update_frontier_v_push_if_out_nbr_Rmat, + Tests_MG_UpdateVFrontierFromOutgoingE_Rmat, ::testing::Combine(::testing::Values(Prims_Usecase{true}), ::testing::Values(cugraph::test::Rmat_Usecase( 10, 16, 0.57, 0.19, 0.19, 0, false, false, 0, true)))); @@ -391,7 +391,7 @@ INSTANTIATE_TEST_SUITE_P( vertex & edge type combination) by command line arguments and do not include more than one Rmat_Usecase that differ only in scale or edge factor (to avoid running same benchmarks more than once) */ - Tests_MG_update_frontier_v_push_if_out_nbr_Rmat, + Tests_MG_UpdateVFrontierFromOutgoingE_Rmat, ::testing::Combine(::testing::Values(Prims_Usecase{false}), ::testing::Values(cugraph::test::Rmat_Usecase( 20, 32, 0.57, 0.19, 0.19, 0, false, false, 0, true))));