Skip to content

Commit

Permalink
Update per_v_transform_reduce_incoming|outgoing_e to support edge mas…
Browse files Browse the repository at this point in the history
…king (#4085)

per_v_transform_reduce_(incoming|outgoing_e) now supports edge masking.

Authors:
  - Seunghwa Kang (https://github.com/seunghwak)

Approvers:
  - Chuck Hastings (https://github.com/ChuckHastings)
  - Naim (https://github.com/naimnv)
  - Joseph Nke (https://github.com/jnke2016)

URL: #4085
  • Loading branch information
seunghwak authored Jan 25, 2024
1 parent 9a261ff commit ff76a38
Show file tree
Hide file tree
Showing 7 changed files with 381 additions and 191 deletions.
4 changes: 2 additions & 2 deletions cpp/src/centrality/eigenvector_centrality_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 Down Expand Up @@ -117,7 +117,7 @@ rmm::device_uvector<weight_t> eigenvector_centrality(
edge_src_centralities.view(),
edge_dst_dummy_property_t{}.view(),
edge_dummy_property_t{}.view(),
[] __device__(vertex_t, vertex_t, auto src_val, auto, auto) { return src_val * 1.0; },
[] __device__(vertex_t, vertex_t, auto src_val, auto, auto) { return src_val; },
weight_t{0},
reduce_op::plus<weight_t>{},
centralities.begin());
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/link_analysis/pagerank_impl.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 @@ -288,7 +288,7 @@ centrality_algorithm_metadata_t pagerank(
edge_dst_dummy_property_t{}.view(),
edge_dummy_property_t{}.view(),
[alpha] __device__(vertex_t, vertex_t, auto src_val, auto, auto) {
return src_val * 1.0 * alpha;
return src_val * alpha;
},
unvarying_part,
reduce_op::plus<result_t>{},
Expand Down
60 changes: 60 additions & 0 deletions cpp/src/prims/detail/prim_functors.cuh
Original file line number Diff line number Diff line change
@@ -0,0 +1,60 @@
/*
* Copyright (c) 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.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#pragma once

#include <cugraph/edge_partition_device_view.cuh>

namespace cugraph {

namespace detail {

template <typename GraphViewType,
typename EdgePartitionSrcValueInputWrapper,
typename EdgePartitionDstValueInputWrapper,
typename EdgePartitionEdgeValueInputWrapper,
typename EdgeOp>
struct call_e_op_t {
edge_partition_device_view_t<typename GraphViewType::vertex_type,
typename GraphViewType::edge_type,
GraphViewType::is_multi_gpu> const& edge_partition{};
EdgePartitionSrcValueInputWrapper const& edge_partition_src_value_input{};
EdgePartitionDstValueInputWrapper const& edge_partition_dst_value_input{};
EdgePartitionEdgeValueInputWrapper const& edge_partition_e_value_input{};
EdgeOp const& e_op{};
typename GraphViewType::vertex_type major{};
typename GraphViewType::vertex_type major_offset{};
typename GraphViewType::vertex_type const* indices{nullptr};
typename GraphViewType::edge_type edge_offset{};

__device__ auto operator()(typename GraphViewType::edge_type i) const
{
auto minor = indices[i];
auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor);
auto src = GraphViewType::is_storage_transposed ? minor : major;
auto dst = GraphViewType::is_storage_transposed ? major : minor;
auto src_offset = GraphViewType::is_storage_transposed ? minor_offset : major_offset;
auto dst_offset = GraphViewType::is_storage_transposed ? major_offset : minor_offset;
return e_op(src,
dst,
edge_partition_src_value_input.get(src_offset),
edge_partition_dst_value_input.get(dst_offset),
edge_partition_e_value_input.get(edge_offset + i));
}
};

} // namespace detail

} // namespace cugraph
Loading

0 comments on commit ff76a38

Please sign in to comment.