Skip to content

Commit

Permalink
Rename sort_and_shuffle to groupby_gpuid_and_shuffle (#1392)
Browse files Browse the repository at this point in the history
Rename to better reflect what this function should do than how it is currently implemented (which can change in the future for better performance).

Authors:
  - Seunghwa Kang (@seunghwak)

Approvers:
  - Chuck Hastings (@ChuckHastings)
  - Alex Fender (@afender)

URL: #1392
  • Loading branch information
seunghwak authored Feb 22, 2021
1 parent 3f13ffc commit 369beee
Show file tree
Hide file tree
Showing 9 changed files with 33 additions and 31 deletions.
Original file line number Diff line number Diff line change
Expand Up @@ -234,7 +234,7 @@ void copy_v_transform_reduce_key_aggregated_out_nbr(

rmm::device_uvector<vertex_t> rx_unique_keys(0, handle.get_stream());
std::vector<size_t> rx_value_counts{};
std::tie(rx_unique_keys, rx_value_counts) = sort_and_shuffle_values(
std::tie(rx_unique_keys, rx_value_counts) = groupby_gpuid_and_shuffle_values(
comm,
unique_keys.begin(),
unique_keys.end(),
Expand Down Expand Up @@ -372,7 +372,7 @@ void copy_v_transform_reduce_key_aggregated_out_nbr(
rmm::device_uvector<weight_t> rx_key_aggregated_edge_weights(0, handle.get_stream());
std::forward_as_tuple(
std::tie(rx_major_vertices, rx_minor_keys, rx_key_aggregated_edge_weights), std::ignore) =
sort_and_shuffle_values(
groupby_gpuid_and_shuffle_values(
sub_comm,
triplet_first,
triplet_first + tmp_major_vertices.size(),
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -241,7 +241,7 @@ transform_reduce_by_adj_matrix_row_col_key_e(
rmm::device_uvector<vertex_t> rx_unique_keys(0, handle.get_stream());
auto rx_value_for_unique_key_buffer = allocate_dataframe_buffer<T>(0, handle.get_stream());
std::tie(rx_unique_keys, rx_value_for_unique_key_buffer, std::ignore) =
sort_and_shuffle_kv_pairs(
groupby_gpuid_and_shuffle_kv_pairs(
comm,
unique_keys.begin(),
unique_keys.end(),
Expand Down
3 changes: 2 additions & 1 deletion cpp/include/utilities/cython.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -451,7 +451,8 @@ std::unique_ptr<cy_multi_edgelists_t> call_egonet(raft::handle_t const& handle,
template <typename vertex_t, typename edge_t, typename weight_t>
std::unique_ptr<major_minor_weights_t<vertex_t, weight_t>> call_shuffle(
raft::handle_t const& handle,
vertex_t* edgelist_major_vertices, // [IN / OUT]: sort_and_shuffle_values() sorts in-place
vertex_t*
edgelist_major_vertices, // [IN / OUT]: groupby_gpuid_and_shuffle_values() sorts in-place
vertex_t* edgelist_minor_vertices, // [IN / OUT]
weight_t* edgelist_weights, // [IN / OUT]
edge_t num_edgelist_edges,
Expand Down
22 changes: 11 additions & 11 deletions cpp/include/utilities/shuffle_comm.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -232,11 +232,11 @@ auto shuffle_values(raft::comms::comms_t const &comm,
}

template <typename ValueIterator, typename ValueToGPUIdOp>
auto sort_and_shuffle_values(raft::comms::comms_t const &comm,
ValueIterator tx_value_first /* [INOUT */,
ValueIterator tx_value_last /* [INOUT */,
ValueToGPUIdOp value_to_gpu_id_op,
cudaStream_t stream)
auto groupby_gpuid_and_shuffle_values(raft::comms::comms_t const &comm,
ValueIterator tx_value_first /* [INOUT */,
ValueIterator tx_value_last /* [INOUT */,
ValueToGPUIdOp value_to_gpu_id_op,
cudaStream_t stream)
{
auto const comm_size = comm.get_size();

Expand Down Expand Up @@ -275,12 +275,12 @@ auto sort_and_shuffle_values(raft::comms::comms_t const &comm,
}

template <typename VertexIterator, typename ValueIterator, typename KeyToGPUIdOp>
auto sort_and_shuffle_kv_pairs(raft::comms::comms_t const &comm,
VertexIterator tx_key_first /* [INOUT */,
VertexIterator tx_key_last /* [INOUT */,
ValueIterator tx_value_first /* [INOUT */,
KeyToGPUIdOp key_to_gpu_id_op,
cudaStream_t stream)
auto groupby_gpuid_and_shuffle_kv_pairs(raft::comms::comms_t const &comm,
VertexIterator tx_key_first /* [INOUT */,
VertexIterator tx_key_last /* [INOUT */,
ValueIterator tx_value_first /* [INOUT */,
KeyToGPUIdOp key_to_gpu_id_op,
cudaStream_t stream)
{
auto d_tx_value_counts = detail::sort_and_count(
comm, tx_key_first, tx_key_last, tx_value_first, key_to_gpu_id_op, stream);
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/experimental/coarsen_graph.cu
Original file line number Diff line number Diff line change
Expand Up @@ -329,7 +329,7 @@ coarsen_graph(
std::forward_as_tuple(
std::tie(rx_edgelist_major_vertices, rx_edgelist_minor_vertices, rx_edgelist_weights),
std::ignore) =
sort_and_shuffle_values(
groupby_gpuid_and_shuffle_values(
handle.get_comms(),
edge_first,
edge_first + coarsened_edgelist_major_vertices.size(),
Expand Down Expand Up @@ -371,7 +371,7 @@ coarsen_graph(
handle.get_stream());

rmm::device_uvector<vertex_t> rx_unique_labels(0, handle.get_stream());
std::tie(rx_unique_labels, std::ignore) = sort_and_shuffle_values(
std::tie(rx_unique_labels, std::ignore) = groupby_gpuid_and_shuffle_values(
handle.get_comms(),
unique_labels.begin(),
unique_labels.end(),
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/experimental/relabel.cu
Original file line number Diff line number Diff line change
Expand Up @@ -106,7 +106,7 @@ void relabel(raft::handle_t const& handle,
thrust::make_tuple(label_pair_old_labels.begin(), label_pair_new_labels.begin()));
std::forward_as_tuple(std::tie(rx_label_pair_old_labels, rx_label_pair_new_labels),
std::ignore) =
sort_and_shuffle_values(
groupby_gpuid_and_shuffle_values(
handle.get_comms(),
pair_first,
pair_first + num_label_pairs,
Expand Down Expand Up @@ -142,7 +142,7 @@ void relabel(raft::handle_t const& handle,
{
rmm::device_uvector<vertex_t> rx_unique_old_labels(0, handle.get_stream());
std::vector<size_t> rx_value_counts{};
std::tie(rx_unique_old_labels, rx_value_counts) = sort_and_shuffle_values(
std::tie(rx_unique_old_labels, rx_value_counts) = groupby_gpuid_and_shuffle_values(
handle.get_comms(),
unique_old_labels.begin(),
unique_old_labels.end(),
Expand Down
18 changes: 9 additions & 9 deletions cpp/src/experimental/renumber_edgelist.cu
Original file line number Diff line number Diff line change
Expand Up @@ -151,14 +151,14 @@ rmm::device_uvector<vertex_t> compute_renumber_map(
auto pair_first = thrust::make_zip_iterator(thrust::make_tuple(labels.begin(), counts.begin()));
rmm::device_uvector<vertex_t> rx_labels(0, handle.get_stream());
rmm::device_uvector<edge_t> rx_counts(0, handle.get_stream());
std::forward_as_tuple(std::tie(rx_labels, rx_counts), std::ignore) = sort_and_shuffle_values(
comm,
pair_first,
pair_first + labels.size(),
[key_func = detail::compute_gpu_id_from_vertex_t<vertex_t>{comm_size}] __device__(auto val) {
return key_func(thrust::get<0>(val));
},
handle.get_stream());
std::forward_as_tuple(std::tie(rx_labels, rx_counts), std::ignore) =
groupby_gpuid_and_shuffle_values(
comm,
pair_first,
pair_first + labels.size(),
[key_func = detail::compute_gpu_id_from_vertex_t<vertex_t>{comm_size}] __device__(
auto val) { return key_func(thrust::get<0>(val)); },
handle.get_stream());

labels.resize(rx_labels.size(), handle.get_stream());
counts.resize(labels.size(), handle.get_stream());
Expand Down Expand Up @@ -309,7 +309,7 @@ void expensive_check_edgelist(
handle.get_stream());

rmm::device_uvector<vertex_t> rx_unique_edge_vertices(0, handle.get_stream());
std::tie(rx_unique_edge_vertices, std::ignore) = sort_and_shuffle_values(
std::tie(rx_unique_edge_vertices, std::ignore) = groupby_gpuid_and_shuffle_values(
handle.get_comms(),
unique_edge_vertices.begin(),
unique_edge_vertices.end(),
Expand Down
5 changes: 3 additions & 2 deletions cpp/src/utilities/cython.cu
Original file line number Diff line number Diff line change
Expand Up @@ -749,7 +749,8 @@ void call_sssp(raft::handle_t const& handle,
template <typename vertex_t, typename edge_t, typename weight_t>
std::unique_ptr<major_minor_weights_t<vertex_t, weight_t>> call_shuffle(
raft::handle_t const& handle,
vertex_t* edgelist_major_vertices, // [IN / OUT]: sort_and_shuffle_values() sorts in-place
vertex_t*
edgelist_major_vertices, // [IN / OUT]: groupby_gpuid_and_shuffle_values() sorts in-place
vertex_t* edgelist_minor_vertices, // [IN / OUT]
weight_t* edgelist_weights, // [IN / OUT]
edge_t num_edgelist_edges,
Expand All @@ -770,7 +771,7 @@ std::unique_ptr<major_minor_weights_t<vertex_t, weight_t>> call_shuffle(
std::forward_as_tuple(
std::tie(ptr_ret->get_major(), ptr_ret->get_minor(), ptr_ret->get_weights()),
std::ignore) =
cugraph::experimental::sort_and_shuffle_values(
cugraph::experimental::groupby_gpuid_and_shuffle_values(
comm, // handle.get_comms(),
zip_edge,
zip_edge + num_edgelist_edges,
Expand Down
2 changes: 1 addition & 1 deletion python/cugraph/structure/graph_primtypes.pxd
Original file line number Diff line number Diff line change
Expand Up @@ -200,7 +200,7 @@ cdef extern from "utilities/cython.hpp" namespace "cugraph::cython":
vertex_t get_part_matrix_partition_minor_first()
vertex_t get_part_matrix_partition_minor_last()

# 4. `sort_and_shuffle_values()` wrapper:
# 4. `groupby_gpuid_and_shuffle_values()` wrapper:
#
cdef extern from "utilities/cython.hpp" namespace "cugraph::cython":

Expand Down

0 comments on commit 369beee

Please sign in to comment.