Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
8 changes: 7 additions & 1 deletion cpp/tests/utilities/test_graphs.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -537,13 +537,19 @@ std::tuple<cugraph::graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_
construct_graph(raft::handle_t const& handle,
input_usecase_t const& input_usecase,
bool test_weighted,
bool renumber = true)
bool renumber = true,
bool drop_self_loops = false,
bool drop_multi_edges = false)
{
auto [d_src_v, d_dst_v, d_weights_v, d_vertices_v, num_vertices, is_symmetric] =
input_usecase
.template construct_edgelist<vertex_t, edge_t, weight_t, store_transposed, multi_gpu>(
handle, test_weighted);

if (drop_self_loops) { remove_self_loops(handle, d_src_v, d_dst_v, d_weights_v); }

if (drop_multi_edges) { sort_and_remove_multi_edges(handle, d_src_v, d_dst_v, d_weights_v); }

return cugraph::
create_graph_from_edgelist<vertex_t, edge_t, weight_t, store_transposed, multi_gpu>(
handle,
Expand Down
129 changes: 129 additions & 0 deletions cpp/tests/utilities/thrust_wrapper.cu
Original file line number Diff line number Diff line change
Expand Up @@ -21,8 +21,10 @@
#include <rmm/exec_policy.hpp>

#include <thrust/copy.h>
#include <thrust/remove.h>
#include <thrust/shuffle.h>
#include <thrust/sort.h>
#include <thrust/unique.h>

namespace cugraph {
namespace test {
Expand Down Expand Up @@ -164,5 +166,132 @@ template rmm::device_uvector<int32_t> randomly_select(raft::handle_t const& hand
template rmm::device_uvector<int64_t> randomly_select(raft::handle_t const& handle,
rmm::device_uvector<int64_t> const& input,
size_t count);

template <typename vertex_t, typename weight_t>
void remove_self_loops(raft::handle_t const& handle,
rmm::device_uvector<vertex_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<vertex_t>& d_dst_v /* [INOUT] */,
std::optional<rmm::device_uvector<weight_t>>& d_weight_v /* [INOUT] */)
{
if (d_weight_v) {
auto edge_first = thrust::make_zip_iterator(
thrust::make_tuple(d_src_v.begin(), d_dst_v.begin(), (*d_weight_v).begin()));
d_src_v.resize(
thrust::distance(edge_first,
thrust::remove_if(
handle.get_thrust_policy(),
edge_first,
edge_first + d_src_v.size(),
[] __device__(auto e) { return thrust::get<0>(e) == thrust::get<1>(e); })),
handle.get_stream());
d_dst_v.resize(d_src_v.size(), handle.get_stream());
(*d_weight_v).resize(d_src_v.size(), handle.get_stream());
} else {
auto edge_first =
thrust::make_zip_iterator(thrust::make_tuple(d_src_v.begin(), d_dst_v.begin()));
d_src_v.resize(
thrust::distance(edge_first,
thrust::remove_if(
handle.get_thrust_policy(),
edge_first,
edge_first + d_src_v.size(),
[] __device__(auto e) { return thrust::get<0>(e) == thrust::get<1>(e); })),
handle.get_stream());
d_dst_v.resize(d_src_v.size(), handle.get_stream());
}

d_src_v.shrink_to_fit(handle.get_stream());
d_dst_v.shrink_to_fit(handle.get_stream());
if (d_weight_v) { (*d_weight_v).shrink_to_fit(handle.get_stream()); }
}

template void remove_self_loops(
raft::handle_t const& handle,
rmm::device_uvector<int32_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<int32_t>& d_dst_v /* [INOUT] */,
std::optional<rmm::device_uvector<float>>& d_weight_v /* [INOUT] */);

template void remove_self_loops(
raft::handle_t const& handle,
rmm::device_uvector<int32_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<int32_t>& d_dst_v /* [INOUT] */,
std::optional<rmm::device_uvector<double>>& d_weight_v /* [INOUT] */);

template void remove_self_loops(
raft::handle_t const& handle,
rmm::device_uvector<int64_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<int64_t>& d_dst_v /* [INOUT] */,
std::optional<rmm::device_uvector<float>>& d_weight_v /* [INOUT] */);

template void remove_self_loops(
raft::handle_t const& handle,
rmm::device_uvector<int64_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<int64_t>& d_dst_v /* [INOUT] */,
std::optional<rmm::device_uvector<double>>& d_weight_v /* [INOUT] */);

template <typename vertex_t, typename weight_t>
void sort_and_remove_multi_edges(
raft::handle_t const& handle,
rmm::device_uvector<vertex_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<vertex_t>& d_dst_v /* [INOUT] */,
std::optional<rmm::device_uvector<weight_t>>& d_weight_v /* [INOUT] */)
{
if (d_weight_v) {
auto edge_first = thrust::make_zip_iterator(
thrust::make_tuple(d_src_v.begin(), d_dst_v.begin(), (*d_weight_v).begin()));
thrust::sort(handle.get_thrust_policy(), edge_first, edge_first + d_src_v.size());
d_src_v.resize(
thrust::distance(edge_first,
thrust::unique(handle.get_thrust_policy(),
edge_first,
edge_first + d_src_v.size(),
[] __device__(auto lhs, auto rhs) {
return (thrust::get<0>(lhs) == thrust::get<0>(rhs)) &&
(thrust::get<1>(lhs) == thrust::get<1>(rhs));
})),
handle.get_stream());
d_dst_v.resize(d_src_v.size(), handle.get_stream());
(*d_weight_v).resize(d_src_v.size(), handle.get_stream());
} else {
auto edge_first =
thrust::make_zip_iterator(thrust::make_tuple(d_src_v.begin(), d_dst_v.begin()));
thrust::sort(handle.get_thrust_policy(), edge_first, edge_first + d_src_v.size());
d_src_v.resize(
thrust::distance(
edge_first,
thrust::unique(handle.get_thrust_policy(), edge_first, edge_first + d_src_v.size())),
handle.get_stream());
d_dst_v.resize(d_src_v.size(), handle.get_stream());
}

d_src_v.shrink_to_fit(handle.get_stream());
d_dst_v.shrink_to_fit(handle.get_stream());
if (d_weight_v) { (*d_weight_v).shrink_to_fit(handle.get_stream()); }
}

template void sort_and_remove_multi_edges(
raft::handle_t const& handle,
rmm::device_uvector<int32_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<int32_t>& d_dst_v /* [INOUT] */,
std::optional<rmm::device_uvector<float>>& d_weight_v /* [INOUT] */);

template void sort_and_remove_multi_edges(
raft::handle_t const& handle,
rmm::device_uvector<int32_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<int32_t>& d_dst_v /* [INOUT] */,
std::optional<rmm::device_uvector<double>>& d_weight_v /* [INOUT] */);

template void sort_and_remove_multi_edges(
raft::handle_t const& handle,
rmm::device_uvector<int64_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<int64_t>& d_dst_v /* [INOUT] */,
std::optional<rmm::device_uvector<float>>& d_weight_v /* [INOUT] */);

template void sort_and_remove_multi_edges(
raft::handle_t const& handle,
rmm::device_uvector<int64_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<int64_t>& d_dst_v /* [INOUT] */,
std::optional<rmm::device_uvector<double>>& d_weight_v /* [INOUT] */);

} // namespace test
} // namespace cugraph
19 changes: 16 additions & 3 deletions cpp/tests/utilities/thrust_wrapper.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -29,19 +29,32 @@ std::tuple<key_buffer_type, value_buffer_type> sort_by_key(raft::handle_t const&

template <typename vertex_t>
void translate_vertex_ids(raft::handle_t const& handle,
rmm::device_uvector<vertex_t>& d_src_v,
rmm::device_uvector<vertex_t>& d_dst_v,
rmm::device_uvector<vertex_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<vertex_t>& d_dst_v /* [INOUT] */,
vertex_t vertex_id_offset);

template <typename vertex_t>
void populate_vertex_ids(raft::handle_t const& handle,
rmm::device_uvector<vertex_t>& d_vertices_v,
rmm::device_uvector<vertex_t>& d_vertices_v /* [INOUT] */,
vertex_t vertex_id_offset);

template <typename T>
rmm::device_uvector<T> randomly_select(raft::handle_t const& handle,
rmm::device_uvector<T> const& input,
size_t count);

template <typename vertex_t, typename weight_t>
void remove_self_loops(raft::handle_t const& handle,
rmm::device_uvector<vertex_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<vertex_t>& d_dst_v /* [INOUT] */,
std::optional<rmm::device_uvector<weight_t>>& d_weight_v /* [INOUT] */);

template <typename vertex_t, typename weight_t>
void sort_and_remove_multi_edges(
raft::handle_t const& handle,
rmm::device_uvector<vertex_t>& d_src_v /* [INOUT] */,
rmm::device_uvector<vertex_t>& d_dst_v /* [INOUT] */,
std::optional<rmm::device_uvector<weight_t>>& d_weight_v /* [INOUT] */);

} // namespace test
} // namespace cugraph