diff --git a/cpp/include/cugraph/detail/shuffle_wrappers.hpp b/cpp/include/cugraph/detail/shuffle_wrappers.hpp index 6592150e74..7dffcce298 100644 --- a/cpp/include/cugraph/detail/shuffle_wrappers.hpp +++ b/cpp/include/cugraph/detail/shuffle_wrappers.hpp @@ -142,30 +142,6 @@ shuffle_ext_vertex_value_pairs_to_local_gpu_by_vertex_partitioning( rmm::device_uvector&& vertices, rmm::device_uvector&& values); -/** - * @brief Shuffle external (i.e. before renumbering) vertex & values pairs to their local GPU based - * on vertex partitioning. - * - * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. - * @tparam value_t Type of values. - * - * @param[in] handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, - * and handles to various CUDA libraries) to run graph algorithms. - * @param[in] vertices Vertices to shuffle. - * @param[in] values_0 First values to shuffle. - * @param[in] values_1 Second values to shuffle. - * - * @return Tuple of vectors storing shuffled vertex & value pairs. - */ -template -std:: - tuple, rmm::device_uvector, rmm::device_uvector> - shuffle_ext_vertex_values_pairs_to_local_gpu_by_vertex_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& vertices, - rmm::device_uvector&& values_0, - rmm::device_uvector&& values_1); - /** * @brief Permute a range. * diff --git a/cpp/include/cugraph/detail/utility_wrappers.hpp b/cpp/include/cugraph/detail/utility_wrappers.hpp index 681ab31907..e0307bc372 100644 --- a/cpp/include/cugraph/detail/utility_wrappers.hpp +++ b/cpp/include/cugraph/detail/utility_wrappers.hpp @@ -68,7 +68,7 @@ void scalar_fill(raft::handle_t const& handle, value_t* d_value, size_t size, va /** * @brief Sort a device span * - * @tparam value_t type of the value to operate on + * @tparam value_t type of the value to operate on. Must be either int32_t or int64_t. * * @param [in] handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, * and handles to various CUDA libraries) to run graph algorithms. @@ -76,12 +76,12 @@ void scalar_fill(raft::handle_t const& handle, value_t* d_value, size_t size, va * */ template -void sort(raft::handle_t const& handle, raft::device_span values); +void sort_ints(raft::handle_t const& handle, raft::device_span values); /** * @brief Keep unique element from a device span * - * @tparam value_t type of the value to operate on + * @tparam value_t type of the value to operate on. Must be either int32_t or int64_t. * * @param [in] handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, * and handles to various CUDA libraries) to run graph algorithms. @@ -90,23 +90,22 @@ void sort(raft::handle_t const& handle, raft::device_span values); * */ template -size_t unique(raft::handle_t const& handle, raft::device_span values); +size_t unique_ints(raft::handle_t const& handle, raft::device_span values); /** * @brief Increment the values of a device span by a constant value * - * @tparam value_t type of the value to operate on + * @tparam value_t type of the value to operate on. Must be either int32_t or int64_t. * - * @param [in] handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, - * and handles to various CUDA libraries) to run graph algorithms. * @param[out] values device span to update * @param[in] value value to be added to each element of the buffer + * @param[in] stream_view stream view * */ template -void transform_increment(rmm::cuda_stream_view const& stream_view, - raft::device_span values, - value_t value); +void transform_increment_ints(raft::device_span values, + value_t value, + rmm::cuda_stream_view const& stream_view); /** * @brief Fill a buffer with a sequence of values @@ -116,7 +115,7 @@ void transform_increment(rmm::cuda_stream_view const& stream_view, * * Similar to the function std::iota, wraps the function thrust::sequence * - * @tparam value_t type of the value to operate on + * @tparam value_t type of the value to operate on. * * @param[in] stream_view stream view * @param[out] d_value device array to fill diff --git a/cpp/src/c_api/neighbor_sampling.cpp b/cpp/src/c_api/neighbor_sampling.cpp index 2988c992fc..842241be81 100644 --- a/cpp/src/c_api/neighbor_sampling.cpp +++ b/cpp/src/c_api/neighbor_sampling.cpp @@ -885,10 +885,10 @@ struct neighbor_sampling_functor : public cugraph::c_api::abstract_functor { // Compute the global start_vertex_label_offsets - cugraph::detail::transform_increment( - handle_.get_stream(), + cugraph::detail::transform_increment_ints( raft::device_span{(*start_vertex_labels).data(), (*start_vertex_labels).size()}, - (label_t)global_labels[handle_.get_comms().get_rank()]); + (label_t)global_labels[handle_.get_comms().get_rank()], + handle_.get_stream()); } if constexpr (multi_gpu) { @@ -902,11 +902,11 @@ struct neighbor_sampling_functor : public cugraph::c_api::abstract_functor { // Get unique labels // sort the start_vertex_labels - cugraph::detail::sort( + cugraph::detail::sort_ints( handle_.get_stream(), raft::device_span{unique_labels.data(), unique_labels.size()}); - auto num_unique_labels = cugraph::detail::unique( + auto num_unique_labels = cugraph::detail::unique_ints( handle_.get_stream(), raft::device_span{unique_labels.data(), unique_labels.size()}); diff --git a/cpp/src/detail/utility_wrappers_32.cu b/cpp/src/detail/utility_wrappers_32.cu index c59006985a..879a1adf33 100644 --- a/cpp/src/detail/utility_wrappers_32.cu +++ b/cpp/src/detail/utility_wrappers_32.cu @@ -63,10 +63,9 @@ template void scalar_fill(raft::handle_t const& handle, size_t* d_value, size_t template void scalar_fill(raft::handle_t const& handle, float* d_value, size_t size, float value); -template void sort(raft::handle_t const& handle, raft::device_span d_span); +template void sort_ints(raft::handle_t const& handle, raft::device_span values); -template size_t unique(raft::handle_t const& handle, raft::device_span d_span); -template size_t unique(raft::handle_t const& handle, raft::device_span d_span); +template size_t unique_ints(raft::handle_t const& handle, raft::device_span values); template void sequence_fill(rmm::cuda_stream_view const& stream_view, int32_t* d_value, @@ -78,13 +77,9 @@ template void sequence_fill(rmm::cuda_stream_view const& stream_view, size_t size, uint32_t start_value); -template void transform_increment(rmm::cuda_stream_view const& stream_view, - raft::device_span d_span, - int32_t value); - -template void transform_increment(rmm::cuda_stream_view const& stream_view, - raft::device_span d_span, - uint32_t value); +template void transform_increment_ints(raft::device_span values, + int32_t value, + rmm::cuda_stream_view const& stream_view); template void stride_fill(rmm::cuda_stream_view const& stream_view, int32_t* d_value, diff --git a/cpp/src/detail/utility_wrappers_64.cu b/cpp/src/detail/utility_wrappers_64.cu index f78c709024..742cb18d71 100644 --- a/cpp/src/detail/utility_wrappers_64.cu +++ b/cpp/src/detail/utility_wrappers_64.cu @@ -61,10 +61,9 @@ template void scalar_fill(raft::handle_t const& handle, template void scalar_fill(raft::handle_t const& handle, double* d_value, size_t size, double value); -template void sort(raft::handle_t const& handle, raft::device_span d_span); +template void sort_ints(raft::handle_t const& handle, raft::device_span values); -template size_t unique(raft::handle_t const& handle, raft::device_span d_span); -template size_t unique(raft::handle_t const& handle, raft::device_span d_span); +template size_t unique_ints(raft::handle_t const& handle, raft::device_span values); template void sequence_fill(rmm::cuda_stream_view const& stream_view, int64_t* d_value, @@ -76,13 +75,9 @@ template void sequence_fill(rmm::cuda_stream_view const& stream_view, size_t size, uint64_t start_value); -template void transform_increment(rmm::cuda_stream_view const& stream_view, - raft::device_span d_span, - int64_t value); - -template void transform_increment(rmm::cuda_stream_view const& stream_view, - raft::device_span d_span, - uint64_t value); +template void transform_increment_ints(raft::device_span values, + int64_t value, + rmm::cuda_stream_view const& stream_view); template void stride_fill(rmm::cuda_stream_view const& stream_view, int64_t* d_value, diff --git a/cpp/src/detail/utility_wrappers_impl.cuh b/cpp/src/detail/utility_wrappers_impl.cuh index 7522f3a9c5..93bd14c4d0 100644 --- a/cpp/src/detail/utility_wrappers_impl.cuh +++ b/cpp/src/detail/utility_wrappers_impl.cuh @@ -65,17 +65,17 @@ void scalar_fill(raft::handle_t const& handle, value_t* d_value, size_t size, va } template -void sort(raft::handle_t const& handle, raft::device_span d_span) +void sort_ints(raft::handle_t const& handle, raft::device_span values) { - thrust::sort(handle.get_thrust_policy(), d_span.begin(), d_span.end()); + thrust::sort(handle.get_thrust_policy(), values.begin(), values.end()); } template -size_t unique(raft::handle_t const& handle, raft::device_span d_span) +size_t unique_ints(raft::handle_t const& handle, raft::device_span values) { auto unique_element_last = - thrust::unique(handle.get_thrust_policy(), d_span.begin(), d_span.end()); - return thrust::distance(d_span.begin(), unique_element_last); + thrust::unique(handle.get_thrust_policy(), values.begin(), values.end()); + return thrust::distance(values.begin(), unique_element_last); } template @@ -88,14 +88,14 @@ void sequence_fill(rmm::cuda_stream_view const& stream_view, } template -void transform_increment(rmm::cuda_stream_view const& stream_view, - raft::device_span d_span, - value_t incr) +void transform_increment_ints(raft::device_span values, + value_t incr, + rmm::cuda_stream_view const& stream_view) { thrust::transform(rmm::exec_policy(stream_view), - d_span.begin(), - d_span.end(), - d_span.begin(), + values.begin(), + values.end(), + values.begin(), cuda::proclaim_return_type([incr] __device__(value_t value) { return static_cast(value + incr); }));