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

Jump to bottom

Add view_concat for edge_minor_property_view_t and update transform_reduce_e_by_dst_key to support reduce_op on tuple types #2879

Merged
merged 4 commits into from
Nov 3, 2022
Merged

Add view_concat for edge_minor_property_view_t and update transform_reduce_e_by_dst_key to support reduce_op on tuple types #2879

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
4425c46
Add view_concat for edge_minor_property_view_t and change transform_r…
Nov 2, 2022
a7f2a47
Remove empty comment line
Nov 2, 2022
de81a19
Rename variable
Nov 2, 2022
191258b
Fix formatting
Nov 2, 2022
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
26 changes: 26 additions & 0 deletions cpp/include/cugraph/edge_src_dst_property.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -581,4 +581,30 @@ auto view_concat(detail::edge_major_property_view_t<vertex_t, Ts> const&... view
}
}

template <typename vertex_t, typename... Ts>
auto view_concat(detail::edge_minor_property_view_t<vertex_t, Ts> const&... views)
{
using concat_value_iterator = decltype(
thrust::make_zip_iterator(thrust_tuple_cat(detail::to_thrust_tuple(views.value_first())...)));

concat_value_iterator edge_partition_concat_value_first{};

auto first_view = detail::get_first_of_pack(views...);

edge_partition_concat_value_first =
thrust::make_zip_iterator(thrust_tuple_cat(detail::to_thrust_tuple(views.value_first())...));

if (first_view.key_chunk_size()) {
return detail::edge_minor_property_view_t<vertex_t, concat_value_iterator>(
*(first_view.keys()),
*(first_view.key_chunk_start_offsets()),
*(first_view.key_chunk_size()),
edge_partition_concat_value_first,
first_view.minor_range_first());
} else {
return detail::edge_minor_property_view_t<vertex_t, concat_value_iterator>(
edge_partition_concat_value_first, first_view.minor_range_first());
}
}

} // namespace cugraph
50 changes: 25 additions & 25 deletions cpp/src/prims/transform_reduce_e_by_src_dst_key.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,7 +52,7 @@ template <bool edge_partition_src_key,
typename EdgePartitionDstValueInputWrapper,
typename EdgePartitionSrcDstKeyInputWrapper,
typename EdgeOp,
typename T>
typename ValueIterator>
__device__ void update_buffer_element(
edge_partition_device_view_t<typename GraphViewType::vertex_type,
typename GraphViewType::edge_type,
Expand All @@ -66,7 +66,7 @@ __device__ void update_buffer_element(
EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input,
EdgeOp e_op,
typename GraphViewType::vertex_type* key,
T* value)
ValueIterator value_first)
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

A bit of nitpicking,

but in our naming convention, value_first implies the first element of a multi-element array. But here, value_first is actually a reference to the one buffer element we need to update. So, better rename this to something like just value or value_iter.

Copy link
Contributor Author

@naimnv naimnv Nov 2, 2022 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Renamed.
@seunghwak Yes, it compiles locally.

{
using vertex_t = typename GraphViewType::vertex_type;

Expand All @@ -80,17 +80,17 @@ __device__ void update_buffer_element(
*key = edge_partition_src_dst_key_input.get(
((GraphViewType::is_storage_transposed != edge_partition_src_key) ? major_offset
: minor_offset));
*value = evaluate_edge_op<GraphViewType,
vertex_t,
EdgePartitionSrcValueInputWrapper,
EdgePartitionDstValueInputWrapper,
EdgeOp>()
.compute(src,
dst,
weight,
edge_partition_src_value_input.get(src_offset),
edge_partition_dst_value_input.get(dst_offset),
e_op);
*value_first = evaluate_edge_op<GraphViewType,
vertex_t,
EdgePartitionSrcValueInputWrapper,
EdgePartitionDstValueInputWrapper,
EdgeOp>()
.compute(src,
dst,
weight,
edge_partition_src_value_input.get(src_offset),
edge_partition_dst_value_input.get(dst_offset),
e_op);
}

template <bool edge_partition_src_key,
Expand All @@ -99,7 +99,7 @@ template <bool edge_partition_src_key,
typename EdgePartitionDstValueInputWrapper,
typename EdgePartitionSrcDstKeyInputWrapper,
typename EdgeOp,
typename T>
typename ValueIterator>
__global__ void transform_reduce_by_src_dst_key_hypersparse(
edge_partition_device_view_t<typename GraphViewType::vertex_type,
typename GraphViewType::edge_type,
Expand All @@ -110,7 +110,7 @@ __global__ void transform_reduce_by_src_dst_key_hypersparse(
EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input,
EdgeOp e_op,
typename GraphViewType::vertex_type* keys,
T* values)
ValueIterator value_first)
{
using vertex_t = typename GraphViewType::vertex_type;
using edge_t = typename GraphViewType::edge_type;
Expand Down Expand Up @@ -145,7 +145,7 @@ __global__ void transform_reduce_by_src_dst_key_hypersparse(
edge_partition_src_dst_key_input,
e_op,
keys + local_offset + i,
values + local_offset + i);
value_first + local_offset + i);
}

