Skip to content

Commit

Permalink
added major_idx_from_major_nocheck
Browse files Browse the repository at this point in the history
  • Loading branch information
@seunghwak
seunghwak committed Jan 16, 2024
1 parent dfbc33a commit e6f6784
Show file tree
Hide file tree
Showing 3 changed files with 38 additions and 89 deletions.
19 changes: 19 additions & 0 deletions cpp/include/cugraph/edge_partition_device_view.cuh
View file
Original file line number Diff line number Diff line change
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 @@ -461,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
67 changes: 9 additions & 58 deletions cpp/src/prims/detail/nbr_intersection.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -154,24 +154,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 +312,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 +329,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
41 changes: 10 additions & 31 deletions cpp/src/prims/transform_e.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -426,28 +426,15 @@ void transform_e(raft::handle_t const& handle,
edge_first + edge_partition_offsets[i + 1],
[edge_partition,
edge_partition_e_mask] __device__(thrust::tuple<vertex_t, vertex_t> edge) {
auto major = thrust::get<0>(edge);
auto minor = thrust::get<1>(edge);
vertex_t major_idx{};
auto major_hypersparse_first = edge_partition.major_hypersparse_first();
if (major_hypersparse_first) {
if (major < *major_hypersparse_first) {
major_idx = edge_partition.major_offset_from_major_nocheck(major);
} else {
auto major_hypersparse_idx =
edge_partition.major_hypersparse_idx_from_major_nocheck(major);
if (!major_hypersparse_idx) { return true; }
major_idx =
edge_partition.major_offset_from_major_nocheck(*major_hypersparse_first) +
*major_hypersparse_idx;
}
} else {
major_idx = edge_partition.major_offset_from_major_nocheck(major);
}
auto major = thrust::get<0>(edge);
auto minor = thrust::get<1>(edge);
auto major_idx = edge_partition.major_idx_from_major_nocheck(major);
if (!major_idx) { return true; }
vertex_t const* indices{nullptr};
edge_t edge_offset{};
edge_t local_degree{};
thrust::tie(indices, edge_offset, local_degree) = edge_partition.local_edges(major_idx);
thrust::tie(indices, edge_offset, local_degree) =
edge_partition.local_edges(*major_idx);
auto lower_it =
thrust::lower_bound(thrust::seq, indices, indices + local_degree, minor);
if (*lower_it != minor) { return true; }
Expand Down Expand Up @@ -494,24 +481,16 @@ void transform_e(raft::handle_t const& handle,
auto major = thrust::get<0>(edge);
auto minor = thrust::get<1>(edge);

auto major_hypersparse_first = edge_partition.major_hypersparse_first();
auto major_offset = edge_partition.major_offset_from_major_nocheck(major);
vertex_t major_idx{major_offset};

if ((major_hypersparse_first) && (major >= *major_hypersparse_first)) {
auto major_hypersparse_idx =
edge_partition.major_hypersparse_idx_from_major_nocheck(major);
assert(major_hypersparse_idx);
major_idx = edge_partition.major_offset_from_major_nocheck(*major_hypersparse_first) +
*major_hypersparse_idx;
}
auto major_offset = edge_partition.major_offset_from_major_nocheck(major);
auto major_idx = edge_partition.major_idx_from_major_nocheck(major);
assert(major_idx);

auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor);

vertex_t const* indices{nullptr};
edge_t edge_offset{};
edge_t local_degree{};
thrust::tie(indices, edge_offset, local_degree) = edge_partition.local_edges(major_idx);
thrust::tie(indices, edge_offset, local_degree) = edge_partition.local_edges(*major_idx);
auto lower_it = thrust::lower_bound(thrust::seq, indices, indices + local_degree, minor);
auto upper_it = thrust::upper_bound(thrust::seq, lower_it, indices + local_degree, minor);

Expand Down

0 comments on commit e6f6784

Please sign in to comment.