From 0092777d1140b0d1af5f96df87ed9ac07a8594de Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Mon, 28 Feb 2022 17:23:50 -0800 Subject: [PATCH 1/9] renamed row_col_properties_t to edge_partition_src_dst_properties --- ...uh => edge_partition_src_dst_property.cuh} | 189 ++++++++++-------- 1 file changed, 103 insertions(+), 86 deletions(-) rename cpp/include/cugraph/prims/{row_col_properties.cuh => edge_partition_src_dst_property.cuh} (71%) diff --git a/cpp/include/cugraph/prims/row_col_properties.cuh b/cpp/include/cugraph/prims/edge_partition_src_dst_property.cuh similarity index 71% rename from cpp/include/cugraph/prims/row_col_properties.cuh rename to cpp/include/cugraph/prims/edge_partition_src_dst_property.cuh index 37620c15e40..713cca77ed7 100644 --- a/cpp/include/cugraph/prims/row_col_properties.cuh +++ b/cpp/include/cugraph/prims/edge_partition_src_dst_property.cuh @@ -37,31 +37,31 @@ namespace cugraph { namespace detail { template -class major_properties_device_view_t { +class edge_partition_major_property_device_view_t { public: using value_type = typename thrust::iterator_traits::value_type; - major_properties_device_view_t() = default; + edge_partition_major_property_device_view_t() = default; - major_properties_device_view_t( + edge_partition_major_property_device_view_t( ValueIterator value_first) // for single-GPU only and for advanced users : value_first_(value_first) { set_local_adj_matrix_partition_idx(size_t{0}); } - major_properties_device_view_t(ValueIterator value_first, - vertex_t const* matrix_partition_major_value_start_offsets) + edge_partition_major_property_device_view_t( + ValueIterator value_first, vertex_t const* matrix_partition_major_value_start_offsets) : value_first_(value_first), matrix_partition_major_value_start_offsets_(matrix_partition_major_value_start_offsets) { set_local_adj_matrix_partition_idx(size_t{0}); } - major_properties_device_view_t(vertex_t const* key_first, - ValueIterator value_first, - vertex_t const* matrix_partition_key_offsets, - vertex_t const* matrix_partition_major_firsts) + edge_partition_major_property_device_view_t(vertex_t const* key_first, + ValueIterator value_first, + vertex_t const* matrix_partition_key_offsets, + vertex_t const* matrix_partition_major_firsts) : key_first_(key_first), value_first_(value_first), matrix_partition_key_offsets_(matrix_partition_key_offsets), @@ -154,18 +154,20 @@ class major_properties_device_view_t { }; template -class minor_properties_device_view_t { +class edge_partition_minor_property_device_view_t { public: using value_type = typename thrust::iterator_traits::value_type; - minor_properties_device_view_t() = default; + edge_partition_minor_property_device_view_t() = default; - minor_properties_device_view_t(ValueIterator value_first) : value_first_(value_first) {} + edge_partition_minor_property_device_view_t(ValueIterator value_first) : value_first_(value_first) + { + } - minor_properties_device_view_t(vertex_t const* key_first, - vertex_t const* key_last, - vertex_t minor_first, - ValueIterator value_first) + edge_partition_minor_property_device_view_t(vertex_t const* key_first, + vertex_t const* key_last, + vertex_t minor_first, + ValueIterator value_first) : key_first_(key_first), key_last_(key_last), minor_first_(minor_first), @@ -197,28 +199,32 @@ class minor_properties_device_view_t { }; template -class major_properties_t { +class edge_partition_major_property_t { public: - major_properties_t() : buffer_(allocate_dataframe_buffer(0, rmm::cuda_stream_view{})) {} + edge_partition_major_property_t() + : buffer_(allocate_dataframe_buffer(0, rmm::cuda_stream_view{})) + { + } - major_properties_t(raft::handle_t const& handle, vertex_t buffer_size) + edge_partition_major_property_t(raft::handle_t const& handle, vertex_t buffer_size) : buffer_(allocate_dataframe_buffer(buffer_size, handle.get_stream())) { } - major_properties_t(raft::handle_t const& handle, - vertex_t buffer_size, - std::vector&& matrix_partition_major_value_start_offsets) + edge_partition_major_property_t( + raft::handle_t const& handle, + vertex_t buffer_size, + std::vector&& matrix_partition_major_value_start_offsets) : buffer_(allocate_dataframe_buffer(buffer_size, handle.get_stream())), matrix_partition_major_value_start_offsets_( std::move(matrix_partition_major_value_start_offsets)) { } - major_properties_t(raft::handle_t const& handle, - vertex_t const* key_first, - std::vector&& matrix_partition_key_offsets, - std::vector&& matrix_partition_major_firsts) + edge_partition_major_property_t(raft::handle_t const& handle, + vertex_t const* key_first, + std::vector&& matrix_partition_key_offsets, + std::vector&& matrix_partition_major_firsts) : key_first_(key_first), buffer_( allocate_dataframe_buffer(matrix_partition_key_offsets.back(), handle.get_stream())), @@ -246,16 +252,17 @@ class major_properties_t { { auto value_first = get_dataframe_buffer_cbegin(buffer_); if (key_first_) { - return major_properties_device_view_t( + return edge_partition_major_property_device_view_t( *key_first_, value_first, (*matrix_partition_key_offsets_).data(), (*matrix_partition_major_firsts_).data()); } else if (matrix_partition_major_value_start_offsets_) { - return major_properties_device_view_t( + return edge_partition_major_property_device_view_t( value_first, (*matrix_partition_major_value_start_offsets_).data()); } else { - return major_properties_device_view_t(value_first); + return edge_partition_major_property_device_view_t( + value_first); } } @@ -263,16 +270,17 @@ class major_properties_t { { auto value_first = get_dataframe_buffer_begin(buffer_); if (key_first_) { - return major_properties_device_view_t( + return edge_partition_major_property_device_view_t( *key_first_, value_first, (*matrix_partition_key_offsets_).data(), (*matrix_partition_major_firsts_).data()); } else if (matrix_partition_major_value_start_offsets_) { - return major_properties_device_view_t( + return edge_partition_major_property_device_view_t( value_first, (*matrix_partition_major_value_start_offsets_).data()); } else { - return major_properties_device_view_t(value_first); + return edge_partition_major_property_device_view_t( + value_first); } } @@ -288,19 +296,22 @@ class major_properties_t { }; template -class minor_properties_t { +class edge_partition_minor_property_t { public: - minor_properties_t() : buffer_(allocate_dataframe_buffer(0, rmm::cuda_stream_view{})) {} + edge_partition_minor_property_t() + : buffer_(allocate_dataframe_buffer(0, rmm::cuda_stream_view{})) + { + } - minor_properties_t(raft::handle_t const& handle, vertex_t buffer_size) + edge_partition_minor_property_t(raft::handle_t const& handle, vertex_t buffer_size) : buffer_(allocate_dataframe_buffer(buffer_size, handle.get_stream())) { } - minor_properties_t(raft::handle_t const& handle, - vertex_t const* key_first, - vertex_t const* key_last, - vertex_t minor_first) + edge_partition_minor_property_t(raft::handle_t const& handle, + vertex_t const* key_first, + vertex_t const* key_last, + vertex_t minor_first) : key_first_(key_first), key_last_(key_last), minor_first_(minor_first), @@ -323,10 +334,11 @@ class minor_properties_t { { auto value_first = get_dataframe_buffer_cbegin(buffer_); if (key_first_) { - return minor_properties_device_view_t( + return edge_partition_minor_property_device_view_t( *key_first_, *key_last_, *minor_first_, value_first); } else { - return minor_properties_device_view_t(value_first); + return edge_partition_minor_property_device_view_t( + value_first); } } @@ -334,10 +346,11 @@ class minor_properties_t { { auto value_first = get_dataframe_buffer_begin(buffer_); if (key_first_) { - return minor_properties_device_view_t( + return edge_partition_minor_property_device_view_t( *key_first_, *key_last_, *minor_first_, value_first); } else { - return minor_properties_device_view_t(value_first); + return edge_partition_minor_property_device_view_t( + value_first); } } @@ -374,15 +387,15 @@ decltype(auto) get_first_of_pack(T&& t, Ts&&...) } // namespace detail template -class row_properties_t { +class edge_partition_src_property_t { public: using value_type = T; static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); - row_properties_t() = default; + edge_partition_src_property_t() = default; - row_properties_t(raft::handle_t const& handle, GraphViewType const& graph_view) + edge_partition_src_property_t(raft::handle_t const& handle, GraphViewType const& graph_view) { using vertex_t = typename GraphViewType::vertex_type; @@ -391,7 +404,7 @@ class row_properties_t { if constexpr (GraphViewType::is_multi_gpu) { if constexpr (GraphViewType::is_adj_matrix_transposed) { auto key_last = graph_view.get_local_sorted_unique_edge_row_end(); - properties_ = detail::minor_properties_t( + property_ = detail::edge_partition_minor_property_t( handle, *key_first, *key_last, graph_view.get_local_adj_matrix_partition_row_first()); } else { std::vector matrix_partition_major_firsts( @@ -400,7 +413,7 @@ class row_properties_t { matrix_partition_major_firsts[i] = graph_view.get_local_adj_matrix_partition_row_first(i); } - properties_ = detail::major_properties_t( + property_ = detail::edge_partition_major_property_t( handle, *key_first, *(graph_view.get_local_sorted_unique_edge_row_offsets()), @@ -411,7 +424,7 @@ class row_properties_t { } } else { if constexpr (GraphViewType::is_adj_matrix_transposed) { - properties_ = detail::minor_properties_t( + property_ = detail::edge_partition_minor_property_t( handle, graph_view.get_number_of_local_adj_matrix_partition_rows()); } else { if constexpr (GraphViewType::is_multi_gpu) { @@ -421,45 +434,46 @@ class row_properties_t { matrix_partition_major_value_start_offsets[i] = graph_view.get_local_adj_matrix_partition_row_value_start_offset(i); } - properties_ = detail::major_properties_t( + property_ = detail::edge_partition_major_property_t( handle, graph_view.get_number_of_local_adj_matrix_partition_rows(), std::move(matrix_partition_major_value_start_offsets)); } else { - properties_ = detail::major_properties_t( + property_ = detail::edge_partition_major_property_t( handle, graph_view.get_number_of_local_adj_matrix_partition_rows()); } } } } - void fill(T value, rmm::cuda_stream_view stream) { properties_.fill(value, stream); } + void fill(T value, rmm::cuda_stream_view stream) { property_.fill(value, stream); } - auto key_first() { return properties_.key_first(); } - auto key_last() { return properties_.key_last(); } + auto key_first() { return property_.key_first(); } + auto key_last() { return property_.key_last(); } - auto value_data() { return properties_.value_data(); } + auto value_data() { return property_.value_data(); } - auto device_view() const { return properties_.device_view(); } - auto mutable_device_view() { return properties_.mutable_device_view(); } + auto device_view() const { return property_.device_view(); } + auto mutable_device_view() { return property_.mutable_device_view(); } private: - std::conditional_t, - detail::major_properties_t> - properties_{}; + std::conditional_t< + GraphViewType::is_adj_matrix_transposed, + detail::edge_partition_minor_property_t, + detail::edge_partition_major_property_t> + property_{}; }; template -class col_properties_t { +class edge_partition_dst_property_t { public: using value_type = T; static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); - col_properties_t() = default; + edge_partition_dst_property_t() = default; - col_properties_t(raft::handle_t const& handle, GraphViewType const& graph_view) + edge_partition_dst_property_t(raft::handle_t const& handle, GraphViewType const& graph_view) { using vertex_t = typename GraphViewType::vertex_type; @@ -473,14 +487,14 @@ class col_properties_t { matrix_partition_major_firsts[i] = graph_view.get_local_adj_matrix_partition_col_first(i); } - properties_ = detail::major_properties_t( + property_ = detail::edge_partition_major_property_t( handle, *key_first, *(graph_view.get_local_sorted_unique_edge_col_offsets()), std::move(matrix_partition_major_firsts)); } else { auto key_last = graph_view.get_local_sorted_unique_edge_col_end(); - properties_ = detail::minor_properties_t( + property_ = detail::edge_partition_minor_property_t( handle, *key_first, *key_last, graph_view.get_local_adj_matrix_partition_col_first()); } } else { @@ -495,40 +509,41 @@ class col_properties_t { matrix_partition_major_value_start_offsets[i] = graph_view.get_local_adj_matrix_partition_col_value_start_offset(i); } - properties_ = detail::major_properties_t( + property_ = detail::edge_partition_major_property_t( handle, graph_view.get_number_of_local_adj_matrix_partition_cols(), std::move(matrix_partition_major_value_start_offsets)); } else { - properties_ = detail::major_properties_t( + property_ = detail::edge_partition_major_property_t( handle, graph_view.get_number_of_local_adj_matrix_partition_cols()); } } else { - properties_ = detail::minor_properties_t( + property_ = detail::edge_partition_minor_property_t( handle, graph_view.get_number_of_local_adj_matrix_partition_cols()); } } } - void fill(T value, rmm::cuda_stream_view stream) { properties_.fill(value, stream); } + void fill(T value, rmm::cuda_stream_view stream) { property_.fill(value, stream); } - auto key_first() { return properties_.key_first(); } - auto key_last() { return properties_.key_last(); } + auto key_first() { return property_.key_first(); } + auto key_last() { return property_.key_last(); } - auto value_data() { return properties_.value_data(); } + auto value_data() { return property_.value_data(); } - auto device_view() const { return properties_.device_view(); } - auto mutable_device_view() { return properties_.mutable_device_view(); } + auto device_view() const { return property_.device_view(); } + auto mutable_device_view() { return property_.mutable_device_view(); } private: - std::conditional_t, - detail::minor_properties_t> - properties_{}; + std::conditional_t< + GraphViewType::is_adj_matrix_transposed, + detail::edge_partition_major_property_t, + detail::edge_partition_minor_property_t> + property_{}; }; template -class dummy_properties_device_view_t { +class dummy_property_device_view_t { public: using value_type = thrust::nullopt_t; @@ -538,30 +553,32 @@ class dummy_properties_device_view_t { }; template -class dummy_properties_t { +class dummy_property_t { public: using value_type = thrust::nullopt_t; - auto device_view() const { return dummy_properties_device_view_t{}; } + auto device_view() const { return dummy_property_device_view_t{}; } }; template -auto device_view_concat(detail::major_properties_device_view_t const&... device_views) +auto device_view_concat( + detail::edge_partition_major_property_device_view_t const&... device_views) { auto concat_first = thrust::make_zip_iterator( thrust_tuple_cat(detail::to_thrust_tuple(device_views.value_data())...)); auto first = detail::get_first_of_pack(device_views...); if (first.key_data()) { - return detail::major_properties_device_view_t( + return detail::edge_partition_major_property_device_view_t( *(first.key_data()), concat_first, *(first.matrix_partition_key_offsets()), *(first.matrix_partition_major_firsts())); } else if (first.matrix_partition_major_value_start_offsets()) { - return detail::major_properties_device_view_t( + return detail::edge_partition_major_property_device_view_t( concat_first, *(first.matrix_partition_major_value_start_offsets())); } else { - return detail::major_properties_device_view_t(concat_first); + return detail::edge_partition_major_property_device_view_t( + concat_first); } } From d06aa484f2356d7db976b81ad1bfc7915ae9c2fe Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Tue, 1 Mar 2022 17:49:05 -0800 Subject: [PATCH 2/9] primitives API improvements (rename copy_to_adj_matrix_row_col to update_edge_partition_src_dst_property) --- .../prims/copy_to_adj_matrix_row_col.cuh | 369 +++++++++--------- 1 file changed, 192 insertions(+), 177 deletions(-) diff --git a/cpp/include/cugraph/prims/copy_to_adj_matrix_row_col.cuh b/cpp/include/cugraph/prims/copy_to_adj_matrix_row_col.cuh index 508294c9e89..b1819f661e5 100644 --- a/cpp/include/cugraph/prims/copy_to_adj_matrix_row_col.cuh +++ b/cpp/include/cugraph/prims/copy_to_adj_matrix_row_col.cuh @@ -45,12 +45,13 @@ namespace cugraph { namespace detail { template -void copy_to_matrix_major(raft::handle_t const& handle, - GraphViewType const& graph_view, - VertexValueInputIterator vertex_value_input_first, - MatrixMajorValueOutputWrapper& matrix_major_value_output) + typename VertexPropertyInputIterator, + typename EdgePartitionMajorPropertyOutputWrapper> +void update_edge_partition_major_property( + raft::handle_t const& handle, + GraphViewType const& graph_view, + VertexPropertyInputIterator vertex_property_input_first, + EdgePartitionMajorPropertyOutputWrapper& edge_partition_major_property_output) { if constexpr (GraphViewType::is_multi_gpu) { using vertex_t = typename GraphViewType::vertex_type; @@ -64,7 +65,7 @@ void copy_to_matrix_major(raft::handle_t const& handle, auto const col_comm_rank = col_comm.get_rank(); auto const col_comm_size = col_comm.get_size(); - if (matrix_major_value_output.key_first()) { + if (edge_partition_major_property_output.key_first()) { auto key_offsets = GraphViewType::is_adj_matrix_transposed ? *(graph_view.get_local_sorted_unique_edge_col_offsets()) : *(graph_view.get_local_sorted_unique_edge_row_offsets()); @@ -75,26 +76,26 @@ void copy_to_matrix_major(raft::handle_t const& handle, max_rx_size, graph_view.get_vertex_partition_size(i * row_comm_size + row_comm_rank)); } auto rx_value_buffer = allocate_dataframe_buffer< - typename std::iterator_traits::value_type>(max_rx_size, - handle.get_stream()); + typename std::iterator_traits::value_type>( + max_rx_size, handle.get_stream()); auto rx_value_first = get_dataframe_buffer_begin(rx_value_buffer); for (int i = 0; i < col_comm_size; ++i) { device_bcast(col_comm, - vertex_value_input_first, + vertex_property_input_first, rx_value_first, graph_view.get_vertex_partition_size(i * row_comm_size + row_comm_rank), i, handle.get_stream()); auto v_offset_first = thrust::make_transform_iterator( - *(matrix_major_value_output.key_first()) + key_offsets[i], + *(edge_partition_major_property_output.key_first()) + key_offsets[i], [v_first = graph_view.get_vertex_partition_first( i * row_comm_size + row_comm_rank)] __device__(auto v) { return v - v_first; }); thrust::gather(handle.get_thrust_policy(), v_offset_first, v_offset_first + (key_offsets[i + 1] - key_offsets[i]), rx_value_first, - matrix_major_value_output.value_data() + key_offsets[i]); + edge_partition_major_property_output.value_data() + key_offsets[i]); } } else { std::vector rx_counts(col_comm_size, size_t{0}); @@ -104,34 +105,35 @@ void copy_to_matrix_major(raft::handle_t const& handle, displacements[i] = (i == 0) ? 0 : displacements[i - 1] + rx_counts[i - 1]; } device_allgatherv(col_comm, - vertex_value_input_first, - matrix_major_value_output.value_data(), + vertex_property_input_first, + edge_partition_major_property_output.value_data(), rx_counts, displacements, handle.get_stream()); } } else { - assert(!(matrix_major_value_output.key_first())); + assert(!(edge_partition_major_property_output.key_first())); assert(graph_view.get_number_of_local_vertices() == GraphViewType::is_adj_matrix_transposed ? graph_view.get_number_of_local_adj_matrix_partition_cols() : graph_view.get_number_of_local_adj_matrix_partition_rows()); thrust::copy(handle.get_thrust_policy(), - vertex_value_input_first, - vertex_value_input_first + graph_view.get_number_of_local_vertices(), - matrix_major_value_output.value_data()); + vertex_property_input_first, + vertex_property_input_first + graph_view.get_number_of_local_vertices(), + edge_partition_major_property_output.value_data()); } } template -void copy_to_matrix_major(raft::handle_t const& handle, - GraphViewType const& graph_view, - VertexIterator vertex_first, - VertexIterator vertex_last, - VertexValueInputIterator vertex_value_input_first, - MatrixMajorValueOutputWrapper& matrix_major_value_output) + typename VertexPropertyInputIterator, + typename EdgePartitionMajorPropertyOutputWrapper> +void update_edge_partition_major_property( + raft::handle_t const& handle, + GraphViewType const& graph_view, + VertexIterator vertex_first, + VertexIterator vertex_last, + VertexPropertyInputIterator vertex_property_input_first, + EdgePartitionMajorPropertyOutputWrapper& edge_partition_major_property_output) { using vertex_t = typename GraphViewType::vertex_type; using edge_t = typename GraphViewType::edge_type; @@ -157,8 +159,8 @@ void copy_to_matrix_major(raft::handle_t const& handle, }); rmm::device_uvector rx_vertices(max_rx_size, handle.get_stream()); auto rx_tmp_buffer = allocate_dataframe_buffer< - typename std::iterator_traits::value_type>(max_rx_size, - handle.get_stream()); + typename std::iterator_traits::value_type>(max_rx_size, + handle.get_stream()); auto rx_value_first = get_dataframe_buffer_begin(rx_tmp_buffer); auto key_offsets = GraphViewType::is_adj_matrix_transposed @@ -183,7 +185,7 @@ void copy_to_matrix_major(raft::handle_t const& handle, thrust::gather(handle.get_thrust_policy(), map_first, map_first + thrust::distance(vertex_first, vertex_last), - vertex_value_input_first, + vertex_property_input_first, rx_value_first); } @@ -193,7 +195,7 @@ void copy_to_matrix_major(raft::handle_t const& handle, col_comm, vertex_first, rx_vertices.begin(), rx_counts[i], i, handle.get_stream()); device_bcast(col_comm, rx_value_first, rx_value_first, rx_counts[i], i, handle.get_stream()); - if (matrix_major_value_output.key_first()) { + if (edge_partition_major_property_output.key_first()) { thrust::for_each( handle.get_thrust_policy(), thrust::make_counting_iterator(vertex_t{0}), @@ -201,9 +203,10 @@ void copy_to_matrix_major(raft::handle_t const& handle, [rx_vertex_first = rx_vertices.begin(), rx_vertex_last = rx_vertices.end(), rx_value_first, - output_key_first = *(matrix_major_value_output.key_first()) + (*key_offsets)[i], - output_value_first = - matrix_major_value_output.value_data() + (*key_offsets)[i]] __device__(auto i) { + output_key_first = + *(edge_partition_major_property_output.key_first()) + (*key_offsets)[i], + output_value_first = edge_partition_major_property_output.value_data() + + (*key_offsets)[i]] __device__(auto i) { auto major = *(output_key_first + i); auto it = thrust::lower_bound(thrust::seq, rx_vertex_first, rx_vertex_last, major); if ((it != rx_vertex_last) && (*it == major)) { @@ -218,35 +221,36 @@ void copy_to_matrix_major(raft::handle_t const& handle, }); // FIXME: this scatter is unnecessary if NCCL directly takes a permutation iterator (and // directly scatters from the internal buffer) - thrust::scatter( - handle.get_thrust_policy(), - rx_value_first, - rx_value_first + rx_counts[i], - map_first, - matrix_major_value_output.value_data() + matrix_partition.get_major_value_start_offset()); + thrust::scatter(handle.get_thrust_policy(), + rx_value_first, + rx_value_first + rx_counts[i], + map_first, + edge_partition_major_property_output.value_data() + + matrix_partition.get_major_value_start_offset()); } } } else { - assert(!(matrix_major_value_output.key_first())); + assert(!(edge_partition_major_property_output.key_first())); assert(graph_view.get_number_of_local_vertices() == GraphViewType::is_adj_matrix_transposed ? graph_view.get_number_of_local_adj_matrix_partition_cols() : graph_view.get_number_of_local_adj_matrix_partition_rows()); - auto val_first = thrust::make_permutation_iterator(vertex_value_input_first, vertex_first); + auto val_first = thrust::make_permutation_iterator(vertex_property_input_first, vertex_first); thrust::scatter(handle.get_thrust_policy(), val_first, val_first + thrust::distance(vertex_first, vertex_last), vertex_first, - matrix_major_value_output.value_data()); + edge_partition_major_property_output.value_data()); } } template -void copy_to_matrix_minor(raft::handle_t const& handle, - GraphViewType const& graph_view, - VertexValueInputIterator vertex_value_input_first, - MatrixMinorValueOutputWrapper& matrix_minor_value_output) + typename VertexPropertyInputIterator, + typename EdgePartitionMinorPropertyOutputWrapper> +void update_edge_partition_minor_property( + raft::handle_t const& handle, + GraphViewType const& graph_view, + VertexPropertyInputIterator vertex_property_input_first, + EdgePartitionMinorPropertyOutputWrapper& edge_partition_minor_property_output) { if constexpr (GraphViewType::is_multi_gpu) { using vertex_t = typename GraphViewType::vertex_type; @@ -260,7 +264,7 @@ void copy_to_matrix_minor(raft::handle_t const& handle, auto const col_comm_rank = col_comm.get_rank(); auto const col_comm_size = col_comm.get_size(); - if (matrix_minor_value_output.key_first()) { + if (edge_partition_minor_property_output.key_first()) { auto key_offsets = GraphViewType::is_adj_matrix_transposed ? *(graph_view.get_local_sorted_unique_edge_row_offsets()) : *(graph_view.get_local_sorted_unique_edge_col_offsets()); @@ -271,26 +275,26 @@ void copy_to_matrix_minor(raft::handle_t const& handle, max_rx_size, graph_view.get_vertex_partition_size(col_comm_rank * row_comm_size + i)); } auto rx_value_buffer = allocate_dataframe_buffer< - typename std::iterator_traits::value_type>(max_rx_size, - handle.get_stream()); + typename std::iterator_traits::value_type>( + max_rx_size, handle.get_stream()); auto rx_value_first = get_dataframe_buffer_begin(rx_value_buffer); for (int i = 0; i < row_comm_size; ++i) { device_bcast(row_comm, - vertex_value_input_first, + vertex_property_input_first, rx_value_first, graph_view.get_vertex_partition_size(col_comm_rank * row_comm_size + i), i, handle.get_stream()); auto v_offset_first = thrust::make_transform_iterator( - *(matrix_minor_value_output.key_first()) + key_offsets[i], + *(edge_partition_minor_property_output.key_first()) + key_offsets[i], [v_first = graph_view.get_vertex_partition_first( col_comm_rank * row_comm_size + i)] __device__(auto v) { return v - v_first; }); thrust::gather(handle.get_thrust_policy(), v_offset_first, v_offset_first + (key_offsets[i + 1] - key_offsets[i]), rx_value_first, - matrix_minor_value_output.value_data() + key_offsets[i]); + edge_partition_minor_property_output.value_data() + key_offsets[i]); } } else { std::vector rx_counts(row_comm_size, size_t{0}); @@ -300,34 +304,35 @@ void copy_to_matrix_minor(raft::handle_t const& handle, displacements[i] = (i == 0) ? 0 : displacements[i - 1] + rx_counts[i - 1]; } device_allgatherv(row_comm, - vertex_value_input_first, - matrix_minor_value_output.value_data(), + vertex_property_input_first, + edge_partition_minor_property_output.value_data(), rx_counts, displacements, handle.get_stream()); } } else { - assert(!(matrix_minor_value_output.key_first())); + assert(!(edge_partition_minor_property_output.key_first())); assert(graph_view.get_number_of_local_vertices() == GraphViewType::is_adj_matrix_transposed ? graph_view.get_number_of_local_adj_matrix_partition_rows() : graph_view.get_number_of_local_adj_matrix_partition_cols()); thrust::copy(handle.get_thrust_policy(), - vertex_value_input_first, - vertex_value_input_first + graph_view.get_number_of_local_vertices(), - matrix_minor_value_output.value_data()); + vertex_property_input_first, + vertex_property_input_first + graph_view.get_number_of_local_vertices(), + edge_partition_minor_property_output.value_data()); } } template -void copy_to_matrix_minor(raft::handle_t const& handle, - GraphViewType const& graph_view, - VertexIterator vertex_first, - VertexIterator vertex_last, - VertexValueInputIterator vertex_value_input_first, - MatrixMinorValueOutputWrapper& matrix_minor_value_output) + typename VertexPropertyInputIterator, + typename EdgePartitionMinorPropertyOutputWrapper> +void update_edge_partition_minor_property( + raft::handle_t const& handle, + GraphViewType const& graph_view, + VertexIterator vertex_first, + VertexIterator vertex_last, + VertexPropertyInputIterator vertex_property_input_first, + EdgePartitionMinorPropertyOutputWrapper& edge_partition_minor_property_output) { using vertex_t = typename GraphViewType::vertex_type; using edge_t = typename GraphViewType::edge_type; @@ -353,8 +358,8 @@ void copy_to_matrix_minor(raft::handle_t const& handle, }); rmm::device_uvector rx_vertices(max_rx_size, handle.get_stream()); auto rx_tmp_buffer = allocate_dataframe_buffer< - typename std::iterator_traits::value_type>(max_rx_size, - handle.get_stream()); + typename std::iterator_traits::value_type>(max_rx_size, + handle.get_stream()); auto rx_value_first = get_dataframe_buffer_begin(rx_tmp_buffer); auto key_offsets = GraphViewType::is_adj_matrix_transposed @@ -378,7 +383,7 @@ void copy_to_matrix_minor(raft::handle_t const& handle, thrust::gather(handle.get_thrust_policy(), map_first, map_first + thrust::distance(vertex_first, vertex_last), - vertex_value_input_first, + vertex_property_input_first, rx_value_first); } @@ -388,7 +393,7 @@ void copy_to_matrix_minor(raft::handle_t const& handle, row_comm, vertex_first, rx_vertices.begin(), rx_counts[i], i, handle.get_stream()); device_bcast(row_comm, rx_value_first, rx_value_first, rx_counts[i], i, handle.get_stream()); - if (matrix_minor_value_output.key_first()) { + if (edge_partition_minor_property_output.key_first()) { thrust::for_each( handle.get_thrust_policy(), thrust::make_counting_iterator(vertex_t{0}), @@ -396,9 +401,10 @@ void copy_to_matrix_minor(raft::handle_t const& handle, [rx_vertex_first = rx_vertices.begin(), rx_vertex_last = rx_vertices.end(), rx_value_first, - output_key_first = *(matrix_minor_value_output.key_first()) + (*key_offsets)[i], - output_value_first = - matrix_minor_value_output.value_data() + (*key_offsets)[i]] __device__(auto i) { + output_key_first = + *(edge_partition_minor_property_output.key_first()) + (*key_offsets)[i], + output_value_first = edge_partition_minor_property_output.value_data() + + (*key_offsets)[i]] __device__(auto i) { auto minor = *(output_key_first + i); auto it = thrust::lower_bound(thrust::seq, rx_vertex_first, rx_vertex_last, minor); if ((it != rx_vertex_last) && (*it == minor)) { @@ -417,191 +423,200 @@ void copy_to_matrix_minor(raft::handle_t const& handle, rx_value_first, rx_value_first + rx_counts[i], map_first, - matrix_minor_value_output.value_data()); + edge_partition_minor_property_output.value_data()); } } } else { - assert(!(matrix_minor_value_output.key_first())); + assert(!(edge_partition_minor_property_output.key_first())); assert(graph_view.get_number_of_local_vertices() == graph_view.get_number_of_local_adj_matrix_partition_rows()); - auto val_first = thrust::make_permutation_iterator(vertex_value_input_first, vertex_first); + auto val_first = thrust::make_permutation_iterator(vertex_property_input_first, vertex_first); thrust::scatter(handle.get_thrust_policy(), val_first, val_first + thrust::distance(vertex_first, vertex_last), vertex_first, - matrix_minor_value_output.value_data()); + edge_partition_minor_property_output.value_data()); } } } // namespace detail /** - * @brief Copy vertex property values to the corresponding graph adjacency matrix row property - * variables. + * @brief Update graph edge partition source property values from the input vertex property values. * - * This version fills the entire set of graph adjacency matrix row property values. + * This version updates graph edge partition property values for the entire edge partition source + * ranges (assigned to this process in multi-GPU). * * @tparam GraphViewType Type of the passed non-owning graph object. - * @tparam VertexValueInputIterator Type of the iterator for vertex properties. + * @tparam VertexPropertyInputIterator Type of the iterator for vertex property values. * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Non-owning graph object. - * @param vertex_value_input_first Iterator pointing to the vertex properties for the first - * (inclusive) vertex (assigned to this process in multi-GPU). `vertex_value_input_last` (exclusive) - * is deduced as @p vertex_value_input_first + @p graph_view.get_number_of_local_vertices(). - * @param adj_matrix_row_value_output Wrapper used to access data storage to copy row properties - * (for the rows assigned to this process in multi-GPU). + * @param vertex_property_input_first Iterator pointing to the vertex property value for the first + * (inclusive) vertex (of the vertex partition assigned to this process in multi-GPU). + * `vertex_property_input_last` (exclusive) is deduced as @p vertex_property_input_first + @p + * graph_view.get_number_of_local_vertices(). + * @param edge_partition_source_property_output Wrapper used to store edge partition source property + * values (for the edge partitions assigned to this process in multi-GPU). */ -template -void copy_to_adj_matrix_row( +template +void update_edge_partition_src_property( raft::handle_t const& handle, GraphViewType const& graph_view, - VertexValueInputIterator vertex_value_input_first, - row_properties_t::value_type>& - adj_matrix_row_value_output) + VertexPropertyInputIterator vertex_property_input_first, + edge_partition_src_property_t< + GraphViewType, + typename std::iterator_traits::value_type>& + edge_partition_src_property_output) { if constexpr (GraphViewType::is_adj_matrix_transposed) { - copy_to_matrix_minor(handle, graph_view, vertex_value_input_first, adj_matrix_row_value_output); + update_edge_partition_minor_property( + handle, graph_view, vertex_property_input_first, edge_partition_src_property_output); } else { - copy_to_matrix_major(handle, graph_view, vertex_value_input_first, adj_matrix_row_value_output); + update_edge_partition_major_property( + handle, graph_view, vertex_property_input_first, edge_partition_src_property_output); } } /** - * @brief Copy vertex property values to the corresponding graph adjacency matrix row property - * variables. + * @brief Update graph edge partition source property values from the input vertex property values. * - * This version fills only a subset of graph adjacency matrix row property values. [@p vertex_first, - * @p vertex_last) specifies the vertices with new values to be copied to graph adjacency matrix row - * property variables. + * This version updates only a subset of graph edge partition source property values. [@p + * vertex_first, @p vertex_last) specifies the vertices with new property values to be updated. * * @tparam GraphViewType Type of the passed non-owning graph object. * @tparam VertexIterator Type of the iterator for vertex identifiers. - * @tparam VertexValueInputIterator Type of the iterator for vertex properties. + * @tparam VertexPropertyInputIterator Type of the iterator for vertex property values. * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Non-owning graph object. - * @param vertex_first Iterator pointing to the first (inclusive) vertex with new values to be - * copied. v in [vertex_first, vertex_last) should be distinct (and should belong to this process in - * multi-GPU), otherwise undefined behavior - * @param vertex_last Iterator pointing to the last (exclusive) vertex with new values to be copied. - * @param vertex_value_input_first Iterator pointing to the vertex properties for the first - * (inclusive) vertex (assigned to this process in multi-GPU). `vertex_value_input_last` (exclusive) - * is deduced as @p vertex_value_input_first + @p graph_view.get_number_of_local_vertices(). - * @param adj_matrix_row_value_output Wrapper used to access data storage to copy row properties - * (for the rows assigned to this process in multi-GPU). + * @param vertex_first Iterator pointing to the first (inclusive) vertex with a new value to be + * updated. v in [vertex_first, vertex_last) should be distinct (and should belong to the vertex + * partition assigned to this process in multi-GPU), otherwise undefined behavior. + * @param vertex_last Iterator pointing to the last (exclusive) vertex with a new value. + * @param vertex_property_input_first Iterator pointing to the vertex property value for the first + * (inclusive) vertex (of the vertex partition assigned to this process in multi-GPU). + * `vertex_property_input_last` (exclusive) is deduced as @p vertex_property_input_first + @p + * graph_view.get_number_of_local_vertices(). + * @param edge_partition_source_property_output Wrapper used to store edge partition source property + * values (for the edge partitions assigned to this process in multi-GPU). */ -template -void copy_to_adj_matrix_row( +template +void update_edge_partition_src_property( raft::handle_t const& handle, GraphViewType const& graph_view, VertexIterator vertex_first, VertexIterator vertex_last, - VertexValueInputIterator vertex_value_input_first, - row_properties_t::value_type>& - adj_matrix_row_value_output) + VertexPropertyInputIterator vertex_property_input_first, + edge_partition_src_property_t< + GraphViewType, + typename std::iterator_traits::value_type>& + edge_partition_src_property_output) { if constexpr (GraphViewType::is_adj_matrix_transposed) { - copy_to_matrix_minor(handle, - graph_view, - vertex_first, - vertex_last, - vertex_value_input_first, - adj_matrix_row_value_output); + detail::update_edge_partition_minor_property(handle, + graph_view, + vertex_first, + vertex_last, + vertex_property_input_first, + edge_partition_src_property_output); } else { - copy_to_matrix_major(handle, - graph_view, - vertex_first, - vertex_last, - vertex_value_input_first, - adj_matrix_row_value_output); + detail::update_edge_partition_major_property(handle, + graph_view, + vertex_first, + vertex_last, + vertex_property_input_first, + edge_partition_src_property_output); } } /** - * @brief Copy vertex property values to the corresponding graph adjacency matrix column property - * variables. + * @brief Update graph edge partition destination property values from the input vertex property + * values. * - * This version fills the entire set of graph adjacency matrix column property values. + * This version updates graph edge partition property values for the entire edge partition + * destination ranges (assigned to this process in multi-GPU). * * @tparam GraphViewType Type of the passed non-owning graph object. - * @tparam VertexValueInputIterator Type of the iterator for vertex properties. + * @tparam VertexPropertyInputIterator Type of the iterator for vertex property values. * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Non-owning graph object. - * @param vertex_value_input_first Iterator pointing to the vertex properties for the first - * (inclusive) vertex (assigned to this process in multi-GPU). `vertex_value_input_last` (exclusive) - * is deduced as @p vertex_value_input_first + @p graph_view.get_number_of_local_vertices(). - * @param adj_matrix_col_value_output Wrapper used to access data storage to copy column properties - * (for the columns assigned to this process in multi-GPU). + * @param vertex_property_input_first Iterator pointing to the vertex property value for the first + * (inclusive) vertex (of the vertex partition assigned to this process in multi-GPU). + * `vertex_property_input_last` (exclusive) is deduced as @p vertex_property_input_first + @p + * graph_view.get_number_of_local_vertices(). + * @param edge_partition_dst_property_output Wrapper used to store edge partition source property + * values (for the edge partitions assigned to this process in multi-GPU). */ -template -void copy_to_adj_matrix_col( +template +void update_edge_partition_dst_property( raft::handle_t const& handle, GraphViewType const& graph_view, - VertexValueInputIterator vertex_value_input_first, + VertexPropertyInputIterator vertex_property_input_first, col_properties_t::value_type>& - adj_matrix_col_value_output) + typename std::iterator_traits::value_type>& + edge_partition_dst_property_output) { if constexpr (GraphViewType::is_adj_matrix_transposed) { - copy_to_matrix_major(handle, graph_view, vertex_value_input_first, adj_matrix_col_value_output); + detail::update_edge_partition_major_property( + handle, graph_view, vertex_property_input_first, edge_partition_dst_property_output); } else { - copy_to_matrix_minor(handle, graph_view, vertex_value_input_first, adj_matrix_col_value_output); + detail::update_edge_partition_minor_property( + handle, graph_view, vertex_property_input_first, edge_partition_dst_property_output); } } /** - * @brief Copy vertex property values to the corresponding graph adjacency matrix column property - * variables. + * @brief Update graph edge partition destination property values from the input vertex property + * values. * - * This version fills only a subset of graph adjacency matrix column property values. [@p - * vertex_first, @p vertex_last) specifies the vertices with new values to be copied to graph - * adjacency matrix column property variables. + * This version updates only a subset of graph edge partition destination property values. [@p + * vertex_first, @p vertex_last) specifies the vertices with new property values to be updated. * * @tparam GraphViewType Type of the passed non-owning graph object. * @tparam VertexIterator Type of the iterator for vertex identifiers. - * @tparam VertexValueInputIterator Type of the iterator for vertex properties. + * @tparam VertexPropertyInputIterator Type of the iterator for vertex property values. * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Non-owning graph object. - * @param vertex_first Iterator pointing to the first (inclusive) vertex with new values to be - * copied. v in [vertex_first, vertex_last) should be distinct (and should belong to this process in - * multi-GPU), otherwise undefined behavior - * @param vertex_last Iterator pointing to the last (exclusive) vertex with new values to be copied. - * @param vertex_value_input_first Iterator pointing to the vertex properties for the first - * (inclusive) vertex (assigned to this process in multi-GPU). `vertex_value_input_last` (exclusive) - * is deduced as @p vertex_value_input_first + @p graph_view.get_number_of_local_vertices(). - * @param adj_matrix_col_value_output Wrapper used to access data storage to copy column properties - * (for the columns assigned to this process in multi-GPU). + * @param vertex_first Iterator pointing to the first (inclusive) vertex with a new value to be + * updated. v in [vertex_first, vertex_last) should be distinct (and should belong to the vertex + * partition assigned to this process in multi-GPU), otherwise undefined behavior. + * @param vertex_last Iterator pointing to the last (exclusive) vertex with a new value. + * @param vertex_property_input_first Iterator pointing to the vertex property value for the first + * (inclusive) vertex (of the vertex partition assigned to this process in multi-GPU). + * `vertex_property_input_last` (exclusive) is deduced as @p vertex_property_input_first + @p + * graph_view.get_number_of_local_vertices(). + * @param edge_partition_dst_property_output Wrapper used to store edge partition source property + * values (for the edge partitions assigned to this process in multi-GPU). (for the columns assigned + * to this process in multi-GPU). */ -template -void copy_to_adj_matrix_col( +template +void update_edge_partition_dst( raft::handle_t const& handle, GraphViewType const& graph_view, VertexIterator vertex_first, VertexIterator vertex_last, - VertexValueInputIterator vertex_value_input_first, + VertexPropertyInputIterator vertex_property_input_first, col_properties_t::value_type>& - adj_matrix_col_value_output) + typename std::iterator_traits::value_type>& + edge_partition_dst_property_output) { if constexpr (GraphViewType::is_adj_matrix_transposed) { - copy_to_matrix_major(handle, - graph_view, - vertex_first, - vertex_last, - vertex_value_input_first, - adj_matrix_col_value_output); + detail::update_edge_partition_major_property(handle, + graph_view, + vertex_first, + vertex_last, + vertex_property_input_first, + edge_partition_dst_property_output); } else { - copy_to_matrix_minor(handle, - graph_view, - vertex_first, - vertex_last, - vertex_value_input_first, - adj_matrix_col_value_output); + detail::update_edge_partition_minor_property(handle, + graph_view, + vertex_first, + vertex_last, + vertex_property_input_first, + edge_partition_dst_property_output); } } From 939281f1e1c81e949b428aa5ab6ec206ce4ac53b Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Tue, 1 Mar 2022 17:49:37 -0800 Subject: [PATCH 3/9] file renaming --- ...rix_row_col.cuh => update_edge_partition_src_dst_property.cuh} | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename cpp/include/cugraph/prims/{copy_to_adj_matrix_row_col.cuh => update_edge_partition_src_dst_property.cuh} (100%) diff --git a/cpp/include/cugraph/prims/copy_to_adj_matrix_row_col.cuh b/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh similarity index 100% rename from cpp/include/cugraph/prims/copy_to_adj_matrix_row_col.cuh rename to cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh From 2d2f4be92378cabfb5ca69f2dbd5acaef85ef315 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Tue, 8 Mar 2022 13:24:35 -0800 Subject: [PATCH 4/9] fix some compile errors --- .../update_edge_partition_src_dst_property.cuh | 14 ++++++++------ cpp/src/cores/core_number_impl.cuh | 12 ++++++------ cpp/src/link_analysis/hits_impl.cuh | 8 ++++---- 3 files changed, 18 insertions(+), 16 deletions(-) diff --git a/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh b/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh index b1819f661e5..32a65c90a59 100644 --- a/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh +++ b/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh @@ -18,7 +18,7 @@ #include #include #include -#include +#include #include #include #include @@ -554,8 +554,9 @@ void update_edge_partition_dst_property( raft::handle_t const& handle, GraphViewType const& graph_view, VertexPropertyInputIterator vertex_property_input_first, - col_properties_t::value_type>& + edge_partition_dst_property_t< + GraphViewType, + typename std::iterator_traits::value_type>& edge_partition_dst_property_output) { if constexpr (GraphViewType::is_adj_matrix_transposed) { @@ -593,14 +594,15 @@ void update_edge_partition_dst_property( * to this process in multi-GPU). */ template -void update_edge_partition_dst( +void update_edge_partition_dst_property( raft::handle_t const& handle, GraphViewType const& graph_view, VertexIterator vertex_first, VertexIterator vertex_last, VertexPropertyInputIterator vertex_property_input_first, - col_properties_t::value_type>& + edge_partition_dst_property_t< + GraphViewType, + typename std::iterator_traits::value_type>& edge_partition_dst_property_output) { if constexpr (GraphViewType::is_adj_matrix_transposed) { diff --git a/cpp/src/cores/core_number_impl.cuh b/cpp/src/cores/core_number_impl.cuh index 049a238a3cd..ad355eaf835 100644 --- a/cpp/src/cores/core_number_impl.cuh +++ b/cpp/src/cores/core_number_impl.cuh @@ -17,9 +17,9 @@ #include #include -#include +#include #include -#include +#include #include #include #include @@ -145,9 +145,9 @@ void core_number(raft::handle_t const& handle, VertexFrontier(Bucket::num_buckets)> vertex_frontier(handle); - col_properties_t, edge_t> + edge_partition_dst_property_t, edge_t> dst_core_numbers(handle, graph_view); - copy_to_adj_matrix_col(handle, graph_view, core_numbers, dst_core_numbers); + update_edge_partition_dst_property(handle, graph_view, core_numbers, dst_core_numbers); auto k = std::max(k_first, size_t{2}); // degree 0|1 vertices belong to 0|1-core if (graph_view.is_symmetric() && (degree_type == k_core_degree_type_t::INOUT) && @@ -199,7 +199,7 @@ void core_number(raft::handle_t const& handle, vertex_frontier, static_cast(Bucket::cur), std::vector{static_cast(Bucket::next)}, - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), dst_core_numbers.device_view(), [k, delta] __device__(vertex_t src, vertex_t dst, auto, auto dst_val) { return dst_val >= k ? thrust::optional{delta} : thrust::nullopt; @@ -224,7 +224,7 @@ void core_number(raft::handle_t const& handle, CUGRAPH_FAIL("unimplemented."); } - copy_to_adj_matrix_col( + update_edge_partition_dst_property( handle, graph_view, vertex_frontier.get_bucket(static_cast(Bucket::next)).begin(), diff --git a/cpp/src/link_analysis/hits_impl.cuh b/cpp/src/link_analysis/hits_impl.cuh index d9925453be7..6cb62c08531 100644 --- a/cpp/src/link_analysis/hits_impl.cuh +++ b/cpp/src/link_analysis/hits_impl.cuh @@ -17,12 +17,12 @@ #include #include -#include #include #include +#include #include -#include #include +#include #include #include @@ -90,8 +90,8 @@ std::tuple hits(raft::handle_t const& handle, } // Property wrappers - row_properties_t prev_src_hubs(handle, graph_view); - col_properties_t curr_dst_auth(handle, graph_view); + edge_partition_src_property_t prev_src_hubs(handle, graph_view); + edge_partition_dst_property_t curr_dst_auth(handle, graph_view); rmm::device_uvector temp_hubs(graph_view.get_number_of_local_vertices(), handle.get_stream()); From 92350d11a4f876aa83cee02220c32503dae1ab2d Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 10 Mar 2022 17:52:43 -0800 Subject: [PATCH 5/9] fix compile errors due to API changes --- .../copy_v_transform_reduce_in_out_nbr.cuh | 35 ++++---- ...ransform_reduce_key_aggregated_out_nbr.cuh | 9 +- cpp/include/cugraph/prims/count_if_e.cuh | 12 +-- .../prims/edge_partition_src_dst_property.cuh | 47 +++++----- cpp/include/cugraph/prims/extract_if_e.cuh | 14 +-- ...orm_reduce_by_adj_matrix_row_col_key_e.cuh | 32 +++---- .../cugraph/prims/transform_reduce_e.cuh | 13 +-- .../update_frontier_v_push_if_out_nbr.cuh | 12 +-- cpp/src/centrality/katz_centrality_impl.cuh | 12 +-- cpp/src/community/louvain.cuh | 85 ++++++++++--------- .../weakly_connected_components_impl.cuh | 17 ++-- cpp/src/link_analysis/hits_impl.cuh | 10 +-- cpp/src/link_analysis/pagerank_impl.cuh | 16 ++-- cpp/src/structure/coarsen_graph_impl.cuh | 23 ++--- cpp/src/structure/graph_view_impl.cuh | 26 +++--- cpp/src/traversal/bfs_impl.cuh | 35 ++++---- cpp/src/traversal/sssp_impl.cuh | 23 ++--- .../mg_copy_v_transform_reduce_inout_nbr.cu | 14 +-- cpp/tests/prims/mg_count_if_e.cu | 14 +-- cpp/tests/prims/mg_extract_if_e.cu | 52 ++++++------ cpp/tests/prims/mg_transform_reduce_e.cu | 14 +-- .../mg_update_frontier_v_push_if_out_nbr.cu | 52 ++++++------ 22 files changed, 294 insertions(+), 273 deletions(-) diff --git a/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh b/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh index 5fb3b6544f7..c314e7253d8 100644 --- a/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh +++ b/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh @@ -17,9 +17,10 @@ #include #include +#include +#include #include #include -#include #include #include #include @@ -483,14 +484,14 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); [[maybe_unused]] std::conditional_t, - col_properties_t> + edge_partition_src_property_t, + edge_partition_dst_property_t> minor_tmp_buffer(handle); // relevant only when (GraphViewType::is_multi_gpu && !update_major if constexpr (GraphViewType::is_multi_gpu && !update_major) { if constexpr (GraphViewType::is_adj_matrix_transposed) { - minor_tmp_buffer = row_properties_t(handle, graph_view); + minor_tmp_buffer = edge_partition_src_property_t(handle, graph_view); } else { - minor_tmp_buffer = col_properties_t(handle, graph_view); + minor_tmp_buffer = edge_partition_dst_property_t(handle, graph_view); } } @@ -761,14 +762,14 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, * @param graph_view Non-owning graph object. * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::row_properties_t::device_view() (if @p e_op needs to access row properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access row properties). Use - * copy_to_adj_matrix_row to fill the wrapper. + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) + * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use + * update_edge_partition_src_property to fill the wrapper. * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::col_properties_t::device_view() (if @p e_op needs to access column properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access column properties). Use - * copy_to_adj_matrix_col to fill the wrapper. + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column + * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column + * properties). Use update_edge_partition_dst_property to fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) * and returns a value to be reduced. @@ -821,14 +822,14 @@ void copy_v_transform_reduce_in_nbr(raft::handle_t const& handle, * @param graph_view Non-owning graph object. * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::row_properties_t::device_view() (if @p e_op needs to access row properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access row properties). Use - * copy_to_adj_matrix_row to fill the wrapper. + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) + * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use + * update_edge_partition_src_property to fill the wrapper. * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::col_properties_t::device_view() (if @p e_op needs to access column properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access column properties). Use - * copy_to_adj_matrix_col to fill the wrapper. + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column + * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column + * properties). Use update_edge_partition_dst_property to fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) * and returns a value to be reduced. diff --git a/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh b/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh index 46a0a6a91ae..32f1407597d 100644 --- a/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh +++ b/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh @@ -167,12 +167,13 @@ struct reduce_with_init_t { * @param graph_view Non-owning graph object. * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::row_properties_t::device_view() (if @p e_op needs to access row properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access row properties). Use - * copy_to_adj_matrix_row to fill the wrapper. + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) + * or cugraph::dummy_:property_t:device_view() (if @p e_op does not access row properties). Use + * update_edge_partition_src_property to fill the wrapper. * @param adj_matrix_col_key_input Device-copyable wrapper used to access column keys (for the * columns assigned to this process in multi-GPU). Use either - * cugraph::col_properties_t::device_view(). Use copy_to_adj_matrix_col to fill the wrapper. + * cugraph::edge_partition_dst_property_t::device_view(). Use update_edge_partition_dst_property to + * fill the wrapper. * @param map_unique_key_first Iterator pointing to the first (inclusive) key in (key, value) pairs * (assigned to this process in multi-GPU, `cugraph::detail::compute_gpu_id_from_vertex_t` is used * to map keys to processes). (Key, value) pairs may be provided by diff --git a/cpp/include/cugraph/prims/count_if_e.cuh b/cpp/include/cugraph/prims/count_if_e.cuh index a715003e7b9..1cabe0519ec 100644 --- a/cpp/include/cugraph/prims/count_if_e.cuh +++ b/cpp/include/cugraph/prims/count_if_e.cuh @@ -41,14 +41,14 @@ namespace cugraph { * @param graph_view Non-owning graph object. * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::row_properties_t::device_view() (if @p e_op needs to access row properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access row properties). Use - * copy_to_adj_matrix_row to fill the wrapper. + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) + * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use + * update_edge_partition_src_property to fill the wrapper. * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::col_properties_t::device_view() (if @p e_op needs to access column properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access column properties). Use - * copy_to_adj_matrix_col to fill the wrapper. + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column + * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column + * properties). Use update_edge_partition_dst_property to fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) * and returns true if this edge should be included in the returned count. diff --git a/cpp/include/cugraph/prims/edge_partition_src_dst_property.cuh b/cpp/include/cugraph/prims/edge_partition_src_dst_property.cuh index b011aabe85f..2648a7ec61b 100644 --- a/cpp/include/cugraph/prims/edge_partition_src_dst_property.cuh +++ b/cpp/include/cugraph/prims/edge_partition_src_dst_property.cuh @@ -213,13 +213,8 @@ class edge_partition_minor_property_device_view_t { template class edge_partition_major_property_t { public: -<<<<<<< HEAD:cpp/include/cugraph/prims/edge_partition_src_dst_property.cuh - edge_partition_major_property_t() - : buffer_(allocate_dataframe_buffer(0, rmm::cuda_stream_view{})) -======= - major_properties_t(raft::handle_t const& handle) + edge_partition_major_property_t(raft::handle_t const& handle) : buffer_(allocate_dataframe_buffer(size_t{0}, handle.get_stream())) ->>>>>>> ab72ed53d4de1fed46bdb81c3c1b6e54b41770e7:cpp/include/cugraph/prims/row_col_properties.cuh { } @@ -328,7 +323,7 @@ class edge_partition_major_property_t { template class edge_partition_minor_property_t { public: - minor_properties_t(raft::handle_t const& handle) + edge_partition_minor_property_t(raft::handle_t const& handle) : buffer_(allocate_dataframe_buffer(size_t{0}, handle.get_stream())) { } @@ -433,10 +428,10 @@ class edge_partition_src_property_t { static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); - edge_partition_src_properties_t(raft::handle_t const& handle) : properties_(handle) {} + edge_partition_src_property_t(raft::handle_t const& handle) : property_(handle) {} - edge_partition_src_properties_t(raft::handle_t const& handle, GraphViewType const& graph_view) - : properties_(handle) + edge_partition_src_property_t(raft::handle_t const& handle, GraphViewType const& graph_view) + : property_(handle) { using vertex_t = typename GraphViewType::vertex_type; @@ -487,9 +482,9 @@ class edge_partition_src_property_t { } } - void clear(raft::handle_t const& handle) { properties_.clear(handle); } + void clear(raft::handle_t const& handle) { property_.clear(handle); } - void fill(T value, rmm::cuda_stream_view stream) { properties_.fill(value, stream); } + void fill(T value, rmm::cuda_stream_view stream) { property_.fill(value, stream); } auto key_first() { return property_.key_first(); } auto key_last() { return property_.key_last(); } @@ -500,10 +495,11 @@ class edge_partition_src_property_t { auto mutable_device_view() { return property_.mutable_device_view(); } private: - std::conditional_t, - detail::major_properties_t> - properties_; + std::conditional_t< + GraphViewType::is_adj_matrix_transposed, + detail::edge_partition_minor_property_t, + detail::edge_partition_major_property_t> + property_; }; template @@ -513,10 +509,10 @@ class edge_partition_dst_property_t { static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); - edge_partition_dst_properties_t(raft::handle_t const& handle) : properties_(handle) {} + edge_partition_dst_property_t(raft::handle_t const& handle) : property_(handle) {} - edge_partition_dst_properties_t(raft::handle_t const& handle, GraphViewType const& graph_view) - : properties_(handle) + edge_partition_dst_property_t(raft::handle_t const& handle, GraphViewType const& graph_view) + : property_(handle) { using vertex_t = typename GraphViewType::vertex_type; @@ -567,9 +563,9 @@ class edge_partition_dst_property_t { } } - void clear(raft::handle_t const& handle) { properties_.clear(handle); } + void clear(raft::handle_t const& handle) { property_.clear(handle); } - void fill(T value, rmm::cuda_stream_view stream) { properties_.fill(value, stream); } + void fill(T value, rmm::cuda_stream_view stream) { property_.fill(value, stream); } auto key_first() { return property_.key_first(); } auto key_last() { return property_.key_last(); } @@ -580,10 +576,11 @@ class edge_partition_dst_property_t { auto mutable_device_view() { return property_.mutable_device_view(); } private: - std::conditional_t, - detail::minor_properties_t> - properties_; + std::conditional_t< + GraphViewType::is_adj_matrix_transposed, + detail::edge_partition_major_property_t, + detail::edge_partition_minor_property_t> + property_; }; template diff --git a/cpp/include/cugraph/prims/extract_if_e.cuh b/cpp/include/cugraph/prims/extract_if_e.cuh index 1c4ed54b220..d5b6f2b3a71 100644 --- a/cpp/include/cugraph/prims/extract_if_e.cuh +++ b/cpp/include/cugraph/prims/extract_if_e.cuh @@ -18,9 +18,9 @@ #include #include #include +#include #include #include -#include #include #include @@ -101,14 +101,14 @@ struct call_e_op_t { * @param graph_view Non-owning graph object. * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::row_properties_t::device_view() (if @p e_op needs to access row properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access row properties). Use - * copy_to_adj_matrix_row to fill the wrapper. + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) + * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use + * update_edge_partition_src_property to fill the wrapper. * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::col_properties_t::device_view() (if @p e_op needs to access column properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access column properties). Use - * copy_to_adj_matrix_col to fill the wrapper. + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column + * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column + * properties). Use update_edge_partition_dst_property to fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) * and returns a boolean value to designate whether to include this edge in the returned edge list diff --git a/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh b/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh index c81cf2d133e..67487ec4b6f 100644 --- a/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh @@ -571,17 +571,18 @@ transform_reduce_by_adj_matrix_row_col_key_e( * @param graph_view Non-owning graph object. * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::row_properties_t::device_view() (if @p e_op needs to access row properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access row properties). Use - * copy_to_adj_matrix_row to fill the wrapper. + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) + * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use + * update_edge_partition_src_property to fill the wrapper. * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::col_properties_t::device_view() (if @p e_op needs to access column properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access column properties). Use - * copy_to_adj_matrix_col to fill the wrapper. + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column + * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column + * properties). Use update_edge_partition_dst_property to fill the wrapper. * @param adj_matrix_row_key_input Device-copyable wrapper used to access row keys(for the rows - * assigned to this process in multi-GPU). Use either cugraph::row_properties_t::device_view(). Use - * copy_to_adj_matrix_row to fill the wrapper. + * assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_src_property_t::device_view(). Use update_edge_partition_src_property to + * fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) * and returns a transformed value to be reduced. @@ -640,17 +641,18 @@ auto transform_reduce_by_adj_matrix_row_key_e( * @param graph_view Non-owning graph object. * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::row_properties_t::device_view() (if @p e_op needs to access row properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access row properties). Use - * copy_to_adj_matrix_row to fill the wrapper. + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) + * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use + * update_edge_partition_src_property to fill the wrapper. * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::col_properties_t::device_view() (if @p e_op needs to access column properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access column properties). Use - * copy_to_adj_matrix_col to fill the wrapper. + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column + * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column + * properties). Use update_edge_partition_dst_property to fill the wrapper. * @param adj_matrix_col_key_input Device-copyable wrapper used to access column keys(for the * columns assigned to this process in multi-GPU). Use either - * cugraph::col_properties_t::device_view(). Use copy_to_adj_matrix_col to fill the wrapper. + * cugraph::edge_partition_dst_property_t::device_view(). Use update_edge_partition_dst_property to + * fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) * and returns a transformed value to be reduced. diff --git a/cpp/include/cugraph/prims/transform_reduce_e.cuh b/cpp/include/cugraph/prims/transform_reduce_e.cuh index 23e79f10b27..89d783dd845 100644 --- a/cpp/include/cugraph/prims/transform_reduce_e.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_e.cuh @@ -19,6 +19,7 @@ #include #include #include +#include #include #include @@ -376,14 +377,14 @@ __global__ void for_all_major_for_all_nbr_high_degree( * @param graph_view Non-owning graph object. * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::row_properties_t::device_view() (if @p e_op needs to access row properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access row properties). Use - * copy_to_adj_matrix_row to fill the wrapper. + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) + * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use + * update_edge_partition_src_property to fill the wrapper. * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::col_properties_t::device_view() (if @p e_op needs to access column properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access column properties). Use - * copy_to_adj_matrix_col to fill the wrapper. + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column + * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column + * properties). Use update_edge_partition_dst_property to fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) * and returns a value to be reduced. diff --git a/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh b/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh index f2cad3166e8..a55f5cfe86a 100644 --- a/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh +++ b/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh @@ -961,14 +961,14 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( * vertices for the next iteration. * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::row_properties_t::device_view() (if @p e_op needs to access row properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access row properties). Use - * copy_to_adj_matrix_row to fill the wrapper. + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) + * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use + * update_edge_partition_src_property to fill the wrapper. * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::col_properties_t::device_view() (if @p e_op needs to access column properties) or - * cugraph::dummy_properties_t::device_view() (if @p e_op does not access column properties). Use - * copy_to_adj_matrix_col to fill the wrapper. + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column + * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column + * properties). Use update_edge_partition_dst_property to fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) * and returns a value to be reduced the @p reduce_op. diff --git a/cpp/src/centrality/katz_centrality_impl.cuh b/cpp/src/centrality/katz_centrality_impl.cuh index b5dbe5a7af7..57bda084c1f 100644 --- a/cpp/src/centrality/katz_centrality_impl.cuh +++ b/cpp/src/centrality/katz_centrality_impl.cuh @@ -17,11 +17,11 @@ #include #include -#include #include #include -#include +#include #include +#include #include #include @@ -93,22 +93,22 @@ void katz_centrality(raft::handle_t const& handle, // old katz centrality values rmm::device_uvector tmp_katz_centralities( pull_graph_view.get_number_of_local_vertices(), handle.get_stream()); - row_properties_t adj_matrix_row_katz_centralities(handle, - pull_graph_view); + edge_partition_src_property_t adj_matrix_row_katz_centralities( + handle, pull_graph_view); auto new_katz_centralities = katz_centralities; auto old_katz_centralities = tmp_katz_centralities.data(); size_t iter{0}; while (true) { std::swap(new_katz_centralities, old_katz_centralities); - copy_to_adj_matrix_row( + update_edge_partition_src_property( handle, pull_graph_view, old_katz_centralities, adj_matrix_row_katz_centralities); copy_v_transform_reduce_in_nbr( handle, pull_graph_view, adj_matrix_row_katz_centralities.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), [alpha] __device__(vertex_t, vertex_t, weight_t w, auto src_val, auto) { return static_cast(alpha * src_val * w); }, diff --git a/cpp/src/community/louvain.cuh b/cpp/src/community/louvain.cuh index 50360d3e260..e42102d6182 100644 --- a/cpp/src/community/louvain.cuh +++ b/cpp/src/community/louvain.cuh @@ -20,13 +20,13 @@ #include #include -#include #include #include -#include +#include #include #include #include +#include #include #include @@ -162,8 +162,8 @@ class Louvain { weight_t total_edge_weight = transform_reduce_e( handle_, current_graph_view_, - dummy_properties_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), [] __device__(auto, auto, weight_t wt, auto, auto) { return wt; }, weight_t{0}); @@ -261,11 +261,11 @@ class Louvain { current_graph_view_, graph_view_t::is_multi_gpu ? src_clusters_cache_.device_view() - : detail::major_properties_device_view_t( + : detail::edge_partition_major_property_device_view_t( next_clusters_v_.begin()), graph_view_t::is_multi_gpu ? dst_clusters_cache_.device_view() - : detail::minor_properties_device_view_t( + : detail::edge_partition_minor_property_device_view_t( next_clusters_v_.begin()), [] __device__(auto, auto, weight_t wt, auto src_cluster, auto nbr_cluster) { if (src_cluster == nbr_cluster) { @@ -324,8 +324,8 @@ class Louvain { if constexpr (graph_view_t::is_multi_gpu) { src_vertex_weights_cache_ = - row_properties_t(handle_, current_graph_view_); - copy_to_adj_matrix_row( + edge_partition_src_property_t(handle_, current_graph_view_); + update_edge_partition_src_property( handle_, current_graph_view_, vertex_weights_v_.begin(), src_vertex_weights_cache_); vertex_weights_v_.resize(0, handle_.get_stream()); vertex_weights_v_.shrink_to_fit(handle_.get_stream()); @@ -347,11 +347,13 @@ class Louvain { handle_.get_stream()); if constexpr (graph_view_t::is_multi_gpu) { - src_clusters_cache_ = row_properties_t(handle_, current_graph_view_); - copy_to_adj_matrix_row( + src_clusters_cache_ = + edge_partition_src_property_t(handle_, current_graph_view_); + update_edge_partition_src_property( handle_, current_graph_view_, next_clusters_v_.begin(), src_clusters_cache_); - dst_clusters_cache_ = col_properties_t(handle_, current_graph_view_); - copy_to_adj_matrix_col( + dst_clusters_cache_ = + edge_partition_dst_property_t(handle_, current_graph_view_); + update_edge_partition_dst_property( handle_, current_graph_view_, next_clusters_v_.begin(), dst_clusters_cache_); } @@ -397,11 +399,11 @@ class Louvain { current_graph_view_, graph_view_t::is_multi_gpu ? src_clusters_cache_.device_view() - : detail::major_properties_device_view_t( + : detail::edge_partition_major_property_device_view_t( next_clusters_v_.data()), graph_view_t::is_multi_gpu ? dst_clusters_cache_.device_view() - : detail::minor_properties_device_view_t( + : detail::edge_partition_minor_property_device_view_t( next_clusters_v_.data()), [] __device__(auto src, auto dst, auto wt, auto src_cluster, auto nbr_cluster) { weight_t sum{0}; @@ -427,7 +429,7 @@ class Louvain { bool up_down) { rmm::device_uvector vertex_cluster_weights_v(0, handle_.get_stream()); - row_properties_t src_cluster_weights(handle_); + edge_partition_src_property_t src_cluster_weights(handle_); if constexpr (graph_view_t::is_multi_gpu) { cugraph::detail::compute_gpu_id_from_vertex_t vertex_to_gpu_id_op{ handle_.get_comms().get_size()}; @@ -441,8 +443,9 @@ class Louvain { vertex_to_gpu_id_op, handle_.get_stream()); - src_cluster_weights = row_properties_t(handle_, current_graph_view_); - copy_to_adj_matrix_row( + src_cluster_weights = + edge_partition_src_property_t(handle_, current_graph_view_); + update_edge_partition_src_property( handle_, current_graph_view_, vertex_cluster_weights_v.begin(), src_cluster_weights); vertex_cluster_weights_v.resize(0, handle_.get_stream()); vertex_cluster_weights_v.shrink_to_fit(handle_.get_stream()); @@ -468,17 +471,17 @@ class Louvain { auto [old_cluster_sum_v, cluster_subtract_v] = compute_cluster_sum_and_subtract(); - row_properties_t> + edge_partition_src_property_t> src_old_cluster_sum_subtract_pairs(handle_); if constexpr (graph_view_t::is_multi_gpu) { src_old_cluster_sum_subtract_pairs = - row_properties_t>(handle_, - current_graph_view_); - copy_to_adj_matrix_row(handle_, - current_graph_view_, - thrust::make_zip_iterator(thrust::make_tuple( - old_cluster_sum_v.begin(), cluster_subtract_v.begin())), - src_old_cluster_sum_subtract_pairs); + edge_partition_src_property_t>( + handle_, current_graph_view_); + update_edge_partition_src_property(handle_, + current_graph_view_, + thrust::make_zip_iterator(thrust::make_tuple( + old_cluster_sum_v.begin(), cluster_subtract_v.begin())), + src_old_cluster_sum_subtract_pairs); old_cluster_sum_v.resize(0, handle_.get_stream()); old_cluster_sum_v.shrink_to_fit(handle_.get_stream()); cluster_subtract_v.resize(0, handle_.get_stream()); @@ -497,15 +500,15 @@ class Louvain { src_cluster_weights.device_view(), src_old_cluster_sum_subtract_pairs.device_view()) : device_view_concat( - detail::major_properties_device_view_t( + detail::edge_partition_major_property_device_view_t( vertex_weights_v_.data()), - detail::major_properties_device_view_t( + detail::edge_partition_major_property_device_view_t( next_clusters_v_.data()), - detail::major_properties_device_view_t( + detail::edge_partition_major_property_device_view_t( vertex_cluster_weights_v.data()), - detail::major_properties_device_view_t( - cluster_old_sum_subtract_pair_first)); + detail::edge_partition_major_property_device_view_t< + vertex_t, + decltype(cluster_old_sum_subtract_pair_first)>(cluster_old_sum_subtract_pair_first)); copy_v_transform_reduce_key_aggregated_out_nbr( handle_, @@ -513,7 +516,7 @@ class Louvain { zipped_src_device_view, graph_view_t::is_multi_gpu ? dst_clusters_cache_.device_view() - : detail::minor_properties_device_view_t( + : detail::edge_partition_minor_property_device_view_t( next_clusters_v_.data()), cluster_keys_v_.begin(), cluster_keys_v_.end(), @@ -531,9 +534,9 @@ class Louvain { detail::cluster_update_op_t{up_down}); if constexpr (graph_view_t::is_multi_gpu) { - copy_to_adj_matrix_row( + update_edge_partition_src_property( handle_, current_graph_view_, next_clusters_v_.begin(), src_clusters_cache_); - copy_to_adj_matrix_col( + update_edge_partition_dst_property( handle_, current_graph_view_, next_clusters_v_.begin(), dst_clusters_cache_); } @@ -541,11 +544,11 @@ class Louvain { cugraph::transform_reduce_by_adj_matrix_row_key_e( handle_, current_graph_view_, - dummy_properties_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), graph_view_t::is_multi_gpu ? src_clusters_cache_.device_view() - : detail::major_properties_device_view_t( + : detail::edge_partition_major_property_device_view_t( next_clusters_v_.data()), detail::return_edge_weight_t{}, weight_t{0}); @@ -609,12 +612,14 @@ class Louvain { rmm::device_uvector cluster_weights_v_; rmm::device_uvector vertex_weights_v_; - row_properties_t + edge_partition_src_property_t src_vertex_weights_cache_; // src cache for vertex_weights_v_ rmm::device_uvector next_clusters_v_; - row_properties_t src_clusters_cache_; // src cache for next_clusters_v_ - col_properties_t dst_clusters_cache_; // dst cache for next_clusters_v_ + edge_partition_src_property_t + src_clusters_cache_; // src cache for next_clusters_v_ + edge_partition_dst_property_t + dst_clusters_cache_; // dst cache for next_clusters_v_ #ifdef TIMING HighResTimer hr_timer_; diff --git a/cpp/src/components/weakly_connected_components_impl.cuh b/cpp/src/components/weakly_connected_components_impl.cuh index c5382096044..b9aa32075b8 100644 --- a/cpp/src/components/weakly_connected_components_impl.cuh +++ b/cpp/src/components/weakly_connected_components_impl.cuh @@ -19,8 +19,8 @@ #include #include #include -#include -#include +#include +#include #include #include #include @@ -460,8 +460,8 @@ void weakly_connected_components_impl(raft::handle_t const& handle, auto adj_matrix_col_components = GraphViewType::is_multi_gpu - ? col_properties_t(handle, level_graph_view) - : col_properties_t(handle); + ? edge_partition_dst_property_t(handle, level_graph_view) + : edge_partition_dst_property_t(handle); if constexpr (GraphViewType::is_multi_gpu) { adj_matrix_col_components.fill(invalid_component_id::value, handle.get_stream()); } @@ -505,7 +505,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, } if constexpr (GraphViewType::is_multi_gpu) { - copy_to_adj_matrix_col( + update_edge_partition_dst_property( handle, level_graph_view, thrust::get<0>(vertex_frontier.get_bucket(static_cast(Bucket::cur)) @@ -538,12 +538,13 @@ void weakly_connected_components_impl(raft::handle_t const& handle, GraphViewType::is_multi_gpu ? std::vector{static_cast(Bucket::next), static_cast(Bucket::conflict)} : std::vector{static_cast(Bucket::next)}, - dummy_properties_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), [col_components = GraphViewType::is_multi_gpu ? adj_matrix_col_components.mutable_device_view() - : detail::minor_properties_device_view_t(level_components), + : detail::edge_partition_minor_property_device_view_t( + level_components), col_first = level_graph_view.get_local_adj_matrix_partition_col_first(), edge_buffer_first = get_dataframe_buffer_begin(edge_buffer), num_edge_inserts = diff --git a/cpp/src/link_analysis/hits_impl.cuh b/cpp/src/link_analysis/hits_impl.cuh index 6cb62c08531..553e5e1fbf6 100644 --- a/cpp/src/link_analysis/hits_impl.cuh +++ b/cpp/src/link_analysis/hits_impl.cuh @@ -100,7 +100,7 @@ std::tuple hits(raft::handle_t const& handle, // Initialize hubs from user input if provided if (has_initial_hubs_guess) { - copy_to_adj_matrix_row(handle, graph_view, prev_hubs, prev_src_hubs); + update_edge_partition_src_property(handle, graph_view, prev_hubs, prev_src_hubs); } else { prev_src_hubs.fill(result_t{1.0} / num_vertices, handle.get_stream()); thrust::fill(handle.get_thrust_policy(), @@ -114,18 +114,18 @@ std::tuple hits(raft::handle_t const& handle, handle, graph_view, prev_src_hubs.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), [] __device__(auto, auto, auto, auto prev_src_hub_value, auto) { return prev_src_hub_value; }, result_t{0}, authorities); - copy_to_adj_matrix_col(handle, graph_view, authorities, curr_dst_auth); + update_edge_partition_dst_property(handle, graph_view, authorities, curr_dst_auth); // Update current source hubs property copy_v_transform_reduce_out_nbr( handle, graph_view, - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), curr_dst_auth.device_view(), [] __device__(auto src, auto dst, auto, auto, auto curr_dst_auth_value) { return curr_dst_auth_value; @@ -152,7 +152,7 @@ std::tuple hits(raft::handle_t const& handle, break; } - copy_to_adj_matrix_row(handle, graph_view, curr_hubs, prev_src_hubs); + update_edge_partition_src_property(handle, graph_view, curr_hubs, prev_src_hubs); // Swap pointers for the next iteration // After this swap call, prev_hubs has the latest value of hubs diff --git a/cpp/src/link_analysis/pagerank_impl.cuh b/cpp/src/link_analysis/pagerank_impl.cuh index b6023d21bf2..e518f19c1ab 100644 --- a/cpp/src/link_analysis/pagerank_impl.cuh +++ b/cpp/src/link_analysis/pagerank_impl.cuh @@ -17,13 +17,13 @@ #include #include -#include #include #include #include +#include #include -#include #include +#include #include #include @@ -103,8 +103,8 @@ void pagerank( auto num_nonpositive_edge_weights = count_if_e(handle, pull_graph_view, - dummy_properties_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), [] __device__(vertex_t, vertex_t, weight_t w, auto, auto) { return w <= 0.0; }); CUGRAPH_EXPECTS(num_nonpositive_edge_weights == 0, "Invalid input argument: input graph should have postive edge weights."); @@ -189,7 +189,8 @@ void pagerank( // old PageRank values rmm::device_uvector old_pageranks(pull_graph_view.get_number_of_local_vertices(), handle.get_stream()); - row_properties_t adj_matrix_row_pageranks(handle, pull_graph_view); + edge_partition_src_property_t adj_matrix_row_pageranks(handle, + pull_graph_view); size_t iter{0}; while (true) { thrust::copy(handle.get_thrust_policy(), @@ -223,7 +224,8 @@ void pagerank( return pagerank / divisor; }); - copy_to_adj_matrix_row(handle, pull_graph_view, pageranks, adj_matrix_row_pageranks); + update_edge_partition_src_property( + handle, pull_graph_view, pageranks, adj_matrix_row_pageranks); auto unvarying_part = aggregate_personalization_vector_size == 0 ? (dangling_sum * alpha + static_cast(1.0 - alpha)) / @@ -234,7 +236,7 @@ void pagerank( handle, pull_graph_view, adj_matrix_row_pageranks.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), [alpha] __device__(vertex_t, vertex_t, weight_t w, auto src_val, auto) { return src_val * w * alpha; }, diff --git a/cpp/src/structure/coarsen_graph_impl.cuh b/cpp/src/structure/coarsen_graph_impl.cuh index c6f2f05cf67..bca9d836001 100644 --- a/cpp/src/structure/coarsen_graph_impl.cuh +++ b/cpp/src/structure/coarsen_graph_impl.cuh @@ -21,8 +21,8 @@ #include #include #include -#include -#include +#include +#include #include #include @@ -252,17 +252,18 @@ coarsen_graph( bool lower_triangular_only = graph_view.is_symmetric(); - std::conditional_t< - store_transposed, - row_properties_t, - vertex_t>, - col_properties_t, - vertex_t>> + std::conditional_t, + vertex_t>, + edge_partition_dst_property_t< + graph_view_t, + vertex_t>> adj_matrix_minor_labels(handle, graph_view); if constexpr (store_transposed) { - copy_to_adj_matrix_row(handle, graph_view, labels, adj_matrix_minor_labels); + update_edge_partition_src_property(handle, graph_view, labels, adj_matrix_minor_labels); } else { - copy_to_adj_matrix_col(handle, graph_view, labels, adj_matrix_minor_labels); + update_edge_partition_dst_property(handle, graph_view, labels, adj_matrix_minor_labels); } std::vector> coarsened_edgelist_majors{}; @@ -521,7 +522,7 @@ coarsen_graph( matrix_partition_device_view_t( graph_view.get_matrix_partition_view()), labels, - detail::minor_properties_device_view_t(labels), + detail::edge_partition_minor_property_device_view_t(labels), graph_view.get_local_adj_matrix_partition_segment_offsets(0), lower_triangular_only); diff --git a/cpp/src/structure/graph_view_impl.cuh b/cpp/src/structure/graph_view_impl.cuh index 1eef31ae806..7c5187d132e 100644 --- a/cpp/src/structure/graph_view_impl.cuh +++ b/cpp/src/structure/graph_view_impl.cuh @@ -22,7 +22,7 @@ #include #include #include -#include +#include #include #include #include @@ -193,8 +193,8 @@ rmm::device_uvector compute_minor_degrees( copy_v_transform_reduce_out_nbr( handle, graph_view, - dummy_properties_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), [] __device__(vertex_t, vertex_t, weight_t, auto, auto) { return edge_t{1}; }, edge_t{0}, minor_degrees.data()); @@ -202,8 +202,8 @@ rmm::device_uvector compute_minor_degrees( copy_v_transform_reduce_in_nbr( handle, graph_view, - dummy_properties_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), [] __device__(vertex_t, vertex_t, weight_t, auto, auto) { return edge_t{1}; }, edge_t{0}, minor_degrees.data()); @@ -228,8 +228,8 @@ rmm::device_uvector compute_weight_sums( copy_v_transform_reduce_in_nbr( handle, graph_view, - dummy_properties_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), [] __device__(vertex_t, vertex_t, weight_t w, auto, auto) { return w; }, weight_t{0.0}, weight_sums.data()); @@ -237,8 +237,8 @@ rmm::device_uvector compute_weight_sums( copy_v_transform_reduce_out_nbr( handle, graph_view, - dummy_properties_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), [] __device__(vertex_t, vertex_t, weight_t w, auto, auto) { return w; }, weight_t{0.0}, weight_sums.data()); @@ -874,8 +874,8 @@ graph_view_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), [] __device__(vertex_t src, vertex_t dst, auto src_val, auto dst_val) { return src == dst ? edge_t{1} : edge_t{0}; }, @@ -898,8 +898,8 @@ edge_t graph_view_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), [] __device__(vertex_t src, vertex_t dst, auto src_val, auto dst_val) { return src == dst ? edge_t{1} : edge_t{0}; }, diff --git a/cpp/src/traversal/bfs_impl.cuh b/cpp/src/traversal/bfs_impl.cuh index 995534b277c..dc94dccee87 100644 --- a/cpp/src/traversal/bfs_impl.cuh +++ b/cpp/src/traversal/bfs_impl.cuh @@ -17,10 +17,10 @@ #include #include -#include #include +#include #include -#include +#include #include #include #include @@ -47,9 +47,10 @@ namespace { template struct e_op_t { - std:: - conditional_t, uint32_t*> - visited_flags{nullptr}; + std::conditional_t, + uint32_t*> + visited_flags{nullptr}; uint32_t const* prev_visited_flags{ nullptr}; // relevant only if multi_gpu is false (this affects only local-computing with 0 // impact in communication volume, so this may improve performance in small-scale but @@ -183,9 +184,10 @@ void bfs(raft::handle_t const& handle, handle.get_stream()); // relevant only if GraphViewType::is_multi_gpu is false auto dst_visited_flags = GraphViewType::is_multi_gpu - ? col_properties_t(handle, push_graph_view) - : col_properties_t(handle); // relevant only if GraphViewType::is_multi_gpu is true + ? edge_partition_dst_property_t(handle, push_graph_view) + : edge_partition_dst_property_t( + handle); // relevant only if GraphViewType::is_multi_gpu is true if constexpr (GraphViewType::is_multi_gpu) { dst_visited_flags.fill(uint8_t{0}, handle.get_stream()); } @@ -197,12 +199,13 @@ void bfs(raft::handle_t const& handle, CUGRAPH_FAIL("unimplemented."); } else { if (GraphViewType::is_multi_gpu) { - copy_to_adj_matrix_col(handle, - push_graph_view, - vertex_frontier.get_bucket(static_cast(Bucket::cur)).begin(), - vertex_frontier.get_bucket(static_cast(Bucket::cur)).end(), - thrust::make_constant_iterator(uint8_t{1}), - dst_visited_flags); + update_edge_partition_dst_property( + handle, + push_graph_view, + vertex_frontier.get_bucket(static_cast(Bucket::cur)).begin(), + vertex_frontier.get_bucket(static_cast(Bucket::cur)).end(), + thrust::make_constant_iterator(uint8_t{1}), + dst_visited_flags); } else { thrust::copy(handle.get_thrust_policy(), visited_flags.begin(), @@ -225,8 +228,8 @@ void bfs(raft::handle_t const& handle, vertex_frontier, static_cast(Bucket::cur), std::vector{static_cast(Bucket::next)}, - dummy_properties_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), #if 1 e_op, #else diff --git a/cpp/src/traversal/sssp_impl.cuh b/cpp/src/traversal/sssp_impl.cuh index 48745d6dae4..7d64ade6f69 100644 --- a/cpp/src/traversal/sssp_impl.cuh +++ b/cpp/src/traversal/sssp_impl.cuh @@ -17,11 +17,11 @@ #include #include -#include #include +#include #include -#include #include +#include #include #include #include @@ -79,8 +79,8 @@ void sssp(raft::handle_t const& handle, auto num_negative_edge_weights = count_if_e(handle, push_graph_view, - dummy_properties_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), [] __device__(vertex_t, vertex_t, weight_t w, auto, auto) { return w < 0.0; }); CUGRAPH_EXPECTS(num_negative_edge_weights == 0, "Invalid input argument: input graph should have non-negative edge weights."); @@ -111,8 +111,8 @@ void sssp(raft::handle_t const& handle, thrust::tie(average_vertex_degree, average_edge_weight) = transform_reduce_e( handle, push_graph_view, - dummy_properties_t{}.device_view(), - dummy_properties_t{}.device_view(), + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), [] __device__(vertex_t, vertex_t, weight_t w, auto, auto) { return thrust::make_tuple(weight_t{1.0}, w); }, @@ -134,8 +134,9 @@ void sssp(raft::handle_t const& handle, // 5. SSSP iteration auto adj_matrix_row_distances = - GraphViewType::is_multi_gpu ? row_properties_t(handle, push_graph_view) - : row_properties_t(handle); + GraphViewType::is_multi_gpu + ? edge_partition_src_property_t(handle, push_graph_view) + : edge_partition_src_property_t(handle); if (GraphViewType::is_multi_gpu) { adj_matrix_row_distances.fill(std::numeric_limits::max(), handle.get_stream()); } @@ -147,7 +148,7 @@ void sssp(raft::handle_t const& handle, auto near_far_threshold = delta; while (true) { if (GraphViewType::is_multi_gpu) { - copy_to_adj_matrix_row( + update_edge_partition_src_property( handle, push_graph_view, vertex_frontier.get_bucket(static_cast(Bucket::cur_near)).begin(), @@ -167,8 +168,8 @@ void sssp(raft::handle_t const& handle, std::vector{static_cast(Bucket::next_near), static_cast(Bucket::far)}, GraphViewType::is_multi_gpu ? adj_matrix_row_distances.device_view() - : detail::major_properties_device_view_t(distances), - dummy_properties_t{}.device_view(), + : detail::edge_partition_major_property_device_view_t(distances), + dummy_property_t{}.device_view(), [vertex_partition, distances, cutoff] __device__( vertex_t src, vertex_t dst, weight_t w, auto src_val, auto) { auto push = true; diff --git a/cpp/tests/prims/mg_copy_v_transform_reduce_inout_nbr.cu b/cpp/tests/prims/mg_copy_v_transform_reduce_inout_nbr.cu index 60f2e2857cf..a037a41e951 100644 --- a/cpp/tests/prims/mg_copy_v_transform_reduce_inout_nbr.cu +++ b/cpp/tests/prims/mg_copy_v_transform_reduce_inout_nbr.cu @@ -28,9 +28,9 @@ #include #include #include -#include #include -#include +#include +#include #include #include @@ -132,8 +132,9 @@ struct generate_impl { graph_view_type const& graph_view, property_buffer_type& property) { - auto output_property = cugraph::col_properties_t(handle, graph_view); - copy_to_adj_matrix_col( + auto output_property = + cugraph::edge_partition_dst_property_t(handle, graph_view); + update_edge_partition_dst_property( handle, graph_view, cugraph::get_dataframe_buffer_begin(property), output_property); return output_property; } @@ -143,8 +144,9 @@ struct generate_impl { graph_view_type const& graph_view, property_buffer_type& property) { - auto output_property = cugraph::row_properties_t(handle, graph_view); - copy_to_adj_matrix_row( + auto output_property = + cugraph::edge_partition_src_property_t(handle, graph_view); + update_edge_partition_src_property( handle, graph_view, cugraph::get_dataframe_buffer_begin(property), output_property); return output_property; } diff --git a/cpp/tests/prims/mg_count_if_e.cu b/cpp/tests/prims/mg_count_if_e.cu index d76df6a550c..04c0ceb9011 100644 --- a/cpp/tests/prims/mg_count_if_e.cu +++ b/cpp/tests/prims/mg_count_if_e.cu @@ -28,9 +28,9 @@ #include #include #include -#include #include -#include +#include +#include #include #include @@ -130,8 +130,9 @@ struct generate_impl { graph_view_type const& graph_view, property_buffer_type& property) { - auto output_property = cugraph::col_properties_t(handle, graph_view); - copy_to_adj_matrix_col( + auto output_property = + cugraph::edge_partition_dst_property_t(handle, graph_view); + update_edge_partition_dst_property( handle, graph_view, cugraph::get_dataframe_buffer_begin(property), output_property); return output_property; } @@ -141,8 +142,9 @@ struct generate_impl { graph_view_type const& graph_view, property_buffer_type& property) { - auto output_property = cugraph::row_properties_t(handle, graph_view); - copy_to_adj_matrix_row( + auto output_property = + cugraph::edge_partition_src_property_t(handle, graph_view); + update_edge_partition_src_property( handle, graph_view, cugraph::get_dataframe_buffer_begin(property), output_property); return output_property; } diff --git a/cpp/tests/prims/mg_extract_if_e.cu b/cpp/tests/prims/mg_extract_if_e.cu index dd3f7078098..ef8d0c1fbda 100644 --- a/cpp/tests/prims/mg_extract_if_e.cu +++ b/cpp/tests/prims/mg_extract_if_e.cu @@ -26,9 +26,9 @@ #include #include -#include #include #include +#include #include #include @@ -191,19 +191,19 @@ class Tests_MG_ExtractIfE cugraph::get_dataframe_buffer_begin(mg_property_buffer), property_transform_t{hash_bin_count}); - cugraph::row_properties_t mg_src_properties(handle, - mg_graph_view); - cugraph::col_properties_t mg_dst_properties(handle, - mg_graph_view); + cugraph::edge_partition_src_property_t mg_src_properties( + handle, mg_graph_view); + cugraph::edge_partition_dst_property_t mg_dst_properties( + handle, mg_graph_view); - copy_to_adj_matrix_row(handle, - mg_graph_view, - cugraph::get_dataframe_buffer_cbegin(mg_property_buffer), - mg_src_properties); - copy_to_adj_matrix_col(handle, - mg_graph_view, - cugraph::get_dataframe_buffer_cbegin(mg_property_buffer), - mg_dst_properties); + update_edge_partition_src_property(handle, + mg_graph_view, + cugraph::get_dataframe_buffer_cbegin(mg_property_buffer), + mg_src_properties); + update_edge_partition_dst_property(handle, + mg_graph_view, + cugraph::get_dataframe_buffer_cbegin(mg_property_buffer), + mg_dst_properties); if (cugraph::test::g_perf) { RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement @@ -280,19 +280,19 @@ class Tests_MG_ExtractIfE cugraph::get_dataframe_buffer_begin(sg_property_buffer), property_transform_t{hash_bin_count}); - cugraph::row_properties_t sg_src_properties( - handle, sg_graph_view); - cugraph::col_properties_t sg_dst_properties( - handle, sg_graph_view); - - copy_to_adj_matrix_row(handle, - sg_graph_view, - cugraph::get_dataframe_buffer_cbegin(sg_property_buffer), - sg_src_properties); - copy_to_adj_matrix_col(handle, - sg_graph_view, - cugraph::get_dataframe_buffer_cbegin(sg_property_buffer), - sg_dst_properties); + cugraph::edge_partition_src_property_t + sg_src_properties(handle, sg_graph_view); + cugraph::edge_partition_dst_property_t + sg_dst_properties(handle, sg_graph_view); + + update_edge_partition_src_property(handle, + sg_graph_view, + cugraph::get_dataframe_buffer_cbegin(sg_property_buffer), + sg_src_properties); + update_edge_partition_dst_property(handle, + sg_graph_view, + cugraph::get_dataframe_buffer_cbegin(sg_property_buffer), + sg_dst_properties); auto [sg_edgelist_srcs, sg_edgelist_dsts, sg_edgelist_weights] = extract_if_e(handle, diff --git a/cpp/tests/prims/mg_transform_reduce_e.cu b/cpp/tests/prims/mg_transform_reduce_e.cu index 8c07af604ca..2a24a1d6249 100644 --- a/cpp/tests/prims/mg_transform_reduce_e.cu +++ b/cpp/tests/prims/mg_transform_reduce_e.cu @@ -28,9 +28,9 @@ #include #include #include -#include -#include +#include #include +#include #include #include @@ -130,8 +130,9 @@ struct generate_impl { graph_view_type const& graph_view, property_buffer_type& property) { - auto output_property = cugraph::col_properties_t(handle, graph_view); - copy_to_adj_matrix_col( + auto output_property = + cugraph::edge_partition_dst_property_t(handle, graph_view); + update_edge_partition_dst_property( handle, graph_view, cugraph::get_dataframe_buffer_begin(property), output_property); return output_property; } @@ -141,8 +142,9 @@ struct generate_impl { graph_view_type const& graph_view, property_buffer_type& property) { - auto output_property = cugraph::row_properties_t(handle, graph_view); - copy_to_adj_matrix_row( + auto output_property = + cugraph::edge_partition_src_property_t(handle, graph_view); + update_edge_partition_src_property( handle, graph_view, cugraph::get_dataframe_buffer_begin(property), output_property); return output_property; } diff --git a/cpp/tests/prims/mg_update_frontier_v_push_if_out_nbr.cu b/cpp/tests/prims/mg_update_frontier_v_push_if_out_nbr.cu index a004e34dc29..94d4d494557 100644 --- a/cpp/tests/prims/mg_update_frontier_v_push_if_out_nbr.cu +++ b/cpp/tests/prims/mg_update_frontier_v_push_if_out_nbr.cu @@ -28,8 +28,8 @@ #include #include #include -#include -#include +#include +#include #include #include @@ -162,19 +162,19 @@ class Tests_MG_UpdateFrontierVPushIfOutNbr sources.end(), mg_graph_view.get_local_vertex_first()); - cugraph::row_properties_t mg_src_properties(handle, - mg_graph_view); - cugraph::col_properties_t mg_dst_properties(handle, - mg_graph_view); + cugraph::edge_partition_src_property_t mg_src_properties( + handle, mg_graph_view); + cugraph::edge_partition_dst_property_t mg_dst_properties( + handle, mg_graph_view); - copy_to_adj_matrix_row(handle, - mg_graph_view, - cugraph::get_dataframe_buffer_cbegin(mg_property_buffer), - mg_src_properties); - copy_to_adj_matrix_col(handle, - mg_graph_view, - cugraph::get_dataframe_buffer_cbegin(mg_property_buffer), - mg_dst_properties); + update_edge_partition_src_property(handle, + mg_graph_view, + cugraph::get_dataframe_buffer_cbegin(mg_property_buffer), + mg_src_properties); + update_edge_partition_dst_property(handle, + mg_graph_view, + cugraph::get_dataframe_buffer_cbegin(mg_property_buffer), + mg_dst_properties); enum class Bucket { cur, next, num_buckets }; cugraph::VertexFrontier(Bucket::num_buckets)> @@ -255,18 +255,18 @@ class Tests_MG_UpdateFrontierVPushIfOutNbr cugraph::get_dataframe_buffer_begin(sg_property_buffer), property_transform_t{hash_bin_count}); - cugraph::row_properties_t sg_src_properties( - handle, sg_graph_view); - cugraph::col_properties_t sg_dst_properties( - handle, sg_graph_view); - copy_to_adj_matrix_row(handle, - sg_graph_view, - cugraph::get_dataframe_buffer_cbegin(sg_property_buffer), - sg_src_properties); - copy_to_adj_matrix_col(handle, - sg_graph_view, - cugraph::get_dataframe_buffer_cbegin(sg_property_buffer), - sg_dst_properties); + cugraph::edge_partition_src_property_t + sg_src_properties(handle, sg_graph_view); + cugraph::edge_partition_dst_property_t + sg_dst_properties(handle, sg_graph_view); + update_edge_partition_src_property(handle, + sg_graph_view, + cugraph::get_dataframe_buffer_cbegin(sg_property_buffer), + sg_src_properties); + update_edge_partition_dst_property(handle, + sg_graph_view, + cugraph::get_dataframe_buffer_cbegin(sg_property_buffer), + sg_dst_properties); cugraph:: VertexFrontier(Bucket::num_buckets)> sg_vertex_frontier(handle); From 27eee7ff88103ababf84ebed978c9bfdf82bd351 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Fri, 11 Mar 2022 15:40:53 -0800 Subject: [PATCH 6/9] replace additional rows/cols in prims to srcs/dsts --- .../copy_v_transform_reduce_in_out_nbr.cuh | 194 ++++++----- ...ransform_reduce_key_aggregated_out_nbr.cuh | 57 +-- cpp/include/cugraph/prims/count_if_e.cuh | 48 +-- cpp/include/cugraph/prims/extract_if_e.cuh | 85 ++--- .../cugraph/prims/property_op_utils.cuh | 16 +- ...orm_reduce_by_adj_matrix_row_col_key_e.cuh | 329 +++++++++--------- .../cugraph/prims/transform_reduce_e.cuh | 116 +++--- ...update_edge_partition_src_dst_property.cuh | 21 +- .../update_frontier_v_push_if_out_nbr.cuh | 124 +++---- 9 files changed, 499 insertions(+), 491 deletions(-) diff --git a/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh b/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh index c314e7253d8..7d62d773976 100644 --- a/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh +++ b/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh @@ -47,8 +47,8 @@ int32_t constexpr copy_v_transform_reduce_nbr_for_all_block_size = 512; template matrix_partition, typename GraphViewType::vertex_type major_hypersparse_first, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, ResultValueOutputIteratorOrWrapper result_value_output, EdgeOp e_op, T init /* relevent only if update_major == true */) @@ -91,8 +91,8 @@ __global__ void for_all_major_for_all_nbr_hypersparse( thrust::tie(indices, weights, local_degree) = matrix_partition.get_local_edges(static_cast(major_idx)); auto transform_op = [&matrix_partition, - &adj_matrix_row_value_input, - &adj_matrix_col_value_input, + &edge_partition_src_value_input, + &edge_partition_dst_value_input, &e_op, major, indices, @@ -111,14 +111,14 @@ __global__ void for_all_major_for_all_nbr_hypersparse( : minor_offset; return evaluate_edge_op() .compute(row, col, weight, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); }; @@ -162,8 +162,8 @@ __global__ void for_all_major_for_all_nbr_hypersparse( template matrix_partition, typename GraphViewType::vertex_type major_first, typename GraphViewType::vertex_type major_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, ResultValueOutputIteratorOrWrapper result_value_output, EdgeOp e_op, T init /* relevent only if update_major == true */) @@ -201,8 +201,8 @@ __global__ void for_all_major_for_all_nbr_low_degree( thrust::tie(indices, weights, local_degree) = matrix_partition.get_local_edges(static_cast(major_offset)); auto transform_op = [&matrix_partition, - &adj_matrix_row_value_input, - &adj_matrix_col_value_input, + &edge_partition_src_value_input, + &edge_partition_dst_value_input, &e_op, major_offset, indices, @@ -224,14 +224,14 @@ __global__ void for_all_major_for_all_nbr_low_degree( : minor_offset; return evaluate_edge_op() .compute(row, col, weight, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); }; @@ -275,8 +275,8 @@ __global__ void for_all_major_for_all_nbr_low_degree( template matrix_partition, typename GraphViewType::vertex_type major_first, typename GraphViewType::vertex_type major_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, ResultValueOutputIteratorOrWrapper result_value_output, EdgeOp e_op, T init /* relevent only if update_major == true */) @@ -340,14 +340,14 @@ __global__ void for_all_major_for_all_nbr_mid_degree( : minor_offset; auto e_op_result = evaluate_edge_op() .compute(row, col, weight, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); if constexpr (update_major) { e_op_result_sum = edge_property_add(e_op_result_sum, e_op_result); @@ -371,8 +371,8 @@ __global__ void for_all_major_for_all_nbr_mid_degree( template matrix_partition, typename GraphViewType::vertex_type major_first, typename GraphViewType::vertex_type major_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, ResultValueOutputIteratorOrWrapper result_value_output, EdgeOp e_op, T init /* relevent only if update_major == true */) @@ -433,14 +433,14 @@ __global__ void for_all_major_for_all_nbr_high_degree( : minor_offset; auto e_op_result = evaluate_edge_op() .compute(row, col, weight, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); if constexpr (update_major) { e_op_result_sum = edge_property_add(e_op_result_sum, e_op_result); @@ -463,15 +463,15 @@ __global__ void for_all_major_for_all_nbr_high_degree( template void copy_v_transform_reduce_nbr(raft::handle_t const& handle, GraphViewType const& graph_view, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgeOp e_op, T init, VertexValueOutputIterator vertex_value_output_first) @@ -536,8 +536,8 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, } } - auto matrix_partition_row_value_input = adj_matrix_row_value_input; - auto matrix_partition_col_value_input = adj_matrix_col_value_input; + auto matrix_partition_row_value_input = edge_partition_src_value_input; + auto matrix_partition_col_value_input = edge_partition_dst_value_input; if constexpr (GraphViewType::is_adj_matrix_transposed) { matrix_partition_col_value_input.set_local_adj_matrix_partition_idx(i); } else { @@ -750,29 +750,29 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, * and thrust::copy() (update vertex properties part, take transform_reduce output as copy input). * * @tparam GraphViewType Type of the passed non-owning graph object. - * @tparam AdjMatrixRowValueInputWrapper Type of the wrapper for graph adjacency matrix row input - * properties. - * @tparam AdjMatrixColValueInputWrapper Type of the wrapper for graph adjacency matrix column input - * properties. + * @tparam EdgePartitionSrcValueInputWrapper Type of the wrapper for edge partition source property + * values. + * @tparam EdgePartitionDstValueInputWrapper Type of the wrapper for edge partition destination + * property values. * @tparam EdgeOp Type of the quaternary (or quinary) edge operator. * @tparam T Type of the initial value for reduction over the incoming edges. * @tparam VertexValueOutputIterator Type of the iterator for vertex output property variables. * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Non-owning graph object. - * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties - * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) - * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use - * update_edge_partition_src_property to fill the wrapper. - * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties - * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column - * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column - * properties). Use update_edge_partition_dst_property to fill the wrapper. + * @param edge_partition_src_value_input Device-copyable wrapper used to access source input + * property values (for the edge sources assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access source property + * values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access source property + * values). Use update_edge_partition_src_property to fill the wrapper. + * @param edge_partition_dst_value_input Device-copyable wrapper used to access destination input + * property values (for the edge destinations assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access destination + * property values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access + * destination property values). Use update_edge_partition_dst_property to fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge - * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) - * and returns a value to be reduced. + * weight), property values for the source, and property values for the destination and returns a + * value to be reduced. * @param init Initial value to be added to the reduced @p e_op return values for each vertex. * @param vertex_value_output_first Iterator pointing to the vertex property variables for the first * (inclusive) vertex (assigned to tihs process in multi-GPU). `vertex_value_output_last` @@ -780,23 +780,24 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, * graph_view.get_number_of_local_vertices(). */ template -void copy_v_transform_reduce_in_nbr(raft::handle_t const& handle, - GraphViewType const& graph_view, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, - EdgeOp e_op, - T init, - VertexValueOutputIterator vertex_value_output_first) +void copy_v_transform_reduce_in_nbr( + raft::handle_t const& handle, + GraphViewType const& graph_view, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgeOp e_op, + T init, + VertexValueOutputIterator vertex_value_output_first) { detail::copy_v_transform_reduce_nbr(handle, graph_view, - adj_matrix_row_value_input, - adj_matrix_col_value_input, + edge_partition_src_value_input, + edge_partition_dst_value_input, e_op, init, vertex_value_output_first); @@ -810,29 +811,29 @@ void copy_v_transform_reduce_in_nbr(raft::handle_t const& handle, * input). * * @tparam GraphViewType Type of the passed non-owning graph object. - * @tparam AdjMatrixRowValueInputWrapper Type of the wrapper for graph adjacency matrix row input - * properties. - * @tparam AdjMatrixColValueInputWrapper Type of the wrapper for graph adjacency matrix column input - * properties. + * @tparam EdgePartitionSrcValueInputWrapper Type of the wrapper for edge partition source property + * values. + * @tparam EdgePartitionDstValueInputWrapper Type of the wrapper for edge partition destination + * property values. * @tparam EdgeOp Type of the quaternary (or quinary) edge operator. * @tparam T Type of the initial value for reduction over the outgoing edges. * @tparam VertexValueOutputIterator Type of the iterator for vertex output property variables. * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Non-owning graph object. - * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties - * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) - * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use - * update_edge_partition_src_property to fill the wrapper. - * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties - * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column - * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column - * properties). Use update_edge_partition_dst_property to fill the wrapper. + * @param edge_partition_src_value_input Device-copyable wrapper used to access source input + * property values (for the edge sources assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access source property + * values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access source property + * values). Use update_edge_partition_src_property to fill the wrapper. + * @param edge_partition_dst_value_input Device-copyable wrapper used to access destination input + * property values (for the edge destinations assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access destination + * property values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access + * destination property values). Use update_edge_partition_dst_property to fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge - * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) - * and returns a value to be reduced. + * weight), property values for the source, and property values for the destination and returns a + * value to be reduced. * @param init Initial value to be added to the reduced @p e_op return values for each vertex. * @param vertex_value_output_first Iterator pointing to the vertex property variables for the * first (inclusive) vertex (assigned to tihs process in multi-GPU). `vertex_value_output_last` @@ -840,23 +841,24 @@ void copy_v_transform_reduce_in_nbr(raft::handle_t const& handle, * graph_view.get_number_of_local_vertices(). */ template -void copy_v_transform_reduce_out_nbr(raft::handle_t const& handle, - GraphViewType const& graph_view, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, - EdgeOp e_op, - T init, - VertexValueOutputIterator vertex_value_output_first) +void copy_v_transform_reduce_out_nbr( + raft::handle_t const& handle, + GraphViewType const& graph_view, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgeOp e_op, + T init, + VertexValueOutputIterator vertex_value_output_first) { detail::copy_v_transform_reduce_nbr(handle, graph_view, - adj_matrix_row_value_input, - adj_matrix_col_value_input, + edge_partition_src_value_input, + edge_partition_dst_value_input, e_op, init, vertex_value_output_first); diff --git a/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh b/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh index 32f1407597d..a536ffc4940 100644 --- a/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh +++ b/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh @@ -39,14 +39,14 @@ namespace cugraph { namespace detail { // a workaround for cudaErrorInvalidDeviceFunction error when device lambda is used -template +template struct minor_to_key_t { - using vertex_t = typename AdjMatrixColKeyInputWrapper::value_type; - AdjMatrixColKeyInputWrapper adj_matrix_col_key_input{}; + using vertex_t = typename EdgePartitionDstKeyInputWrapper::value_type; + EdgePartitionDstKeyInputWrapper edge_partition_dst_key_input{}; vertex_t minor_first{}; __device__ vertex_t operator()(vertex_t minor) const { - return adj_matrix_col_key_input.get(minor - minor_first); + return edge_partition_dst_key_input.get(minor - minor_first); } }; @@ -81,12 +81,12 @@ struct pair_to_binary_partition_id_t { // a workaround for cudaErrorInvalidDeviceFunction error when device lambda is used template struct call_key_aggregated_e_op_t { - AdjMatrixRowValueInputWrapper matrix_partition_row_value_input{}; + EdgePartitionSrcValueInputWrapper matrix_partition_row_value_input{}; KeyAggregatedEdgeOp key_aggregated_e_op{}; MatrixPartitionDeviceView matrix_partition{}; StaticMapDeviceView kv_map{}; @@ -152,11 +152,12 @@ struct reduce_with_init_t { * support two level reduction for every vertex. * * @tparam GraphViewType Type of the passed non-owning graph object. - * @tparam AdjMatrixRowValueInputWrapper Type of the wrapper for graph adjacency matrix row input - * properties. - * @tparam AdjMatrixColKeyInputWrapper Type of the wrapper for graph adjacency matrix column keys. - * @tparam VertexIterator Type of the iterator for graph adjacency matrix column key values for - * aggregation (key type should coincide with vertex type). + * @tparam EdgePartitionSrcValueInputWrapper Type of the wrapper for edge partition source property + * values. + * @tparam EdgePartitionDstKeyInputWrapper Type of the wrapper for edge partition destination key + * values. + * @tparam VertexIterator Type of the iterator for keys in (key, value) pairs (key type should + * coincide with vertex type). * @tparam ValueIterator Type of the iterator for values in (key, value) pairs. * @tparam KeyAggregatedEdgeOp Type of the quinary key-aggregated edge operator. * @tparam ReduceOp Type of the binary reduction operator. @@ -165,26 +166,26 @@ struct reduce_with_init_t { * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Non-owning graph object. - * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties - * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) - * or cugraph::dummy_:property_t:device_view() (if @p e_op does not access row properties). Use - * update_edge_partition_src_property to fill the wrapper. - * @param adj_matrix_col_key_input Device-copyable wrapper used to access column keys (for the - * columns assigned to this process in multi-GPU). Use either + * @param edge_partition_src_value_input Device-copyable wrapper used to access source input + * property values (for the edge sources assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access source property + * values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access source property + * values). Use update_edge_partition_src_property to fill the wrapper. + * @param edge_partition_dst_key_input Device-copyable wrapper used to access destination input key + * values (for the edge destinations assigned to this process in multi-GPU). Use * cugraph::edge_partition_dst_property_t::device_view(). Use update_edge_partition_dst_property to * fill the wrapper. * @param map_unique_key_first Iterator pointing to the first (inclusive) key in (key, value) pairs * (assigned to this process in multi-GPU, `cugraph::detail::compute_gpu_id_from_vertex_t` is used * to map keys to processes). (Key, value) pairs may be provided by - * transform_reduce_by_adj_matrix_row_key_e() or transform_reduce_by_adj_matrix_col_key_e(). + * transform_reduce_by_src_key_e() or transform_reduce_by_dst_key_e(). * @param map_unique_key_last Iterator pointing to the last (exclusive) key in (key, value) pairs * (assigned to this process in multi-GPU). * @param map_value_first Iterator pointing to the first (inclusive) value in (key, value) pairs * (assigned to this process in multi-GPU). `map_value_last` (exclusive) is deduced as @p * map_value_first + thrust::distance(@p map_unique_key_first, @p map_unique_key_last). * @param key_aggregated_e_op Quinary operator takes edge source, key, aggregated edge weight, *(@p - * adj_matrix_row_value_input_first + i), and value for the key stored in the input (key, value) + * edge_partition_src_value_input_first + i), and value for the key stored in the input (key, value) * pairs provided by @p map_unique_key_first, @p map_unique_key_last, and @p map_value_first * (aggregated over the entire set of processes in multi-GPU). * @param reduce_op Binary operator takes two input arguments and reduce the two variables to one. @@ -195,8 +196,8 @@ struct reduce_with_init_t { * graph_view.get_number_of_local_vertices(). */ template {adj_matrix_col_key_input, - matrix_partition.get_minor_first()}); + detail::minor_to_key_t{ + edge_partition_dst_key_input, matrix_partition.get_minor_first()}); // to limit memory footprint ((1 << 20) is a tuning parameter) auto approx_edges_to_sort_per_iteration = @@ -578,7 +579,7 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( auto tmp_e_op_result_buffer = allocate_dataframe_buffer(tmp_majors.size(), handle.get_stream()); - auto matrix_partition_row_value_input = adj_matrix_row_value_input; + auto matrix_partition_row_value_input = edge_partition_src_value_input; matrix_partition_row_value_input.set_local_adj_matrix_partition_idx(i); auto triplet_first = thrust::make_zip_iterator(thrust::make_tuple( @@ -589,7 +590,7 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( get_dataframe_buffer_begin(tmp_e_op_result_buffer), detail::call_key_aggregated_e_op_t{ diff --git a/cpp/include/cugraph/prims/count_if_e.cuh b/cpp/include/cugraph/prims/count_if_e.cuh index 1cabe0519ec..030a194ea71 100644 --- a/cpp/include/cugraph/prims/count_if_e.cuh +++ b/cpp/include/cugraph/prims/count_if_e.cuh @@ -31,38 +31,38 @@ namespace cugraph { * This function is inspired by thrust::count_if(). * * @tparam GraphViewType Type of the passed non-owning graph object. - * @tparam AdjMatrixRowValueInputWrapper Type of the wrapper for graph adjacency matrix row input - * properties. - * @tparam AdjMatrixColValueInputWrapper Type of the wrapper for graph adjacency matrix column input - * properties. + * @tparam EdgePartitionSrcValueInputWrapper Type of the wrapper for edge partition source property + * values. + * @tparam EdgePartitionDstValueInputWrapper Type of the wrapper for edge partition destination + * property values. * @tparam EdgeOp Type of the quaternary (or quinary) edge operator. * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Non-owning graph object. - * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties - * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) - * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use - * update_edge_partition_src_property to fill the wrapper. - * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties - * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column - * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column - * properties). Use update_edge_partition_dst_property to fill the wrapper. + * @param edge_partition_src_value_input Device-copyable wrapper used to access source input + * property values (for the edge sources assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access source property + * values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access source property + * values). Use update_edge_partition_src_property to fill the wrapper. + * @param edge_partition_dst_value_input Device-copyable wrapper used to access destination input + * property values (for the edge destinations assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access destination + * property values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access + * destination property values). Use update_edge_partition_dst_property to fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge - * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) - * and returns true if this edge should be included in the returned count. + * weight), property values for the source, and property values for the destination and returns if + * this edge should be included in the returned count. * @return GraphViewType::edge_type Number of times @p e_op returned true. */ template typename GraphViewType::edge_type count_if_e( raft::handle_t const& handle, GraphViewType const& graph_view, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgeOp e_op) { using vertex_t = typename GraphViewType::vertex_type; @@ -70,12 +70,12 @@ typename GraphViewType::edge_type count_if_e( return transform_reduce_e(handle, graph_view, - adj_matrix_row_value_input, - adj_matrix_col_value_input, + edge_partition_src_value_input, + edge_partition_dst_value_input, cast_edge_op_bool_to_integer{e_op}, edge_t{0}); diff --git a/cpp/include/cugraph/prims/extract_if_e.cuh b/cpp/include/cugraph/prims/extract_if_e.cuh index d5b6f2b3a71..18b0925ea7b 100644 --- a/cpp/include/cugraph/prims/extract_if_e.cuh +++ b/cpp/include/cugraph/prims/extract_if_e.cuh @@ -37,8 +37,8 @@ namespace cugraph { namespace detail { template struct call_e_op_t { matrix_partition_device_view_t matrix_partition{}; - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input{}; - AdjMatrixColValueInputWrapper adj_matrix_col_value_input{}; + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input{}; + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input{}; EdgeOp e_op{}; template @@ -70,14 +70,14 @@ struct call_e_op_t { auto col_offset = GraphViewType::is_adj_matrix_transposed ? major_offset : minor_offset; return !evaluate_edge_op() .compute(row, col, weight, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); } }; @@ -91,44 +91,44 @@ struct call_e_op_t { * This function is inspired by thrust::copy_if & thrust::remove_if(). * * @tparam GraphViewType Type of the passed non-owning graph object. - * @tparam AdjMatrixRowValueInputWrapper Type of the wrapper for graph adjacency matrix row input - * properties. - * @tparam AdjMatrixColValueInputWrapper Type of the wrapper for graph adjacency matrix column input - * properties. + * @tparam EdgePartitionSrcValueInputWrapper Type of the wrapper for edge partition source property + * values. + * @tparam EdgePartitionDstValueInputWrapper Type of the wrapper for edge partition destination + * property values. * @tparam EdgeOp Type of the quaternary (or quinary) edge operator. * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Non-owning graph object. - * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties - * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) - * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use - * update_edge_partition_src_property to fill the wrapper. - * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties - * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column - * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column - * properties). Use update_edge_partition_dst_property to fill the wrapper. + * @param edge_partition_src_value_input Device-copyable wrapper used to access source input + * property values (for the edge sources assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access source property + * values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access source property + * values). Use update_edge_partition_src_property to fill the wrapper. + * @param edge_partition_dst_value_input Device-copyable wrapper used to access destination input + * property values (for the edge destinations assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access destination + * property values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access + * destination property values). Use update_edge_partition_dst_property to fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge - * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) - * and returns a boolean value to designate whether to include this edge in the returned edge list - * (if true is returned) or not (if false is returned). + * weight), property values for the source, and property values for the destination and returns a + * boolean value to designate whether to include this edge in the returned edge list (if true is + * returned) or not (if false is returned). * @return std::tuple, * rmm::device_uvector, * std::optional>> Tuple storing an * extracted edge list (sources, destinations, and optional weights). */ template std::tuple, rmm::device_uvector, std::optional>> extract_if_e(raft::handle_t const& handle, GraphViewType const& graph_view, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgeOp e_op) { using vertex_t = typename GraphViewType::vertex_type; @@ -157,8 +157,8 @@ extract_if_e(raft::handle_t const& handle, matrix_partition_device_view_t( graph_view.get_matrix_partition_view(i)); - auto matrix_partition_row_value_input = adj_matrix_row_value_input; - auto matrix_partition_col_value_input = adj_matrix_col_value_input; + auto matrix_partition_row_value_input = edge_partition_src_value_input; + auto matrix_partition_col_value_input = edge_partition_dst_value_input; if constexpr (GraphViewType::is_adj_matrix_transposed) { matrix_partition_col_value_input.set_local_adj_matrix_partition_idx(i); } else { @@ -182,8 +182,8 @@ extract_if_e(raft::handle_t const& handle, edge_first + cur_size, edge_first + cur_size + edgelist_edge_counts[i], detail::call_e_op_t{matrix_partition, matrix_partition_row_value_input, matrix_partition_col_value_input, @@ -193,15 +193,16 @@ extract_if_e(raft::handle_t const& handle, thrust::make_tuple(edgelist_majors.begin(), edgelist_minors.begin())); cur_size += static_cast(thrust::distance( edge_first + cur_size, - thrust::remove_if( - handle.get_thrust_policy(), - edge_first + cur_size, - edge_first + cur_size + edgelist_edge_counts[i], - detail::call_e_op_t{ - matrix_partition, adj_matrix_row_value_input, adj_matrix_col_value_input, e_op}))); + thrust::remove_if(handle.get_thrust_policy(), + edge_first + cur_size, + edge_first + cur_size + edgelist_edge_counts[i], + detail::call_e_op_t{matrix_partition, + edge_partition_src_value_input, + edge_partition_dst_value_input, + e_op}))); } } diff --git a/cpp/include/cugraph/prims/property_op_utils.cuh b/cpp/include/cugraph/prims/property_op_utils.cuh index 1168617ae63..61be94c4414 100644 --- a/cpp/include/cugraph/prims/property_op_utils.cuh +++ b/cpp/include/cugraph/prims/property_op_utils.cuh @@ -96,14 +96,14 @@ struct edge_op_result_type< template struct evaluate_edge_op { using vertex_type = typename GraphViewType::vertex_type; using weight_type = typename GraphViewType::weight_type; - using row_value_type = typename AdjMatrixRowValueInputWrapper::value_type; - using col_value_type = typename AdjMatrixColValueInputWrapper::value_type; + using row_value_type = typename EdgePartitionSrcValueInputWrapper::value_type; + using col_value_type = typename EdgePartitionDstValueInputWrapper::value_type; using result_type = typename detail:: edge_op_result_type:: type; @@ -139,16 +139,16 @@ struct evaluate_edge_op { template struct cast_edge_op_bool_to_integer { static_assert(std::is_integral::value); using vertex_type = typename GraphViewType::vertex_type; using weight_type = typename GraphViewType::weight_type; - using row_value_type = typename AdjMatrixRowValueInputWrapper::value_type; - using col_value_type = typename AdjMatrixColValueInputWrapper::value_type; + using row_value_type = typename EdgePartitionSrcValueInputWrapper::value_type; + using col_value_type = typename EdgePartitionDstValueInputWrapper::value_type; EdgeOp e_op{}; diff --git a/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh b/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh index 67487ec4b6f..bd47562b49e 100644 --- a/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh @@ -32,13 +32,13 @@ namespace cugraph { namespace detail { // FIXME: block size requires tuning -int32_t constexpr transform_reduce_by_adj_matrix_row_col_key_e_for_all_block_size = 128; +int32_t constexpr transform_reduce_by_src_dst_key_e_for_all_block_size = 128; -template __device__ void update_buffer_element( @@ -49,9 +49,9 @@ __device__ void update_buffer_element( typename GraphViewType::vertex_type major, typename GraphViewType::vertex_type minor, typename GraphViewType::weight_type weight, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, - AdjMatrixRowColKeyInputWrapper adj_matrix_row_col_key_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input, EdgeOp e_op, typename GraphViewType::vertex_type* key, T* value) @@ -65,26 +65,27 @@ __device__ void update_buffer_element( auto row_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : major_offset; auto col_offset = GraphViewType::is_adj_matrix_transposed ? major_offset : minor_offset; - *key = adj_matrix_row_col_key_input.get(( - (GraphViewType::is_adj_matrix_transposed != adj_matrix_row_key) ? major_offset : minor_offset)); + *key = edge_partition_src_dst_key_input.get( + ((GraphViewType::is_adj_matrix_transposed != edge_partition_src_key) ? major_offset + : minor_offset)); *value = evaluate_edge_op() .compute(row, col, weight, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); } -template __global__ void for_all_major_for_all_nbr_hypersparse( @@ -93,9 +94,9 @@ __global__ void for_all_major_for_all_nbr_hypersparse( typename GraphViewType::weight_type, GraphViewType::is_multi_gpu> matrix_partition, typename GraphViewType::vertex_type major_hypersparse_first, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, - AdjMatrixRowColKeyInputWrapper adj_matrix_row_col_key_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input, EdgeOp e_op, typename GraphViewType::vertex_type* keys, T* values) @@ -123,14 +124,14 @@ __global__ void for_all_major_for_all_nbr_hypersparse( matrix_partition.get_local_edges(static_cast(major_idx)); auto local_offset = matrix_partition.get_local_offset(major_idx); for (edge_t i = 0; i < local_degree; ++i) { - update_buffer_element( + update_buffer_element( matrix_partition, major, indices[i], weights ? (*weights)[i] : weight_t{1.0}, - adj_matrix_row_value_input, - adj_matrix_col_value_input, - adj_matrix_row_col_key_input, + edge_partition_src_value_input, + edge_partition_dst_value_input, + edge_partition_src_dst_key_input, e_op, keys + local_offset + i, values + local_offset + i); @@ -140,11 +141,11 @@ __global__ void for_all_major_for_all_nbr_hypersparse( } } -template __global__ void for_all_major_for_all_nbr_low_degree( @@ -154,9 +155,9 @@ __global__ void for_all_major_for_all_nbr_low_degree( GraphViewType::is_multi_gpu> matrix_partition, typename GraphViewType::vertex_type major_first, typename GraphViewType::vertex_type major_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, - AdjMatrixRowColKeyInputWrapper adj_matrix_row_col_key_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input, EdgeOp e_op, typename GraphViewType::vertex_type* keys, T* values) @@ -180,14 +181,14 @@ __global__ void for_all_major_for_all_nbr_low_degree( matrix_partition.get_local_edges(static_cast(major_offset)); auto local_offset = matrix_partition.get_local_offset(major_offset); for (edge_t i = 0; i < local_degree; ++i) { - update_buffer_element( + update_buffer_element( matrix_partition, major, indices[i], weights ? (*weights)[i] : weight_t{1.0}, - adj_matrix_row_value_input, - adj_matrix_col_value_input, - adj_matrix_row_col_key_input, + edge_partition_src_value_input, + edge_partition_dst_value_input, + edge_partition_src_dst_key_input, e_op, keys + local_offset + i, values + local_offset + i); @@ -197,11 +198,11 @@ __global__ void for_all_major_for_all_nbr_low_degree( } } -template __global__ void for_all_major_for_all_nbr_mid_degree( @@ -211,9 +212,9 @@ __global__ void for_all_major_for_all_nbr_mid_degree( GraphViewType::is_multi_gpu> matrix_partition, typename GraphViewType::vertex_type major_first, typename GraphViewType::vertex_type major_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, - AdjMatrixRowColKeyInputWrapper adj_matrix_row_col_key_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input, EdgeOp e_op, typename GraphViewType::vertex_type* keys, T* values) @@ -223,8 +224,7 @@ __global__ void for_all_major_for_all_nbr_mid_degree( using weight_t = typename GraphViewType::weight_type; auto const tid = threadIdx.x + blockIdx.x * blockDim.x; - static_assert( - transform_reduce_by_adj_matrix_row_col_key_e_for_all_block_size % raft::warp_size() == 0); + static_assert(transform_reduce_by_src_dst_key_e_for_all_block_size % raft::warp_size() == 0); auto const lane_id = tid % raft::warp_size(); auto major_start_offset = static_cast(major_first - matrix_partition.get_major_first()); size_t idx = static_cast(tid / raft::warp_size()); @@ -240,14 +240,14 @@ __global__ void for_all_major_for_all_nbr_mid_degree( matrix_partition.get_local_edges(static_cast(major_offset)); auto local_offset = matrix_partition.get_local_offset(major_offset); for (edge_t i = lane_id; i < local_degree; i += raft::warp_size()) { - update_buffer_element( + update_buffer_element( matrix_partition, major, indices[i], weights ? (*weights)[i] : weight_t{1.0}, - adj_matrix_row_value_input, - adj_matrix_col_value_input, - adj_matrix_row_col_key_input, + edge_partition_src_value_input, + edge_partition_dst_value_input, + edge_partition_src_dst_key_input, e_op, keys + local_offset + i, values + local_offset + i); @@ -257,11 +257,11 @@ __global__ void for_all_major_for_all_nbr_mid_degree( } } -template __global__ void for_all_major_for_all_nbr_high_degree( @@ -271,9 +271,9 @@ __global__ void for_all_major_for_all_nbr_high_degree( GraphViewType::is_multi_gpu> matrix_partition, typename GraphViewType::vertex_type major_first, typename GraphViewType::vertex_type major_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, - AdjMatrixRowColKeyInputWrapper adj_matrix_row_col_key_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input, EdgeOp e_op, typename GraphViewType::vertex_type* keys, T* values) @@ -296,14 +296,14 @@ __global__ void for_all_major_for_all_nbr_high_degree( matrix_partition.get_local_edges(static_cast(major_offset)); auto local_offset = matrix_partition.get_local_offset(major_offset); for (edge_t i = threadIdx.x; i < local_degree; i += blockDim.x) { - update_buffer_element( + update_buffer_element( matrix_partition, major, indices[i], weights ? (*weights)[i] : weight_t{1.0}, - adj_matrix_row_value_input, - adj_matrix_col_value_input, - adj_matrix_row_col_key_input, + edge_partition_src_value_input, + edge_partition_dst_value_input, + edge_partition_src_dst_key_input, e_op, keys + local_offset + i, values + local_offset + i); @@ -340,26 +340,26 @@ std::tuple, BufferType> reduce_to_unique_kv_pairs( return std::make_tuple(std::move(unique_keys), std::move(value_for_unique_key_buffer)); } -template std::tuple, decltype(allocate_dataframe_buffer(0, cudaStream_t{nullptr}))> -transform_reduce_by_adj_matrix_row_col_key_e( +transform_reduce_by_src_dst_key_e( raft::handle_t const& handle, GraphViewType const& graph_view, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, - AdjMatrixRowColKeyInputWrapper adj_matrix_row_col_key_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgePartitionSrcDstKeyInputWrapper edge_partition_src_dst_key_input, EdgeOp e_op, T init) { static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); - static_assert(std::is_same::value); using vertex_t = typename GraphViewType::vertex_type; @@ -389,16 +389,16 @@ transform_reduce_by_adj_matrix_row_col_key_e( auto tmp_value_buffer = allocate_dataframe_buffer(tmp_keys.size(), handle.get_stream()); if (graph_view.get_vertex_partition_size(comm_root_rank) > 0) { - auto matrix_partition_row_value_input = adj_matrix_row_value_input; - auto matrix_partition_col_value_input = adj_matrix_col_value_input; + auto matrix_partition_row_value_input = edge_partition_src_value_input; + auto matrix_partition_col_value_input = edge_partition_dst_value_input; if constexpr (GraphViewType::is_adj_matrix_transposed) { matrix_partition_col_value_input.set_local_adj_matrix_partition_idx(i); } else { matrix_partition_row_value_input.set_local_adj_matrix_partition_idx(i); } - auto matrix_partition_row_col_key_input = adj_matrix_row_col_key_input; - if constexpr ((adj_matrix_row_key && !GraphViewType::is_adj_matrix_transposed) || - (!adj_matrix_row_key && GraphViewType::is_adj_matrix_transposed)) { + auto matrix_partition_row_col_key_input = edge_partition_src_dst_key_input; + if constexpr ((edge_partition_src_key && !GraphViewType::is_adj_matrix_transposed) || + (!edge_partition_src_key && GraphViewType::is_adj_matrix_transposed)) { matrix_partition_row_col_key_input.set_local_adj_matrix_partition_idx(i); } @@ -411,9 +411,9 @@ transform_reduce_by_adj_matrix_row_col_key_e( if ((*segment_offsets)[1] > 0) { raft::grid_1d_block_t update_grid( (*segment_offsets)[1], - detail::transform_reduce_by_adj_matrix_row_col_key_e_for_all_block_size, + detail::transform_reduce_by_src_dst_key_e_for_all_block_size, handle.get_device_properties().maxGridSize[0]); - detail::for_all_major_for_all_nbr_high_degree + detail::for_all_major_for_all_nbr_high_degree <<>>( matrix_partition, matrix_partition.get_major_first(), @@ -428,9 +428,9 @@ transform_reduce_by_adj_matrix_row_col_key_e( if ((*segment_offsets)[2] - (*segment_offsets)[1] > 0) { raft::grid_1d_warp_t update_grid( (*segment_offsets)[2] - (*segment_offsets)[1], - detail::transform_reduce_by_adj_matrix_row_col_key_e_for_all_block_size, + detail::transform_reduce_by_src_dst_key_e_for_all_block_size, handle.get_device_properties().maxGridSize[0]); - detail::for_all_major_for_all_nbr_mid_degree + detail::for_all_major_for_all_nbr_mid_degree <<>>( matrix_partition, matrix_partition.get_major_first() + (*segment_offsets)[1], @@ -445,9 +445,9 @@ transform_reduce_by_adj_matrix_row_col_key_e( if ((*segment_offsets)[3] - (*segment_offsets)[2] > 0) { raft::grid_1d_thread_t update_grid( (*segment_offsets)[3] - (*segment_offsets)[2], - detail::transform_reduce_by_adj_matrix_row_col_key_e_for_all_block_size, + detail::transform_reduce_by_src_dst_key_e_for_all_block_size, handle.get_device_properties().maxGridSize[0]); - detail::for_all_major_for_all_nbr_low_degree + detail::for_all_major_for_all_nbr_low_degree <<>>( matrix_partition, matrix_partition.get_major_first() + (*segment_offsets)[2], @@ -463,9 +463,9 @@ transform_reduce_by_adj_matrix_row_col_key_e( (*(matrix_partition.get_dcs_nzd_vertex_count()) > 0)) { raft::grid_1d_thread_t update_grid( *(matrix_partition.get_dcs_nzd_vertex_count()), - detail::transform_reduce_by_adj_matrix_row_col_key_e_for_all_block_size, + detail::transform_reduce_by_src_dst_key_e_for_all_block_size, handle.get_device_properties().maxGridSize[0]); - detail::for_all_major_for_all_nbr_hypersparse + detail::for_all_major_for_all_nbr_hypersparse <<>>( matrix_partition, matrix_partition.get_major_first() + (*segment_offsets)[3], @@ -479,10 +479,10 @@ transform_reduce_by_adj_matrix_row_col_key_e( } else { raft::grid_1d_thread_t update_grid( matrix_partition.get_major_size(), - detail::transform_reduce_by_adj_matrix_row_col_key_e_for_all_block_size, + detail::transform_reduce_by_src_dst_key_e_for_all_block_size, handle.get_device_properties().maxGridSize[0]); - detail::for_all_major_for_all_nbr_low_degree + detail::for_all_major_for_all_nbr_low_degree <<>>( matrix_partition, matrix_partition.get_major_first(), @@ -556,68 +556,68 @@ transform_reduce_by_adj_matrix_row_col_key_e( * @brief Iterate over the entire set of edges and reduce @p 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 graph adjacency matrix rows. + * edges are determined by the edge sources. * * @tparam GraphViewType Type of the passed non-owning graph object. - * @tparam AdjMatrixRowValueInputWrapper Type of the wrapper for graph adjacency matrix row input - * properties. - * @tparam AdjMatrixColValueInputWrapper Type of the wrapper for graph adjacency matrix column input - * properties. - * @tparam AdjMatrixRowKeyInputWrapper Type of the wrapper for graph adjacency matrix row keys. + * @tparam EdgePartitionSrcValueInputWrapper Type of the wrapper for edge partition source property + * values. + * @tparam EdgePartitionDstValueInputWrapper Type of the wrapper for edge partition destination + * property values. + * @tparam EdgePartitionSrcKeyInputWrapper Type of the wrapper for edge partition source key values. * @tparam EdgeOp Type of the quaternary (or quinary) edge operator. * @tparam T Type of the values in (key, value) pairs. * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Non-owning graph object. - * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties - * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) - * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use - * update_edge_partition_src_property to fill the wrapper. - * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties - * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column - * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column - * properties). Use update_edge_partition_dst_property to fill the wrapper. - * @param adj_matrix_row_key_input Device-copyable wrapper used to access row keys(for the rows - * assigned to this process in multi-GPU). Use either + * @param edge_partition_src_value_input Device-copyable wrapper used to access source input + * property values (for the edge sources assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access source property + * values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access source property + * values). Use update_edge_partition_src_property to fill the wrapper. + * @param edge_partition_dst_value_input Device-copyable wrapper used to access destination input + * property values (for the edge destinations assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access destination + * property values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access + * destination property values). Use update_edge_partition_dst_property to fill the wrapper. + * @param edge_partition_src_key_input Device-copyable wrapper used to access source input key + * values (for the edge sources assigned to this process in multi-GPU). Use * cugraph::edge_partition_src_property_t::device_view(). Use update_edge_partition_src_property to * fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge - * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) - * and returns a transformed value to be reduced. - * @param init Initial value to be added to the value in each transform-reduced (key, value) pair. + * weight), property values for the source, and property values for the destination and returns a + * transformed value to be reduced to (source key, value) pairs. + * @param init Initial value to be added to the value in each transform-reduced (source key, value) + * pair. * @return std::tuple Tuple of rmm::device_uvector and * rmm::device_uvector (if T is arithmetic scalar) or a tuple of rmm::device_uvector objects (if * T is a thrust::tuple type of arithmetic scalar types, one rmm::device_uvector object per scalar * type). */ template -auto transform_reduce_by_adj_matrix_row_key_e( - raft::handle_t const& handle, - GraphViewType const& graph_view, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, - AdjMatrixRowKeyInputWrapper adj_matrix_row_key_input, - EdgeOp e_op, - T init) +auto transform_reduce_by_src_key_e(raft::handle_t const& handle, + GraphViewType const& graph_view, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgePartitionSrcKeyInputWrapper edge_partition_src_key_input, + EdgeOp e_op, + T init) { static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); - static_assert(std::is_same::value); - return detail::transform_reduce_by_adj_matrix_row_col_key_e(handle, - graph_view, - adj_matrix_row_value_input, - adj_matrix_col_value_input, - adj_matrix_row_key_input, - e_op, - init); + return detail::transform_reduce_by_src_dst_key_e(handle, + graph_view, + edge_partition_src_value_input, + edge_partition_dst_value_input, + edge_partition_src_key_input, + e_op, + init); } // FIXME: EdgeOp & VertexOp in update_frontier_v_push_if_out_nbr concatenates push inidicator or @@ -626,68 +626,69 @@ auto transform_reduce_by_adj_matrix_row_key_e( * @brief Iterate over the entire set of edges and reduce @p 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 graph adjacency matrix columns. + * edges are determined by the edge destinations. * * @tparam GraphViewType Type of the passed non-owning graph object. - * @tparam AdjMatrixRowValueInputWrapper Type of the wrapper for graph adjacency matrix row input - * properties. - * @tparam AdjMatrixColValueInputWrapper Type of the wrapper for graph adjacency matrix column input - * properties. - * @tparam AdjMatrixColKeyInputWrapper Type of the wrapper for graph adjacency matrix column keys. + * @tparam EdgePartitionSrcValueInputWrapper Type of the wrapper for edge partition source property + * values. + * @tparam EdgePartitionDstValueInputWrapper Type of the wrapper for edge partition destination + * property values. + * @tparam EdgePartitionDstKeyInputWrapper Type of the wrapper for edge partition destination key + * values. * @tparam EdgeOp Type of the quaternary (or quinary) edge operator. * @tparam T Type of the values in (key, value) pairs. * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Non-owning graph object. - * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties - * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) - * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use - * update_edge_partition_src_property to fill the wrapper. - * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties - * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column - * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column - * properties). Use update_edge_partition_dst_property to fill the wrapper. - * @param adj_matrix_col_key_input Device-copyable wrapper used to access column keys(for the - * columns assigned to this process in multi-GPU). Use either + * @param edge_partition_src_value_input Device-copyable wrapper used to access source input + * property values (for the edge sources assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access source property + * values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access source property + * values). Use update_edge_partition_src_property to fill the wrapper. + * @param edge_partition_dst_value_input Device-copyable wrapper used to access destination input + * property values (for the edge destinations assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access destination + * property values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access + * destination property values). Use update_edge_partition_dst_property to fill the wrapper. + * @param edge_partition_dst_key_input Device-copyable wrapper used to access destination input key + * values (for the edge destinations assigned to this process in multi-GPU). Use * cugraph::edge_partition_dst_property_t::device_view(). Use update_edge_partition_dst_property to * fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge - * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) - * and returns a transformed value to be reduced. - * @param init Initial value to be added to the value in each transform-reduced (key, value) pair. + * weight), property values for the source, and property values for the destination and returns a + * transformed value to be reduced to (destination key, value) pairs. + * @param init Initial value to be added to the value in each transform-reduced (destination key, + * value) pair. * @return std::tuple Tuple of rmm::device_uvector and * rmm::device_uvector (if T is arithmetic scalar) or a tuple of rmm::device_uvector objects (if * T is a thrust::tuple type of arithmetic scalar types, one rmm::device_uvector object per scalar * type). */ template -auto transform_reduce_by_adj_matrix_col_key_e( - raft::handle_t const& handle, - GraphViewType const& graph_view, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, - AdjMatrixColKeyInputWrapper adj_matrix_col_key_input, - EdgeOp e_op, - T init) +auto transform_reduce_by_dst_key_e(raft::handle_t const& handle, + GraphViewType const& graph_view, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, + EdgePartitionDstKeyInputWrapper edge_partition_dst_key_input, + EdgeOp e_op, + T init) { static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); - static_assert(std::is_same::value); - return detail::transform_reduce_by_adj_matrix_row_col_key_e(handle, - graph_view, - adj_matrix_row_value_input, - adj_matrix_col_value_input, - adj_matrix_col_key_input, - e_op, - init); + return detail::transform_reduce_by_src_dst_key_e(handle, + graph_view, + edge_partition_src_value_input, + edge_partition_dst_value_input, + edge_partition_dst_key_input, + e_op, + init); } } // namespace cugraph diff --git a/cpp/include/cugraph/prims/transform_reduce_e.cuh b/cpp/include/cugraph/prims/transform_reduce_e.cuh index 89d783dd845..f976e908d92 100644 --- a/cpp/include/cugraph/prims/transform_reduce_e.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_e.cuh @@ -40,8 +40,8 @@ namespace detail { int32_t constexpr transform_reduce_e_for_all_block_size = 128; template __global__ void for_all_major_for_all_nbr_hypersparse( @@ -50,8 +50,8 @@ __global__ void for_all_major_for_all_nbr_hypersparse( typename GraphViewType::weight_type, GraphViewType::is_multi_gpu> matrix_partition, typename GraphViewType::vertex_type major_hypersparse_first, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, ResultIterator result_iter /* size 1 */, EdgeOp e_op) { @@ -86,8 +86,8 @@ __global__ void for_all_major_for_all_nbr_hypersparse( thrust::make_counting_iterator(edge_t{0}), thrust::make_counting_iterator(local_degree), [&matrix_partition, - &adj_matrix_row_value_input, - &adj_matrix_col_value_input, + &edge_partition_src_value_input, + &edge_partition_dst_value_input, &e_op, major, indices, @@ -106,14 +106,14 @@ __global__ void for_all_major_for_all_nbr_hypersparse( : minor_offset; return evaluate_edge_op() .compute(row, col, weight, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); }, e_op_result_t{}, @@ -128,8 +128,8 @@ __global__ void for_all_major_for_all_nbr_hypersparse( } template __global__ void for_all_major_for_all_nbr_low_degree( @@ -139,8 +139,8 @@ __global__ void for_all_major_for_all_nbr_low_degree( GraphViewType::is_multi_gpu> matrix_partition, typename GraphViewType::vertex_type major_first, typename GraphViewType::vertex_type major_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, ResultIterator result_iter /* size 1 */, EdgeOp e_op) { @@ -169,8 +169,8 @@ __global__ void for_all_major_for_all_nbr_low_degree( thrust::make_counting_iterator(edge_t{0}), thrust::make_counting_iterator(local_degree), [&matrix_partition, - &adj_matrix_row_value_input, - &adj_matrix_col_value_input, + &edge_partition_src_value_input, + &edge_partition_dst_value_input, &e_op, major_offset, indices, @@ -192,14 +192,14 @@ __global__ void for_all_major_for_all_nbr_low_degree( : minor_offset; return evaluate_edge_op() .compute(row, col, weight, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); }, e_op_result_t{}, @@ -214,8 +214,8 @@ __global__ void for_all_major_for_all_nbr_low_degree( } template __global__ void for_all_major_for_all_nbr_mid_degree( @@ -225,8 +225,8 @@ __global__ void for_all_major_for_all_nbr_mid_degree( GraphViewType::is_multi_gpu> matrix_partition, typename GraphViewType::vertex_type major_first, typename GraphViewType::vertex_type major_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, ResultIterator result_iter /* size 1 */, EdgeOp e_op) { @@ -269,14 +269,14 @@ __global__ void for_all_major_for_all_nbr_mid_degree( : minor_offset; auto e_op_result = evaluate_edge_op() .compute(row, col, weight, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); e_op_result_sum = edge_property_add(e_op_result_sum, e_op_result); } @@ -288,8 +288,8 @@ __global__ void for_all_major_for_all_nbr_mid_degree( } template __global__ void for_all_major_for_all_nbr_high_degree( @@ -299,8 +299,8 @@ __global__ void for_all_major_for_all_nbr_high_degree( GraphViewType::is_multi_gpu> matrix_partition, typename GraphViewType::vertex_type major_first, typename GraphViewType::vertex_type major_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, ResultIterator result_iter /* size 1 */, EdgeOp e_op) { @@ -340,14 +340,14 @@ __global__ void for_all_major_for_all_nbr_high_degree( : minor_offset; auto e_op_result = evaluate_edge_op() .compute(row, col, weight, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); e_op_result_sum = edge_property_add(e_op_result_sum, e_op_result); } @@ -366,40 +366,40 @@ __global__ void for_all_major_for_all_nbr_high_degree( * This function is inspired by thrust::transform_reduce(). * * @tparam GraphViewType Type of the passed non-owning graph object. - * @tparam AdjMatrixRowValueInputWrapper Type of the wrapper for graph adjacency matrix row input - * properties. - * @tparam AdjMatrixColValueInputWrapper Type of the wrapper for graph adjacency matrix column input - * properties. + * @tparam EdgePartitionSrcValueInputWrapper Type of the wrapper for edge partition source property + * values. + * @tparam EdgePartitionDstValueInputWrapper Type of the wrapper for edge partition destination + * property values. * @tparam EdgeOp Type of the quaternary (or quinary) edge operator. * @tparam T Type of the initial value. * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Non-owning graph object. - * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties - * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) - * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use - * update_edge_partition_src_property to fill the wrapper. - * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties - * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column - * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column - * properties). Use update_edge_partition_dst_property to fill the wrapper. + * @param edge_partition_src_value_input Device-copyable wrapper used to access source input + * property values (for the edge sources assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access source property + * values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access source property + * values). Use update_edge_partition_src_property to fill the wrapper. + * @param edge_partition_dst_value_input Device-copyable wrapper used to access destination input + * property values (for the edge destinations assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access destination + * property values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access + * destination property values). Use update_edge_partition_dst_property to fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge - * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) - * and returns a value to be reduced. + * weight), property values for the source, and property values for the destination and returns a + * value to be reduced. * @param init Initial value to be added to the transform-reduced input vertex properties. * @return T Reduction of the @p edge_op outputs. */ template T transform_reduce_e(raft::handle_t const& handle, GraphViewType const& graph_view, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgeOp e_op, T init) { @@ -422,8 +422,8 @@ T transform_reduce_e(raft::handle_t const& handle, matrix_partition_device_view_t( graph_view.get_matrix_partition_view(i)); - auto matrix_partition_row_value_input = adj_matrix_row_value_input; - auto matrix_partition_col_value_input = adj_matrix_col_value_input; + auto matrix_partition_row_value_input = edge_partition_src_value_input; + auto matrix_partition_col_value_input = edge_partition_dst_value_input; if constexpr (GraphViewType::is_adj_matrix_transposed) { matrix_partition_col_value_input.set_local_adj_matrix_partition_idx(i); } else { diff --git a/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh b/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh index 32a65c90a59..0d39e1aa5ae 100644 --- a/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh +++ b/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh @@ -456,8 +456,9 @@ void update_edge_partition_minor_property( * (inclusive) vertex (of the vertex partition assigned to this process in multi-GPU). * `vertex_property_input_last` (exclusive) is deduced as @p vertex_property_input_first + @p * graph_view.get_number_of_local_vertices(). - * @param edge_partition_source_property_output Wrapper used to store edge partition source property - * values (for the edge partitions assigned to this process in multi-GPU). + * @param edge_partition_src_property_output Device-copyable wrapper used to store source property + * values (for the edge sources assigned to this process in multi-GPU). Use + * cugraph::edge_partition_src_property_t::device_view(). */ template void update_edge_partition_src_property( @@ -498,8 +499,9 @@ void update_edge_partition_src_property( * (inclusive) vertex (of the vertex partition assigned to this process in multi-GPU). * `vertex_property_input_last` (exclusive) is deduced as @p vertex_property_input_first + @p * graph_view.get_number_of_local_vertices(). - * @param edge_partition_source_property_output Wrapper used to store edge partition source property - * values (for the edge partitions assigned to this process in multi-GPU). + * @param edge_partition_src_property_output Device-copyable wrapper used to store source property + * values (for the edge sources assigned to this process in multi-GPU). Use + * cugraph::edge_partition_src_property_t::device_view(). */ template void update_edge_partition_src_property( @@ -546,8 +548,9 @@ void update_edge_partition_src_property( * (inclusive) vertex (of the vertex partition assigned to this process in multi-GPU). * `vertex_property_input_last` (exclusive) is deduced as @p vertex_property_input_first + @p * graph_view.get_number_of_local_vertices(). - * @param edge_partition_dst_property_output Wrapper used to store edge partition source property - * values (for the edge partitions assigned to this process in multi-GPU). + * @param edge_partition_dst_property_output Device-copyable wrapper used to store destination + * property values (for the edge destinations assigned to this process in multi-GPU). Use + * cugraph::edge_partition_dst_property_t::device_view(). */ template void update_edge_partition_dst_property( @@ -589,9 +592,9 @@ void update_edge_partition_dst_property( * (inclusive) vertex (of the vertex partition assigned to this process in multi-GPU). * `vertex_property_input_last` (exclusive) is deduced as @p vertex_property_input_first + @p * graph_view.get_number_of_local_vertices(). - * @param edge_partition_dst_property_output Wrapper used to store edge partition source property - * values (for the edge partitions assigned to this process in multi-GPU). (for the columns assigned - * to this process in multi-GPU). + * @param edge_partition_dst_property_output Device-copyable wrapper used to store destination + * property values (for the edge destinations assigned to this process in multi-GPU). Use + * cugraph::edge_partition_dst_property_t::device_view(). */ template void update_edge_partition_dst_property( diff --git a/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh b/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh index a55f5cfe86a..fc00a140577 100644 --- a/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh +++ b/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh @@ -211,8 +211,8 @@ __device__ void push_buffer_element(vertex_t col, template @@ -224,8 +224,8 @@ __global__ void for_all_frontier_row_for_all_nbr_hypersparse( typename GraphViewType::vertex_type major_hypersparse_first, KeyIterator key_first, KeyIterator key_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, BufferKeyOutputIterator buffer_key_output_first, BufferPayloadOutputIterator buffer_payload_output_first, size_t* buffer_idx_ptr, @@ -241,8 +241,8 @@ __global__ void for_all_frontier_row_for_all_nbr_hypersparse( typename optional_payload_buffer_value_type_t::value; using e_op_result_t = typename evaluate_edge_op::result_type; static_assert(!GraphViewType::is_adj_matrix_transposed, @@ -343,14 +343,14 @@ __global__ void for_all_frontier_row_for_all_nbr_hypersparse( auto col_offset = matrix_partition.get_minor_offset_from_minor_nocheck(col); e_op_result = evaluate_edge_op() .compute(key, col, weights ? (*weights)[local_edge_offset] : weight_t{1.0}, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); } auto ballot_e_op = @@ -381,8 +381,8 @@ __global__ void for_all_frontier_row_for_all_nbr_hypersparse( template @@ -393,8 +393,8 @@ __global__ void for_all_frontier_row_for_all_nbr_low_degree( GraphViewType::is_multi_gpu> matrix_partition, KeyIterator key_first, KeyIterator key_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, BufferKeyOutputIterator buffer_key_output_first, BufferPayloadOutputIterator buffer_payload_output_first, size_t* buffer_idx_ptr, @@ -410,8 +410,8 @@ __global__ void for_all_frontier_row_for_all_nbr_low_degree( typename optional_payload_buffer_value_type_t::value; using e_op_result_t = typename evaluate_edge_op::result_type; static_assert(!GraphViewType::is_adj_matrix_transposed, @@ -503,14 +503,14 @@ __global__ void for_all_frontier_row_for_all_nbr_low_degree( auto col_offset = matrix_partition.get_minor_offset_from_minor_nocheck(col); e_op_result = evaluate_edge_op() .compute(key, col, weights ? (*weights)[local_edge_offset] : weight_t{1.0}, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); } auto ballot = __ballot_sync(uint32_t{0xffffffff}, e_op_result ? uint32_t{1} : uint32_t{0}); @@ -541,8 +541,8 @@ __global__ void for_all_frontier_row_for_all_nbr_low_degree( template @@ -553,8 +553,8 @@ __global__ void for_all_frontier_row_for_all_nbr_mid_degree( GraphViewType::is_multi_gpu> matrix_partition, KeyIterator key_first, KeyIterator key_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, BufferKeyOutputIterator buffer_key_output_first, BufferPayloadOutputIterator buffer_payload_output_first, size_t* buffer_idx_ptr, @@ -570,8 +570,8 @@ __global__ void for_all_frontier_row_for_all_nbr_mid_degree( typename optional_payload_buffer_value_type_t::value; using e_op_result_t = typename evaluate_edge_op::result_type; static_assert(!GraphViewType::is_adj_matrix_transposed, @@ -610,14 +610,14 @@ __global__ void for_all_frontier_row_for_all_nbr_mid_degree( auto col_offset = matrix_partition.get_minor_offset_from_minor_nocheck(col); e_op_result = evaluate_edge_op() .compute(key, col, weights ? (*weights)[i] : weight_t{1.0}, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); } auto ballot = __ballot_sync(uint32_t{0xffffffff}, e_op_result ? uint32_t{1} : uint32_t{0}); @@ -648,8 +648,8 @@ __global__ void for_all_frontier_row_for_all_nbr_mid_degree( template @@ -660,8 +660,8 @@ __global__ void for_all_frontier_row_for_all_nbr_high_degree( GraphViewType::is_multi_gpu> matrix_partition, KeyIterator key_first, KeyIterator key_last, - AdjMatrixRowValueInputWrapper adj_matrix_row_value_input, - AdjMatrixColValueInputWrapper adj_matrix_col_value_input, + EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, + EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, BufferKeyOutputIterator buffer_key_output_first, BufferPayloadOutputIterator buffer_payload_output_first, size_t* buffer_idx_ptr, @@ -677,8 +677,8 @@ __global__ void for_all_frontier_row_for_all_nbr_high_degree( typename optional_payload_buffer_value_type_t::value; using e_op_result_t = typename evaluate_edge_op::result_type; static_assert(!GraphViewType::is_adj_matrix_transposed, @@ -718,14 +718,14 @@ __global__ void for_all_frontier_row_for_all_nbr_high_degree( auto col_offset = matrix_partition.get_minor_offset_from_minor_nocheck(col); e_op_result = evaluate_edge_op() .compute(key, col, weights ? (*weights)[i] : weight_t{1.0}, - adj_matrix_row_value_input.get(row_offset), - adj_matrix_col_value_input.get(col_offset), + edge_partition_src_value_input.get(row_offset), + edge_partition_dst_value_input.get(col_offset), e_op); } BlockScan(temp_storage) @@ -941,10 +941,10 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( * @tparam GraphViewType Type of the passed non-owning graph object. * @tparam VertexFrontierType Type of the vertex frontier class which abstracts vertex frontier * managements. - * @tparam AdjMatrixRowValueInputWrapper Type of the wrapper for graph adjacency matrix row input - * properties. - * @tparam AdjMatrixColValueInputWrapper Type of the wrapper for graph adjacency matrix column input - * properties. + * @tparam EdgePartitionSrcValueInputWrapper Type of the wrapper for edge partition source property + * values. + * @tparam EdgePartitionDstValueInputWrapper Type of the wrapper for edge partition destination + * property values. * @tparam EdgeOp Type of the quaternary (or quinary) edge operator. * @tparam ReduceOp Type of the binary reduction operator. * @tparam VertexValueInputIterator Type of the iterator for vertex properties. @@ -959,19 +959,19 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( * current iteration. * @param next_frontier_bucket_indices Indices of the VertexFrontier buckets to store new frontier * vertices for the next iteration. - * @param adj_matrix_row_value_input Device-copyable wrapper used to access row input properties - * (for the rows assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access row properties) - * or cugraph::dummy_property_t::device_view() (if @p e_op does not access row properties). Use - * update_edge_partition_src_property to fill the wrapper. - * @param adj_matrix_col_value_input Device-copyable wrapper used to access column input properties - * (for the columns assigned to this process in multi-GPU). Use either - * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access column - * properties) or cugraph::dummy_property_t::device_view() (if @p e_op does not access column - * properties). Use update_edge_partition_dst_property to fill the wrapper. + * @param edge_partition_src_value_input Device-copyable wrapper used to access source input + * property values (for the edge sources assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_src_property_t::device_view() (if @p e_op needs to access source property + * values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access source property + * values). Use update_edge_partition_src_property to fill the wrapper. + * @param edge_partition_dst_value_input Device-copyable wrapper used to access destination input + * property values (for the edge destinations assigned to this process in multi-GPU). Use either + * cugraph::edge_partition_dst_property_t::device_view() (if @p e_op needs to access destination + * property values) or cugraph::dummy_property_t::device_view() (if @p e_op does not access + * destination property values). Use update_edge_partition_dst_property to fill the wrapper. * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge - * weight), properties for the row (i.e. source), and properties for the column (i.e. destination) - * and returns a value to be reduced the @p reduce_op. + * weight), property values for the source, and property values for the destination and returns a + * value to be reduced the @p reduce_op. * @param reduce_op Binary operator takes two input arguments and reduce the two variables to one. * @param vertex_value_input_first Iterator pointing to the vertex properties for the first * (inclusive) vertex (assigned to this process in multi-GPU). `vertex_value_input_last` (exclusive) @@ -988,8 +988,8 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( */ template Date: Fri, 11 Mar 2022 15:48:35 -0800 Subject: [PATCH 7/9] renumae transform_reduce_by_adj_matrix_row|col_key_e to transform_redcue_by_src|dst_key_e --- ... => transform_reduce_by_src_dst_key_e.cuh} | 0 cpp/src/community/louvain.cuh | 25 +++++++++---------- 2 files changed, 12 insertions(+), 13 deletions(-) rename cpp/include/cugraph/prims/{transform_reduce_by_adj_matrix_row_col_key_e.cuh => transform_reduce_by_src_dst_key_e.cuh} (100%) diff --git a/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh b/cpp/include/cugraph/prims/transform_reduce_by_src_dst_key_e.cuh similarity index 100% rename from cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh rename to cpp/include/cugraph/prims/transform_reduce_by_src_dst_key_e.cuh diff --git a/cpp/src/community/louvain.cuh b/cpp/src/community/louvain.cuh index e42102d6182..cbc4b83182d 100644 --- a/cpp/src/community/louvain.cuh +++ b/cpp/src/community/louvain.cuh @@ -23,7 +23,7 @@ #include #include #include -#include +#include #include #include #include @@ -540,18 +540,17 @@ class Louvain { handle_, current_graph_view_, next_clusters_v_.begin(), dst_clusters_cache_); } - std::tie(cluster_keys_v_, cluster_weights_v_) = - cugraph::transform_reduce_by_adj_matrix_row_key_e( - handle_, - current_graph_view_, - dummy_property_t{}.device_view(), - dummy_property_t{}.device_view(), - graph_view_t::is_multi_gpu - ? src_clusters_cache_.device_view() - : detail::edge_partition_major_property_device_view_t( - next_clusters_v_.data()), - detail::return_edge_weight_t{}, - weight_t{0}); + std::tie(cluster_keys_v_, cluster_weights_v_) = cugraph::transform_reduce_by_src_key_e( + handle_, + current_graph_view_, + dummy_property_t{}.device_view(), + dummy_property_t{}.device_view(), + graph_view_t::is_multi_gpu + ? src_clusters_cache_.device_view() + : detail::edge_partition_major_property_device_view_t( + next_clusters_v_.data()), + detail::return_edge_weight_t{}, + weight_t{0}); } void shrink_graph() From c6a8ba501f553a945e046e94aad8be00fd006017 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Fri, 11 Mar 2022 17:14:15 -0800 Subject: [PATCH 8/9] further replace row/col with src/dst in prims --- .../copy_v_transform_reduce_in_out_nbr.cuh | 92 +++---- ...ransform_reduce_key_aggregated_out_nbr.cuh | 10 +- cpp/include/cugraph/prims/extract_if_e.cuh | 28 +-- .../cugraph/prims/property_op_utils.cuh | 118 ++++----- .../transform_reduce_by_src_dst_key_e.cuh | 58 ++--- .../cugraph/prims/transform_reduce_e.cuh | 92 +++---- .../update_frontier_v_push_if_out_nbr.cuh | 232 +++++++++--------- 7 files changed, 315 insertions(+), 315 deletions(-) diff --git a/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh b/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh index 7d62d773976..a7698324a07 100644 --- a/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh +++ b/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh @@ -101,12 +101,12 @@ __global__ void for_all_major_for_all_nbr_hypersparse( auto minor = indices[i]; auto weight = weights ? (*weights)[i] : weight_t{1.0}; auto minor_offset = matrix_partition.get_minor_offset_from_minor_nocheck(minor); - auto row = GraphViewType::is_adj_matrix_transposed ? minor : major; - auto col = GraphViewType::is_adj_matrix_transposed ? major : minor; - auto row_offset = GraphViewType::is_adj_matrix_transposed + auto src = GraphViewType::is_adj_matrix_transposed ? minor : major; + auto dst = GraphViewType::is_adj_matrix_transposed ? major : minor; + auto src_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : static_cast(major_offset); - auto col_offset = GraphViewType::is_adj_matrix_transposed + auto dst_offset = GraphViewType::is_adj_matrix_transposed ? static_cast(major_offset) : minor_offset; return evaluate_edge_op() - .compute(row, - col, + .compute(src, + dst, weight, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); }; @@ -210,16 +210,16 @@ __global__ void for_all_major_for_all_nbr_low_degree( auto minor = indices[i]; auto weight = weights ? (*weights)[i] : weight_t{1.0}; auto minor_offset = matrix_partition.get_minor_offset_from_minor_nocheck(minor); - auto row = GraphViewType::is_adj_matrix_transposed + auto src = GraphViewType::is_adj_matrix_transposed ? minor : matrix_partition.get_major_from_major_offset_nocheck(major_offset); - auto col = GraphViewType::is_adj_matrix_transposed + auto dst = GraphViewType::is_adj_matrix_transposed ? matrix_partition.get_major_from_major_offset_nocheck(major_offset) : minor; - auto row_offset = GraphViewType::is_adj_matrix_transposed + auto src_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : static_cast(major_offset); - auto col_offset = GraphViewType::is_adj_matrix_transposed + auto dst_offset = GraphViewType::is_adj_matrix_transposed ? static_cast(major_offset) : minor_offset; return evaluate_edge_op() - .compute(row, - col, + .compute(src, + dst, weight, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); }; @@ -326,16 +326,16 @@ __global__ void for_all_major_for_all_nbr_mid_degree( auto minor = indices[i]; auto weight = weights ? (*weights)[i] : weight_t{1.0}; auto minor_offset = matrix_partition.get_minor_offset_from_minor_nocheck(minor); - auto row = GraphViewType::is_adj_matrix_transposed + auto src = GraphViewType::is_adj_matrix_transposed ? minor : matrix_partition.get_major_from_major_offset_nocheck(major_offset); - auto col = GraphViewType::is_adj_matrix_transposed + auto dst = GraphViewType::is_adj_matrix_transposed ? matrix_partition.get_major_from_major_offset_nocheck(major_offset) : minor; - auto row_offset = GraphViewType::is_adj_matrix_transposed + auto src_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : static_cast(major_offset); - auto col_offset = GraphViewType::is_adj_matrix_transposed + auto dst_offset = GraphViewType::is_adj_matrix_transposed ? static_cast(major_offset) : minor_offset; auto e_op_result = evaluate_edge_op() - .compute(row, - col, + .compute(src, + dst, weight, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); if constexpr (update_major) { e_op_result_sum = edge_property_add(e_op_result_sum, e_op_result); @@ -419,16 +419,16 @@ __global__ void for_all_major_for_all_nbr_high_degree( auto minor = indices[i]; auto weight = weights ? (*weights)[i] : weight_t{1.0}; auto minor_offset = matrix_partition.get_minor_offset_from_minor_nocheck(minor); - auto row = GraphViewType::is_adj_matrix_transposed + auto src = GraphViewType::is_adj_matrix_transposed ? minor : matrix_partition.get_major_from_major_offset_nocheck(major_offset); - auto col = GraphViewType::is_adj_matrix_transposed + auto dst = GraphViewType::is_adj_matrix_transposed ? matrix_partition.get_major_from_major_offset_nocheck(major_offset) : minor; - auto row_offset = GraphViewType::is_adj_matrix_transposed + auto src_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : static_cast(major_offset); - auto col_offset = GraphViewType::is_adj_matrix_transposed + auto dst_offset = GraphViewType::is_adj_matrix_transposed ? static_cast(major_offset) : minor_offset; auto e_op_result = evaluate_edge_op() - .compute(row, - col, + .compute(src, + dst, weight, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); if constexpr (update_major) { e_op_result_sum = edge_property_add(e_op_result_sum, e_op_result); @@ -536,12 +536,12 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, } } - auto matrix_partition_row_value_input = edge_partition_src_value_input; - auto matrix_partition_col_value_input = edge_partition_dst_value_input; + auto matrix_partition_src_value_input = edge_partition_src_value_input; + auto matrix_partition_dst_value_input = edge_partition_dst_value_input; if constexpr (GraphViewType::is_adj_matrix_transposed) { - matrix_partition_col_value_input.set_local_adj_matrix_partition_idx(i); + matrix_partition_dst_value_input.set_local_adj_matrix_partition_idx(i); } else { - matrix_partition_row_value_input.set_local_adj_matrix_partition_idx(i); + matrix_partition_src_value_input.set_local_adj_matrix_partition_idx(i); } std::conditional_t>>( matrix_partition, matrix_partition.get_major_first() + (*segment_offsets)[3], - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, segment_output_buffer, e_op, major_init); @@ -650,8 +650,8 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, matrix_partition, matrix_partition.get_major_first(), matrix_partition.get_major_last(), - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, output_buffer, e_op, major_init); diff --git a/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh b/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh index a536ffc4940..64eeb029e8e 100644 --- a/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh +++ b/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh @@ -86,7 +86,7 @@ template struct call_key_aggregated_e_op_t { - EdgePartitionSrcValueInputWrapper matrix_partition_row_value_input{}; + EdgePartitionSrcValueInputWrapper matrix_partition_src_value_input{}; KeyAggregatedEdgeOp key_aggregated_e_op{}; MatrixPartitionDeviceView matrix_partition{}; StaticMapDeviceView kv_map{}; @@ -99,7 +99,7 @@ struct call_key_aggregated_e_op_t { return key_aggregated_e_op(major, key, w, - matrix_partition_row_value_input.get( + matrix_partition_src_value_input.get( matrix_partition.get_major_offset_from_major_nocheck(major)), kv_map.find(key)->second.load(cuda::std::memory_order_relaxed)); } @@ -579,8 +579,8 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( auto tmp_e_op_result_buffer = allocate_dataframe_buffer(tmp_majors.size(), handle.get_stream()); - auto matrix_partition_row_value_input = edge_partition_src_value_input; - matrix_partition_row_value_input.set_local_adj_matrix_partition_idx(i); + auto matrix_partition_src_value_input = edge_partition_src_value_input; + matrix_partition_src_value_input.set_local_adj_matrix_partition_idx(i); auto triplet_first = thrust::make_zip_iterator(thrust::make_tuple( tmp_majors.begin(), tmp_minor_keys.begin(), tmp_key_aggregated_edge_weights.begin())); @@ -594,7 +594,7 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( KeyAggregatedEdgeOp, decltype(matrix_partition), decltype(kv_map.get_device_view())>{ - matrix_partition_row_value_input, + matrix_partition_src_value_input, key_aggregated_e_op, matrix_partition, GraphViewType::is_multi_gpu ? multi_gpu_kv_map_ptr->get_device_view() diff --git a/cpp/include/cugraph/prims/extract_if_e.cuh b/cpp/include/cugraph/prims/extract_if_e.cuh index 18b0925ea7b..bce5089c13e 100644 --- a/cpp/include/cugraph/prims/extract_if_e.cuh +++ b/cpp/include/cugraph/prims/extract_if_e.cuh @@ -64,20 +64,20 @@ struct call_e_op_t { if constexpr (thrust::tuple_size::value == 3) { weight = thrust::get<2>(e); } auto major_offset = matrix_partition.get_major_offset_from_major_nocheck(major); auto minor_offset = matrix_partition.get_minor_offset_from_minor_nocheck(minor); - auto row = GraphViewType::is_adj_matrix_transposed ? minor : major; - auto col = GraphViewType::is_adj_matrix_transposed ? major : minor; - auto row_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : major_offset; - auto col_offset = GraphViewType::is_adj_matrix_transposed ? major_offset : minor_offset; + auto src = GraphViewType::is_adj_matrix_transposed ? minor : major; + auto dst = GraphViewType::is_adj_matrix_transposed ? major : minor; + auto src_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : major_offset; + auto dst_offset = GraphViewType::is_adj_matrix_transposed ? major_offset : minor_offset; return !evaluate_edge_op() - .compute(row, - col, + .compute(src, + dst, weight, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); } }; @@ -157,12 +157,12 @@ extract_if_e(raft::handle_t const& handle, matrix_partition_device_view_t( graph_view.get_matrix_partition_view(i)); - auto matrix_partition_row_value_input = edge_partition_src_value_input; - auto matrix_partition_col_value_input = edge_partition_dst_value_input; + auto matrix_partition_src_value_input = edge_partition_src_value_input; + auto matrix_partition_dst_value_input = edge_partition_dst_value_input; if constexpr (GraphViewType::is_adj_matrix_transposed) { - matrix_partition_col_value_input.set_local_adj_matrix_partition_idx(i); + matrix_partition_dst_value_input.set_local_adj_matrix_partition_idx(i); } else { - matrix_partition_row_value_input.set_local_adj_matrix_partition_idx(i); + matrix_partition_src_value_input.set_local_adj_matrix_partition_idx(i); } detail::decompress_matrix_partition_to_edgelist( @@ -185,8 +185,8 @@ extract_if_e(raft::handle_t const& handle, EdgePartitionSrcValueInputWrapper, EdgePartitionDstValueInputWrapper, EdgeOp>{matrix_partition, - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, e_op}))); } else { auto edge_first = thrust::make_zip_iterator( diff --git a/cpp/include/cugraph/prims/property_op_utils.cuh b/cpp/include/cugraph/prims/property_op_utils.cuh index 61be94c4414..c50a1fde93f 100644 --- a/cpp/include/cugraph/prims/property_op_utils.cuh +++ b/cpp/include/cugraph/prims/property_op_utils.cuh @@ -48,8 +48,8 @@ struct is_valid_edge_op< template struct edge_op_result_type; @@ -57,39 +57,39 @@ struct edge_op_result_type; template struct edge_op_result_type< key_t, vertex_t, weight_t, - row_value_t, - col_value_t, + src_value_t, + dst_value_t, EdgeOp, std::enable_if_t>:: + typename std::invoke_result>:: valid>> { using type = - typename std::invoke_result::type; + typename std::invoke_result::type; }; template struct edge_op_result_type< key_t, vertex_t, weight_t, - row_value_t, - col_value_t, + src_value_t, + dst_value_t, EdgeOp, std::enable_if_t>::valid>> { - using type = typename std::invoke_result::type; + typename std::invoke_result>::valid>> { + using type = typename std::invoke_result::type; }; } // namespace detail @@ -102,38 +102,38 @@ template :: + edge_op_result_type:: type; - template - __device__ - std::enable_if_t>::valid, - typename std::invoke_result::type> - compute(K r, V c, W w, R rv, C cv, E e) + template + __device__ std::enable_if_t< + detail::is_valid_edge_op>::valid, + typename std::invoke_result::type> + compute(K s, V d, W w, SV sv, DV dv, E e) { - return e(r, c, w, rv, cv); + return e(s, d, w, sv, dv); } - template + template __device__ - std::enable_if_t>::valid, - typename std::invoke_result::type> - compute(K r, V c, W w, R rv, C cv, E e) + std::enable_if_t>::valid, + typename std::invoke_result::type> + compute(K s, V d, W w, SV sv, DV dv, E e) { - return e(r, c, rv, cv); + return e(s, d, sv, dv); } }; @@ -147,35 +147,35 @@ struct cast_edge_op_bool_to_integer { static_assert(std::is_integral::value); using vertex_type = typename GraphViewType::vertex_type; using weight_type = typename GraphViewType::weight_type; - using row_value_type = typename EdgePartitionSrcValueInputWrapper::value_type; - using col_value_type = typename EdgePartitionDstValueInputWrapper::value_type; + using src_value_type = typename EdgePartitionSrcValueInputWrapper::value_type; + using dst_value_type = typename EdgePartitionDstValueInputWrapper::value_type; EdgeOp e_op{}; - template - __device__ - std::enable_if_t>::valid, - T> - operator()(K r, V c, W w, R rv, C cv) + template + __device__ std:: + enable_if_t>::valid, T> + operator()(K s, V d, W w, SV sv, DV dv) { - return e_op(r, c, w, rv, cv) ? T{1} : T{0}; + return e_op(s, d, w, sv, dv) ? T{1} : T{0}; } - template + template __device__ - std::enable_if_t>::valid, T> - operator()(K r, V c, R rv, C cv) + std::enable_if_t>::valid, + T> + operator()(K s, V d, SV sv, DV dv) { - return e_op(r, c, rv, cv) ? T{1} : T{0}; + return e_op(s, d, sv, dv) ? T{1} : T{0}; } }; diff --git a/cpp/include/cugraph/prims/transform_reduce_by_src_dst_key_e.cuh b/cpp/include/cugraph/prims/transform_reduce_by_src_dst_key_e.cuh index bd47562b49e..9a6bde282f3 100644 --- a/cpp/include/cugraph/prims/transform_reduce_by_src_dst_key_e.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_by_src_dst_key_e.cuh @@ -60,10 +60,10 @@ __device__ void update_buffer_element( auto major_offset = matrix_partition.get_major_offset_from_major_nocheck(major); auto minor_offset = matrix_partition.get_minor_offset_from_minor_nocheck(minor); - auto row = GraphViewType::is_adj_matrix_transposed ? minor : major; - auto col = GraphViewType::is_adj_matrix_transposed ? major : minor; - auto row_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : major_offset; - auto col_offset = GraphViewType::is_adj_matrix_transposed ? major_offset : minor_offset; + auto src = GraphViewType::is_adj_matrix_transposed ? minor : major; + auto dst = GraphViewType::is_adj_matrix_transposed ? major : minor; + auto src_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : major_offset; + auto dst_offset = GraphViewType::is_adj_matrix_transposed ? major_offset : minor_offset; *key = edge_partition_src_dst_key_input.get( ((GraphViewType::is_adj_matrix_transposed != edge_partition_src_key) ? major_offset @@ -73,11 +73,11 @@ __device__ void update_buffer_element( EdgePartitionSrcValueInputWrapper, EdgePartitionDstValueInputWrapper, EdgeOp>() - .compute(row, - col, + .compute(src, + dst, weight, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); } @@ -389,17 +389,17 @@ transform_reduce_by_src_dst_key_e( auto tmp_value_buffer = allocate_dataframe_buffer(tmp_keys.size(), handle.get_stream()); if (graph_view.get_vertex_partition_size(comm_root_rank) > 0) { - auto matrix_partition_row_value_input = edge_partition_src_value_input; - auto matrix_partition_col_value_input = edge_partition_dst_value_input; + auto matrix_partition_src_value_input = edge_partition_src_value_input; + auto matrix_partition_dst_value_input = edge_partition_dst_value_input; if constexpr (GraphViewType::is_adj_matrix_transposed) { - matrix_partition_col_value_input.set_local_adj_matrix_partition_idx(i); + matrix_partition_dst_value_input.set_local_adj_matrix_partition_idx(i); } else { - matrix_partition_row_value_input.set_local_adj_matrix_partition_idx(i); + matrix_partition_src_value_input.set_local_adj_matrix_partition_idx(i); } - auto matrix_partition_row_col_key_input = edge_partition_src_dst_key_input; + auto matrix_partition_src_dst_key_input = edge_partition_src_dst_key_input; if constexpr ((edge_partition_src_key && !GraphViewType::is_adj_matrix_transposed) || (!edge_partition_src_key && GraphViewType::is_adj_matrix_transposed)) { - matrix_partition_row_col_key_input.set_local_adj_matrix_partition_idx(i); + matrix_partition_src_dst_key_input.set_local_adj_matrix_partition_idx(i); } auto segment_offsets = graph_view.get_local_adj_matrix_partition_segment_offsets(i); @@ -418,9 +418,9 @@ transform_reduce_by_src_dst_key_e( matrix_partition, matrix_partition.get_major_first(), matrix_partition.get_major_first() + (*segment_offsets)[1], - matrix_partition_row_value_input, - matrix_partition_col_value_input, - matrix_partition_row_col_key_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, + matrix_partition_src_dst_key_input, e_op, tmp_keys.data(), get_dataframe_buffer_begin(tmp_value_buffer)); @@ -435,9 +435,9 @@ transform_reduce_by_src_dst_key_e( matrix_partition, matrix_partition.get_major_first() + (*segment_offsets)[1], matrix_partition.get_major_first() + (*segment_offsets)[2], - matrix_partition_row_value_input, - matrix_partition_col_value_input, - matrix_partition_row_col_key_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, + matrix_partition_src_dst_key_input, e_op, tmp_keys.data(), get_dataframe_buffer_begin(tmp_value_buffer)); @@ -452,9 +452,9 @@ transform_reduce_by_src_dst_key_e( matrix_partition, matrix_partition.get_major_first() + (*segment_offsets)[2], matrix_partition.get_major_first() + (*segment_offsets)[3], - matrix_partition_row_value_input, - matrix_partition_col_value_input, - matrix_partition_row_col_key_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, + matrix_partition_src_dst_key_input, e_op, tmp_keys.data(), get_dataframe_buffer_begin(tmp_value_buffer)); @@ -469,9 +469,9 @@ transform_reduce_by_src_dst_key_e( <<>>( matrix_partition, matrix_partition.get_major_first() + (*segment_offsets)[3], - matrix_partition_row_value_input, - matrix_partition_col_value_input, - matrix_partition_row_col_key_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, + matrix_partition_src_dst_key_input, e_op, tmp_keys.data(), get_dataframe_buffer_begin(tmp_value_buffer)); @@ -487,9 +487,9 @@ transform_reduce_by_src_dst_key_e( matrix_partition, matrix_partition.get_major_first(), matrix_partition.get_major_last(), - matrix_partition_row_value_input, - matrix_partition_col_value_input, - matrix_partition_row_col_key_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, + matrix_partition_src_dst_key_input, e_op, tmp_keys.data(), get_dataframe_buffer_begin(tmp_value_buffer)); diff --git a/cpp/include/cugraph/prims/transform_reduce_e.cuh b/cpp/include/cugraph/prims/transform_reduce_e.cuh index f976e908d92..e8a42e9a0d7 100644 --- a/cpp/include/cugraph/prims/transform_reduce_e.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_e.cuh @@ -96,12 +96,12 @@ __global__ void for_all_major_for_all_nbr_hypersparse( auto minor = indices[i]; auto weight = weights ? (*weights)[i] : weight_t{1.0}; auto minor_offset = matrix_partition.get_minor_offset_from_minor_nocheck(minor); - auto row = GraphViewType::is_adj_matrix_transposed ? minor : major; - auto col = GraphViewType::is_adj_matrix_transposed ? major : minor; - auto row_offset = GraphViewType::is_adj_matrix_transposed + auto src = GraphViewType::is_adj_matrix_transposed ? minor : major; + auto dst = GraphViewType::is_adj_matrix_transposed ? major : minor; + auto src_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : static_cast(major_offset); - auto col_offset = GraphViewType::is_adj_matrix_transposed + auto dst_offset = GraphViewType::is_adj_matrix_transposed ? static_cast(major_offset) : minor_offset; return evaluate_edge_op() - .compute(row, - col, + .compute(src, + dst, weight, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); }, e_op_result_t{}, @@ -178,16 +178,16 @@ __global__ void for_all_major_for_all_nbr_low_degree( auto minor = indices[i]; auto weight = weights ? (*weights)[i] : weight_t{1.0}; auto minor_offset = matrix_partition.get_minor_offset_from_minor_nocheck(minor); - auto row = GraphViewType::is_adj_matrix_transposed + auto src = GraphViewType::is_adj_matrix_transposed ? minor : matrix_partition.get_major_from_major_offset_nocheck(major_offset); - auto col = GraphViewType::is_adj_matrix_transposed + auto dst = GraphViewType::is_adj_matrix_transposed ? matrix_partition.get_major_from_major_offset_nocheck(major_offset) : minor; - auto row_offset = GraphViewType::is_adj_matrix_transposed + auto src_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : static_cast(major_offset); - auto col_offset = GraphViewType::is_adj_matrix_transposed + auto dst_offset = GraphViewType::is_adj_matrix_transposed ? static_cast(major_offset) : minor_offset; return evaluate_edge_op() - .compute(row, - col, + .compute(src, + dst, weight, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); }, e_op_result_t{}, @@ -255,16 +255,16 @@ __global__ void for_all_major_for_all_nbr_mid_degree( auto minor = indices[i]; auto weight = weights ? (*weights)[i] : weight_t{1.0}; auto minor_offset = matrix_partition.get_minor_offset_from_minor_nocheck(minor); - auto row = GraphViewType::is_adj_matrix_transposed + auto src = GraphViewType::is_adj_matrix_transposed ? minor : matrix_partition.get_major_from_major_offset_nocheck(major_offset); - auto col = GraphViewType::is_adj_matrix_transposed + auto dst = GraphViewType::is_adj_matrix_transposed ? matrix_partition.get_major_from_major_offset_nocheck(major_offset) : minor; - auto row_offset = GraphViewType::is_adj_matrix_transposed + auto src_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : static_cast(major_offset); - auto col_offset = GraphViewType::is_adj_matrix_transposed + auto dst_offset = GraphViewType::is_adj_matrix_transposed ? static_cast(major_offset) : minor_offset; auto e_op_result = evaluate_edge_op() - .compute(row, - col, + .compute(src, + dst, weight, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); e_op_result_sum = edge_property_add(e_op_result_sum, e_op_result); } @@ -326,16 +326,16 @@ __global__ void for_all_major_for_all_nbr_high_degree( auto minor = indices[i]; auto weight = weights ? (*weights)[i] : weight_t{1.0}; auto minor_offset = matrix_partition.get_minor_offset_from_minor_nocheck(minor); - auto row = GraphViewType::is_adj_matrix_transposed + auto src = GraphViewType::is_adj_matrix_transposed ? minor : matrix_partition.get_major_from_major_offset_nocheck(major_offset); - auto col = GraphViewType::is_adj_matrix_transposed + auto dst = GraphViewType::is_adj_matrix_transposed ? matrix_partition.get_major_from_major_offset_nocheck(major_offset) : minor; - auto row_offset = GraphViewType::is_adj_matrix_transposed + auto src_offset = GraphViewType::is_adj_matrix_transposed ? minor_offset : static_cast(major_offset); - auto col_offset = GraphViewType::is_adj_matrix_transposed + auto dst_offset = GraphViewType::is_adj_matrix_transposed ? static_cast(major_offset) : minor_offset; auto e_op_result = evaluate_edge_op() - .compute(row, - col, + .compute(src, + dst, weight, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); e_op_result_sum = edge_property_add(e_op_result_sum, e_op_result); } @@ -422,12 +422,12 @@ T transform_reduce_e(raft::handle_t const& handle, matrix_partition_device_view_t( graph_view.get_matrix_partition_view(i)); - auto matrix_partition_row_value_input = edge_partition_src_value_input; - auto matrix_partition_col_value_input = edge_partition_dst_value_input; + auto matrix_partition_src_value_input = edge_partition_src_value_input; + auto matrix_partition_dst_value_input = edge_partition_dst_value_input; if constexpr (GraphViewType::is_adj_matrix_transposed) { - matrix_partition_col_value_input.set_local_adj_matrix_partition_idx(i); + matrix_partition_dst_value_input.set_local_adj_matrix_partition_idx(i); } else { - matrix_partition_row_value_input.set_local_adj_matrix_partition_idx(i); + matrix_partition_src_value_input.set_local_adj_matrix_partition_idx(i); } auto segment_offsets = graph_view.get_local_adj_matrix_partition_segment_offsets(i); @@ -445,8 +445,8 @@ T transform_reduce_e(raft::handle_t const& handle, matrix_partition, matrix_partition.get_major_first(), matrix_partition.get_major_first() + (*segment_offsets)[1], - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, get_dataframe_buffer_begin(result_buffer), e_op); } @@ -459,8 +459,8 @@ T transform_reduce_e(raft::handle_t const& handle, matrix_partition, matrix_partition.get_major_first() + (*segment_offsets)[1], matrix_partition.get_major_first() + (*segment_offsets)[2], - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, get_dataframe_buffer_begin(result_buffer), e_op); } @@ -473,8 +473,8 @@ T transform_reduce_e(raft::handle_t const& handle, matrix_partition, matrix_partition.get_major_first() + (*segment_offsets)[2], matrix_partition.get_major_first() + (*segment_offsets)[3], - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, get_dataframe_buffer_begin(result_buffer), e_op); } @@ -487,8 +487,8 @@ T transform_reduce_e(raft::handle_t const& handle, <<>>( matrix_partition, matrix_partition.get_major_first() + (*segment_offsets)[3], - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, get_dataframe_buffer_begin(result_buffer), e_op); } @@ -503,8 +503,8 @@ T transform_reduce_e(raft::handle_t const& handle, matrix_partition, matrix_partition.get_major_first(), matrix_partition.get_major_last(), - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, get_dataframe_buffer_begin(result_buffer), e_op); } diff --git a/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh b/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh index fc00a140577..a879d2f50af 100644 --- a/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh +++ b/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh @@ -184,7 +184,7 @@ template -__device__ void push_buffer_element(vertex_t col, +__device__ void push_buffer_element(vertex_t dst, e_op_result_t e_op_result, BufferKeyOutputIterator buffer_key_output_first, BufferPayloadOutputIterator buffer_payload_output_first, @@ -197,14 +197,14 @@ __device__ void push_buffer_element(vertex_t col, assert(e_op_result.has_value()); if constexpr (std::is_same_v && std::is_same_v) { - *(buffer_key_output_first + buffer_idx) = col; + *(buffer_key_output_first + buffer_idx) = dst; } else if constexpr (std::is_same_v && !std::is_same_v) { - *(buffer_key_output_first + buffer_idx) = col; + *(buffer_key_output_first + buffer_idx) = dst; *(buffer_payload_output_first + buffer_idx) = *e_op_result; } else if constexpr (!std::is_same_v && std::is_same_v) { - *(buffer_key_output_first + buffer_idx) = thrust::make_tuple(col, *e_op_result); + *(buffer_key_output_first + buffer_idx) = thrust::make_tuple(dst, *e_op_result); } else { - *(buffer_key_output_first + buffer_idx) = thrust::make_tuple(col, thrust::get<0>(*e_op_result)); + *(buffer_key_output_first + buffer_idx) = thrust::make_tuple(dst, thrust::get<0>(*e_op_result)); *(buffer_payload_output_first + buffer_idx) = thrust::get<1>(*e_op_result); } } @@ -216,7 +216,7 @@ template -__global__ void for_all_frontier_row_for_all_nbr_hypersparse( +__global__ void for_all_frontier_src_for_all_nbr_hypersparse( matrix_partition_device_view_t(major_hypersparse_first - matrix_partition.get_major_first()); auto idx = static_cast(tid); @@ -284,17 +284,17 @@ __global__ void for_all_frontier_row_for_all_nbr_hypersparse( edge_t local_degree{0}; if (lane_id < static_cast(max_key_idx - min_key_idx)) { auto key = *(key_first + idx); - vertex_t row{}; + vertex_t src{}; if constexpr (std::is_same_v) { - row = key; + src = key; } else { - row = thrust::get<0>(key); + src = thrust::get<0>(key); } - auto row_hypersparse_idx = matrix_partition.get_major_hypersparse_idx_from_major_nocheck(row); - if (row_hypersparse_idx) { - auto row_idx = row_start_offset + *row_hypersparse_idx; - local_degree = matrix_partition.get_local_degree(row_idx); - warp_key_local_edge_offsets[threadIdx.x] = matrix_partition.get_local_offset(row_idx); + auto src_hypersparse_idx = matrix_partition.get_major_hypersparse_idx_from_major_nocheck(src); + if (src_hypersparse_idx) { + auto src_idx = src_start_offset + *src_hypersparse_idx; + local_degree = matrix_partition.get_local_degree(src_idx); + warp_key_local_edge_offsets[threadIdx.x] = matrix_partition.get_local_offset(src_idx); } else { local_degree = edge_t{0}; warp_key_local_edge_offsets[threadIdx.x] = edge_t{0}; // dummy @@ -314,7 +314,7 @@ __global__ void for_all_frontier_row_for_all_nbr_hypersparse( for (size_t i = lane_id; i < rounded_up_num_edges_this_warp; i += raft::warp_size()) { e_op_result_t e_op_result{}; - vertex_t col{}; + vertex_t dst{}; if (i < static_cast(num_edges_this_warp)) { auto key_idx_this_warp = static_cast(thrust::distance( @@ -332,25 +332,25 @@ __global__ void for_all_frontier_row_for_all_nbr_hypersparse( : warp_local_degree_inclusive_sums[warp_id * raft::warp_size() + key_idx_this_warp - 1])); auto key = *(key_first + (min_key_idx + key_idx_this_warp)); - vertex_t row{}; + vertex_t src{}; if constexpr (std::is_same_v) { - row = key; + src = key; } else { - row = thrust::get<0>(key); + src = thrust::get<0>(key); } - col = indices[local_edge_offset]; - auto row_offset = matrix_partition.get_major_offset_from_major_nocheck(row); - auto col_offset = matrix_partition.get_minor_offset_from_minor_nocheck(col); + dst = indices[local_edge_offset]; + auto src_offset = matrix_partition.get_major_offset_from_major_nocheck(src); + auto dst_offset = matrix_partition.get_minor_offset_from_minor_nocheck(dst); e_op_result = evaluate_edge_op() .compute(key, - col, + dst, weights ? (*weights)[local_edge_offset] : weight_t{1.0}, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); } auto ballot_e_op = @@ -367,7 +367,7 @@ __global__ void for_all_frontier_row_for_all_nbr_hypersparse( if (e_op_result) { auto buffer_warp_offset = static_cast(__popc(ballot_e_op & ~(uint32_t{0xffffffff} << lane_id))); - push_buffer_element(col, + push_buffer_element(dst, e_op_result, buffer_key_output_first, buffer_payload_output_first, @@ -386,7 +386,7 @@ template -__global__ void for_all_frontier_row_for_all_nbr_low_degree( +__global__ void for_all_frontier_src_for_all_nbr_low_degree( matrix_partition_device_view_t(max_key_idx - min_key_idx)) { auto key = *(key_first + idx); - vertex_t row{}; + vertex_t src{}; if constexpr (std::is_same_v) { - row = key; + src = key; } else { - row = thrust::get<0>(key); + src = thrust::get<0>(key); } - auto row_offset = matrix_partition.get_major_offset_from_major_nocheck(row); - local_degree = matrix_partition.get_local_degree(row_offset); - warp_key_local_edge_offsets[threadIdx.x] = matrix_partition.get_local_offset(row_offset); + auto src_offset = matrix_partition.get_major_offset_from_major_nocheck(src); + local_degree = matrix_partition.get_local_degree(src_offset); + warp_key_local_edge_offsets[threadIdx.x] = matrix_partition.get_local_offset(src_offset); } WarpScan(temp_storage) .InclusiveSum(local_degree, warp_local_degree_inclusive_sums[threadIdx.x]); @@ -474,7 +474,7 @@ __global__ void for_all_frontier_row_for_all_nbr_low_degree( raft::warp_size(); for (size_t i = lane_id; i < rounded_up_num_edges_this_warp; i += raft::warp_size()) { e_op_result_t e_op_result{}; - vertex_t col{}; + vertex_t dst{}; if (i < static_cast(num_edges_this_warp)) { auto key_idx_this_warp = static_cast(thrust::distance( @@ -492,25 +492,25 @@ __global__ void for_all_frontier_row_for_all_nbr_low_degree( : warp_local_degree_inclusive_sums[warp_id * raft::warp_size() + key_idx_this_warp - 1])); auto key = *(key_first + (min_key_idx + key_idx_this_warp)); - vertex_t row{}; + vertex_t src{}; if constexpr (std::is_same_v) { - row = key; + src = key; } else { - row = thrust::get<0>(key); + src = thrust::get<0>(key); } - col = indices[local_edge_offset]; - auto row_offset = matrix_partition.get_major_offset_from_major_nocheck(row); - auto col_offset = matrix_partition.get_minor_offset_from_minor_nocheck(col); + dst = indices[local_edge_offset]; + auto src_offset = matrix_partition.get_major_offset_from_major_nocheck(src); + auto dst_offset = matrix_partition.get_minor_offset_from_minor_nocheck(dst); e_op_result = evaluate_edge_op() .compute(key, - col, + dst, weights ? (*weights)[local_edge_offset] : weight_t{1.0}, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); } auto ballot = __ballot_sync(uint32_t{0xffffffff}, e_op_result ? uint32_t{1} : uint32_t{0}); @@ -526,7 +526,7 @@ __global__ void for_all_frontier_row_for_all_nbr_low_degree( if (e_op_result) { auto buffer_warp_offset = static_cast(__popc(ballot & ~(uint32_t{0xffffffff} << lane_id))); - push_buffer_element(col, + push_buffer_element(dst, e_op_result, buffer_key_output_first, buffer_payload_output_first, @@ -546,7 +546,7 @@ template -__global__ void for_all_frontier_row_for_all_nbr_mid_degree( +__global__ void for_all_frontier_src_for_all_nbr_mid_degree( matrix_partition_device_view_t(thrust::distance(key_first, key_last))) { auto key = *(key_first + idx); - vertex_t row{}; + vertex_t src{}; if constexpr (std::is_same_v) { - row = key; + src = key; } else { - row = thrust::get<0>(key); + src = thrust::get<0>(key); } - auto row_offset = matrix_partition.get_major_offset_from_major_nocheck(row); + auto src_offset = matrix_partition.get_major_offset_from_major_nocheck(src); vertex_t const* indices{nullptr}; thrust::optional weights{thrust::nullopt}; edge_t local_out_degree{}; - thrust::tie(indices, weights, local_out_degree) = matrix_partition.get_local_edges(row_offset); + thrust::tie(indices, weights, local_out_degree) = matrix_partition.get_local_edges(src_offset); auto rounded_up_local_out_degree = ((static_cast(local_out_degree) + (raft::warp_size() - 1)) / raft::warp_size()) * raft::warp_size(); for (size_t i = lane_id; i < rounded_up_local_out_degree; i += raft::warp_size()) { e_op_result_t e_op_result{}; - vertex_t col{}; + vertex_t dst{}; if (i < static_cast(local_out_degree)) { - col = indices[i]; - auto col_offset = matrix_partition.get_minor_offset_from_minor_nocheck(col); + dst = indices[i]; + auto dst_offset = matrix_partition.get_minor_offset_from_minor_nocheck(dst); e_op_result = evaluate_edge_op() .compute(key, - col, + dst, weights ? (*weights)[i] : weight_t{1.0}, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); } auto ballot = __ballot_sync(uint32_t{0xffffffff}, e_op_result ? uint32_t{1} : uint32_t{0}); @@ -633,7 +633,7 @@ __global__ void for_all_frontier_row_for_all_nbr_mid_degree( if (e_op_result) { auto buffer_warp_offset = static_cast(__popc(ballot & ~(uint32_t{0xffffffff} << lane_id))); - push_buffer_element(col, + push_buffer_element(dst, e_op_result, buffer_key_output_first, buffer_payload_output_first, @@ -653,7 +653,7 @@ template -__global__ void for_all_frontier_row_for_all_nbr_high_degree( +__global__ void for_all_frontier_src_for_all_nbr_high_degree( matrix_partition_device_view_t(thrust::distance(key_first, key_last))) { auto key = *(key_first + idx); - vertex_t row{}; + vertex_t src{}; if constexpr (std::is_same_v) { - row = key; + src = key; } else { - row = thrust::get<0>(key); + src = thrust::get<0>(key); } - auto row_offset = matrix_partition.get_major_offset_from_major_nocheck(row); + auto src_offset = matrix_partition.get_major_offset_from_major_nocheck(src); vertex_t const* indices{nullptr}; thrust::optional weights{thrust::nullopt}; edge_t local_out_degree{}; - thrust::tie(indices, weights, local_out_degree) = matrix_partition.get_local_edges(row_offset); + thrust::tie(indices, weights, local_out_degree) = matrix_partition.get_local_edges(src_offset); auto rounded_up_local_out_degree = ((static_cast(local_out_degree) + (update_frontier_v_push_if_out_nbr_for_all_block_size - 1)) / @@ -710,22 +710,22 @@ __global__ void for_all_frontier_row_for_all_nbr_high_degree( update_frontier_v_push_if_out_nbr_for_all_block_size; for (size_t i = threadIdx.x; i < rounded_up_local_out_degree; i += blockDim.x) { e_op_result_t e_op_result{}; - vertex_t col{}; + vertex_t dst{}; edge_t buffer_block_offset{0}; if (i < static_cast(local_out_degree)) { - col = indices[i]; - auto col_offset = matrix_partition.get_minor_offset_from_minor_nocheck(col); + dst = indices[i]; + auto dst_offset = matrix_partition.get_minor_offset_from_minor_nocheck(dst); e_op_result = evaluate_edge_op() .compute(key, - col, + dst, weights ? (*weights)[i] : weight_t{1.0}, - edge_partition_src_value_input.get(row_offset), - edge_partition_dst_value_input.get(col_offset), + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), e_op); } BlockScan(temp_storage) @@ -741,7 +741,7 @@ __global__ void for_all_frontier_row_for_all_nbr_high_degree( } __syncthreads(); if (e_op_result) { - push_buffer_element(col, + push_buffer_element(dst, e_op_result, buffer_key_output_first, buffer_payload_output_first, @@ -1082,17 +1082,17 @@ void update_frontier_v_push_if_out_nbr( get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer)); } - vertex_t const* matrix_partition_frontier_row_first{nullptr}; - vertex_t const* matrix_partition_frontier_row_last{nullptr}; + vertex_t const* matrix_partition_frontier_src_first{nullptr}; + vertex_t const* matrix_partition_frontier_src_last{nullptr}; if constexpr (std::is_same_v) { - matrix_partition_frontier_row_first = + matrix_partition_frontier_src_first = get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer); - matrix_partition_frontier_row_last = + matrix_partition_frontier_src_last = get_dataframe_buffer_end(matrix_partition_frontier_key_buffer); } else { - matrix_partition_frontier_row_first = thrust::get<0>( + matrix_partition_frontier_src_first = thrust::get<0>( get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer).get_iterator_tuple()); - matrix_partition_frontier_row_last = thrust::get<0>( + matrix_partition_frontier_src_last = thrust::get<0>( get_dataframe_buffer_end(matrix_partition_frontier_key_buffer).get_iterator_tuple()); } @@ -1106,24 +1106,24 @@ void update_frontier_v_push_if_out_nbr( auto max_pushes = use_dcs ? thrust::transform_reduce( execution_policy, - matrix_partition_frontier_row_first, - matrix_partition_frontier_row_last, + matrix_partition_frontier_src_first, + matrix_partition_frontier_src_last, [matrix_partition, major_hypersparse_first = matrix_partition.get_major_first() + (*segment_offsets) - [detail::num_sparse_segments_per_vertex_partition]] __device__(auto row) { - if (row < major_hypersparse_first) { - auto row_offset = matrix_partition.get_major_offset_from_major_nocheck(row); - return matrix_partition.get_local_degree(row_offset); + [detail::num_sparse_segments_per_vertex_partition]] __device__(auto src) { + if (src < major_hypersparse_first) { + auto src_offset = matrix_partition.get_major_offset_from_major_nocheck(src); + return matrix_partition.get_local_degree(src_offset); } else { - auto row_hypersparse_idx = - matrix_partition.get_major_hypersparse_idx_from_major_nocheck(row); - return row_hypersparse_idx + auto src_hypersparse_idx = + matrix_partition.get_major_hypersparse_idx_from_major_nocheck(src); + return src_hypersparse_idx ? matrix_partition.get_local_degree( matrix_partition.get_major_offset_from_major_nocheck( major_hypersparse_first) + - *row_hypersparse_idx) + *src_hypersparse_idx) : edge_t{0}; } }, @@ -1131,11 +1131,11 @@ void update_frontier_v_push_if_out_nbr( thrust::plus()) : thrust::transform_reduce( execution_policy, - matrix_partition_frontier_row_first, - matrix_partition_frontier_row_last, - [matrix_partition] __device__(auto row) { - auto row_offset = matrix_partition.get_major_offset_from_major_nocheck(row); - return matrix_partition.get_local_degree(row_offset); + matrix_partition_frontier_src_first, + matrix_partition_frontier_src_last, + [matrix_partition] __device__(auto src) { + auto src_offset = matrix_partition.get_major_offset_from_major_nocheck(src); + return matrix_partition.get_local_degree(src_offset); }, edge_t{0}, thrust::plus()); @@ -1146,9 +1146,9 @@ void update_frontier_v_push_if_out_nbr( resize_dataframe_buffer(payload_buffer, new_buffer_size, handle.get_stream()); } - auto matrix_partition_row_value_input = edge_partition_src_value_input; - auto matrix_partition_col_value_input = edge_partition_dst_value_input; - matrix_partition_row_value_input.set_local_adj_matrix_partition_idx(i); + auto matrix_partition_src_value_input = edge_partition_src_value_input; + auto matrix_partition_dst_value_input = edge_partition_dst_value_input; + matrix_partition_src_value_input.set_local_adj_matrix_partition_idx(i); if (segment_offsets) { static_assert(detail::num_sparse_segments_per_vertex_partition == 3); @@ -1162,8 +1162,8 @@ void update_frontier_v_push_if_out_nbr( d_thresholds.data(), h_thresholds.data(), h_thresholds.size(), handle.get_stream()); rmm::device_uvector d_offsets(d_thresholds.size(), handle.get_stream()); thrust::lower_bound(handle.get_thrust_policy(), - matrix_partition_frontier_row_first, - matrix_partition_frontier_row_last, + matrix_partition_frontier_src_first, + matrix_partition_frontier_src_last, d_thresholds.begin(), d_thresholds.end(), d_offsets.begin()); @@ -1179,13 +1179,13 @@ void update_frontier_v_push_if_out_nbr( h_offsets[0], detail::update_frontier_v_push_if_out_nbr_for_all_block_size, handle.get_device_properties().maxGridSize[0]); - detail::for_all_frontier_row_for_all_nbr_high_degree + detail::for_all_frontier_src_for_all_nbr_high_degree <<>>( matrix_partition, get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer), get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer) + h_offsets[0], - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, get_dataframe_buffer_begin(key_buffer), detail::get_optional_payload_buffer_begin(payload_buffer), buffer_idx.data(), @@ -1196,13 +1196,13 @@ void update_frontier_v_push_if_out_nbr( h_offsets[1] - h_offsets[0], detail::update_frontier_v_push_if_out_nbr_for_all_block_size, handle.get_device_properties().maxGridSize[0]); - detail::for_all_frontier_row_for_all_nbr_mid_degree + detail::for_all_frontier_src_for_all_nbr_mid_degree <<>>( matrix_partition, get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer) + h_offsets[0], get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer) + h_offsets[1], - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, get_dataframe_buffer_begin(key_buffer), detail::get_optional_payload_buffer_begin(payload_buffer), buffer_idx.data(), @@ -1213,13 +1213,13 @@ void update_frontier_v_push_if_out_nbr( h_offsets[2] - h_offsets[1], detail::update_frontier_v_push_if_out_nbr_for_all_block_size, handle.get_device_properties().maxGridSize[0]); - detail::for_all_frontier_row_for_all_nbr_low_degree + detail::for_all_frontier_src_for_all_nbr_low_degree <<>>( matrix_partition, get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer) + h_offsets[1], get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer) + h_offsets[2], - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, get_dataframe_buffer_begin(key_buffer), detail::get_optional_payload_buffer_begin(payload_buffer), buffer_idx.data(), @@ -1230,14 +1230,14 @@ void update_frontier_v_push_if_out_nbr( h_offsets[3] - h_offsets[2], detail::update_frontier_v_push_if_out_nbr_for_all_block_size, handle.get_device_properties().maxGridSize[0]); - detail::for_all_frontier_row_for_all_nbr_hypersparse + detail::for_all_frontier_src_for_all_nbr_hypersparse <<>>( matrix_partition, matrix_partition.get_major_first() + (*segment_offsets)[3], get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer) + h_offsets[2], get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer) + h_offsets[3], - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, get_dataframe_buffer_begin(key_buffer), detail::get_optional_payload_buffer_begin(payload_buffer), buffer_idx.data(), @@ -1250,13 +1250,13 @@ void update_frontier_v_push_if_out_nbr( detail::update_frontier_v_push_if_out_nbr_for_all_block_size, handle.get_device_properties().maxGridSize[0]); - detail::for_all_frontier_row_for_all_nbr_low_degree + detail::for_all_frontier_src_for_all_nbr_low_degree <<>>( matrix_partition, get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer), get_dataframe_buffer_end(matrix_partition_frontier_key_buffer), - matrix_partition_row_value_input, - matrix_partition_col_value_input, + matrix_partition_src_value_input, + matrix_partition_dst_value_input, get_dataframe_buffer_begin(key_buffer), detail::get_optional_payload_buffer_begin(payload_buffer), buffer_idx.data(), @@ -1291,15 +1291,15 @@ void update_frontier_v_push_if_out_nbr( d_vertex_lasts.data(), h_vertex_lasts.data(), h_vertex_lasts.size(), handle.get_stream()); rmm::device_uvector d_tx_buffer_last_boundaries(d_vertex_lasts.size(), handle.get_stream()); - vertex_t const* row_first{nullptr}; + vertex_t const* src_first{nullptr}; if constexpr (std::is_same_v) { - row_first = get_dataframe_buffer_begin(key_buffer); + src_first = get_dataframe_buffer_begin(key_buffer); } else { - row_first = thrust::get<0>(get_dataframe_buffer_begin(key_buffer).get_iterator_tuple()); + src_first = thrust::get<0>(get_dataframe_buffer_begin(key_buffer).get_iterator_tuple()); } thrust::lower_bound(handle.get_thrust_policy(), - row_first, - row_first + num_buffer_elements, + src_first, + src_first + num_buffer_elements, d_vertex_lasts.begin(), d_vertex_lasts.end(), d_tx_buffer_last_boundaries.begin()); From c74009b6609c3039c002373ed72d4fc996a0f5c7 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Fri, 11 Mar 2022 17:19:21 -0800 Subject: [PATCH 9/9] copyright year --- cpp/include/cugraph/prims/count_if_e.cuh | 2 +- cpp/include/cugraph/prims/transform_reduce_e.cuh | 2 +- .../cugraph/prims/update_edge_partition_src_dst_property.cuh | 2 +- cpp/src/centrality/katz_centrality_impl.cuh | 2 +- cpp/src/cores/core_number_impl.cuh | 2 +- cpp/src/link_analysis/hits_impl.cuh | 2 +- cpp/src/link_analysis/pagerank_impl.cuh | 2 +- 7 files changed, 7 insertions(+), 7 deletions(-) diff --git a/cpp/include/cugraph/prims/count_if_e.cuh b/cpp/include/cugraph/prims/count_if_e.cuh index 030a194ea71..e7caaedc2fa 100644 --- a/cpp/include/cugraph/prims/count_if_e.cuh +++ b/cpp/include/cugraph/prims/count_if_e.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/include/cugraph/prims/transform_reduce_e.cuh b/cpp/include/cugraph/prims/transform_reduce_e.cuh index e8a42e9a0d7..0582cad1522 100644 --- a/cpp/include/cugraph/prims/transform_reduce_e.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_e.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh b/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh index 0d39e1aa5ae..e306aa1f7df 100644 --- a/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh +++ b/cpp/include/cugraph/prims/update_edge_partition_src_dst_property.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/src/centrality/katz_centrality_impl.cuh b/cpp/src/centrality/katz_centrality_impl.cuh index 57bda084c1f..4e5dd1f8d0a 100644 --- a/cpp/src/centrality/katz_centrality_impl.cuh +++ b/cpp/src/centrality/katz_centrality_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/src/cores/core_number_impl.cuh b/cpp/src/cores/core_number_impl.cuh index ad355eaf835..0e861a65b8b 100644 --- a/cpp/src/cores/core_number_impl.cuh +++ b/cpp/src/cores/core_number_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/src/link_analysis/hits_impl.cuh b/cpp/src/link_analysis/hits_impl.cuh index 553e5e1fbf6..162d1f17ac9 100644 --- a/cpp/src/link_analysis/hits_impl.cuh +++ b/cpp/src/link_analysis/hits_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/src/link_analysis/pagerank_impl.cuh b/cpp/src/link_analysis/pagerank_impl.cuh index e518f19c1ab..3e77a8bc026 100644 --- a/cpp/src/link_analysis/pagerank_impl.cuh +++ b/cpp/src/link_analysis/pagerank_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License.