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

Forward-merge branch-24.02 to branch-24.04 #4114

Merged
merged 1 commit into from
Jan 25, 2024
Merged

Forward-merge branch-24.02 to branch-24.04 #4114

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
View file
Open in desktop
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
View file
Open in desktop
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
View file
Open in desktop
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
View file
Open in desktop
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
3 changes: 2 additions & 1 deletion cpp/src/prims/per_v_pair_transform_dst_nbr_intersection.cuh
View file
Open in desktop
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