Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Improve MG graph creation #2044

Merged
merged 17 commits into from
Feb 9, 2022
Merged

Improve MG graph creation #2044

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
17 commits
Select commit Hold shift + click to select a range
6102cd1
update group_by_and_count to not use reduce_by_key (which is expensiv…
seunghwak Jan 27, 2022
5146b19
add time measurements (should be undone)
seunghwak Jan 27, 2022
5073304
Merge branch 'branch-22.04' of github.com:rapidsai/cugraph into enh_g…
seunghwak Feb 1, 2022
7c02fcf
improve weak scaling behavior of renumber
seunghwak Feb 1, 2022
0744160
move is_first_in_run_t to graph_utils.cuh
seunghwak Feb 1, 2022
077008b
avoid using device lambdas
seunghwak Feb 1, 2022
6a0dfa1
fix compile errors
seunghwak Feb 1, 2022
41645aa
code cleanup
seunghwak Feb 2, 2022
a0b009e
fix overflow bug with 2^31 or more vertices
seunghwak Feb 2, 2022
b5cb9c4
Merge branch 'branch-22.04' of github.com:rapidsai/cugraph into enh_g…
seunghwak Feb 3, 2022
9a70472
delete temporary performance measurement code
seunghwak Feb 3, 2022
214ada9
delete additional temporary performance measurement code
seunghwak Feb 3, 2022
3a605b5
remove temporary performance measurement code
seunghwak Feb 3, 2022
c537e2d
Merge branch 'branch-22.04' of github.com:rapidsai/cugraph into enh_g…
seunghwak Feb 3, 2022
4852ce4
clang-format & copyright year
seunghwak Feb 3, 2022
ac8a6b7
added static_assert
seunghwak Feb 9, 2022
f17db5f
Merge branch 'branch-22.04' of github.com:rapidsai/cugraph into enh_g…
seunghwak Feb 9, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
11 changes: 10 additions & 1 deletion cpp/include/cugraph/detail/graph_utils.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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.
Expand Down Expand Up @@ -77,5 +77,14 @@ struct compute_partition_id_from_edge_t {
}
};

template <typename vertex_t>
struct is_first_in_run_t {
vertex_t const* vertices{nullptr};
__device__ bool operator()(size_t i) const
{
return (i == 0) || (vertices[i - 1] != vertices[i]);
}
};

} // namespace detail
} // namespace cugraph
10 changes: 0 additions & 10 deletions cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -99,16 +99,6 @@ struct call_key_aggregated_e_op_t {
}
};

// a workaround for cudaErrorInvalidDeviceFunction error when device lambda is used
template <typename vertex_t>
struct is_first_in_run_t {
vertex_t const* major_vertices{nullptr};
__device__ bool operator()(size_t i) const
{
return ((i == 0) || (major_vertices[i] != major_vertices[i - 1])) ? true : false;
}
};

