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.

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

Implement has_edge() & compute_multiplicity() #4096

Merged
merged 21 commits into from
Jan 25, 2024
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
21 commits
Select commit Hold shift + click to select a range
dfbc33a
add an empty line between two functions
seunghwak Jan 16, 2024
e6f6784
added major_idx_from_major_nocheck
seunghwak Jan 16, 2024
c9c3a2b
add initial implementation of has_edge() and compute_multiplicity
seunghwak Jan 17, 2024
e827457
Merge branch 'branch-24.02' of https://github.com/rapidsai/cugraph in…
seunghwak Jan 17, 2024
06d4f77
move count_invalid_vertex_pairs to error_check_utils.cuh
seunghwak Jan 17, 2024
4b4fb46
refactor has_edge() and compute_multiplicity()
seunghwak Jan 17, 2024
7f25cfc
Merge branch 'branch-24.02' of https://github.com/rapidsai/cugraph in…
seunghwak Jan 17, 2024
a7d0fff
clang-format and copyright year
seunghwak Jan 17, 2024
becf133
to_host, to_device specialization for std::vector<bool>
seunghwak Jan 17, 2024
5d3ed2a
remove repetitive tests
seunghwak Jan 17, 2024
e8d0ccc
fix compile error
seunghwak Jan 17, 2024
737f438
add bool specialization for device_gatherv and device_allgatherv
seunghwak Jan 17, 2024
00789bf
add tests for has_edge() and compute_multiplicity
seunghwak Jan 17, 2024
0b45356
copyright year
seunghwak Jan 17, 2024
49f46d0
Merge branch 'branch-24.02' of https://github.com/rapidsai/cugraph in…
seunghwak Jan 17, 2024
7cfe60a
bug fix
seunghwak Jan 18, 2024
c81655d
Merge branch 'branch-24.02' of https://github.com/rapidsai/cugraph in…
seunghwak Jan 18, 2024
43f8485
bug fix in tests
seunghwak Jan 19, 2024
8ff673a
test bug fix
seunghwak Jan 20, 2024
7b0d132
Merge branch 'branch-24.02' of https://github.com/rapidsai/cugraph in…
seunghwak Jan 20, 2024
9ee207f
Merge branch 'branch-24.02' of https://github.com/rapidsai/cugraph in…
seunghwak Jan 24, 2024
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
22 changes: 21 additions & 1 deletion cpp/include/cugraph/edge_partition_device_view.cuh
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020-2023, NVIDIA CORPORATION.
* Copyright (c) 2020-2024, 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 @@ -298,6 +298,20 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
return major_range_first_ + major_offset;
}

__device__ thrust::optional<vertex_t> major_idx_from_major_nocheck(vertex_t major) const noexcept
{
if (major_hypersparse_first_ && (major >= *major_hypersparse_first_)) {
auto major_hypersparse_idx =
detail::major_hypersparse_idx_from_major_nocheck_impl(*dcs_nzd_vertices_, major);
return major_hypersparse_idx
? thrust::make_optional((*major_hypersparse_first_ - major_range_first_) +
*major_hypersparse_idx)
: thrust::nullopt;
} else {
return major - major_range_first_;
}
}

