transform_reduce_e_by_src_dst_key.cuh#

namespace cugraph

Functions

template<typename GraphViewType, typename EdgeSrcValueInputWrapper, typename EdgeDstValueInputWrapper, typename EdgeValueInputWrapper, typename EdgeSrcKeyInputWrapper, typename EdgeOp, typename ReduceOp, typename T> auto transform_reduce_e_by_src_key( raft::handle_t const &handle, GraphViewType const &graph_view, EdgeSrcValueInputWrapper edge_src_value_input, EdgeDstValueInputWrapper edge_dst_value_input, EdgeValueInputWrapper edge_value_input, EdgeSrcKeyInputWrapper edge_src_key_input, EdgeOp e_op, T init, ReduceOp reduce_op, bool do_expensive_check = false )

Iterate over the entire set of edges and reduce edge_op outputs to (key, value) pairs.

This function is inspired by thrust::transform_reduce() and thrust::reduce_by_key(). Keys for edges are determined by the edge sources.

Template Parameters:

GraphViewType – Type of the passed non-owning graph object.
EdgeSrcValueInputWrapper – Type of the wrapper for edge source property values.
EdgeDstValueInputWrapper – Type of the wrapper for edge destination property values.
EdgeSrcKeyInputWrapper – Type of the wrapper for edge source key values.
EdgeValueInputWrapper – Type of the wrapper for edge property values.
EdgeOp – Type of the quinary edge operator.
T – Type of the values in (key, value) pairs.

Parameters:

handle – RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and handles to various CUDA libraries) to run graph algorithms.
graph_view – Non-owning graph object.
edge_src_value_input – Wrapper used to access source input property values (for the edge sources assigned to this process in multi-GPU). Use either cugraph::edge_src_property_t::view() (if e_op needs to access source property values) or cugraph::edge_src_dummy_property_t::view() (if e_op does not access source property values). Use update_edge_src_property to fill the wrapper.
edge_dst_value_input – Wrapper used to access destination input property values (for the edge destinations assigned to this process in multi-GPU). Use either cugraph::edge_dst_property_t::view() (if e_op needs to access destination property values) or cugraph::edge_dst_dummy_property_t::view() (if e_op does not access destination property values). Use update_edge_dst_property to fill the wrapper.
edge_src_key_input – Wrapper used to access source input ke values (for the edge sources assigned to this process in multi-GPU). Use cugraph::edge_src_property_t::view(). Use update_edge_src_property to fill the wrapper.
edge_value_input – Wrapper used to access edge input property values (for the edges assigned to this process in multi-GPU). Use either cugraph::edge_property_t::view() (if e_op needs to access edge property values) or cugraph::edge_dummy_property_t::view() (if e_op does not access edge property values).
e_op – Quinary operator takes edge source, edge destination, property values for the source, destination, and edge and returns a value to be reduced to (source key, value) pairs.
init – Initial value to be added to the value in each transform-reduced (source key, value) pair.
reduce_op – Binary operator that takes two input arguments and reduce the two values to one. There are pre-defined reduction operators in src/prims/reduce_op.cuh. It is recommended to use the pre-defined reduction operators whenever possible as the current (and future) implementations of graph primitives may check whether ReduceOp is a known type (or has known member variables) to take a more optimized code path. See the documentation in the reduce_op.cuh file for instructions on writing custom reduction operators.
do_expensive_check – A flag to run expensive checks for input arguments (if set to true).

Returns:

std::tuple Tuple of rmm::device_uvector<typename GraphView::vertex_type> and rmm::device_uvector<T> (if T is arithmetic scalar) or a tuple of rmm::device_uvector objects (if T is a cuda::std::tuple type of arithmetic scalar types, one rmm::device_uvector object per scalar type).

template<typename GraphViewType, typename EdgeSrcValueInputWrapper, typename EdgeDstValueInputWrapper, typename EdgeValueInputWrapper, typename EdgeDstKeyInputWrapper, typename EdgeOp, typename ReduceOp, typename T> auto transform_reduce_e_by_dst_key( raft::handle_t const &handle, GraphViewType const &graph_view, EdgeSrcValueInputWrapper edge_src_value_input, EdgeDstValueInputWrapper edge_dst_value_input, EdgeValueInputWrapper edge_value_input, EdgeDstKeyInputWrapper edge_dst_key_input, EdgeOp e_op, T init, ReduceOp reduce_op, bool do_expensive_check = false )

Iterate over the entire set of edges and reduce edge_op outputs to (key, value) pairs.