// a workaround for cudaErrorInvalidDeviceFunction error when device lambda is used
template <typename vertex_t>
struct is_valid_vertex_t {
Expand Down
68 changes: 34 additions & 34 deletions cpp/include/cugraph/utilities/shuffle_comm.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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.
Expand All @@ -22,6 +22,7 @@
#include <rmm/device_uvector.hpp>
#include <rmm/exec_policy.hpp>

#include <thrust/binary_search.h>
#include <thrust/distance.h>
#include <thrust/fill.h>
#include <thrust/reduce.h>
Expand All @@ -36,6 +37,23 @@ namespace cugraph {

namespace detail {

template <typename GroupIdIterator>
struct compute_group_id_count_pair_t {
GroupIdIterator group_id_first{};
GroupIdIterator group_id_last{};

__device__ thrust::tuple<int, size_t> operator()(size_t i) const
{
static_assert(
std::is_same_v<typename thrust::iterator_traits<GroupIdIterator>::value_type, int>);
auto lower_it =
thrust::lower_bound(thrust::seq, group_id_first, group_id_last, static_cast<int>(i));
auto upper_it = thrust::upper_bound(thrust::seq, lower_it, group_id_last, static_cast<int>(i));
return thrust::make_tuple(static_cast<int>(i),
static_cast<size_t>(thrust::distance(lower_it, upper_it)));
Comment on lines +49 to +53
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Instead of casting i to int can you cast it to
typename std::iterator_traits<GroupIdIterator>::value_type while you are using it for binary_search?
It could be cast to int while returning in the tuple.

Copy link
Contributor Author

@seunghwak seunghwak Feb 9, 2022 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yeah...

Actually, typename std::iterator_traits<GroupIdIterator>::value_type should be int here (Group IDs or GPU IDs or process IDs are all assumed to be int). Using GroupIdIterator intead of int type pointer because I am using thrust::transform_iterator to compute group ID on the fly.

I'll add

static_assert(std::is_same_v(typename std::iterator_traits<GroupIdIterator>::value_type, int)); to make this clearer.

And I may modify operator()(size_t i) to operator()(int i) to avoid type casting... but I have mixed thoughts about this (I assume thrust::tabulate computes the ranges of i from thrust::distance() and i can potentially overflow int. Shouldn't happen unless we have more than 2^31-1 GPUs and that's the assumption for using 'int' here for GPU IDs. I guess explicitly casting this makes this intention more clearer than relying on implicit casting).

What do you think?

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

yeah, that works.

}
};

// inline to suppress a complaint about ODR violation
inline std::tuple<std::vector<size_t>,
std::vector<size_t>,
Expand Down Expand Up @@ -128,23 +146,14 @@ rmm::device_uvector<size_t> groupby_and_count(ValueIterator tx_value_first /* [I
[value_to_group_id_op] __device__(auto value) { return value_to_group_id_op(value); });
rmm::device_uvector<int> d_tx_dst_ranks(num_groups, stream_view);
rmm::device_uvector<size_t> d_tx_value_counts(d_tx_dst_ranks.size(), stream_view);
auto last =
thrust::reduce_by_key(rmm::exec_policy(stream_view),
group_id_first,
group_id_first + thrust::distance(tx_value_first, tx_value_last),
thrust::make_constant_iterator(size_t{1}),
d_tx_dst_ranks.begin(),
d_tx_value_counts.begin());
if (thrust::distance(d_tx_dst_ranks.begin(), thrust::get<0>(last)) < num_groups) {
rmm::device_uvector<size_t> d_counts(num_groups, stream_view);
thrust::fill(rmm::exec_policy(stream_view), d_counts.begin(), d_counts.end(), size_t{0});
thrust::scatter(rmm::exec_policy(stream_view),
d_tx_value_counts.begin(),
thrust::get<1>(last),
d_tx_dst_ranks.begin(),
d_counts.begin());
d_tx_value_counts = std::move(d_counts);
}
auto rank_count_pair_first = thrust::make_zip_iterator(
thrust::make_tuple(d_tx_dst_ranks.begin(), d_tx_value_counts.begin()));
thrust::tabulate(
rmm::exec_policy(stream_view),
rank_count_pair_first,
rank_count_pair_first + num_groups,
detail::compute_group_id_count_pair_t<decltype(group_id_first)>{
group_id_first, group_id_first + thrust::distance(tx_value_first, tx_value_last)});

return d_tx_value_counts;
}
Expand All @@ -169,22 +178,13 @@ rmm::device_uvector<size_t> groupby_and_count(VertexIterator tx_key_first /* [IN
tx_key_first, [key_to_group_id_op] __device__(auto key) { return key_to_group_id_op(key); });
rmm::device_uvector<int> d_tx_dst_ranks(num_groups, stream_view);
rmm::device_uvector<size_t> d_tx_value_counts(d_tx_dst_ranks.size(), stream_view);
auto last = thrust::reduce_by_key(rmm::exec_policy(stream_view),
group_id_first,
group_id_first + thrust::distance(tx_key_first, tx_key_last),
thrust::make_constant_iterator(size_t{1}),
d_tx_dst_ranks.begin(),
d_tx_value_counts.begin());
if (thrust::distance(d_tx_dst_ranks.begin(), thrust::get<0>(last)) < num_groups) {
rmm::device_uvector<size_t> d_counts(num_groups, stream_view);
thrust::fill(rmm::exec_policy(stream_view), d_counts.begin(), d_counts.end(), size_t{0});
thrust::scatter(rmm::exec_policy(stream_view),
d_tx_value_counts.begin(),
thrust::get<1>(last),
d_tx_dst_ranks.begin(),
d_counts.begin());
d_tx_value_counts = std::move(d_counts);
}
auto rank_count_pair_first = thrust::make_zip_iterator(
thrust::make_tuple(d_tx_dst_ranks.begin(), d_tx_value_counts.begin()));
thrust::tabulate(rmm::exec_policy(stream_view),
rank_count_pair_first,
rank_count_pair_first + num_groups,
detail::compute_group_id_count_pair_t<decltype(group_id_first)>{
group_id_first, group_id_first + thrust::distance(tx_key_first, tx_key_last)});

return d_tx_value_counts;
}
Expand Down
6 changes: 3 additions & 3 deletions cpp/src/generators/generate_rmat_edgelist.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2021, NVIDIA CORPORATION.
* Copyright (c) 2021-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.
Expand Down Expand Up @@ -94,8 +94,8 @@ std::tuple<rmm::device_uvector<vertex_t>, rmm::device_uvector<vertex_t>> generat
}
}
}
src += src_bit_set ? static_cast<vertex_t>(1 << bit) : 0;
dst += dst_bit_set ? static_cast<vertex_t>(1 << bit) : 0;
src += src_bit_set ? static_cast<vertex_t>(vertex_t{1} << bit) : 0;
dst += dst_bit_set ? static_cast<vertex_t>(vertex_t{1} << bit) : 0;
}
return thrust::make_tuple(src, dst);
});
Expand Down
Loading