__device__ vertex_t major_from_major_idx_nocheck(vertex_t major_idx) const noexcept
{
if (major_hypersparse_first_) {
Expand Down Expand Up @@ -339,6 +353,7 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
return dcs_nzd_vertices_ ? thrust::optional<vertex_t const*>{(*dcs_nzd_vertices_).data()}
: thrust::nullopt;
}

__host__ __device__ thrust::optional<vertex_t> dcs_nzd_vertex_count() const
{
return dcs_nzd_vertices_
Expand Down Expand Up @@ -460,6 +475,11 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
return major_offset;
}

__device__ thrust::optional<vertex_t> major_idx_from_major_nocheck(vertex_t major) const noexcept
{
return major_offset_from_major_nocheck(major);
}

__device__ vertex_t major_from_major_idx_nocheck(vertex_t major_idx) const noexcept
{
return major_from_major_offset_nocheck(major_idx);
Expand Down
25 changes: 24 additions & 1 deletion cpp/include/cugraph/graph_view.hpp
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020-2023, NVIDIA CORPORATION.
* Copyright (c) 2020-2024, 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 @@ -631,6 +631,19 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
edge_t count_self_loops(raft::handle_t const& handle) const;
edge_t count_multi_edges(raft::handle_t const& handle) const;

rmm::device_uvector<bool> has_edge(raft::handle_t const& handle,
/* (edge_srcs, edge_dsts) should be pre-shuffled */
raft::device_span<vertex_t const> edge_srcs,
raft::device_span<vertex_t const> edge_dsts,
bool do_expensive_check = false);

rmm::device_uvector<edge_t> compute_multiplicity(
raft::handle_t const& handle,
/* (edge_srcs, edge_dsts) should be pre-shuffled */
raft::device_span<vertex_t const> edge_srcs,
raft::device_span<vertex_t const> edge_dsts,
bool do_expensive_check = false);

template <bool transposed = is_storage_transposed>
std::enable_if_t<transposed, std::optional<raft::device_span<vertex_t const>>>
local_sorted_unique_edge_srcs() const
Expand Down Expand Up @@ -928,6 +941,16 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
edge_t count_self_loops(raft::handle_t const& handle) const;
edge_t count_multi_edges(raft::handle_t const& handle) const;

rmm::device_uvector<bool> has_edge(raft::handle_t const& handle,
raft::device_span<vertex_t const> edge_srcs,
raft::device_span<vertex_t const> edge_dsts,
bool do_expensive_check = false);

rmm::device_uvector<edge_t> compute_multiplicity(raft::handle_t const& handle,
raft::device_span<vertex_t const> edge_srcs,
raft::device_span<vertex_t const> edge_dsts,
bool do_expensive_check = false);

template <bool transposed = is_storage_transposed>
std::enable_if_t<transposed, std::optional<raft::device_span<vertex_t const>>>
local_sorted_unique_edge_srcs() const
Expand Down
3 changes: 2 additions & 1 deletion cpp/src/link_prediction/similarity_impl.cuh
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2023, NVIDIA CORPORATION.
* Copyright (c) 2022-2024, 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 @@ -18,6 +18,7 @@
#include <prims/count_if_e.cuh>
#include <prims/per_v_pair_transform_dst_nbr_intersection.cuh>
#include <prims/update_edge_src_dst_property.cuh>
#include <utilities/error_check_utils.cuh>

#include <cugraph/graph_functions.hpp>
#include <cugraph/graph_view.hpp>
Expand Down
170 changes: 11 additions & 159 deletions cpp/src/prims/detail/nbr_intersection.cuh
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2023, NVIDIA CORPORATION.
* Copyright (c) 2022-2024, 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 @@ -17,6 +17,7 @@

#include <prims/detail/optional_dataframe_buffer.hpp>
#include <prims/kv_store.cuh>
#include <utilities/error_check_utils.cuh>

#include <cugraph/edge_partition_device_view.cuh>
#include <cugraph/edge_partition_edge_property_device_view.cuh>
Expand Down Expand Up @@ -63,35 +64,6 @@ namespace cugraph {

namespace detail {

// check vertices in the pair are valid and first element of the pair is within the local vertex
// partition range
template <typename vertex_t>
struct is_invalid_input_vertex_pair_t {
vertex_t num_vertices{};
raft::device_span<vertex_t const> edge_partition_major_range_firsts{};
raft::device_span<vertex_t const> edge_partition_major_range_lasts{};
vertex_t edge_partition_minor_range_first{};
vertex_t edge_partition_minor_range_last{};

__device__ bool operator()(thrust::tuple<vertex_t, vertex_t> pair) const
{
auto major = thrust::get<0>(pair);
auto minor = thrust::get<1>(pair);
if (!is_valid_vertex(num_vertices, major) || !is_valid_vertex(num_vertices, minor)) {
return true;
}
auto it = thrust::upper_bound(thrust::seq,
edge_partition_major_range_lasts.begin(),
edge_partition_major_range_lasts.end(),
major);
if (it == edge_partition_major_range_lasts.end()) { return true; }
auto edge_partition_idx =
static_cast<size_t>(thrust::distance(edge_partition_major_range_lasts.begin(), it));
if (major < edge_partition_major_range_firsts[edge_partition_idx]) { return true; }
return (minor < edge_partition_minor_range_first) || (minor >= edge_partition_minor_range_last);
}
};

// group index determined by major_comm_rank (primary key) and local edge partition index (secondary
// key)
template <typename vertex_t>
Expand Down Expand Up @@ -154,24 +126,11 @@ struct update_rx_major_local_degree_t {
auto major =
rx_majors[rx_group_firsts[major_comm_rank * minor_comm_size + local_edge_partition_idx] +
offset_in_local_edge_partition];
vertex_t major_idx{0};
edge_t local_degree{0};
if (multi_gpu && (edge_partition.major_hypersparse_first() &&
(major >= *(edge_partition.major_hypersparse_first())))) {
auto major_hypersparse_idx = edge_partition.major_hypersparse_idx_from_major_nocheck(major);
if (major_hypersparse_idx) {
major_idx =
(*(edge_partition.major_hypersparse_first()) - edge_partition.major_range_first()) +
*major_hypersparse_idx;
local_degree = edge_partition.local_degree(major_idx);
}
} else {
major_idx = edge_partition.major_offset_from_major_nocheck(major);
local_degree = edge_partition.local_degree(major_idx);
}
auto major_idx = edge_partition.major_idx_from_major_nocheck(major);
auto local_degree = major_idx ? edge_partition.local_degree(*major_idx) : edge_t{0};

if (edge_partition_e_mask && (local_degree > edge_t{0})) {
auto local_offset = edge_partition.local_offset(major_idx);
auto local_offset = edge_partition.local_offset(*major_idx);
local_degree = static_cast<edge_t>(
count_set_bits((*edge_partition_e_mask).value_first(), local_offset, local_degree));
}
Expand Down Expand Up @@ -325,29 +284,11 @@ struct pick_min_degree_t {
edge_t local_degree0{0};
vertex_t major0 = thrust::get<0>(pair);
if constexpr (std::is_same_v<FirstElementToIdxMap, void*>) {
vertex_t major_idx{0};
if constexpr (multi_gpu) {
if (edge_partition.major_hypersparse_first() &&
(major0 >= *(edge_partition.major_hypersparse_first()))) {
auto major_hypersparse_idx =
edge_partition.major_hypersparse_idx_from_major_nocheck(major0);
if (major_hypersparse_idx) {
major_idx =
(*(edge_partition.major_hypersparse_first()) - edge_partition.major_range_first()) +
*major_hypersparse_idx;
local_degree0 = edge_partition.local_degree(major_idx);
}
} else {
major_idx = edge_partition.major_offset_from_major_nocheck(major0);
local_degree0 = edge_partition.local_degree(major_idx);
}
} else {
major_idx = edge_partition.major_offset_from_major_nocheck(major0);
local_degree0 = edge_partition.local_degree(major_idx);
}
auto major_idx = edge_partition.major_idx_from_major_nocheck(major0);
local_degree0 = major_idx ? edge_partition.local_degree(*major_idx) : edge_t{0};

if (edge_partition_e_mask && (local_degree0 > edge_t{0})) {
auto local_offset = edge_partition.local_offset(major_idx);
auto local_offset = edge_partition.local_offset(*major_idx);
local_degree0 =
count_set_bits((*edge_partition_e_mask).value_first(), local_offset, local_degree0);
}
Expand All @@ -360,29 +301,11 @@ struct pick_min_degree_t {
edge_t local_degree1{0};
vertex_t major1 = thrust::get<1>(pair);
if constexpr (std::is_same_v<SecondElementToIdxMap, void*>) {
vertex_t major_idx{0};
if constexpr (multi_gpu) {
if (edge_partition.major_hypersparse_first() &&
(major1 >= *(edge_partition.major_hypersparse_first()))) {
auto major_hypersparse_idx =
edge_partition.major_hypersparse_idx_from_major_nocheck(major1);
if (major_hypersparse_idx) {
major_idx =
(*(edge_partition.major_hypersparse_first()) - edge_partition.major_range_first()) +
*major_hypersparse_idx;
local_degree1 = edge_partition.local_degree(major_idx);
}
} else {
major_idx = edge_partition.major_offset_from_major_nocheck(major1);
local_degree1 = edge_partition.local_degree(major_idx);
}
} else {
major_idx = edge_partition.major_offset_from_major_nocheck(major1);
local_degree1 = edge_partition.local_degree(major_idx);
}
auto major_idx = edge_partition.major_idx_from_major_nocheck(major1);
local_degree1 = major_idx ? edge_partition.local_degree(*major_idx) : edge_t{0};

if (edge_partition_e_mask && (local_degree1 > edge_t{0})) {
auto local_offset = edge_partition.local_offset(major_idx);
auto local_offset = edge_partition.local_offset(*major_idx);
local_degree1 =
count_set_bits((*edge_partition_e_mask).value_first(), local_offset, local_degree1);
}
Expand Down Expand Up @@ -699,77 +622,6 @@ struct gatherv_indices_t {
}
};

template <typename GraphViewType, typename VertexPairIterator>
size_t count_invalid_vertex_pairs(raft::handle_t const& handle,
GraphViewType const& graph_view,
VertexPairIterator vertex_pair_first,
VertexPairIterator vertex_pair_last)
{
using vertex_t = typename GraphViewType::vertex_type;

std::vector<vertex_t> h_edge_partition_major_range_firsts(
graph_view.number_of_local_edge_partitions());
std::vector<vertex_t> h_edge_partition_major_range_lasts(
h_edge_partition_major_range_firsts.size());
vertex_t edge_partition_minor_range_first{};
vertex_t edge_partition_minor_range_last{};
if constexpr (GraphViewType::is_multi_gpu) {
for (size_t i = 0; i < graph_view.number_of_local_edge_partitions(); i++) {
if constexpr (GraphViewType::is_storage_transposed) {
h_edge_partition_major_range_firsts[i] = graph_view.local_edge_partition_dst_range_first(i);
h_edge_partition_major_range_lasts[i] = graph_view.local_edge_partition_dst_range_last(i);
} else {
h_edge_partition_major_range_firsts[i] = graph_view.local_edge_partition_src_range_first(i);
h_edge_partition_major_range_lasts[i] = graph_view.local_edge_partition_src_range_last(i);
}
}
if constexpr (GraphViewType::is_storage_transposed) {
edge_partition_minor_range_first = graph_view.local_edge_partition_src_range_first();
edge_partition_minor_range_last = graph_view.local_edge_partition_src_range_last();
} else {
edge_partition_minor_range_first = graph_view.local_edge_partition_dst_range_first();
edge_partition_minor_range_last = graph_view.local_edge_partition_dst_range_last();
}
} else {
h_edge_partition_major_range_firsts[0] = vertex_t{0};
h_edge_partition_major_range_lasts[0] = graph_view.number_of_vertices();
edge_partition_minor_range_first = vertex_t{0};
edge_partition_minor_range_last = graph_view.number_of_vertices();
}
rmm::device_uvector<vertex_t> d_edge_partition_major_range_firsts(
h_edge_partition_major_range_firsts.size(), handle.get_stream());
rmm::device_uvector<vertex_t> d_edge_partition_major_range_lasts(
h_edge_partition_major_range_lasts.size(), handle.get_stream());
raft::update_device(d_edge_partition_major_range_firsts.data(),
h_edge_partition_major_range_firsts.data(),
h_edge_partition_major_range_firsts.size(),
handle.get_stream());
raft::update_device(d_edge_partition_major_range_lasts.data(),
h_edge_partition_major_range_lasts.data(),
h_edge_partition_major_range_lasts.size(),
handle.get_stream());

auto num_invalid_pairs = thrust::count_if(
handle.get_thrust_policy(),
vertex_pair_first,
vertex_pair_last,
is_invalid_input_vertex_pair_t<vertex_t>{
graph_view.number_of_vertices(),
raft::device_span<vertex_t const>(d_edge_partition_major_range_firsts.begin(),
d_edge_partition_major_range_firsts.end()),
raft::device_span<vertex_t const>(d_edge_partition_major_range_lasts.begin(),
d_edge_partition_major_range_lasts.end()),
edge_partition_minor_range_first,
edge_partition_minor_range_last});
if constexpr (GraphViewType::is_multi_gpu) {
auto& comm = handle.get_comms();
num_invalid_pairs =
host_scalar_allreduce(comm, num_invalid_pairs, raft::comms::op_t::SUM, handle.get_stream());
}

return num_invalid_pairs;
}

// In multi-GPU, the first element of every vertex pair in [vertex_pair_first, vertex_pair) should
// be within the valid edge partition major range assigned to this process and the second element
// should be within the valid edge partition minor range assigned to this process.
Expand Down
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2023, NVIDIA CORPORATION.
* Copyright (c) 2022-2024, 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 @@ -19,6 +19,7 @@
#include <prims/detail/nbr_intersection.cuh>
#include <prims/property_op_utils.cuh>
#include <utilities/collect_comm.cuh>
#include <utilities/error_check_utils.cuh>

#include <cugraph/edge_partition_device_view.cuh>
#include <cugraph/edge_partition_edge_property_device_view.cuh>
Expand Down
Loading
Loading