This function is inspired by thrust::transform_reduce() and thrust::reduce_by_key(). Keys for edges are determined by the edge destinations.

Template Parameters:

GraphViewType – Type of the passed non-owning graph object.
EdgeSrcValueInputWrapper – Type of the wrapper for edge source property values.
EdgeDstValueInputWrapper – Type of the wrapper for edge destination property values.
EdgeDstKeyInputWrapper – Type of the wrapper for edge destination key values.
EdgeValueInputWrapper – Type of the wrapper for edge property values.
EdgeOp – Type of the quinary edge operator.
T – Type of the values in (key, value) pairs.

Parameters:

handle – RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and handles to various CUDA libraries) to run graph algorithms.
graph_view – Non-owning graph object.
edge_src_value_input – Wrapper used to access source input property values (for the edge sources assigned to this process in multi-GPU). Use either cugraph::edge_src_property_t::view() (if e_op needs to access source property values) or cugraph::edge_src_dummy_property_t::view() (if e_op does not access source property values). Use update_edge_src_property to fill the wrapper.
edge_dst_value_input – Wrapper used to access destination input property values (for the edge destinations assigned to this process in multi-GPU). Use either cugraph::edge_dst_property_t::view() (if e_op needs to access destination property values) or cugraph::edge_dst_dummy_property_t::view() (if e_op does not access destination property values). Use update_edge_dst_property to fill the wrapper.
edge_partition_dst_key_input – Wrapper used to access destination input key values (for the edge destinations assigned to this process in multi-GPU). Use cugraph::edge_dst_property_t::view(). Use update_edge_dst_property to fill the wrapper.
edge_value_input – Wrapper used to access edge input property values (for the edges assigned to this process in multi-GPU). Use either cugraph::edge_property_t::view() (if e_op needs to access edge property values) or cugraph::edge_dummy_property_t::view() (if e_op does not access edge property values).
e_op – Quinary operator takes edge source, edge destination, property values for the source, destination, and edge and returns a value to be reduced to (destination key, value) pairs.
init – Initial value to be added to the value in each transform-reduced (destination key, value) pair.
reduce_op – Binary operator that takes two input arguments and reduce the two values to one. There are pre-defined reduction operators in src/prims/reduce_op.cuh. It is recommended to use the pre-defined reduction operators whenever possible as the current (and future) implementations of graph primitives may check whether ReduceOp is a known type (or has known member variables) to take a more optimized code path. See the documentation in the reduce_op.cuh file for instructions on writing custom reduction operators.
do_expensive_check – A flag to run expensive checks for input arguments (if set to true).

Returns:

namespace detail

Functions

template<bool edge_partition_src_key, typename GraphViewType, typename EdgePartitionSrcValueInputWrapper, typename EdgePartitionDstValueInputWrapper, typename EdgePartitionEdgeValueInputWrapper, typename EdgePartitionSrcDstKeyInputWrapper, typename EdgeOp, typename ValueIterator> __device__ void transform_reduce_e_by_src_dst_key_update_buffer_element (edge_partition_device_view_t< typename GraphViewType::vertex_type, typename GraphViewType::edge_type, GraphViewType::is_multi_gpu > &edge_partition, typename GraphViewType::vertex_type major, typename GraphViewType::vertex_type minor, typename GraphViewType::edge_type edge_offset, EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgePartitionEdgeValueInputWrapper edge_partition_e_value_input, EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input, EdgeOp e_op, typename GraphViewType::vertex_type *key, ValueIterator value)

template<bool edge_partition_src_key, typename GraphViewType, typename EdgePartitionSrcValueInputWrapper, typename EdgePartitionDstValueInputWrapper, typename EdgePartitionEdgeValueInputWrapper, typename EdgePartitionSrcDstKeyInputWrapper, typename EdgePartitionEdgeMaskWrapper, typename EdgeOp, typename ValueIterator> static __global__ void transform_reduce_by_src_dst_key_hypersparse (edge_partition_device_view_t< typename GraphViewType::vertex_type, typename GraphViewType::edge_type, GraphViewType::is_multi_gpu > edge_partition, EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgePartitionEdgeValueInputWrapper edge_partition_e_value_input, EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input, EdgePartitionEdgeMaskWrapper edge_partition_e_mask, cuda::std::optional< raft::device_span< typename GraphViewType::edge_type const > > edge_offsets_with_mask, EdgeOp e_op, typename GraphViewType::vertex_type *keys, ValueIterator value_iter)