idx += gridDim.x * blockDim.x;
Expand All @@ -158,7 +158,7 @@ template <bool edge_partition_src_key,
typename EdgePartitionDstValueInputWrapper,
typename EdgePartitionSrcDstKeyInputWrapper,
typename EdgeOp,
typename T>
typename ValueIterator>
__global__ void transform_reduce_by_src_dst_key_low_degree(
edge_partition_device_view_t<typename GraphViewType::vertex_type,
typename GraphViewType::edge_type,
Expand All @@ -171,7 +171,7 @@ __global__ void transform_reduce_by_src_dst_key_low_degree(
EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input,
EdgeOp e_op,
typename GraphViewType::vertex_type* keys,
T* values)
ValueIterator value_first)
{
using vertex_t = typename GraphViewType::vertex_type;
using edge_t = typename GraphViewType::edge_type;
Expand Down Expand Up @@ -203,7 +203,7 @@ __global__ void transform_reduce_by_src_dst_key_low_degree(
edge_partition_src_dst_key_input,
e_op,
keys + local_offset + i,
values + local_offset + i);
value_first + local_offset + i);
}

idx += gridDim.x * blockDim.x;
Expand All @@ -216,7 +216,7 @@ template <bool edge_partition_src_key,
typename EdgePartitionDstValueInputWrapper,
typename EdgePartitionSrcDstKeyInputWrapper,
typename EdgeOp,
typename T>
typename ValueIterator>
__global__ void transform_reduce_by_src_dst_key_mid_degree(
edge_partition_device_view_t<typename GraphViewType::vertex_type,
typename GraphViewType::edge_type,
Expand All @@ -229,7 +229,7 @@ __global__ void transform_reduce_by_src_dst_key_mid_degree(
EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input,
EdgeOp e_op,
typename GraphViewType::vertex_type* keys,
T* values)
ValueIterator value_first)
{
using vertex_t = typename GraphViewType::vertex_type;
using edge_t = typename GraphViewType::edge_type;
Expand Down Expand Up @@ -263,7 +263,7 @@ __global__ void transform_reduce_by_src_dst_key_mid_degree(
edge_partition_src_dst_key_input,
e_op,
keys + local_offset + i,
values + local_offset + i);
value_first + local_offset + i);
}

idx += gridDim.x * (blockDim.x / raft::warp_size());
Expand All @@ -276,7 +276,7 @@ template <bool edge_partition_src_key,
typename EdgePartitionDstValueInputWrapper,
typename EdgePartitionSrcDstKeyInputWrapper,
typename EdgeOp,
typename T>
typename ValueIterator>
__global__ void transform_reduce_by_src_dst_key_high_degree(
edge_partition_device_view_t<typename GraphViewType::vertex_type,
typename GraphViewType::edge_type,
Expand All @@ -289,7 +289,7 @@ __global__ void transform_reduce_by_src_dst_key_high_degree(
EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input,
EdgeOp e_op,
typename GraphViewType::vertex_type* keys,
T* values)
ValueIterator value_first)
{
using vertex_t = typename GraphViewType::vertex_type;
using edge_t = typename GraphViewType::edge_type;
Expand Down Expand Up @@ -320,7 +320,7 @@ __global__ void transform_reduce_by_src_dst_key_high_degree(
edge_partition_src_dst_key_input,
e_op,
keys + local_offset + i,
values + local_offset + i);
value_first + local_offset + i);
}

idx += gridDim.x;
Expand Down