template<bool edge_partition_src_key, typename GraphViewType, typename EdgePartitionSrcValueInputWrapper, typename EdgePartitionDstValueInputWrapper, typename EdgePartitionEdgeValueInputWrapper, typename EdgePartitionSrcDstKeyInputWrapper, typename EdgePartitionEdgeMaskWrapper, typename EdgeOp, typename ValueIterator> static __global__ void transform_reduce_by_src_dst_key_low_degree (edge_partition_device_view_t< typename GraphViewType::vertex_type, typename GraphViewType::edge_type, GraphViewType::is_multi_gpu > edge_partition, typename GraphViewType::vertex_type major_range_first, typename GraphViewType::vertex_type major_range_last, EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgePartitionEdgeValueInputWrapper edge_partition_e_value_input, EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input, EdgePartitionEdgeMaskWrapper edge_partition_e_mask, cuda::std::optional< raft::device_span< typename GraphViewType::edge_type const > > edge_offsets_with_mask, EdgeOp e_op, typename GraphViewType::vertex_type *keys, ValueIterator value_iter)

template<bool edge_partition_src_key, typename GraphViewType, typename EdgePartitionSrcValueInputWrapper, typename EdgePartitionDstValueInputWrapper, typename EdgePartitionEdgeValueInputWrapper, typename EdgePartitionSrcDstKeyInputWrapper, typename EdgePartitionEdgeMaskWrapper, typename EdgeOp, typename ValueIterator> static __global__ void transform_reduce_by_src_dst_key_mid_degree (edge_partition_device_view_t< typename GraphViewType::vertex_type, typename GraphViewType::edge_type, GraphViewType::is_multi_gpu > edge_partition, typename GraphViewType::vertex_type major_range_first, typename GraphViewType::vertex_type major_range_last, EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgePartitionEdgeValueInputWrapper edge_partition_e_value_input, EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input, EdgePartitionEdgeMaskWrapper edge_partition_e_mask, cuda::std::optional< raft::device_span< typename GraphViewType::edge_type const > > edge_offsets_with_mask, EdgeOp e_op, typename GraphViewType::vertex_type *keys, ValueIterator value_iter)

template<bool edge_partition_src_key, typename GraphViewType, typename EdgePartitionSrcValueInputWrapper, typename EdgePartitionDstValueInputWrapper, typename EdgePartitionEdgeValueInputWrapper, typename EdgePartitionSrcDstKeyInputWrapper, typename EdgePartitionEdgeMaskWrapper, typename EdgeOp, typename ValueIterator> static __global__ void transform_reduce_by_src_dst_key_high_degree (edge_partition_device_view_t< typename GraphViewType::vertex_type, typename GraphViewType::edge_type, GraphViewType::is_multi_gpu > edge_partition, typename GraphViewType::vertex_type major_range_first, typename GraphViewType::vertex_type major_range_last, EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgePartitionEdgeValueInputWrapper edge_partition_e_value_input, EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input, EdgePartitionEdgeMaskWrapper edge_partition_e_mask, cuda::std::optional< raft::device_span< typename GraphViewType::edge_type const > > edge_offsets_with_mask, EdgeOp e_op, typename GraphViewType::vertex_type *keys, ValueIterator value_iter)

template<typename vertex_t, typename value_t, typename BufferType, typename ReduceOp> std::tuple<rmm::device_uvector<vertex_t>, BufferType> reduce_to_unique_kv_pairs( rmm::device_uvector<vertex_t> &&keys, BufferType &&value_buffer, ReduceOp reduce_op, cudaStream_t stream )

template<bool edge_src_key, typename GraphViewType, typename EdgeSrcValueInputWrapper, typename EdgeDstValueInputWrapper, typename EdgeValueInputWrapper, typename EdgeSrcDstKeyInputWrapper, typename EdgeOp, typename ReduceOp, typename T> std::tuple<rmm::device_uvector<typename GraphViewType::vertex_type>, decltype(allocate_dataframe_buffer<T>(0, cudaStream_t{nullptr}))> transform_reduce_e_by_src_dst_key( raft::handle_t const &handle, GraphViewType const &graph_view, EdgeSrcValueInputWrapper edge_src_value_input, EdgeDstValueInputWrapper edge_dst_value_input, EdgeValueInputWrapper edge_value_input, EdgeSrcDstKeyInputWrapper edge_src_dst_key_input, EdgeOp e_op, T init, ReduceOp reduce_op )

Variables

int32_t constexpr transform_reduce_e_by_src_dst_key_kernel_block_size = 128