From 6263b9fd1c8226c2b2f30ccdc4c95716769867e5 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 13 Jul 2023 14:08:50 -0700 Subject: [PATCH 01/17] define sampling output renumbering function API --- cpp/include/cugraph/graph_functions.hpp | 52 +++++++++++++++++++++++++ 1 file changed, 52 insertions(+) diff --git a/cpp/include/cugraph/graph_functions.hpp b/cpp/include/cugraph/graph_functions.hpp index 017b32d0470..4158f2f4088 100644 --- a/cpp/include/cugraph/graph_functions.hpp +++ b/cpp/include/cugraph/graph_functions.hpp @@ -916,4 +916,56 @@ rmm::device_uvector select_random_vertices( bool sort_vertices, bool do_expensive_check = false); +/** + * @brief renumber sampling output + * + * This function renumbers sampling function (e.g. uniform_neighbor_sample) outputs satisfying the + * following requirements. + * + * 1. Say @p edgelist_srcs has N unique vertices. These N unique vertices will be mapped to [0, N). + * 2. Among the N unique vertices, an original vertex with a smaller attached hop number will be + * renumbered to a smaller vertex ID than any other original vertices with a larger attached hop + * number (if @p edgelist_hops.has_value() is true). If a single vertex is attached to multiple hop + * numbers, the minimum hop number is used. + * 3. Say @p edgelist_dsts has M unique vertices that appear only in @p edgelist_dsts (the set of M + * unique vertices does not include any vertices that appear in @p edgelist_srcs). Then, these M + * unique vertices will be mapped to [N, N + M). + * 4. If label_offsets.has_value() is ture, edge lists for different labels will be renumbered + * separately. + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam label_t Type of labels. Needs to be an integral type. + * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * or multi-GPU (true). + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param edgelist_srcs A vector storing original edgelist source vertices. + * @param edgelist_hops An optional pointer to the array storing hops for each edge list source + * vertices (size = @p edgelist_srcs.size()). + * @param edgelist_dsts A vector storing original edgelist destination vertices (size = @p + * edgelist_srcs.size()). + * @param label_offsets An optional tuple of unique labels and the input edge list (@p + * edgelist_srcs, @p edgelist_hops, and @p edgelist_dsts) offsets for the labels (siez = # unique + * labels + 1). + * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`). + * @return Tuple of vectors storing renumbered edge sources (size = @p edgelist_srcs.size()) , + * renumbered edge destinations (size = @p edgelist_dsts.size()), renumber_map to query original + * verties (size = # unique vertices or aggregate # unique vertices for every label), and + * renumber_map offsets (size = std::get<0>(*label_offsets).size() + 1, valid only if @p + * label_offsets.has_value() is true). + */ +template +std::tuple, + rmm::device_uvector, + rmm::device_uvector, + std::optional>> +renumber_sampled_edgelist( + raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + std::optional> edgelist_hops, + rmm::device_uvector&& edgelist_dsts, + std::optional, raft::device_span>> + label_offsets, + bool do_expensive_check = false); + } // namespace cugraph From 6c88e86a52095b0394a9287e6d86d49b3d3428d2 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Mon, 17 Jul 2023 10:24:42 -0700 Subject: [PATCH 02/17] update the API (remove multi_gpu flag) --- cpp/include/cugraph/graph_functions.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/include/cugraph/graph_functions.hpp b/cpp/include/cugraph/graph_functions.hpp index 4158f2f4088..caffef60076 100644 --- a/cpp/include/cugraph/graph_functions.hpp +++ b/cpp/include/cugraph/graph_functions.hpp @@ -933,10 +933,10 @@ rmm::device_uvector select_random_vertices( * 4. If label_offsets.has_value() is ture, edge lists for different labels will be renumbered * separately. * + * This function is single-GPU only (we are not aware of any practical multi-GPU use cases). + * * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. * @tparam label_t Type of labels. Needs to be an integral type. - * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) - * or multi-GPU (true). * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param edgelist_srcs A vector storing original edgelist source vertices. @@ -954,7 +954,7 @@ rmm::device_uvector select_random_vertices( * renumber_map offsets (size = std::get<0>(*label_offsets).size() + 1, valid only if @p * label_offsets.has_value() is true). */ -template +template std::tuple, rmm::device_uvector, rmm::device_uvector, From 055496fb1122e1514ea26d335de80999150ffb56 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Mon, 17 Jul 2023 10:25:21 -0700 Subject: [PATCH 03/17] initial draft implementation --- cpp/CMakeLists.txt | 1 + .../renumber_sampled_edgelist_impl.cuh | 583 ++++++++++++++++++ .../sampling/renumber_sampled_edgelist_sg.cu | 51 ++ 3 files changed, 635 insertions(+) create mode 100644 cpp/src/sampling/renumber_sampled_edgelist_impl.cuh create mode 100644 cpp/src/sampling/renumber_sampled_edgelist_sg.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 996e654734f..b50e3a96927 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -223,6 +223,7 @@ set(CUGRAPH_SOURCES src/sampling/detail/sampling_utils_sg.cu src/sampling/uniform_neighbor_sampling_mg.cpp src/sampling/uniform_neighbor_sampling_sg.cpp + src/sampling/renumber_sampled_edgelist_sg.cu src/cores/core_number_sg.cu src/cores/core_number_mg.cu src/cores/k_core_sg.cu diff --git a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh new file mode 100644 index 00000000000..215fd00c255 --- /dev/null +++ b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh @@ -0,0 +1,583 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include + +#include + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +namespace cugraph { + +namespace { + +template +std::tuple, std::optional>> +compute_renumber_map( + raft::handle_t const& handle, + raft::device_span edgelist_srcs, + std::optional> edgelist_hops, + raft::device_span edgelist_dsts, + std::optional, raft::device_span>> + label_offsets) +{ + std::optional> unique_label_src_pair_label_indices{std::nullopt}; + rmm::device_uvector unique_label_src_pair_vertices(0, handle.get_stream()); + { + rmm::device_uvector srcs(edgelist_srcs.size(), handle.get_stream()); + thrust::copy( + handle.get_thrust_policy(), edgelist_srcs.begin(), edgelist_srcs.end(), srcs.begin()); + + if (label_offsets) { + rmm::device_uvector label_indices(edgelist_srcs.size(), handle.get_stream()); + thrust::transform( + handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(edgelist_srcs.size()), + label_indices.begin(), + [offsets = raft::device_span( + std::get<1>(*label_offsets).data() + 1, + std::get<1>(*label_offsets).size())] __device__(size_t i) { + return static_cast(thrust::distance( + offsets.begin(), thrust::upper_bound(thrust::seq, offsets.begin(), offsets.end(), i))); + }); + + if (edgelist_hops) { + rmm::device_uvector hops((*edgelist_hops).size(), handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), + (*edgelist_hops).begin(), + (*edgelist_hops).end(), + hops.begin()); + auto triplet_first = + thrust::make_zip_iterator(label_indices.begin(), srcs.begin(), hops.begin()); + thrust::sort(handle.get_thrust_policy(), triplet_first, triplet_first + srcs.size()); + auto num_uniques = static_cast( + thrust::distance(triplet_first, + thrust::unique(handle.get_thrust_policy(), + triplet_first, + triplet_first + srcs.size(), + [] __device__(auto lhs, auto rhs) { + return (thrust::get<0>(lhs) == thrust::get<0>(rhs)) && + (thrust::get<1>(lhs) == thrust::get<1>(rhs)); + }))); + label_indices.resize(num_uniques, handle.get_stream()); + srcs.resize(num_uniques, handle.get_stream()); + hops.resize(num_uniques, handle.get_stream()); + label_indices.shrink_to_fit(handle.get_stream()); + srcs.shrink_to_fit(handle.get_stream()); + hops.shrink_to_fit(handle.get_stream()); + + unique_label_src_pair_label_indices = std::move(label_indices); + + auto num_labels = std::get<0>(*label_offsets).size(); + rmm::device_uvector tmp_label_offsets(num_labels + 1, handle.get_stream()); + tmp_label_offsets.set_element_to_zero_async(0, handle.get_stream()); + thrust::upper_bound(handle.get_thrust_policy(), + label_indices.begin(), + label_indices.end(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(num_labels), + tmp_label_offsets.begin() + 1); + + rmm::device_uvector segment_sorted_hops(hops.size(), handle.get_stream()); + rmm::device_uvector segment_sorted_srcs(srcs.size(), handle.get_stream()); + + size_t tmp_storage_bytes{0}; + rmm::device_uvector d_tmp_storage(0, handle.get_stream()); + + cub::DeviceSegmentedSort::SortPairs(static_cast(nullptr), + tmp_storage_bytes, + hops.begin(), + segment_sorted_hops.begin(), + srcs.begin(), + segment_sorted_srcs.begin(), + hops.size(), + num_labels, + tmp_label_offsets.begin(), + tmp_label_offsets.begin() + 1, + handle.get_stream()); + + if (tmp_storage_bytes > d_tmp_storage.size()) { + d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); + } + + cub::DeviceSegmentedSort::SortPairs(d_tmp_storage.data(), + tmp_storage_bytes, + hops.begin(), + segment_sorted_hops.begin(), + srcs.begin(), + segment_sorted_srcs.begin(), + hops.size(), + num_labels, + tmp_label_offsets.begin(), + tmp_label_offsets.begin() + 1, + handle.get_stream()); + + unique_label_src_pair_vertices = std::move(segment_sorted_srcs); + } else { + rmm::device_uvector segment_sorted_srcs(srcs.size(), handle.get_stream()); + + size_t tmp_storage_bytes{0}; + rmm::device_uvector d_tmp_storage(0, handle.get_stream()); + + cub::DeviceSegmentedSort::SortKeys(static_cast(nullptr), + tmp_storage_bytes, + srcs.begin(), + segment_sorted_srcs.begin(), + srcs.size(), + std::get<0>(*label_offsets).size(), + std::get<1>(*label_offsets).begin(), + std::get<1>(*label_offsets).begin() + 1, + handle.get_stream()); + + if (tmp_storage_bytes > d_tmp_storage.size()) { + d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); + } + + cub::DeviceSegmentedSort::SortKeys(d_tmp_storage.data(), + tmp_storage_bytes, + srcs.begin(), + segment_sorted_srcs.begin(), + srcs.size(), + std::get<0>(*label_offsets).size(), + std::get<1>(*label_offsets).begin(), + std::get<1>(*label_offsets).begin() + 1, + handle.get_stream()); + + auto pair_first = + thrust::make_zip_iterator(label_indices.begin(), segment_sorted_srcs.begin()); + auto num_uniques = static_cast(thrust::distance( + pair_first, + thrust::unique( + handle.get_thrust_policy(), pair_first, pair_first + label_indices.size()))); + label_indices.resize(num_uniques, handle.get_stream()); + segment_sorted_srcs.resize(num_uniques, handle.get_stream()); + label_indices.shrink_to_fit(handle.get_stream()); + segment_sorted_srcs.shrink_to_fit(handle.get_stream()); + + unique_label_src_pair_label_indices = std::move(label_indices); + unique_label_src_pair_vertices = std::move(segment_sorted_srcs); + } + } else { + if (edgelist_hops) { + rmm::device_uvector hops((*edgelist_hops).size(), handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), + (*edgelist_hops).begin(), + (*edgelist_hops).end(), + hops.begin()); + + auto pair_first = thrust::make_zip_iterator(hops.begin(), srcs.begin()); + thrust::sort(handle.get_thrust_policy(), pair_first, pair_first + hops.size()); + srcs.resize( + thrust::distance(srcs.begin(), + thrust::unique(handle.get_thrust_policy(), srcs.begin(), srcs.end())), + handle.get_stream()); + } else { + thrust::sort(handle.get_thrust_policy(), srcs.begin(), srcs.end()); + srcs.resize( + thrust::distance(srcs.begin(), + thrust::unique(handle.get_thrust_policy(), srcs.begin(), srcs.end())), + handle.get_stream()); + srcs.shrink_to_fit(handle.get_stream()); + } + + unique_label_src_pair_vertices = std::move(srcs); + } + } + + std::optional> unique_label_dst_pair_label_indices{std::nullopt}; + rmm::device_uvector unique_label_dst_pair_vertices(0, handle.get_stream()); + { + rmm::device_uvector dsts(edgelist_dsts.size(), handle.get_stream()); + thrust::copy( + handle.get_thrust_policy(), edgelist_dsts.begin(), edgelist_dsts.end(), dsts.begin()); + if (label_offsets) { + rmm::device_uvector label_indices(edgelist_dsts.size(), handle.get_stream()); + thrust::transform( + handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(edgelist_dsts.size()), + label_indices.begin(), + [offsets = raft::device_span( + std::get<1>(*label_offsets).data() + 1, + std::get<1>(*label_offsets).size())] __device__(size_t i) { + return static_cast(thrust::distance( + offsets.begin(), thrust::upper_bound(thrust::seq, offsets.begin(), offsets.end(), i))); + }); + + rmm::device_uvector segment_sorted_dsts(dsts.size(), handle.get_stream()); + + size_t tmp_storage_bytes{0}; + rmm::device_uvector d_tmp_storage(0, handle.get_stream()); + + cub::DeviceSegmentedSort::SortKeys(static_cast(nullptr), + tmp_storage_bytes, + dsts.begin(), + segment_sorted_dsts.begin(), + dsts.size(), + std::get<0>(*label_offsets).size(), + std::get<1>(*label_offsets).begin(), + std::get<1>(*label_offsets).begin() + 1, + handle.get_stream()); + + if (tmp_storage_bytes > d_tmp_storage.size()) { + d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); + } + + cub::DeviceSegmentedSort::SortKeys(d_tmp_storage.data(), + tmp_storage_bytes, + dsts.begin(), + segment_sorted_dsts.begin(), + dsts.size(), + std::get<0>(*label_offsets).size(), + std::get<1>(*label_offsets).begin(), + std::get<1>(*label_offsets).begin() + 1, + handle.get_stream()); + + auto pair_first = + thrust::make_zip_iterator(label_indices.begin(), segment_sorted_dsts.begin()); + auto num_uniques = static_cast(thrust::distance( + pair_first, + thrust::unique(handle.get_thrust_policy(), pair_first, pair_first + label_indices.size()))); + label_indices.resize(num_uniques, handle.get_stream()); + segment_sorted_dsts.resize(num_uniques, handle.get_stream()); + label_indices.shrink_to_fit(handle.get_stream()); + segment_sorted_dsts.shrink_to_fit(handle.get_stream()); + + unique_label_dst_pair_label_indices = std::move(label_indices); + unique_label_dst_pair_vertices = std::move(segment_sorted_dsts); + + } else { + thrust::sort(handle.get_thrust_policy(), dsts.begin(), dsts.end()); + dsts.resize( + thrust::distance(dsts.begin(), + thrust::unique(handle.get_thrust_policy(), dsts.begin(), dsts.end())), + handle.get_stream()); + dsts.shrink_to_fit(handle.get_stream()); + } + + unique_label_dst_pair_vertices = std::move(dsts); + } + + if (label_offsets) { + auto label_src_pair_first = thrust::make_zip_iterator( + (*unique_label_src_pair_label_indices).begin(), unique_label_src_pair_vertices.begin()); + auto label_dst_pair_first = thrust::make_zip_iterator( + (*unique_label_dst_pair_label_indices).begin(), unique_label_dst_pair_vertices.begin()); + rmm::device_uvector output_label_indices((*unique_label_dst_pair_label_indices).size(), + handle.get_stream()); + rmm::device_uvector output_vertices((*unique_label_dst_pair_label_indices).size(), + handle.get_stream()); + auto output_label_dst_pair_first = + thrust::make_zip_iterator(output_label_indices.begin(), output_vertices.begin()); + auto output_label_dst_pair_last = + thrust::set_difference(handle.get_thrust_policy(), + label_dst_pair_first, + label_dst_pair_first + (*unique_label_dst_pair_label_indices).size(), + label_src_pair_first, + label_src_pair_first + (*unique_label_src_pair_label_indices).size(), + output_label_dst_pair_first); + output_label_indices.resize( + thrust::distance(output_label_dst_pair_first, output_label_dst_pair_last), + handle.get_stream()); + output_vertices.resize(output_label_indices.size(), handle.get_stream()); + output_label_indices.shrink_to_fit(handle.get_stream()); + output_vertices.shrink_to_fit(handle.get_stream()); + unique_label_dst_pair_label_indices = std::move(output_label_indices); + unique_label_dst_pair_vertices = std::move(output_vertices); + + rmm::device_uvector merged_label_indices( + (*unique_label_src_pair_label_indices).size() + (*unique_label_dst_pair_label_indices).size(), + handle.get_stream()); + rmm::device_uvector merged_vertices(merged_label_indices.size(), handle.get_stream()); + auto label_src_triplet_first = + thrust::make_zip_iterator((*unique_label_src_pair_label_indices).begin(), + thrust::make_constant_iterator(uint8_t{0}), + unique_label_src_pair_vertices.begin()); + auto label_dst_triplet_first = + thrust::make_zip_iterator((*unique_label_dst_pair_label_indices).begin(), + thrust::make_constant_iterator(uint8_t{1}), + unique_label_dst_pair_vertices.begin()); + thrust::merge( + handle.get_thrust_policy(), + label_src_triplet_first, + label_src_triplet_first + (*unique_label_src_pair_label_indices).size(), + label_dst_triplet_first, + label_dst_triplet_first + (*unique_label_dst_pair_label_indices).size(), + thrust::make_zip_iterator( + merged_label_indices.begin(), thrust::make_discard_iterator(), merged_vertices.begin())); + + return std::make_tuple(std::move(merged_vertices), std::move(merged_label_indices)); + } else { + rmm::device_uvector output_vertices(unique_label_dst_pair_vertices.size(), + handle.get_stream()); + auto output_last = thrust::set_difference(handle.get_thrust_policy(), + unique_label_dst_pair_vertices.begin(), + unique_label_dst_pair_vertices.end(), + unique_label_src_pair_vertices.begin(), + unique_label_src_pair_vertices.end(), + output_vertices.begin()); + + auto num_unique_srcs = unique_label_src_pair_vertices.size(); + auto renumber_map = std::move(unique_label_src_pair_vertices); + renumber_map.resize( + renumber_map.size() + thrust::distance(output_vertices.begin(), output_last), + handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), + output_vertices.begin(), + output_last, + renumber_map.begin() + num_unique_srcs); + + return std::make_tuple(std::move(renumber_map), std::nullopt); + } +} + +} // namespace + +template +std::tuple, + rmm::device_uvector, + rmm::device_uvector, + std::optional>> +renumber_sampled_edgelist( + raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + std::optional> edgelist_hops, + rmm::device_uvector&& edgelist_dsts, + std::optional, raft::device_span>> + label_offsets, + bool do_expensive_check) +{ + CUGRAPH_EXPECTS( + edgelist_srcs.size() == edgelist_dsts.size(), + "Invalid input arguments: edgelist_srcs.size() and edgelist_dsts.size() should coincide."); + CUGRAPH_EXPECTS(!edgelist_hops.has_value() || (edgelist_srcs.size() == (*edgelist_hops).size()), + "Invalid input arguments: if edgelist_hops is valid, (*edgelist_hops).size() and " + "edgelist_srcs.size() should coincide."); + CUGRAPH_EXPECTS(!label_offsets.has_value() || + (std::get<1>(*label_offsets).size() == std::get<0>(*label_offsets).size() + 1), + "Invalid input arguments: if label_offsets is valid, " + "std::get<1>(label_offsets).size() (size of the offset array) should be " + "std::get<0>(label_offsets).size() (number of unique labels) + 1."); + + if (do_expensive_check) { + if (label_offsets) { + CUGRAPH_EXPECTS(thrust::is_sorted(handle.get_thrust_policy(), + std::get<1>(*label_offsets).begin(), + std::get<1>(*label_offsets).end()), + "Invalid input arguments: if label_offsets is valid, " + "std::get<1>(*label_offsets) should be sorted."); + size_t back_element{}; + raft::update_host( + &back_element, + std::get<1>(*label_offsets).data() + (std::get<1>(*label_offsets).size() - 1), + size_t{1}, + handle.get_stream()); + handle.get_stream(); + CUGRAPH_EXPECTS(back_element == edgelist_srcs.size(), + "Invalid input arguments: if label_offsets is valid, the last element of " + "std::get<1>(*label_offsets) and edgelist_srcs.size() should coincide."); + } + } + + auto [renumber_map, renumber_map_label_indices] = compute_renumber_map( + handle, + raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), + edgelist_hops, + raft::device_span(edgelist_dsts.data(), edgelist_dsts.size()), + label_offsets); + + std::optional> renumber_map_label_offsets{}; + if (label_offsets) { + auto num_unique_labels = thrust::count_if( + handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator((*renumber_map_label_indices).size()), + detail::is_first_in_run_t{(*renumber_map_label_indices).data()}); + rmm::device_uvector unique_label_indices(num_unique_labels, handle.get_stream()); + rmm::device_uvector vertex_counts(num_unique_labels, handle.get_stream()); + thrust::reduce_by_key(handle.get_thrust_policy(), + (*renumber_map_label_indices).begin(), + (*renumber_map_label_indices).end(), + thrust::make_constant_iterator(size_t{1}), + unique_label_indices.begin(), + vertex_counts.begin()); + + renumber_map_label_offsets = + rmm::device_uvector(std::get<0>(*label_offsets).size() + 1, handle.get_stream()); + thrust::fill(handle.get_thrust_policy(), + (*renumber_map_label_offsets).begin(), + (*renumber_map_label_offsets).end(), + size_t{0}); + thrust::scatter(handle.get_thrust_policy(), + vertex_counts.begin(), + vertex_counts.end(), + unique_label_indices.begin(), + (*renumber_map_label_offsets).begin() + 1); + + thrust::inclusive_scan(handle.get_thrust_policy(), + (*renumber_map_label_offsets).begin(), + (*renumber_map_label_offsets).end(), + (*renumber_map_label_offsets).begin()); + } + + if (label_offsets) { + rmm::device_uvector new_vertices(renumber_map.size(), handle.get_stream()); + thrust::tabulate(handle.get_thrust_policy(), + new_vertices.begin(), + new_vertices.end(), + [label_indices = raft::device_span( + (*renumber_map_label_indices).data(), (*renumber_map_label_indices).size()), + renumber_map_label_offsets = raft::device_span( + (*renumber_map_label_offsets).data(), + (*renumber_map_label_offsets).size())] __device__(size_t i) { + auto label_index = label_indices[i]; + auto label_start_offset = renumber_map_label_offsets[label_index]; + return static_cast(i - label_start_offset); + }); + + auto num_labels = std::get<0>(*label_offsets).size(); + + rmm::device_uvector segment_sorted_renumber_map(renumber_map.size(), + handle.get_stream()); + rmm::device_uvector segment_sorted_new_vertices(new_vertices.size(), + handle.get_stream()); + + size_t tmp_storage_bytes{0}; + rmm::device_uvector d_tmp_storage(0, handle.get_stream()); + + cub::DeviceSegmentedSort::SortPairs(static_cast(nullptr), + tmp_storage_bytes, + renumber_map.begin(), + segment_sorted_renumber_map.begin(), + new_vertices.begin(), + segment_sorted_new_vertices.begin(), + renumber_map.size(), + num_labels, + (*renumber_map_label_offsets).begin(), + (*renumber_map_label_offsets).begin() + 1, + handle.get_stream()); + + if (tmp_storage_bytes > d_tmp_storage.size()) { + d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); + } + + cub::DeviceSegmentedSort::SortPairs(d_tmp_storage.data(), + tmp_storage_bytes, + renumber_map.begin(), + segment_sorted_renumber_map.begin(), + new_vertices.begin(), + segment_sorted_new_vertices.begin(), + renumber_map.size(), + num_labels, + (*renumber_map_label_offsets).begin(), + (*renumber_map_label_offsets).begin() + 1, + handle.get_stream()); + + auto pair_first = + thrust::make_zip_iterator(edgelist_srcs.begin(), (*renumber_map_label_indices).begin()); + thrust::transform( + handle.get_thrust_policy(), + pair_first, + pair_first + edgelist_srcs.size(), + edgelist_srcs.begin(), + [renumber_map_label_offsets = raft::device_span( + (*renumber_map_label_offsets).data(), (*renumber_map_label_offsets).size()), + old_vertices = raft::device_span(segment_sorted_renumber_map.data(), + segment_sorted_renumber_map.size()), + new_vertices = raft::device_span( + segment_sorted_new_vertices.data(), + segment_sorted_new_vertices.size())] __device__(auto pair) { + auto old_vertex = thrust::get<0>(pair); + auto label_index = thrust::get<1>(pair); + auto label_start_offset = renumber_map_label_offsets[label_index]; + auto label_end_offset = renumber_map_label_offsets[label_index + 1]; + auto it = thrust::lower_bound(thrust::seq, + old_vertices.begin() + label_start_offset, + old_vertices.begin() + label_end_offset, + old_vertex); + assert(*it == old_vertex); + return *(new_vertices.begin() + thrust::distance(old_vertices.begin(), it)); + }); + + pair_first = + thrust::make_zip_iterator(edgelist_dsts.begin(), (*renumber_map_label_indices).begin()); + thrust::transform( + handle.get_thrust_policy(), + pair_first, + pair_first + edgelist_dsts.size(), + edgelist_dsts.begin(), + [renumber_map_label_offsets = raft::device_span( + (*renumber_map_label_offsets).data(), (*renumber_map_label_offsets).size()), + old_vertices = raft::device_span(segment_sorted_renumber_map.data(), + segment_sorted_renumber_map.size()), + new_vertices = raft::device_span( + segment_sorted_new_vertices.data(), + segment_sorted_new_vertices.size())] __device__(auto pair) { + auto old_vertex = thrust::get<0>(pair); + auto label_index = thrust::get<1>(pair); + auto label_start_offset = renumber_map_label_offsets[label_index]; + auto label_end_offset = renumber_map_label_offsets[label_index + 1]; + auto it = thrust::lower_bound(thrust::seq, + old_vertices.begin() + label_start_offset, + old_vertices.begin() + label_end_offset, + old_vertex); + assert(*it == old_vertex); + return new_vertices[thrust::distance(old_vertices.begin(), it)]; + }); + + } else { + kv_store_t kv_store(renumber_map.begin(), + renumber_map.end(), + thrust::make_counting_iterator(vertex_t{0}), + std::numeric_limits::max(), + std::numeric_limits::max(), + handle.get_stream()); + auto kv_store_view = kv_store.view(); + + kv_store_view.find( + edgelist_srcs.begin(), edgelist_srcs.end(), edgelist_srcs.begin(), handle.get_stream()); + kv_store_view.find( + edgelist_dsts.begin(), edgelist_dsts.end(), edgelist_dsts.begin(), handle.get_stream()); + } + + return std::make_tuple(std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::move(renumber_map), + std::move(renumber_map_label_offsets)); +} + +} // namespace cugraph diff --git a/cpp/src/sampling/renumber_sampled_edgelist_sg.cu b/cpp/src/sampling/renumber_sampled_edgelist_sg.cu new file mode 100644 index 00000000000..522440108da --- /dev/null +++ b/cpp/src/sampling/renumber_sampled_edgelist_sg.cu @@ -0,0 +1,51 @@ +/* + * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include "renumber_sampled_edgelist_impl.cuh" + +namespace cugraph { + +template +std::tuple, + rmm::device_uvector, + rmm::device_uvector, + std::optional>> +renumber_sampled_edgelist( + raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + std::optional> edgelist_hops, + rmm::device_uvector&& edgelist_dsts, + std::optional, raft::device_span>> + label_offsets, + bool do_expensive_check); + +template +std::tuple, + rmm::device_uvector, + rmm::device_uvector, + std::optional>> +renumber_sampled_edgelist( + raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + std::optional> edgelist_hops, + rmm::device_uvector&& edgelist_dsts, + std::optional, raft::device_span>> + label_offsets, + bool do_expensive_check); + +} // namespace cugraph From 2211d24ced5d1ca163e1c92925b5b9328111d42e Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Tue, 18 Jul 2023 17:45:13 -0700 Subject: [PATCH 04/17] add test code --- cpp/tests/CMakeLists.txt | 5 + .../renumber_sampled_edgelist_test.cu | 432 ++++++++++++++++++ 2 files changed, 437 insertions(+) create mode 100644 cpp/tests/sampling/renumber_sampled_edgelist_test.cu diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 3bcd5546455..a8ed8f7926f 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -352,6 +352,11 @@ ConfigureTest(RANDOM_WALKS_TEST sampling/sg_random_walks_test.cpp) ConfigureTest(UNIFORM_NEIGHBOR_SAMPLING_TEST sampling/sg_uniform_neighbor_sampling.cu) target_link_libraries(UNIFORM_NEIGHBOR_SAMPLING_TEST PRIVATE cuco::cuco) +################################################################################################### +# - RENUMBER SAMPLED EDGE LIST tests -------------------------------------------------------------- +ConfigureTest(RENUMBER_SAMPLED_EDGELIST_TEST sampling/renumber_sampled_edgelist_test.cu) +target_link_libraries(RENUMBER_SAMPLED_EDGELIST_TEST PRIVATE cuco::cuco) + ################################################################################################### # - Renumber tests -------------------------------------------------------------------------------- set(RENUMBERING_TEST_SRCS diff --git a/cpp/tests/sampling/renumber_sampled_edgelist_test.cu b/cpp/tests/sampling/renumber_sampled_edgelist_test.cu new file mode 100644 index 00000000000..797be44cc65 --- /dev/null +++ b/cpp/tests/sampling/renumber_sampled_edgelist_test.cu @@ -0,0 +1,432 @@ +/* + * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include +#include +#include + +#include +#include + +#include + +#include +#include +#include +#include +#include + +struct RenumberSampledEdgelist_Usecase { + size_t num_vertices{}; + size_t num_sampled_edges{}; + size_t num_hops{1}; // enabled if larger than 1 + size_t num_labels{1}; // enabled if larger than 1 + bool check_correctness{true}; +}; + +class Tests_RenumberSampledEdgelist + : public ::testing::TestWithParam { + public: + Tests_RenumberSampledEdgelist() {} + + static void SetUpTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(RenumberSampledEdgelist_Usecase const& usecase) + { + using label_t = int32_t; + + raft::handle_t handle{}; + HighResTimer hr_timer{}; + + raft::random::RngState rng_state(0); + + rmm::device_uvector org_edgelist_srcs(usecase.num_sampled_edges, handle.get_stream()); + rmm::device_uvector org_edgelist_dsts(usecase.num_sampled_edges, handle.get_stream()); + // FIXME: need to check the range ([min,max] or [min, max)?) + cugraph::detail::uniform_random_fill(handle.get_stream(), + org_edgelist_srcs.data(), + org_edgelist_srcs.size(), + vertex_t{0}, + static_cast(usecase.num_vertices), + rng_state); + cugraph::detail::uniform_random_fill(handle.get_stream(), + org_edgelist_dsts.data(), + org_edgelist_dsts.size(), + vertex_t{0}, + static_cast(usecase.num_vertices), + rng_state); + + std::optional> edgelist_hops{std::nullopt}; + if (usecase.num_hops > 1) { + edgelist_hops = rmm::device_uvector(usecase.num_sampled_edges, handle.get_stream()); + // FIXME: need to check the range ([min,max] or [min, max)?) + cugraph::detail::uniform_random_fill(handle.get_stream(), + (*edgelist_hops).data(), + (*edgelist_hops).size(), + int32_t{0}, + static_cast(usecase.num_hops), + rng_state); + } + + std::optional, rmm::device_uvector>> + label_offsets{std::nullopt}; + if (usecase.num_labels > 1) { + rmm::device_uvector labels(usecase.num_labels, handle.get_stream()); + thrust::sequence(handle.get_thrust_policy(), labels.begin(), labels.end(), label_t{0}); + + rmm::device_uvector edgelist_labels(usecase.num_sampled_edges, handle.get_stream()); + cugraph::detail::uniform_random_fill(handle.get_stream(), + edgelist_labels.data(), + edgelist_labels.size(), + label_t{0}, + static_cast(usecase.num_labels), + rng_state); + + rmm::device_uvector offsets(usecase.num_labels + 1, handle.get_stream()); + thrust::fill(handle.get_thrust_policy(), offsets.begin(), offsets.end(), size_t{0}); + + thrust::for_each( + handle.get_thrust_policy(), + edgelist_labels.begin(), + edgelist_labels.end(), + [offsets = + raft::device_span(offsets.data(), offsets.size())] __device__(label_t label) { + cuda::atomic_ref atomic_counter(offsets[label]); + atomic_counter.fetch_add(size_t{1}, cuda::std::memory_order_relaxed); + }); + + thrust::exclusive_scan( + handle.get_thrust_policy(), offsets.begin(), offsets.end(), offsets.begin()); + + label_offsets = std::make_tuple(std::move(labels), std::move(offsets)); + } + + rmm::device_uvector renumbered_edgelist_srcs(org_edgelist_srcs.size(), + handle.get_stream()); + rmm::device_uvector renumbered_edgelist_dsts(org_edgelist_dsts.size(), + handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), + org_edgelist_srcs.begin(), + org_edgelist_srcs.end(), + renumbered_edgelist_srcs.begin()); + thrust::copy(handle.get_thrust_policy(), + org_edgelist_dsts.begin(), + org_edgelist_dsts.end(), + renumbered_edgelist_dsts.begin()); + + rmm::device_uvector renumber_map(0, handle.get_stream()); + std::optional> renumber_map_label_offsets{std::nullopt}; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_timer.start("Renumber sampled edgelist"); + } + + std::tie(renumbered_edgelist_srcs, + renumbered_edgelist_dsts, + renumber_map, + renumber_map_label_offsets) = + cugraph::renumber_sampled_edgelist( + handle, + std::move(renumbered_edgelist_srcs), + edgelist_hops ? std::make_optional>( + (*edgelist_hops).data(), (*edgelist_hops).size()) + : std::nullopt, + std::move(renumbered_edgelist_dsts), + label_offsets + ? std::make_optional< + std::tuple, raft::device_span>>( + std::make_tuple(raft::device_span(std::get<0>(*label_offsets).data(), + std::get<0>(*label_offsets).size()), + raft::device_span(std::get<1>(*label_offsets).data(), + std::get<1>(*label_offsets).size()))) + : std::nullopt); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + if (usecase.check_correctness) { + for (size_t i = 0; i < usecase.num_labels; ++i) { + size_t edgelist_start_offset = + label_offsets ? std::get<1>(*label_offsets).element(i, handle.get_stream()) : size_t{0}; + size_t edgelist_end_offset = label_offsets + ? std::get<1>(*label_offsets).element(i, handle.get_stream()) + : usecase.num_sampled_edges; + auto this_label_org_edgelist_srcs = + raft::device_span(org_edgelist_srcs.data() + edgelist_start_offset, + edgelist_end_offset - edgelist_start_offset); + auto this_label_org_edgelist_dsts = + raft::device_span(org_edgelist_dsts.data() + edgelist_start_offset, + edgelist_end_offset - edgelist_start_offset); + auto this_label_edgelist_hops = edgelist_hops + ? std::make_optional>( + (*edgelist_hops).data() + edgelist_start_offset, + edgelist_end_offset - edgelist_start_offset) + : std::nullopt; + auto this_label_renumbered_edgelist_srcs = + raft::device_span(renumbered_edgelist_srcs.data() + edgelist_start_offset, + edgelist_end_offset - edgelist_start_offset); + auto this_label_renumbered_edgelist_dsts = + raft::device_span(renumbered_edgelist_dsts.data() + edgelist_start_offset, + edgelist_end_offset - edgelist_start_offset); + + size_t renumber_map_start_offset = + renumber_map_label_offsets ? (*renumber_map_label_offsets).element(i, handle.get_stream()) + : size_t{0}; + size_t renumber_map_end_offset = + renumber_map_label_offsets + ? (*renumber_map_label_offsets).element(i + 1, handle.get_stream()) + : renumber_map.size(); + auto this_label_renumber_map = + raft::device_span(renumber_map.data() + renumber_map_start_offset, + renumber_map_end_offset - renumber_map_start_offset); + + // check un-renumbering recovers the original edge list + + auto pair_first = thrust::make_zip_iterator(this_label_org_edgelist_srcs.begin(), + this_label_renumbered_edgelist_srcs.begin()); + auto num_renumber_errors = thrust::count_if( + handle.get_thrust_policy(), + pair_first, + pair_first + this_label_org_edgelist_srcs.size(), + [this_label_renumber_map] __device__(auto pair) { + auto org = thrust::get<0>(pair); + auto renumbered = thrust::get<1>(pair); + return this_label_renumber_map[thrust::get<1>(pair)] != thrust::get<0>(pair); + }); + ASSERT_TRUE(num_renumber_errors == 0) << "Renumber error in edge list sources."; + + pair_first = thrust::make_zip_iterator(this_label_org_edgelist_dsts.begin(), + this_label_renumbered_edgelist_dsts.begin()); + num_renumber_errors = thrust::count_if( + handle.get_thrust_policy(), + pair_first, + pair_first + this_label_org_edgelist_dsts.size(), + [this_label_renumber_map] __device__(auto pair) { + auto org = thrust::get<0>(pair); + auto renumbered = thrust::get<1>(pair); + return this_label_renumber_map[thrust::get<1>(pair)] != thrust::get<0>(pair); + }); + ASSERT_TRUE(num_renumber_errors == 0) << "Renumber error in edge list destinations."; + + // check the invariants in renumber_map (1. vertices appeared in edge list sources should + // have a smaller renumbered vertex ID than the vertices appear only in edge list + // destinations, 2. edge list source vertices with a smaller minimum hop number should have + // a smaller renumbered vertex ID than the edge list source vertices with a larger hop + // number) + + rmm::device_uvector unique_srcs(this_label_org_edgelist_srcs.size(), + handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), + this_label_org_edgelist_srcs.begin(), + this_label_org_edgelist_srcs.end(), + unique_srcs.begin()); + std::optional> unique_src_hops = + this_label_edgelist_hops ? std::make_optional>( + (*this_label_edgelist_hops).size(), handle.get_stream()) + : std::nullopt; + if (this_label_edgelist_hops) { + thrust::copy(handle.get_thrust_policy(), + (*this_label_edgelist_hops).begin(), + (*this_label_edgelist_hops).end(), + (*unique_src_hops).begin()); + + auto pair_first = + thrust::make_zip_iterator(unique_srcs.begin(), (*unique_src_hops).begin()); + thrust::sort(handle.get_thrust_policy(), pair_first, pair_first + unique_srcs.size()); + unique_srcs.resize(thrust::distance(pair_first, + thrust::unique(handle.get_thrust_policy(), + pair_first, + pair_first + unique_srcs.size())), + handle.get_stream()); + (*unique_src_hops).resize(unique_srcs.size(), handle.get_stream()); + } else { + thrust::sort(handle.get_thrust_policy(), unique_srcs.begin(), unique_srcs.end()); + unique_srcs.resize( + thrust::distance( + unique_srcs.begin(), + thrust::unique(handle.get_thrust_policy(), unique_srcs.begin(), unique_srcs.end())), + handle.get_stream()); + } + + rmm::device_uvector unique_dsts(this_label_org_edgelist_dsts.size(), + handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), + this_label_org_edgelist_dsts.begin(), + this_label_org_edgelist_dsts.end(), + unique_dsts.begin()); + thrust::sort(handle.get_thrust_policy(), unique_dsts.begin(), unique_dsts.end()); + unique_dsts.resize( + thrust::distance( + unique_dsts.begin(), + thrust::unique(handle.get_thrust_policy(), unique_dsts.begin(), unique_dsts.end())), + handle.get_stream()); + + unique_dsts.resize( + thrust::distance( + unique_dsts.begin(), + thrust::remove_if(handle.get_thrust_policy(), + unique_dsts.begin(), + unique_dsts.end(), + [sorted_unique_srcs = raft::device_span( + unique_srcs.data(), unique_srcs.size())] __device__(auto dst) { + return thrust::binary_search(thrust::seq, + sorted_unique_srcs.begin(), + sorted_unique_srcs.end(), + dst); + })), + handle.get_stream()); + + rmm::device_uvector sorted_org_vertices(this_label_renumber_map.size(), + handle.get_stream()); + rmm::device_uvector matching_renumbered_vertices(sorted_org_vertices.size(), + handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), + this_label_renumber_map.begin(), + this_label_renumber_map.end(), + sorted_org_vertices.begin()); + thrust::sequence(handle.get_thrust_policy(), + matching_renumbered_vertices.begin(), + matching_renumbered_vertices.end(), + vertex_t{0}); + thrust::sort_by_key(handle.get_thrust_policy(), + sorted_org_vertices.begin(), + sorted_org_vertices.end(), + matching_renumbered_vertices.begin()); + + auto max_src_renumbered_vertex = thrust::transform_reduce( + handle.get_thrust_policy(), + unique_srcs.begin(), + unique_srcs.end(), + [sorted_org_vertices = raft::device_span(sorted_org_vertices.data(), + sorted_org_vertices.size()), + matching_renumbered_vertices = raft::device_span( + matching_renumbered_vertices.data(), + matching_renumbered_vertices.size())] __device__(vertex_t src) { + auto it = thrust::lower_bound( + thrust::seq, sorted_org_vertices.begin(), sorted_org_vertices.end(), src); + return matching_renumbered_vertices[thrust::distance(sorted_org_vertices.begin(), it)]; + }, + std::numeric_limits::lowest(), + thrust::maximum{}); + + auto min_dst_renumbered_vertex = thrust::transform_reduce( + handle.get_thrust_policy(), + unique_dsts.begin(), + unique_dsts.end(), + [sorted_org_vertices = raft::device_span(sorted_org_vertices.data(), + sorted_org_vertices.size()), + matching_renumbered_vertices = raft::device_span( + matching_renumbered_vertices.data(), + matching_renumbered_vertices.size())] __device__(vertex_t dst) { + auto it = thrust::lower_bound( + thrust::seq, sorted_org_vertices.begin(), sorted_org_vertices.end(), dst); + return matching_renumbered_vertices[thrust::distance(sorted_org_vertices.begin(), it)]; + }, + std::numeric_limits::max(), + thrust::minimum{}); + + ASSERT_TRUE(max_src_renumbered_vertex < min_dst_renumbered_vertex) + << "Invariants violated, a source vertex is renumbered to a non-smaller value than a " + "vertex that appear only in the edge list destinations."; + + if (this_label_edgelist_hops) { + thrust::sort_by_key(handle.get_thrust_policy(), + (*unique_src_hops).begin(), + (*unique_src_hops).end(), + unique_srcs.begin()); + rmm::device_uvector min_vertices(usecase.num_hops, handle.get_stream()); + rmm::device_uvector max_vertices(usecase.num_hops, handle.get_stream()); + auto this_label_num_unique_hops = static_cast( + thrust::distance(min_vertices.begin(), + thrust::get<1>(thrust::reduce_by_key(handle.get_thrust_policy(), + (*unique_src_hops).begin(), + (*unique_src_hops).end(), + unique_srcs.begin(), + thrust::make_discard_iterator(), + min_vertices.begin(), + thrust::equal_to{}, + thrust::minimum{})))); + min_vertices.resize(this_label_num_unique_hops, handle.get_stream()); + + thrust::reduce_by_key(handle.get_thrust_policy(), + (*unique_src_hops).begin(), + (*unique_src_hops).end(), + unique_srcs.begin(), + thrust::make_discard_iterator(), + min_vertices.begin(), + thrust::equal_to{}, + thrust::maximum{}); + max_vertices.resize(this_label_num_unique_hops, handle.get_stream()); + + auto num_violations = + thrust::count_if(handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{1}), + thrust::make_counting_iterator(this_label_num_unique_hops), + [min_vertices = raft::device_span(min_vertices.data(), + min_vertices.size()), + max_vertices = raft::device_span( + max_vertices.data(), max_vertices.size())] __device__(size_t i) { + return min_vertices[i] <= max_vertices[i - 1]; + }); + + ASSERT_TRUE(num_violations == 0) + << "Invariant violated, a vertex with a smaller hop is renumbered to a non-smaller " + "value than a vertex with a larger hop."; + } + } + } + } +}; + +TEST_P(Tests_RenumberSampledEdgelist, CheckInt32) +{ + auto param = GetParam(); + run_current_test(param); +} + +TEST_P(Tests_RenumberSampledEdgelist, CheckInt64) +{ + auto param = GetParam(); + run_current_test(param); +} + +INSTANTIATE_TEST_SUITE_P(small_test, + Tests_RenumberSampledEdgelist, + ::testing::Values(RenumberSampledEdgelist_Usecase{16, 64, 1, 1, true}, + RenumberSampledEdgelist_Usecase{16, 64, 3, 1, true}, + RenumberSampledEdgelist_Usecase{16, 512, 1, 8, true}, + RenumberSampledEdgelist_Usecase{16, 512, 3, 8, true})); + +INSTANTIATE_TEST_SUITE_P( + benchmark_test, + Tests_RenumberSampledEdgelist, + ::testing::Values(RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 1, 1, true}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 5, 1, true}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 28, 1, 1 << 20, true}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 28, 5, 1 << 20, true})); + +CUGRAPH_TEST_PROGRAM_MAIN() From 9fe5e13f012d23d9fd950a57401865459a957d98 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Wed, 19 Jul 2023 10:42:49 -0700 Subject: [PATCH 05/17] bug fixes --- .../renumber_sampled_edgelist_impl.cuh | 121 +++++++++++------- .../renumber_sampled_edgelist_test.cu | 30 ++--- 2 files changed, 89 insertions(+), 62 deletions(-) diff --git a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh index 215fd00c255..991b9e3a6c5 100644 --- a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh +++ b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh @@ -51,31 +51,29 @@ compute_renumber_map( raft::device_span edgelist_srcs, std::optional> edgelist_hops, raft::device_span edgelist_dsts, + std::optional> edgelist_label_indices, std::optional, raft::device_span>> label_offsets) { std::optional> unique_label_src_pair_label_indices{std::nullopt}; - rmm::device_uvector unique_label_src_pair_vertices(0, handle.get_stream()); + rmm::device_uvector unique_label_src_pair_vertices( + 0, handle.get_stream()); // sorted by (label, hop, src) + std::optional> sorted_srcs{ + std::nullopt}; // sorted by (label, src), relevant only when edgelist_hops is valid { - rmm::device_uvector srcs(edgelist_srcs.size(), handle.get_stream()); - thrust::copy( - handle.get_thrust_policy(), edgelist_srcs.begin(), edgelist_srcs.end(), srcs.begin()); - if (label_offsets) { - rmm::device_uvector label_indices(edgelist_srcs.size(), handle.get_stream()); - thrust::transform( - handle.get_thrust_policy(), - thrust::make_counting_iterator(size_t{0}), - thrust::make_counting_iterator(edgelist_srcs.size()), - label_indices.begin(), - [offsets = raft::device_span( - std::get<1>(*label_offsets).data() + 1, - std::get<1>(*label_offsets).size())] __device__(size_t i) { - return static_cast(thrust::distance( - offsets.begin(), thrust::upper_bound(thrust::seq, offsets.begin(), offsets.end(), i))); - }); + rmm::device_uvector label_indices((*edgelist_label_indices).size(), + handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), + (*edgelist_label_indices).begin(), + (*edgelist_label_indices).end(), + label_indices.begin()); if (edgelist_hops) { + rmm::device_uvector srcs(edgelist_srcs.size(), handle.get_stream()); + thrust::copy( + handle.get_thrust_policy(), edgelist_srcs.begin(), edgelist_srcs.end(), srcs.begin()); + rmm::device_uvector hops((*edgelist_hops).size(), handle.get_stream()); thrust::copy(handle.get_thrust_policy(), (*edgelist_hops).begin(), @@ -101,6 +99,8 @@ compute_renumber_map( hops.shrink_to_fit(handle.get_stream()); unique_label_src_pair_label_indices = std::move(label_indices); + sorted_srcs = rmm::device_uvector(srcs.size(), handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), srcs.begin(), srcs.end(), (*sorted_srcs).begin()); auto num_labels = std::get<0>(*label_offsets).size(); rmm::device_uvector tmp_label_offsets(num_labels + 1, handle.get_stream()); @@ -148,16 +148,17 @@ compute_renumber_map( unique_label_src_pair_vertices = std::move(segment_sorted_srcs); } else { - rmm::device_uvector segment_sorted_srcs(srcs.size(), handle.get_stream()); + rmm::device_uvector segment_sorted_srcs(edgelist_srcs.size(), + handle.get_stream()); size_t tmp_storage_bytes{0}; rmm::device_uvector d_tmp_storage(0, handle.get_stream()); cub::DeviceSegmentedSort::SortKeys(static_cast(nullptr), tmp_storage_bytes, - srcs.begin(), + edgelist_srcs.begin(), segment_sorted_srcs.begin(), - srcs.size(), + edgelist_srcs.size(), std::get<0>(*label_offsets).size(), std::get<1>(*label_offsets).begin(), std::get<1>(*label_offsets).begin() + 1, @@ -169,9 +170,9 @@ compute_renumber_map( cub::DeviceSegmentedSort::SortKeys(d_tmp_storage.data(), tmp_storage_bytes, - srcs.begin(), + edgelist_srcs.begin(), segment_sorted_srcs.begin(), - srcs.size(), + edgelist_srcs.size(), std::get<0>(*label_offsets).size(), std::get<1>(*label_offsets).begin(), std::get<1>(*label_offsets).begin() + 1, @@ -192,6 +193,10 @@ compute_renumber_map( unique_label_src_pair_vertices = std::move(segment_sorted_srcs); } } else { + rmm::device_uvector srcs(edgelist_srcs.size(), handle.get_stream()); + thrust::copy( + handle.get_thrust_policy(), edgelist_srcs.begin(), edgelist_srcs.end(), srcs.begin()); + if (edgelist_hops) { rmm::device_uvector hops((*edgelist_hops).size(), handle.get_stream()); thrust::copy(handle.get_thrust_policy(), @@ -199,12 +204,20 @@ compute_renumber_map( (*edgelist_hops).end(), hops.begin()); - auto pair_first = thrust::make_zip_iterator(hops.begin(), srcs.begin()); - thrust::sort(handle.get_thrust_policy(), pair_first, pair_first + hops.size()); + auto pair_first = thrust::make_zip_iterator( + srcs.begin(), hops.begin()); // src is a primary key, hop is a secondary key + thrust::sort(handle.get_thrust_policy(), pair_first, pair_first + srcs.size()); srcs.resize( thrust::distance(srcs.begin(), - thrust::unique(handle.get_thrust_policy(), srcs.begin(), srcs.end())), + thrust::get<0>(thrust::unique_by_key( + handle.get_thrust_policy(), srcs.begin(), srcs.end(), hops.begin()))), handle.get_stream()); + hops.resize(srcs.size(), handle.get_stream()); + + sorted_srcs = rmm::device_uvector(srcs.size(), handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), srcs.begin(), srcs.end(), (*sorted_srcs).begin()); + + thrust::sort_by_key(handle.get_thrust_policy(), hops.begin(), hops.end(), srcs.begin()); } else { thrust::sort(handle.get_thrust_policy(), srcs.begin(), srcs.end()); srcs.resize( @@ -225,18 +238,12 @@ compute_renumber_map( thrust::copy( handle.get_thrust_policy(), edgelist_dsts.begin(), edgelist_dsts.end(), dsts.begin()); if (label_offsets) { - rmm::device_uvector label_indices(edgelist_dsts.size(), handle.get_stream()); - thrust::transform( - handle.get_thrust_policy(), - thrust::make_counting_iterator(size_t{0}), - thrust::make_counting_iterator(edgelist_dsts.size()), - label_indices.begin(), - [offsets = raft::device_span( - std::get<1>(*label_offsets).data() + 1, - std::get<1>(*label_offsets).size())] __device__(size_t i) { - return static_cast(thrust::distance( - offsets.begin(), thrust::upper_bound(thrust::seq, offsets.begin(), offsets.end(), i))); - }); + rmm::device_uvector label_indices((*edgelist_label_indices).size(), + handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), + (*edgelist_label_indices).begin(), + (*edgelist_label_indices).end(), + label_indices.begin()); rmm::device_uvector segment_sorted_dsts(dsts.size(), handle.get_stream()); @@ -294,7 +301,8 @@ compute_renumber_map( if (label_offsets) { auto label_src_pair_first = thrust::make_zip_iterator( - (*unique_label_src_pair_label_indices).begin(), unique_label_src_pair_vertices.begin()); + (*unique_label_src_pair_label_indices).begin(), + edgelist_hops ? (*sorted_srcs).begin() : unique_label_src_pair_vertices.begin()); auto label_dst_pair_first = thrust::make_zip_iterator( (*unique_label_dst_pair_label_indices).begin(), unique_label_dst_pair_vertices.begin()); rmm::device_uvector output_label_indices((*unique_label_dst_pair_label_indices).size(), @@ -344,12 +352,13 @@ compute_renumber_map( } else { rmm::device_uvector output_vertices(unique_label_dst_pair_vertices.size(), handle.get_stream()); - auto output_last = thrust::set_difference(handle.get_thrust_policy(), - unique_label_dst_pair_vertices.begin(), - unique_label_dst_pair_vertices.end(), - unique_label_src_pair_vertices.begin(), - unique_label_src_pair_vertices.end(), - output_vertices.begin()); + auto output_last = thrust::set_difference( + handle.get_thrust_policy(), + unique_label_dst_pair_vertices.begin(), + unique_label_dst_pair_vertices.end(), + edgelist_hops ? (*sorted_srcs).begin() : unique_label_src_pair_vertices.begin(), + edgelist_hops ? (*sorted_srcs).end() : unique_label_src_pair_vertices.end(), + output_vertices.begin()); auto num_unique_srcs = unique_label_src_pair_vertices.size(); auto renumber_map = std::move(unique_label_src_pair_vertices); @@ -413,11 +422,30 @@ renumber_sampled_edgelist( } } + std::optional> edgelist_label_indices{std::nullopt}; + if (label_offsets) { + edgelist_label_indices = rmm::device_uvector(edgelist_srcs.size(), handle.get_stream()); + thrust::transform( + handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(edgelist_srcs.size()), + (*edgelist_label_indices).begin(), + [offsets = + raft::device_span(std::get<1>(*label_offsets).data() + 1, + std::get<1>(*label_offsets).size())] __device__(size_t i) { + return static_cast(thrust::distance( + offsets.begin(), thrust::upper_bound(thrust::seq, offsets.begin(), offsets.end(), i))); + }); + } + auto [renumber_map, renumber_map_label_indices] = compute_renumber_map( handle, raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), edgelist_hops, raft::device_span(edgelist_dsts.data(), edgelist_dsts.size()), + edgelist_label_indices ? std::make_optional>( + (*edgelist_label_indices).data(), (*edgelist_label_indices).size()) + : std::nullopt, label_offsets); std::optional> renumber_map_label_offsets{}; @@ -508,7 +536,7 @@ renumber_sampled_edgelist( handle.get_stream()); auto pair_first = - thrust::make_zip_iterator(edgelist_srcs.begin(), (*renumber_map_label_indices).begin()); + thrust::make_zip_iterator(edgelist_srcs.begin(), (*edgelist_label_indices).begin()); thrust::transform( handle.get_thrust_policy(), pair_first, @@ -534,7 +562,7 @@ renumber_sampled_edgelist( }); pair_first = - thrust::make_zip_iterator(edgelist_dsts.begin(), (*renumber_map_label_indices).begin()); + thrust::make_zip_iterator(edgelist_dsts.begin(), (*edgelist_label_indices).begin()); thrust::transform( handle.get_thrust_policy(), pair_first, @@ -558,7 +586,6 @@ renumber_sampled_edgelist( assert(*it == old_vertex); return new_vertices[thrust::distance(old_vertices.begin(), it)]; }); - } else { kv_store_t kv_store(renumber_map.begin(), renumber_map.end(), diff --git a/cpp/tests/sampling/renumber_sampled_edgelist_test.cu b/cpp/tests/sampling/renumber_sampled_edgelist_test.cu index 797be44cc65..32237cfafc3 100644 --- a/cpp/tests/sampling/renumber_sampled_edgelist_test.cu +++ b/cpp/tests/sampling/renumber_sampled_edgelist_test.cu @@ -363,13 +363,13 @@ class Tests_RenumberSampledEdgelist auto this_label_num_unique_hops = static_cast( thrust::distance(min_vertices.begin(), thrust::get<1>(thrust::reduce_by_key(handle.get_thrust_policy(), - (*unique_src_hops).begin(), - (*unique_src_hops).end(), - unique_srcs.begin(), - thrust::make_discard_iterator(), - min_vertices.begin(), - thrust::equal_to{}, - thrust::minimum{})))); + (*unique_src_hops).begin(), + (*unique_src_hops).end(), + unique_srcs.begin(), + thrust::make_discard_iterator(), + min_vertices.begin(), + thrust::equal_to{}, + thrust::minimum{})))); min_vertices.resize(this_label_num_unique_hops, handle.get_stream()); thrust::reduce_by_key(handle.get_thrust_policy(), @@ -416,17 +416,17 @@ TEST_P(Tests_RenumberSampledEdgelist, CheckInt64) INSTANTIATE_TEST_SUITE_P(small_test, Tests_RenumberSampledEdgelist, - ::testing::Values(RenumberSampledEdgelist_Usecase{16, 64, 1, 1, true}, - RenumberSampledEdgelist_Usecase{16, 64, 3, 1, true}, - RenumberSampledEdgelist_Usecase{16, 512, 1, 8, true}, - RenumberSampledEdgelist_Usecase{16, 512, 3, 8, true})); + ::testing::Values(RenumberSampledEdgelist_Usecase{1024, 4096, 1, 1, true}, + RenumberSampledEdgelist_Usecase{1024, 4096, 3, 1, true}, + RenumberSampledEdgelist_Usecase{1024, 32768, 1, 256, true}, + RenumberSampledEdgelist_Usecase{1024, 32768, 3, 256, true})); INSTANTIATE_TEST_SUITE_P( benchmark_test, Tests_RenumberSampledEdgelist, - ::testing::Values(RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 1, 1, true}, - RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 5, 1, true}, - RenumberSampledEdgelist_Usecase{1 << 20, 1 << 28, 1, 1 << 20, true}, - RenumberSampledEdgelist_Usecase{1 << 20, 1 << 28, 5, 1 << 20, true})); + ::testing::Values(RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 1, 1, false}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 5, 1, false}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 28, 1, 1 << 20, false}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 28, 5, 1 << 20, false})); CUGRAPH_TEST_PROGRAM_MAIN() From 9ad3c8f90eca427bf68ec4549b2d80bb81350810 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Wed, 19 Jul 2023 11:44:43 -0700 Subject: [PATCH 06/17] minor tweaks --- .../renumber_sampled_edgelist_impl.cuh | 10 ++++++++ .../renumber_sampled_edgelist_test.cu | 23 +++++++++---------- 2 files changed, 21 insertions(+), 12 deletions(-) diff --git a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh index 991b9e3a6c5..7916148a1c4 100644 --- a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh +++ b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh @@ -390,6 +390,8 @@ renumber_sampled_edgelist( label_offsets, bool do_expensive_check) { + // 1. check input arguments + CUGRAPH_EXPECTS( edgelist_srcs.size() == edgelist_dsts.size(), "Invalid input arguments: edgelist_srcs.size() and edgelist_dsts.size() should coincide."); @@ -422,6 +424,8 @@ renumber_sampled_edgelist( } } + // 2. find label indices for each input edge + std::optional> edgelist_label_indices{std::nullopt}; if (label_offsets) { edgelist_label_indices = rmm::device_uvector(edgelist_srcs.size(), handle.get_stream()); @@ -438,6 +442,8 @@ renumber_sampled_edgelist( }); } + // 3. compute renumber_map + auto [renumber_map, renumber_map_label_indices] = compute_renumber_map( handle, raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), @@ -448,6 +454,8 @@ renumber_sampled_edgelist( : std::nullopt, label_offsets); + // 4. compute renumber map offsets for each label + std::optional> renumber_map_label_offsets{}; if (label_offsets) { auto num_unique_labels = thrust::count_if( @@ -482,6 +490,8 @@ renumber_sampled_edgelist( (*renumber_map_label_offsets).begin()); } + // 5. renumber input edges + if (label_offsets) { rmm::device_uvector new_vertices(renumber_map.size(), handle.get_stream()); thrust::tabulate(handle.get_thrust_policy(), diff --git a/cpp/tests/sampling/renumber_sampled_edgelist_test.cu b/cpp/tests/sampling/renumber_sampled_edgelist_test.cu index 32237cfafc3..e0564207215 100644 --- a/cpp/tests/sampling/renumber_sampled_edgelist_test.cu +++ b/cpp/tests/sampling/renumber_sampled_edgelist_test.cu @@ -62,7 +62,6 @@ class Tests_RenumberSampledEdgelist rmm::device_uvector org_edgelist_srcs(usecase.num_sampled_edges, handle.get_stream()); rmm::device_uvector org_edgelist_dsts(usecase.num_sampled_edges, handle.get_stream()); - // FIXME: need to check the range ([min,max] or [min, max)?) cugraph::detail::uniform_random_fill(handle.get_stream(), org_edgelist_srcs.data(), org_edgelist_srcs.size(), @@ -79,7 +78,6 @@ class Tests_RenumberSampledEdgelist std::optional> edgelist_hops{std::nullopt}; if (usecase.num_hops > 1) { edgelist_hops = rmm::device_uvector(usecase.num_sampled_edges, handle.get_stream()); - // FIXME: need to check the range ([min,max] or [min, max)?) cugraph::detail::uniform_random_fill(handle.get_stream(), (*edgelist_hops).data(), (*edgelist_hops).size(), @@ -414,19 +412,20 @@ TEST_P(Tests_RenumberSampledEdgelist, CheckInt64) run_current_test(param); } -INSTANTIATE_TEST_SUITE_P(small_test, - Tests_RenumberSampledEdgelist, - ::testing::Values(RenumberSampledEdgelist_Usecase{1024, 4096, 1, 1, true}, - RenumberSampledEdgelist_Usecase{1024, 4096, 3, 1, true}, - RenumberSampledEdgelist_Usecase{1024, 32768, 1, 256, true}, - RenumberSampledEdgelist_Usecase{1024, 32768, 3, 256, true})); +INSTANTIATE_TEST_SUITE_P( + small_test, + Tests_RenumberSampledEdgelist, + ::testing::Values(RenumberSampledEdgelist_Usecase{1024, 4096, 1, 1, true}, + RenumberSampledEdgelist_Usecase{1024, 4096, 3, 1, true}, + RenumberSampledEdgelist_Usecase{1024, 32768, 1, 256, true}, + RenumberSampledEdgelist_Usecase{1024, 32768, 3, 256, true})); INSTANTIATE_TEST_SUITE_P( benchmark_test, Tests_RenumberSampledEdgelist, - ::testing::Values(RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 1, 1, false}, - RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 5, 1, false}, - RenumberSampledEdgelist_Usecase{1 << 20, 1 << 28, 1, 1 << 20, false}, - RenumberSampledEdgelist_Usecase{1 << 20, 1 << 28, 5, 1 << 20, false})); + ::testing::Values(RenumberSampledEdgelist_Usecase{1 << 20, 1 << 20, 1, 1, false}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 20, 5, 1, false}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 25, 1, 1 << 20, false}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 25, 5, 1 << 20, false})); CUGRAPH_TEST_PROGRAM_MAIN() From dfbc196d8c00530de88dc009af944e165ebfbb26 Mon Sep 17 00:00:00 2001 From: Charles Hastings Date: Wed, 19 Jul 2023 20:48:35 -0700 Subject: [PATCH 07/17] define API for MFG renumbering in the C API --- cpp/include/cugraph_c/sampling_algorithms.h | 26 +++++++++++ cpp/src/c_api/uniform_neighbor_sampling.cpp | 48 ++++++++++++++++++++- 2 files changed, 73 insertions(+), 1 deletion(-) diff --git a/cpp/include/cugraph_c/sampling_algorithms.h b/cpp/include/cugraph_c/sampling_algorithms.h index 5e792403a88..37124d100dd 100644 --- a/cpp/include/cugraph_c/sampling_algorithms.h +++ b/cpp/include/cugraph_c/sampling_algorithms.h @@ -217,6 +217,14 @@ typedef enum cugraph_prior_sources_behavior_t { cugraph_error_code_t cugraph_sampling_options_create(cugraph_sampling_options_t** options, cugraph_error_t** error); +/** + * @brief Set flag to renumber results + * + * @param options - opaque pointer to the sampling options + * @param value - Boolean value to assign to the option + */ +void cugraph_sampling_set_renumber_results(cugraph_sampling_options_t* options, bool_t value); + /** * @brief Set flag to sample with_replacement * @@ -446,6 +454,24 @@ cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_index( cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_offsets( const cugraph_sample_result_t* result); +/** + * @brief Get the renumber map + * + * @param [in] result The result from a sampling algorithm + * @return type erased array pointing to the renumber map + */ +cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_renumber_map( + const cugraph_sample_result_t* result); + +/** + * @brief Get the renumber map offsets + * + * @param [in] result The result from a sampling algorithm + * @return type erased array pointing to the renumber map offsets + */ +cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_renumber_map_offsets( + const cugraph_sample_result_t* result); + /** * @brief Free a sampling result * diff --git a/cpp/src/c_api/uniform_neighbor_sampling.cpp b/cpp/src/c_api/uniform_neighbor_sampling.cpp index d9dc9ca4d50..6c7bf6296b2 100644 --- a/cpp/src/c_api/uniform_neighbor_sampling.cpp +++ b/cpp/src/c_api/uniform_neighbor_sampling.cpp @@ -37,6 +37,7 @@ struct cugraph_sampling_options_t { bool_t return_hops_{FALSE}; prior_sources_behavior_t prior_sources_behavior_{prior_sources_behavior_t::DEFAULT}; bool_t dedupe_sources_{FALSE}; + bool_t renumber_results_{FALSE}; }; struct cugraph_sample_result_t { @@ -48,6 +49,8 @@ struct cugraph_sample_result_t { cugraph_type_erased_device_array_t* hop_{nullptr}; cugraph_type_erased_device_array_t* label_{nullptr}; cugraph_type_erased_device_array_t* offsets_{nullptr}; + cugraph_type_erased_device_array_t* renumber_map_{nullptr}; + cugraph_type_erased_device_array_t* renumber_map_offsets_{nullptr}; }; } // namespace c_api @@ -226,6 +229,22 @@ struct uniform_neighbor_sampling_functor : public cugraph::c_api::abstract_funct vertex_partition_lasts, do_expensive_check_); + std::optional> renumber_map{std::nullopt}; + std::optional> renumber_map_offsets{std::nullopt}; + +#if 0 + std::tie(src, dst, renumber_map, renumber_map_offsets) = + cugraph::renumber_sampled_edgelist(handle_, + std::move(src), + hop ? std::make_optional(raft::device_span{hop->data(), hop->size()}) : std::nullopt, + std::move(dst), + std::make_optional(std::make_tuple(raft::device_span{edge_label->data(), + edge_label->size()}, + raft::device_span{offsets->data(), + offsets->size()})) + do_expensive_check); +#endif + result_ = new cugraph::c_api::cugraph_sample_result_t{ new cugraph::c_api::cugraph_type_erased_device_array_t(src, graph_->vertex_type_), new cugraph::c_api::cugraph_type_erased_device_array_t(dst, graph_->vertex_type_), @@ -242,7 +261,11 @@ struct uniform_neighbor_sampling_functor : public cugraph::c_api::abstract_funct ? new cugraph::c_api::cugraph_type_erased_device_array_t(edge_label.value(), INT32) : nullptr, (offsets) ? new cugraph::c_api::cugraph_type_erased_device_array_t(offsets.value(), SIZE_T) - : nullptr}; + : nullptr, + (renumber_map) ? new cugraph::c_api::cugraph_type_erased_device_array_t(renumber_map.value(), graph_->vertex_type_) + : nullptr, + (renumber_map_offsets) ? new cugraph::c_api::cugraph_type_erased_device_array_t(renumber_map_offsets.value(), SIZE_T) + : nullptr}; } } }; @@ -263,6 +286,13 @@ extern "C" cugraph_error_code_t cugraph_sampling_options_create( return CUGRAPH_SUCCESS; } +extern "C" void cugraph_sampling_set_renumber_results(cugraph_sampling_options_t* options, + bool_t value) +{ + auto internal_pointer = reinterpret_cast(options); + internal_pointer->renumber_results_ = value; +} + extern "C" void cugraph_sampling_set_with_replacement(cugraph_sampling_options_t* options, bool_t value) { @@ -386,6 +416,22 @@ extern "C" cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_of internal_pointer->offsets_->view()); } +extern "C" cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_renumber_map( + const cugraph_sample_result_t* result) +{ + auto internal_pointer = reinterpret_cast(result); + return reinterpret_cast( + internal_pointer->renumber_map_->view()); +} + +extern "C" cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_renumber_map_offsets( + const cugraph_sample_result_t* result) +{ + auto internal_pointer = reinterpret_cast(result); + return reinterpret_cast( + internal_pointer->renumber_map_offsets_->view()); +} + extern "C" cugraph_error_code_t cugraph_test_uniform_neighborhood_sample_result_create( const cugraph_resource_handle_t* handle, const cugraph_type_erased_device_array_view_t* srcs, From bc5d3e1a66606dc1d3f8c85cdda600011886a4eb Mon Sep 17 00:00:00 2001 From: Charles Hastings Date: Wed, 19 Jul 2023 20:51:00 -0700 Subject: [PATCH 08/17] define API for MFG renumbering in the C API --- cpp/src/c_api/uniform_neighbor_sampling.cpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/cpp/src/c_api/uniform_neighbor_sampling.cpp b/cpp/src/c_api/uniform_neighbor_sampling.cpp index 6c7bf6296b2..fe29b6252bd 100644 --- a/cpp/src/c_api/uniform_neighbor_sampling.cpp +++ b/cpp/src/c_api/uniform_neighbor_sampling.cpp @@ -262,10 +262,12 @@ struct uniform_neighbor_sampling_functor : public cugraph::c_api::abstract_funct : nullptr, (offsets) ? new cugraph::c_api::cugraph_type_erased_device_array_t(offsets.value(), SIZE_T) : nullptr, - (renumber_map) ? new cugraph::c_api::cugraph_type_erased_device_array_t(renumber_map.value(), graph_->vertex_type_) - : nullptr, - (renumber_map_offsets) ? new cugraph::c_api::cugraph_type_erased_device_array_t(renumber_map_offsets.value(), SIZE_T) - : nullptr}; + (renumber_map) ? new cugraph::c_api::cugraph_type_erased_device_array_t( + renumber_map.value(), graph_->vertex_type_) + : nullptr, + (renumber_map_offsets) ? new cugraph::c_api::cugraph_type_erased_device_array_t( + renumber_map_offsets.value(), SIZE_T) + : nullptr}; } } }; From 012d392d12a701ee0a2f949cba9d1051df7cb365 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 20 Jul 2023 10:55:19 -0700 Subject: [PATCH 09/17] memory footprint cut --- .../renumber_sampled_edgelist_impl.cuh | 344 ++++++++++++------ .../renumber_sampled_edgelist_test.cu | 8 +- 2 files changed, 228 insertions(+), 124 deletions(-) diff --git a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh index 7916148a1c4..52397645d02 100644 --- a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh +++ b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh @@ -19,6 +19,7 @@ #include #include +#include #include #include @@ -44,26 +45,31 @@ namespace cugraph { namespace { -template -std::tuple, std::optional>> +template +std::tuple, std::optional>> compute_renumber_map( raft::handle_t const& handle, raft::device_span edgelist_srcs, std::optional> edgelist_hops, raft::device_span edgelist_dsts, - std::optional> edgelist_label_indices, + std::optional> edgelist_label_indices, std::optional, raft::device_span>> label_offsets) { - std::optional> unique_label_src_pair_label_indices{std::nullopt}; + auto approx_edges_to_sort_per_iteration = + static_cast(handle.get_device_properties().multiProcessorCount) * + (1 << 20) /* tuning parameter */; // for segmented sort + + std::optional> unique_label_src_pair_label_indices{ + std::nullopt}; rmm::device_uvector unique_label_src_pair_vertices( 0, handle.get_stream()); // sorted by (label, hop, src) std::optional> sorted_srcs{ std::nullopt}; // sorted by (label, src), relevant only when edgelist_hops is valid { if (label_offsets) { - rmm::device_uvector label_indices((*edgelist_label_indices).size(), - handle.get_stream()); + rmm::device_uvector label_indices((*edgelist_label_indices).size(), + handle.get_stream()); thrust::copy(handle.get_thrust_policy(), (*edgelist_label_indices).begin(), (*edgelist_label_indices).end(), @@ -98,10 +104,6 @@ compute_renumber_map( srcs.shrink_to_fit(handle.get_stream()); hops.shrink_to_fit(handle.get_stream()); - unique_label_src_pair_label_indices = std::move(label_indices); - sorted_srcs = rmm::device_uvector(srcs.size(), handle.get_stream()); - thrust::copy(handle.get_thrust_policy(), srcs.begin(), srcs.end(), (*sorted_srcs).begin()); - auto num_labels = std::get<0>(*label_offsets).size(); rmm::device_uvector tmp_label_offsets(num_labels + 1, handle.get_stream()); tmp_label_offsets.set_element_to_zero_async(0, handle.get_stream()); @@ -112,39 +114,61 @@ compute_renumber_map( thrust::make_counting_iterator(num_labels), tmp_label_offsets.begin() + 1); - rmm::device_uvector segment_sorted_hops(hops.size(), handle.get_stream()); + unique_label_src_pair_label_indices = std::move(label_indices); + sorted_srcs = rmm::device_uvector(srcs.size(), handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), srcs.begin(), srcs.end(), (*sorted_srcs).begin()); + rmm::device_uvector segment_sorted_srcs(srcs.size(), handle.get_stream()); size_t tmp_storage_bytes{0}; rmm::device_uvector d_tmp_storage(0, handle.get_stream()); - cub::DeviceSegmentedSort::SortPairs(static_cast(nullptr), - tmp_storage_bytes, - hops.begin(), - segment_sorted_hops.begin(), - srcs.begin(), - segment_sorted_srcs.begin(), - hops.size(), - num_labels, - tmp_label_offsets.begin(), - tmp_label_offsets.begin() + 1, - handle.get_stream()); - - if (tmp_storage_bytes > d_tmp_storage.size()) { - d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); + auto [h_label_offsets, h_edge_offsets] = detail::compute_offset_aligned_edge_chunks( + handle, + tmp_label_offsets.data(), + static_cast(tmp_label_offsets.size() - 1), + hops.size(), + approx_edges_to_sort_per_iteration); + auto num_chunks = h_label_offsets.size() - 1; + size_t max_chunk_size{0}; + for (size_t i = 0; i < num_chunks; ++i) { + max_chunk_size = std::max(max_chunk_size, + static_cast(h_edge_offsets[i + 1] - h_edge_offsets[i])); + } + rmm::device_uvector segment_sorted_hops(max_chunk_size, handle.get_stream()); + + for (size_t i = 0; i < num_chunks; ++i) { + auto offset_first = + thrust::make_transform_iterator(tmp_label_offsets.data() + h_label_offsets[i], + detail::shift_left_t{h_edge_offsets[i]}); + cub::DeviceSegmentedSort::SortPairs(static_cast(nullptr), + tmp_storage_bytes, + hops.begin() + h_edge_offsets[i], + segment_sorted_hops.begin(), + srcs.begin() + h_edge_offsets[i], + segment_sorted_srcs.begin() + h_edge_offsets[i], + h_edge_offsets[i + 1] - h_edge_offsets[i], + h_label_offsets[i + 1] - h_label_offsets[i], + offset_first, + offset_first + 1, + handle.get_stream()); + + if (tmp_storage_bytes > d_tmp_storage.size()) { + d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); + } + + cub::DeviceSegmentedSort::SortPairs(d_tmp_storage.data(), + tmp_storage_bytes, + hops.begin() + h_edge_offsets[i], + segment_sorted_hops.begin(), + srcs.begin() + h_edge_offsets[i], + segment_sorted_srcs.begin() + h_edge_offsets[i], + h_edge_offsets[i + 1] - h_edge_offsets[i], + h_label_offsets[i + 1] - h_label_offsets[i], + offset_first, + offset_first + 1, + handle.get_stream()); } - - cub::DeviceSegmentedSort::SortPairs(d_tmp_storage.data(), - tmp_storage_bytes, - hops.begin(), - segment_sorted_hops.begin(), - srcs.begin(), - segment_sorted_srcs.begin(), - hops.size(), - num_labels, - tmp_label_offsets.begin(), - tmp_label_offsets.begin() + 1, - handle.get_stream()); unique_label_src_pair_vertices = std::move(segment_sorted_srcs); } else { @@ -154,30 +178,43 @@ compute_renumber_map( size_t tmp_storage_bytes{0}; rmm::device_uvector d_tmp_storage(0, handle.get_stream()); - cub::DeviceSegmentedSort::SortKeys(static_cast(nullptr), - tmp_storage_bytes, - edgelist_srcs.begin(), - segment_sorted_srcs.begin(), - edgelist_srcs.size(), - std::get<0>(*label_offsets).size(), - std::get<1>(*label_offsets).begin(), - std::get<1>(*label_offsets).begin() + 1, - handle.get_stream()); - - if (tmp_storage_bytes > d_tmp_storage.size()) { - d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); + auto [h_label_offsets, h_edge_offsets] = detail::compute_offset_aligned_edge_chunks( + handle, + std::get<1>(*label_offsets).data(), + static_cast(std::get<1>(*label_offsets).size() - 1), + edgelist_srcs.size(), + approx_edges_to_sort_per_iteration); + auto num_chunks = h_label_offsets.size() - 1; + + for (size_t i = 0; i < num_chunks; ++i) { + auto offset_first = + thrust::make_transform_iterator(std::get<1>(*label_offsets).data() + h_label_offsets[i], + detail::shift_left_t{h_edge_offsets[i]}); + cub::DeviceSegmentedSort::SortKeys(static_cast(nullptr), + tmp_storage_bytes, + edgelist_srcs.begin() + h_edge_offsets[i], + segment_sorted_srcs.begin() + h_edge_offsets[i], + h_edge_offsets[i + 1] - h_edge_offsets[i], + h_label_offsets[i + 1] - h_label_offsets[i], + offset_first, + offset_first + 1, + handle.get_stream()); + + if (tmp_storage_bytes > d_tmp_storage.size()) { + d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); + } + + cub::DeviceSegmentedSort::SortKeys(d_tmp_storage.data(), + tmp_storage_bytes, + edgelist_srcs.begin() + h_edge_offsets[i], + segment_sorted_srcs.begin() + h_edge_offsets[i], + h_edge_offsets[i + 1] - h_edge_offsets[i], + h_label_offsets[i + 1] - h_label_offsets[i], + offset_first, + offset_first + 1, + handle.get_stream()); } - cub::DeviceSegmentedSort::SortKeys(d_tmp_storage.data(), - tmp_storage_bytes, - edgelist_srcs.begin(), - segment_sorted_srcs.begin(), - edgelist_srcs.size(), - std::get<0>(*label_offsets).size(), - std::get<1>(*label_offsets).begin(), - std::get<1>(*label_offsets).begin() + 1, - handle.get_stream()); - auto pair_first = thrust::make_zip_iterator(label_indices.begin(), segment_sorted_srcs.begin()); auto num_uniques = static_cast(thrust::distance( @@ -231,15 +268,16 @@ compute_renumber_map( } } - std::optional> unique_label_dst_pair_label_indices{std::nullopt}; + std::optional> unique_label_dst_pair_label_indices{ + std::nullopt}; rmm::device_uvector unique_label_dst_pair_vertices(0, handle.get_stream()); { rmm::device_uvector dsts(edgelist_dsts.size(), handle.get_stream()); thrust::copy( handle.get_thrust_policy(), edgelist_dsts.begin(), edgelist_dsts.end(), dsts.begin()); if (label_offsets) { - rmm::device_uvector label_indices((*edgelist_label_indices).size(), - handle.get_stream()); + rmm::device_uvector label_indices((*edgelist_label_indices).size(), + handle.get_stream()); thrust::copy(handle.get_thrust_policy(), (*edgelist_label_indices).begin(), (*edgelist_label_indices).end(), @@ -250,29 +288,47 @@ compute_renumber_map( size_t tmp_storage_bytes{0}; rmm::device_uvector d_tmp_storage(0, handle.get_stream()); - cub::DeviceSegmentedSort::SortKeys(static_cast(nullptr), - tmp_storage_bytes, - dsts.begin(), - segment_sorted_dsts.begin(), - dsts.size(), - std::get<0>(*label_offsets).size(), - std::get<1>(*label_offsets).begin(), - std::get<1>(*label_offsets).begin() + 1, - handle.get_stream()); + auto [h_label_offsets, h_edge_offsets] = detail::compute_offset_aligned_edge_chunks( + handle, + std::get<1>(*label_offsets).data(), + static_cast(std::get<1>(*label_offsets).size() - 1), + dsts.size(), + approx_edges_to_sort_per_iteration); + auto num_chunks = h_label_offsets.size() - 1; + + for (size_t i = 0; i < num_chunks; ++i) { + auto offset_first = + thrust::make_transform_iterator(std::get<1>(*label_offsets).data() + h_label_offsets[i], + detail::shift_left_t{h_edge_offsets[i]}); + cub::DeviceSegmentedSort::SortKeys(static_cast(nullptr), + tmp_storage_bytes, + dsts.begin() + h_edge_offsets[i], + segment_sorted_dsts.begin() + h_edge_offsets[i], + h_edge_offsets[i + 1] - h_edge_offsets[i], + h_label_offsets[i + 1] - h_label_offsets[i], + offset_first, + offset_first + 1, + handle.get_stream()); - if (tmp_storage_bytes > d_tmp_storage.size()) { - d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); + if (tmp_storage_bytes > d_tmp_storage.size()) { + d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); + } + + cub::DeviceSegmentedSort::SortKeys(d_tmp_storage.data(), + tmp_storage_bytes, + dsts.begin() + h_edge_offsets[i], + segment_sorted_dsts.begin() + h_edge_offsets[i], + h_edge_offsets[i + 1] - h_edge_offsets[i], + h_label_offsets[i + 1] - h_label_offsets[i], + offset_first, + offset_first + 1, + handle.get_stream()); } - cub::DeviceSegmentedSort::SortKeys(d_tmp_storage.data(), - tmp_storage_bytes, - dsts.begin(), - segment_sorted_dsts.begin(), - dsts.size(), - std::get<0>(*label_offsets).size(), - std::get<1>(*label_offsets).begin(), - std::get<1>(*label_offsets).begin() + 1, - handle.get_stream()); + dsts.resize(0, handle.get_stream()); + d_tmp_storage.resize(0, handle.get_stream()); + dsts.shrink_to_fit(handle.get_stream()); + d_tmp_storage.shrink_to_fit(handle.get_stream()); auto pair_first = thrust::make_zip_iterator(label_indices.begin(), segment_sorted_dsts.begin()); @@ -286,7 +342,6 @@ compute_renumber_map( unique_label_dst_pair_label_indices = std::move(label_indices); unique_label_dst_pair_vertices = std::move(segment_sorted_dsts); - } else { thrust::sort(handle.get_thrust_policy(), dsts.begin(), dsts.end()); dsts.resize( @@ -294,9 +349,9 @@ compute_renumber_map( thrust::unique(handle.get_thrust_policy(), dsts.begin(), dsts.end())), handle.get_stream()); dsts.shrink_to_fit(handle.get_stream()); - } - unique_label_dst_pair_vertices = std::move(dsts); + unique_label_dst_pair_vertices = std::move(dsts); + } } if (label_offsets) { @@ -305,8 +360,8 @@ compute_renumber_map( edgelist_hops ? (*sorted_srcs).begin() : unique_label_src_pair_vertices.begin()); auto label_dst_pair_first = thrust::make_zip_iterator( (*unique_label_dst_pair_label_indices).begin(), unique_label_dst_pair_vertices.begin()); - rmm::device_uvector output_label_indices((*unique_label_dst_pair_label_indices).size(), - handle.get_stream()); + rmm::device_uvector output_label_indices( + (*unique_label_dst_pair_label_indices).size(), handle.get_stream()); rmm::device_uvector output_vertices((*unique_label_dst_pair_label_indices).size(), handle.get_stream()); auto output_label_dst_pair_first = @@ -318,6 +373,8 @@ compute_renumber_map( label_src_pair_first, label_src_pair_first + (*unique_label_src_pair_label_indices).size(), output_label_dst_pair_first); + + sorted_srcs = std::nullopt; output_label_indices.resize( thrust::distance(output_label_dst_pair_first, output_label_dst_pair_last), handle.get_stream()); @@ -327,7 +384,7 @@ compute_renumber_map( unique_label_dst_pair_label_indices = std::move(output_label_indices); unique_label_dst_pair_vertices = std::move(output_vertices); - rmm::device_uvector merged_label_indices( + rmm::device_uvector merged_label_indices( (*unique_label_src_pair_label_indices).size() + (*unique_label_dst_pair_label_indices).size(), handle.get_stream()); rmm::device_uvector merged_vertices(merged_label_indices.size(), handle.get_stream()); @@ -360,6 +417,8 @@ compute_renumber_map( edgelist_hops ? (*sorted_srcs).end() : unique_label_src_pair_vertices.end(), output_vertices.begin()); + sorted_srcs = std::nullopt; + auto num_unique_srcs = unique_label_src_pair_vertices.size(); auto renumber_map = std::move(unique_label_src_pair_vertices); renumber_map.resize( @@ -390,8 +449,15 @@ renumber_sampled_edgelist( label_offsets, bool do_expensive_check) { + using label_index_t = uint32_t; + // 1. check input arguments + CUGRAPH_EXPECTS(!label_offsets || (std::get<0>(*label_offsets).size() <= + std::numeric_limits::max()), + "Invalid input arguments: current implementation assumes that the number of " + "unique labels is no larger than std::numeric_limits::max()."); + CUGRAPH_EXPECTS( edgelist_srcs.size() == edgelist_dsts.size(), "Invalid input arguments: edgelist_srcs.size() and edgelist_dsts.size() should coincide."); @@ -426,9 +492,11 @@ renumber_sampled_edgelist( // 2. find label indices for each input edge - std::optional> edgelist_label_indices{std::nullopt}; + // FIXME: how expensive is recomputing label index? + std::optional> edgelist_label_indices{std::nullopt}; if (label_offsets) { - edgelist_label_indices = rmm::device_uvector(edgelist_srcs.size(), handle.get_stream()); + edgelist_label_indices = + rmm::device_uvector(edgelist_srcs.size(), handle.get_stream()); thrust::transform( handle.get_thrust_policy(), thrust::make_counting_iterator(size_t{0}), @@ -449,7 +517,7 @@ renumber_sampled_edgelist( raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), edgelist_hops, raft::device_span(edgelist_dsts.data(), edgelist_dsts.size()), - edgelist_label_indices ? std::make_optional>( + edgelist_label_indices ? std::make_optional>( (*edgelist_label_indices).data(), (*edgelist_label_indices).size()) : std::nullopt, label_offsets); @@ -458,12 +526,28 @@ renumber_sampled_edgelist( std::optional> renumber_map_label_offsets{}; if (label_offsets) { + // FIXME: it seems like count_if also suffers from 32 bit integer overflow +#if 1 // DEBUG + size_t num_unique_labels{0}; + size_t num_scanned{0}; + while (num_scanned < (*renumber_map_label_indices).size()) { + num_unique_labels += thrust::count_if( + handle.get_thrust_policy(), + thrust::make_counting_iterator(num_scanned), + thrust::make_counting_iterator( + std::min((*renumber_map_label_indices).size() - num_scanned, size_t{1024 * 1024 * 1024})), + detail::is_first_in_run_t{(*renumber_map_label_indices).data()}); + num_scanned += + std::min((*renumber_map_label_indices).size() - num_scanned, size_t{1024 * 1024 * 1024}); + } +#else auto num_unique_labels = thrust::count_if( handle.get_thrust_policy(), thrust::make_counting_iterator(size_t{0}), thrust::make_counting_iterator((*renumber_map_label_indices).size()), - detail::is_first_in_run_t{(*renumber_map_label_indices).data()}); - rmm::device_uvector unique_label_indices(num_unique_labels, handle.get_stream()); + detail::is_first_in_run_t{(*renumber_map_label_indices).data()}); +#endif + rmm::device_uvector unique_label_indices(num_unique_labels, handle.get_stream()); rmm::device_uvector vertex_counts(num_unique_labels, handle.get_stream()); thrust::reduce_by_key(handle.get_thrust_policy(), (*renumber_map_label_indices).begin(), @@ -497,7 +581,7 @@ renumber_sampled_edgelist( thrust::tabulate(handle.get_thrust_policy(), new_vertices.begin(), new_vertices.end(), - [label_indices = raft::device_span( + [label_indices = raft::device_span( (*renumber_map_label_indices).data(), (*renumber_map_label_indices).size()), renumber_map_label_offsets = raft::device_span( (*renumber_map_label_offsets).data(), @@ -507,6 +591,9 @@ renumber_sampled_edgelist( return static_cast(i - label_start_offset); }); + (*renumber_map_label_indices).resize(0, handle.get_stream()); + (*renumber_map_label_indices).shrink_to_fit(handle.get_stream()); + auto num_labels = std::get<0>(*label_offsets).size(); rmm::device_uvector segment_sorted_renumber_map(renumber_map.size(), @@ -517,33 +604,50 @@ renumber_sampled_edgelist( size_t tmp_storage_bytes{0}; rmm::device_uvector d_tmp_storage(0, handle.get_stream()); - cub::DeviceSegmentedSort::SortPairs(static_cast(nullptr), - tmp_storage_bytes, - renumber_map.begin(), - segment_sorted_renumber_map.begin(), - new_vertices.begin(), - segment_sorted_new_vertices.begin(), - renumber_map.size(), - num_labels, - (*renumber_map_label_offsets).begin(), - (*renumber_map_label_offsets).begin() + 1, - handle.get_stream()); - - if (tmp_storage_bytes > d_tmp_storage.size()) { - d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); - } + auto approx_edges_to_sort_per_iteration = + static_cast(handle.get_device_properties().multiProcessorCount) * + (1 << 20) /* tuning parameter */; // for segmented sort + + auto [h_label_offsets, h_edge_offsets] = detail::compute_offset_aligned_edge_chunks( + handle, + (*renumber_map_label_offsets).data(), + static_cast((*renumber_map_label_offsets).size() - 1), + renumber_map.size(), + approx_edges_to_sort_per_iteration); + auto num_chunks = h_label_offsets.size() - 1; + + for (size_t i = 0; i < num_chunks; ++i) { + auto offset_first = + thrust::make_transform_iterator((*renumber_map_label_offsets).data() + h_label_offsets[i], + detail::shift_left_t{h_edge_offsets[i]}); + cub::DeviceSegmentedSort::SortPairs(static_cast(nullptr), + tmp_storage_bytes, + renumber_map.begin() + h_edge_offsets[i], + segment_sorted_renumber_map.begin() + h_edge_offsets[i], + new_vertices.begin() + h_edge_offsets[i], + segment_sorted_new_vertices.begin() + h_edge_offsets[i], + h_edge_offsets[i + 1] - h_edge_offsets[i], + h_label_offsets[i + 1] - h_label_offsets[i], + offset_first, + offset_first + 1, + handle.get_stream()); + + if (tmp_storage_bytes > d_tmp_storage.size()) { + d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); + } - cub::DeviceSegmentedSort::SortPairs(d_tmp_storage.data(), - tmp_storage_bytes, - renumber_map.begin(), - segment_sorted_renumber_map.begin(), - new_vertices.begin(), - segment_sorted_new_vertices.begin(), - renumber_map.size(), - num_labels, - (*renumber_map_label_offsets).begin(), - (*renumber_map_label_offsets).begin() + 1, - handle.get_stream()); + cub::DeviceSegmentedSort::SortPairs(d_tmp_storage.data(), + tmp_storage_bytes, + renumber_map.begin() + h_edge_offsets[i], + segment_sorted_renumber_map.begin() + h_edge_offsets[i], + new_vertices.begin() + h_edge_offsets[i], + segment_sorted_new_vertices.begin() + h_edge_offsets[i], + h_edge_offsets[i + 1] - h_edge_offsets[i], + h_label_offsets[i + 1] - h_label_offsets[i], + offset_first, + offset_first + 1, + handle.get_stream()); + } auto pair_first = thrust::make_zip_iterator(edgelist_srcs.begin(), (*edgelist_label_indices).begin()); diff --git a/cpp/tests/sampling/renumber_sampled_edgelist_test.cu b/cpp/tests/sampling/renumber_sampled_edgelist_test.cu index e0564207215..be907846d1b 100644 --- a/cpp/tests/sampling/renumber_sampled_edgelist_test.cu +++ b/cpp/tests/sampling/renumber_sampled_edgelist_test.cu @@ -423,9 +423,9 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( benchmark_test, Tests_RenumberSampledEdgelist, - ::testing::Values(RenumberSampledEdgelist_Usecase{1 << 20, 1 << 20, 1, 1, false}, - RenumberSampledEdgelist_Usecase{1 << 20, 1 << 20, 5, 1, false}, - RenumberSampledEdgelist_Usecase{1 << 20, 1 << 25, 1, 1 << 20, false}, - RenumberSampledEdgelist_Usecase{1 << 20, 1 << 25, 5, 1 << 20, false})); + ::testing::Values(RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 1, 1, false}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 5, 1, false}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 29, 1, 1 << 20, false}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 29, 5, 1 << 20, false})); CUGRAPH_TEST_PROGRAM_MAIN() From 63759b885aadb6683f6ec3ef3bb5d34125a9dddd Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 20 Jul 2023 10:55:49 -0700 Subject: [PATCH 10/17] code improvemnt --- cpp/src/structure/detail/structure_utils.cuh | 13 ++++--------- 1 file changed, 4 insertions(+), 9 deletions(-) diff --git a/cpp/src/structure/detail/structure_utils.cuh b/cpp/src/structure/detail/structure_utils.cuh index 6887acf1af4..e24eb5dc81b 100644 --- a/cpp/src/structure/detail/structure_utils.cuh +++ b/cpp/src/structure/detail/structure_utils.cuh @@ -18,6 +18,7 @@ #include #include +#include #include #include @@ -77,12 +78,6 @@ struct update_edge_t { } }; -template -struct rebase_offset_t { - edge_t base_offset{}; - __device__ edge_t operator()(edge_t offset) const { return offset - base_offset; } -}; - template rmm::device_uvector expand_sparse_offsets(raft::device_span offsets, idx_t base_idx, @@ -360,7 +355,7 @@ void sort_adjacency_list(raft::handle_t const& handle, for (size_t i = 0; i < num_chunks; ++i) { size_t tmp_storage_bytes{0}; auto offset_first = thrust::make_transform_iterator( - offsets.data() + h_vertex_offsets[i], rebase_offset_t{h_edge_offsets[i]}); + offsets.data() + h_vertex_offsets[i], shift_left_t{h_edge_offsets[i]}); cub::DeviceSegmentedSort::SortPairs(static_cast(nullptr), tmp_storage_bytes, index_first + h_edge_offsets[i], @@ -408,7 +403,7 @@ void sort_adjacency_list(raft::handle_t const& handle, for (size_t i = 0; i < num_chunks; ++i) { size_t tmp_storage_bytes{0}; auto offset_first = thrust::make_transform_iterator( - offsets.data() + h_vertex_offsets[i], rebase_offset_t{h_edge_offsets[i]}); + offsets.data() + h_vertex_offsets[i], shift_left_t{h_edge_offsets[i]}); cub::DeviceSegmentedSort::SortPairs(static_cast(nullptr), tmp_storage_bytes, index_first + h_edge_offsets[i], @@ -492,7 +487,7 @@ void sort_adjacency_list(raft::handle_t const& handle, for (size_t i = 0; i < num_chunks; ++i) { size_t tmp_storage_bytes{0}; auto offset_first = thrust::make_transform_iterator(offsets.data() + h_vertex_offsets[i], - rebase_offset_t{h_edge_offsets[i]}); + shift_left_t{h_edge_offsets[i]}); cub::DeviceSegmentedSort::SortKeys(static_cast(nullptr), tmp_storage_bytes, index_first + h_edge_offsets[i], From e21bcd0b2b03c88c49c311d9b7804ca65ae6a2eb Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 20 Jul 2023 16:16:54 -0700 Subject: [PATCH 11/17] bug fix and memory footprint optimization --- .../renumber_sampled_edgelist_impl.cuh | 122 +++++++++--------- 1 file changed, 63 insertions(+), 59 deletions(-) diff --git a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh index 52397645d02..0f57c217580 100644 --- a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh +++ b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh @@ -52,14 +52,30 @@ compute_renumber_map( raft::device_span edgelist_srcs, std::optional> edgelist_hops, raft::device_span edgelist_dsts, - std::optional> edgelist_label_indices, - std::optional, raft::device_span>> + std::optional> label_offsets) { auto approx_edges_to_sort_per_iteration = static_cast(handle.get_device_properties().multiProcessorCount) * (1 << 20) /* tuning parameter */; // for segmented sort + std::optional> edgelist_label_indices{std::nullopt}; + if (label_offsets) { + edgelist_label_indices = + rmm::device_uvector(edgelist_srcs.size(), handle.get_stream()); + thrust::transform( + handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(edgelist_srcs.size()), + (*edgelist_label_indices).begin(), + [offsets = raft::device_span( + (*label_offsets).data() + 1, + (*label_offsets).size() - 1)] __device__(size_t i) { + return static_cast(thrust::distance( + offsets.begin(), thrust::upper_bound(thrust::seq, offsets.begin(), offsets.end(), i))); + }); + } + std::optional> unique_label_src_pair_label_indices{ std::nullopt}; rmm::device_uvector unique_label_src_pair_vertices( @@ -104,7 +120,7 @@ compute_renumber_map( srcs.shrink_to_fit(handle.get_stream()); hops.shrink_to_fit(handle.get_stream()); - auto num_labels = std::get<0>(*label_offsets).size(); + auto num_labels = (*label_offsets).size() - 1; rmm::device_uvector tmp_label_offsets(num_labels + 1, handle.get_stream()); tmp_label_offsets.set_element_to_zero_async(0, handle.get_stream()); thrust::upper_bound(handle.get_thrust_policy(), @@ -120,7 +136,6 @@ compute_renumber_map( rmm::device_uvector segment_sorted_srcs(srcs.size(), handle.get_stream()); - size_t tmp_storage_bytes{0}; rmm::device_uvector d_tmp_storage(0, handle.get_stream()); auto [h_label_offsets, h_edge_offsets] = detail::compute_offset_aligned_edge_chunks( @@ -138,6 +153,8 @@ compute_renumber_map( rmm::device_uvector segment_sorted_hops(max_chunk_size, handle.get_stream()); for (size_t i = 0; i < num_chunks; ++i) { + size_t tmp_storage_bytes{0}; + auto offset_first = thrust::make_transform_iterator(tmp_label_offsets.data() + h_label_offsets[i], detail::shift_left_t{h_edge_offsets[i]}); @@ -169,26 +186,29 @@ compute_renumber_map( offset_first + 1, handle.get_stream()); } + d_tmp_storage.resize(0, handle.get_stream()); + d_tmp_storage.shrink_to_fit(handle.get_stream()); unique_label_src_pair_vertices = std::move(segment_sorted_srcs); } else { rmm::device_uvector segment_sorted_srcs(edgelist_srcs.size(), handle.get_stream()); - size_t tmp_storage_bytes{0}; rmm::device_uvector d_tmp_storage(0, handle.get_stream()); auto [h_label_offsets, h_edge_offsets] = detail::compute_offset_aligned_edge_chunks( handle, - std::get<1>(*label_offsets).data(), - static_cast(std::get<1>(*label_offsets).size() - 1), + (*label_offsets).data(), + static_cast((*label_offsets).size() - 1), edgelist_srcs.size(), approx_edges_to_sort_per_iteration); auto num_chunks = h_label_offsets.size() - 1; for (size_t i = 0; i < num_chunks; ++i) { + size_t tmp_storage_bytes{0}; + auto offset_first = - thrust::make_transform_iterator(std::get<1>(*label_offsets).data() + h_label_offsets[i], + thrust::make_transform_iterator((*label_offsets).data() + h_label_offsets[i], detail::shift_left_t{h_edge_offsets[i]}); cub::DeviceSegmentedSort::SortKeys(static_cast(nullptr), tmp_storage_bytes, @@ -214,6 +234,8 @@ compute_renumber_map( offset_first + 1, handle.get_stream()); } + d_tmp_storage.resize(0, handle.get_stream()); + d_tmp_storage.shrink_to_fit(handle.get_stream()); auto pair_first = thrust::make_zip_iterator(label_indices.begin(), segment_sorted_srcs.begin()); @@ -285,20 +307,21 @@ compute_renumber_map( rmm::device_uvector segment_sorted_dsts(dsts.size(), handle.get_stream()); - size_t tmp_storage_bytes{0}; rmm::device_uvector d_tmp_storage(0, handle.get_stream()); auto [h_label_offsets, h_edge_offsets] = detail::compute_offset_aligned_edge_chunks( handle, - std::get<1>(*label_offsets).data(), - static_cast(std::get<1>(*label_offsets).size() - 1), + (*label_offsets).data(), + static_cast((*label_offsets).size() - 1), dsts.size(), approx_edges_to_sort_per_iteration); auto num_chunks = h_label_offsets.size() - 1; for (size_t i = 0; i < num_chunks; ++i) { + size_t tmp_storage_bytes{0}; + auto offset_first = - thrust::make_transform_iterator(std::get<1>(*label_offsets).data() + h_label_offsets[i], + thrust::make_transform_iterator((*label_offsets).data() + h_label_offsets[i], detail::shift_left_t{h_edge_offsets[i]}); cub::DeviceSegmentedSort::SortKeys(static_cast(nullptr), tmp_storage_bytes, @@ -354,6 +377,8 @@ compute_renumber_map( } } + edgelist_label_indices = std::nullopt; + if (label_offsets) { auto label_src_pair_first = thrust::make_zip_iterator( (*unique_label_src_pair_label_indices).begin(), @@ -490,63 +515,24 @@ renumber_sampled_edgelist( } } - // 2. find label indices for each input edge - - // FIXME: how expensive is recomputing label index? - std::optional> edgelist_label_indices{std::nullopt}; - if (label_offsets) { - edgelist_label_indices = - rmm::device_uvector(edgelist_srcs.size(), handle.get_stream()); - thrust::transform( - handle.get_thrust_policy(), - thrust::make_counting_iterator(size_t{0}), - thrust::make_counting_iterator(edgelist_srcs.size()), - (*edgelist_label_indices).begin(), - [offsets = - raft::device_span(std::get<1>(*label_offsets).data() + 1, - std::get<1>(*label_offsets).size())] __device__(size_t i) { - return static_cast(thrust::distance( - offsets.begin(), thrust::upper_bound(thrust::seq, offsets.begin(), offsets.end(), i))); - }); - } - - // 3. compute renumber_map + // 2. compute renumber_map - auto [renumber_map, renumber_map_label_indices] = compute_renumber_map( + auto [renumber_map, renumber_map_label_indices] = compute_renumber_map( handle, raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), edgelist_hops, raft::device_span(edgelist_dsts.data(), edgelist_dsts.size()), - edgelist_label_indices ? std::make_optional>( - (*edgelist_label_indices).data(), (*edgelist_label_indices).size()) - : std::nullopt, - label_offsets); + label_offsets ? std::make_optional>(std::get<1>(*label_offsets)) : std::nullopt); - // 4. compute renumber map offsets for each label + // 3. compute renumber map offsets for each label std::optional> renumber_map_label_offsets{}; if (label_offsets) { - // FIXME: it seems like count_if also suffers from 32 bit integer overflow -#if 1 // DEBUG - size_t num_unique_labels{0}; - size_t num_scanned{0}; - while (num_scanned < (*renumber_map_label_indices).size()) { - num_unique_labels += thrust::count_if( - handle.get_thrust_policy(), - thrust::make_counting_iterator(num_scanned), - thrust::make_counting_iterator( - std::min((*renumber_map_label_indices).size() - num_scanned, size_t{1024 * 1024 * 1024})), - detail::is_first_in_run_t{(*renumber_map_label_indices).data()}); - num_scanned += - std::min((*renumber_map_label_indices).size() - num_scanned, size_t{1024 * 1024 * 1024}); - } -#else auto num_unique_labels = thrust::count_if( handle.get_thrust_policy(), thrust::make_counting_iterator(size_t{0}), thrust::make_counting_iterator((*renumber_map_label_indices).size()), detail::is_first_in_run_t{(*renumber_map_label_indices).data()}); -#endif rmm::device_uvector unique_label_indices(num_unique_labels, handle.get_stream()); rmm::device_uvector vertex_counts(num_unique_labels, handle.get_stream()); thrust::reduce_by_key(handle.get_thrust_policy(), @@ -574,7 +560,7 @@ renumber_sampled_edgelist( (*renumber_map_label_offsets).begin()); } - // 5. renumber input edges + // 4. renumber input edges if (label_offsets) { rmm::device_uvector new_vertices(renumber_map.size(), handle.get_stream()); @@ -601,7 +587,6 @@ renumber_sampled_edgelist( rmm::device_uvector segment_sorted_new_vertices(new_vertices.size(), handle.get_stream()); - size_t tmp_storage_bytes{0}; rmm::device_uvector d_tmp_storage(0, handle.get_stream()); auto approx_edges_to_sort_per_iteration = @@ -617,6 +602,8 @@ renumber_sampled_edgelist( auto num_chunks = h_label_offsets.size() - 1; for (size_t i = 0; i < num_chunks; ++i) { + size_t tmp_storage_bytes{0}; + auto offset_first = thrust::make_transform_iterator((*renumber_map_label_offsets).data() + h_label_offsets[i], detail::shift_left_t{h_edge_offsets[i]}); @@ -648,9 +635,26 @@ renumber_sampled_edgelist( offset_first + 1, handle.get_stream()); } + new_vertices.resize(0, handle.get_stream()); + d_tmp_storage.resize(0, handle.get_stream()); + new_vertices.shrink_to_fit(handle.get_stream()); + d_tmp_storage.shrink_to_fit(handle.get_stream()); + + rmm::device_uvector edgelist_label_indices(edgelist_srcs.size(), handle.get_stream()); + thrust::transform( + handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(edgelist_srcs.size()), + edgelist_label_indices.begin(), + [offsets = raft::device_span( + std::get<1>(*label_offsets).data() + 1, + std::get<1>(*label_offsets).size() - 1)] __device__(size_t i) { + return static_cast(thrust::distance( + offsets.begin(), thrust::upper_bound(thrust::seq, offsets.begin(), offsets.end(), i))); + }); auto pair_first = - thrust::make_zip_iterator(edgelist_srcs.begin(), (*edgelist_label_indices).begin()); + thrust::make_zip_iterator(edgelist_srcs.begin(), edgelist_label_indices.begin()); thrust::transform( handle.get_thrust_policy(), pair_first, @@ -676,7 +680,7 @@ renumber_sampled_edgelist( }); pair_first = - thrust::make_zip_iterator(edgelist_dsts.begin(), (*edgelist_label_indices).begin()); + thrust::make_zip_iterator(edgelist_dsts.begin(), edgelist_label_indices.begin()); thrust::transform( handle.get_thrust_policy(), pair_first, From 70100812cd24188529754eeaa0d5d75b30f94574 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 20 Jul 2023 16:18:18 -0700 Subject: [PATCH 12/17] bug fix --- .../renumber_sampled_edgelist_test.cu | 96 +++++++++++-------- 1 file changed, 55 insertions(+), 41 deletions(-) diff --git a/cpp/tests/sampling/renumber_sampled_edgelist_test.cu b/cpp/tests/sampling/renumber_sampled_edgelist_test.cu index be907846d1b..49cb674214c 100644 --- a/cpp/tests/sampling/renumber_sampled_edgelist_test.cu +++ b/cpp/tests/sampling/renumber_sampled_edgelist_test.cu @@ -170,9 +170,9 @@ class Tests_RenumberSampledEdgelist for (size_t i = 0; i < usecase.num_labels; ++i) { size_t edgelist_start_offset = label_offsets ? std::get<1>(*label_offsets).element(i, handle.get_stream()) : size_t{0}; - size_t edgelist_end_offset = label_offsets - ? std::get<1>(*label_offsets).element(i, handle.get_stream()) - : usecase.num_sampled_edges; + size_t edgelist_end_offset = + label_offsets ? std::get<1>(*label_offsets).element(i + 1, handle.get_stream()) + : usecase.num_sampled_edges; auto this_label_org_edgelist_srcs = raft::device_span(org_edgelist_srcs.data() + edgelist_start_offset, edgelist_end_offset - edgelist_start_offset); @@ -204,30 +204,29 @@ class Tests_RenumberSampledEdgelist // check un-renumbering recovers the original edge list - auto pair_first = thrust::make_zip_iterator(this_label_org_edgelist_srcs.begin(), + auto pair_first = thrust::make_zip_iterator(this_label_org_edgelist_srcs.begin(), this_label_renumbered_edgelist_srcs.begin()); - auto num_renumber_errors = thrust::count_if( - handle.get_thrust_policy(), - pair_first, - pair_first + this_label_org_edgelist_srcs.size(), - [this_label_renumber_map] __device__(auto pair) { - auto org = thrust::get<0>(pair); - auto renumbered = thrust::get<1>(pair); - return this_label_renumber_map[thrust::get<1>(pair)] != thrust::get<0>(pair); - }); + auto num_renumber_errors = + thrust::count_if(handle.get_thrust_policy(), + pair_first, + pair_first + this_label_org_edgelist_srcs.size(), + [this_label_renumber_map] __device__(auto pair) { + auto org = thrust::get<0>(pair); + auto renumbered = thrust::get<1>(pair); + return this_label_renumber_map[renumbered] != org; + }); ASSERT_TRUE(num_renumber_errors == 0) << "Renumber error in edge list sources."; pair_first = thrust::make_zip_iterator(this_label_org_edgelist_dsts.begin(), this_label_renumbered_edgelist_dsts.begin()); - num_renumber_errors = thrust::count_if( - handle.get_thrust_policy(), - pair_first, - pair_first + this_label_org_edgelist_dsts.size(), - [this_label_renumber_map] __device__(auto pair) { - auto org = thrust::get<0>(pair); - auto renumbered = thrust::get<1>(pair); - return this_label_renumber_map[thrust::get<1>(pair)] != thrust::get<0>(pair); - }); + num_renumber_errors = thrust::count_if(handle.get_thrust_policy(), + pair_first, + pair_first + this_label_org_edgelist_dsts.size(), + [this_label_renumber_map] __device__(auto pair) { + auto org = thrust::get<0>(pair); + auto renumbered = thrust::get<1>(pair); + return this_label_renumber_map[renumbered] != org; + }); ASSERT_TRUE(num_renumber_errors == 0) << "Renumber error in edge list destinations."; // check the invariants in renumber_map (1. vertices appeared in edge list sources should @@ -255,11 +254,13 @@ class Tests_RenumberSampledEdgelist auto pair_first = thrust::make_zip_iterator(unique_srcs.begin(), (*unique_src_hops).begin()); thrust::sort(handle.get_thrust_policy(), pair_first, pair_first + unique_srcs.size()); - unique_srcs.resize(thrust::distance(pair_first, - thrust::unique(handle.get_thrust_policy(), - pair_first, - pair_first + unique_srcs.size())), - handle.get_stream()); + unique_srcs.resize( + thrust::distance(unique_srcs.begin(), + thrust::get<0>(thrust::unique_by_key(handle.get_thrust_policy(), + unique_srcs.begin(), + unique_srcs.end(), + (*unique_src_hops).begin()))), + handle.get_stream()); (*unique_src_hops).resize(unique_srcs.size(), handle.get_stream()); } else { thrust::sort(handle.get_thrust_policy(), unique_srcs.begin(), unique_srcs.end()); @@ -358,12 +359,25 @@ class Tests_RenumberSampledEdgelist unique_srcs.begin()); rmm::device_uvector min_vertices(usecase.num_hops, handle.get_stream()); rmm::device_uvector max_vertices(usecase.num_hops, handle.get_stream()); + auto unique_renumbered_src_first = thrust::make_transform_iterator( + unique_srcs.begin(), + [sorted_org_vertices = raft::device_span(sorted_org_vertices.data(), + sorted_org_vertices.size()), + matching_renumbered_vertices = raft::device_span( + matching_renumbered_vertices.data(), + matching_renumbered_vertices.size())] __device__(vertex_t src) { + auto it = thrust::lower_bound( + thrust::seq, sorted_org_vertices.begin(), sorted_org_vertices.end(), src); + return matching_renumbered_vertices[thrust::distance(sorted_org_vertices.begin(), + it)]; + }); + auto this_label_num_unique_hops = static_cast( thrust::distance(min_vertices.begin(), thrust::get<1>(thrust::reduce_by_key(handle.get_thrust_policy(), (*unique_src_hops).begin(), (*unique_src_hops).end(), - unique_srcs.begin(), + unique_renumbered_src_first, thrust::make_discard_iterator(), min_vertices.begin(), thrust::equal_to{}, @@ -373,9 +387,9 @@ class Tests_RenumberSampledEdgelist thrust::reduce_by_key(handle.get_thrust_policy(), (*unique_src_hops).begin(), (*unique_src_hops).end(), - unique_srcs.begin(), + unique_renumbered_src_first, thrust::make_discard_iterator(), - min_vertices.begin(), + max_vertices.begin(), thrust::equal_to{}, thrust::maximum{}); max_vertices.resize(this_label_num_unique_hops, handle.get_stream()); @@ -412,20 +426,20 @@ TEST_P(Tests_RenumberSampledEdgelist, CheckInt64) run_current_test(param); } -INSTANTIATE_TEST_SUITE_P( - small_test, - Tests_RenumberSampledEdgelist, - ::testing::Values(RenumberSampledEdgelist_Usecase{1024, 4096, 1, 1, true}, - RenumberSampledEdgelist_Usecase{1024, 4096, 3, 1, true}, - RenumberSampledEdgelist_Usecase{1024, 32768, 1, 256, true}, - RenumberSampledEdgelist_Usecase{1024, 32768, 3, 256, true})); +INSTANTIATE_TEST_SUITE_P(small_test, + Tests_RenumberSampledEdgelist, + ::testing::Values(RenumberSampledEdgelist_Usecase{1024, 4096, 1, 1, true}, + RenumberSampledEdgelist_Usecase{1024, 4096, 3, 1, true}, + RenumberSampledEdgelist_Usecase{ + 1024, 32768, 1, 256, true}, + RenumberSampledEdgelist_Usecase{1024, 32768, 3, 256, true})); INSTANTIATE_TEST_SUITE_P( benchmark_test, Tests_RenumberSampledEdgelist, - ::testing::Values(RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 1, 1, false}, - RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 5, 1, false}, - RenumberSampledEdgelist_Usecase{1 << 20, 1 << 29, 1, 1 << 20, false}, - RenumberSampledEdgelist_Usecase{1 << 20, 1 << 29, 5, 1 << 20, false})); + ::testing::Values(RenumberSampledEdgelist_Usecase{1 << 20, 1 << 20, 1, 1, false}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 20, 5, 1, false}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 1, 1 << 20, false}, + RenumberSampledEdgelist_Usecase{1 << 20, 1 << 24, 5, 1 << 20, false})); CUGRAPH_TEST_PROGRAM_MAIN() From d4d540759692019e583aa266cc16ae43cc9e5938 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 20 Jul 2023 16:32:12 -0700 Subject: [PATCH 13/17] clang-format --- .../renumber_sampled_edgelist_impl.cuh | 50 +++++++++---------- .../sampling/renumber_sampled_edgelist_sg.cu | 18 +++---- cpp/src/structure/detail/structure_utils.cuh | 8 +-- .../renumber_sampled_edgelist_test.cu | 14 +++--- 4 files changed, 44 insertions(+), 46 deletions(-) diff --git a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh index 0f57c217580..ae9dd64348a 100644 --- a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh +++ b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh @@ -47,13 +47,11 @@ namespace { template std::tuple, std::optional>> -compute_renumber_map( - raft::handle_t const& handle, - raft::device_span edgelist_srcs, - std::optional> edgelist_hops, - raft::device_span edgelist_dsts, - std::optional> - label_offsets) +compute_renumber_map(raft::handle_t const& handle, + raft::device_span edgelist_srcs, + std::optional> edgelist_hops, + raft::device_span edgelist_dsts, + std::optional> label_offsets) { auto approx_edges_to_sort_per_iteration = static_cast(handle.get_device_properties().multiProcessorCount) * @@ -69,8 +67,7 @@ compute_renumber_map( thrust::make_counting_iterator(edgelist_srcs.size()), (*edgelist_label_indices).begin(), [offsets = raft::device_span( - (*label_offsets).data() + 1, - (*label_offsets).size() - 1)] __device__(size_t i) { + (*label_offsets).data() + 1, (*label_offsets).size() - 1)] __device__(size_t i) { return static_cast(thrust::distance( offsets.begin(), thrust::upper_bound(thrust::seq, offsets.begin(), offsets.end(), i))); }); @@ -309,12 +306,12 @@ compute_renumber_map( rmm::device_uvector d_tmp_storage(0, handle.get_stream()); - auto [h_label_offsets, h_edge_offsets] = detail::compute_offset_aligned_edge_chunks( - handle, - (*label_offsets).data(), - static_cast((*label_offsets).size() - 1), - dsts.size(), - approx_edges_to_sort_per_iteration); + auto [h_label_offsets, h_edge_offsets] = + detail::compute_offset_aligned_edge_chunks(handle, + (*label_offsets).data(), + static_cast((*label_offsets).size() - 1), + dsts.size(), + approx_edges_to_sort_per_iteration); auto num_chunks = h_label_offsets.size() - 1; for (size_t i = 0; i < num_chunks; ++i) { @@ -517,12 +514,15 @@ renumber_sampled_edgelist( // 2. compute renumber_map - auto [renumber_map, renumber_map_label_indices] = compute_renumber_map( - handle, - raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), - edgelist_hops, - raft::device_span(edgelist_dsts.data(), edgelist_dsts.size()), - label_offsets ? std::make_optional>(std::get<1>(*label_offsets)) : std::nullopt); + auto [renumber_map, renumber_map_label_indices] = + compute_renumber_map( + handle, + raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), + edgelist_hops, + raft::device_span(edgelist_dsts.data(), edgelist_dsts.size()), + label_offsets + ? std::make_optional>(std::get<1>(*label_offsets)) + : std::nullopt); // 3. compute renumber map offsets for each label @@ -640,7 +640,8 @@ renumber_sampled_edgelist( new_vertices.shrink_to_fit(handle.get_stream()); d_tmp_storage.shrink_to_fit(handle.get_stream()); - rmm::device_uvector edgelist_label_indices(edgelist_srcs.size(), handle.get_stream()); + rmm::device_uvector edgelist_label_indices(edgelist_srcs.size(), + handle.get_stream()); thrust::transform( handle.get_thrust_policy(), thrust::make_counting_iterator(size_t{0}), @@ -651,7 +652,7 @@ renumber_sampled_edgelist( std::get<1>(*label_offsets).size() - 1)] __device__(size_t i) { return static_cast(thrust::distance( offsets.begin(), thrust::upper_bound(thrust::seq, offsets.begin(), offsets.end(), i))); - }); + }); auto pair_first = thrust::make_zip_iterator(edgelist_srcs.begin(), edgelist_label_indices.begin()); @@ -679,8 +680,7 @@ renumber_sampled_edgelist( return *(new_vertices.begin() + thrust::distance(old_vertices.begin(), it)); }); - pair_first = - thrust::make_zip_iterator(edgelist_dsts.begin(), edgelist_label_indices.begin()); + pair_first = thrust::make_zip_iterator(edgelist_dsts.begin(), edgelist_label_indices.begin()); thrust::transform( handle.get_thrust_policy(), pair_first, diff --git a/cpp/src/sampling/renumber_sampled_edgelist_sg.cu b/cpp/src/sampling/renumber_sampled_edgelist_sg.cu index 522440108da..9ffa3cb67ad 100644 --- a/cpp/src/sampling/renumber_sampled_edgelist_sg.cu +++ b/cpp/src/sampling/renumber_sampled_edgelist_sg.cu @@ -20,11 +20,10 @@ namespace cugraph { -template -std::tuple, - rmm::device_uvector, - rmm::device_uvector, - std::optional>> +template std::tuple, + rmm::device_uvector, + rmm::device_uvector, + std::optional>> renumber_sampled_edgelist( raft::handle_t const& handle, rmm::device_uvector&& edgelist_srcs, @@ -34,11 +33,10 @@ renumber_sampled_edgelist( label_offsets, bool do_expensive_check); -template -std::tuple, - rmm::device_uvector, - rmm::device_uvector, - std::optional>> +template std::tuple, + rmm::device_uvector, + rmm::device_uvector, + std::optional>> renumber_sampled_edgelist( raft::handle_t const& handle, rmm::device_uvector&& edgelist_srcs, diff --git a/cpp/src/structure/detail/structure_utils.cuh b/cpp/src/structure/detail/structure_utils.cuh index e24eb5dc81b..b6c292324fa 100644 --- a/cpp/src/structure/detail/structure_utils.cuh +++ b/cpp/src/structure/detail/structure_utils.cuh @@ -354,8 +354,8 @@ void sort_adjacency_list(raft::handle_t const& handle, if constexpr (std::is_arithmetic_v) { for (size_t i = 0; i < num_chunks; ++i) { size_t tmp_storage_bytes{0}; - auto offset_first = thrust::make_transform_iterator( - offsets.data() + h_vertex_offsets[i], shift_left_t{h_edge_offsets[i]}); + auto offset_first = thrust::make_transform_iterator(offsets.data() + h_vertex_offsets[i], + shift_left_t{h_edge_offsets[i]}); cub::DeviceSegmentedSort::SortPairs(static_cast(nullptr), tmp_storage_bytes, index_first + h_edge_offsets[i], @@ -402,8 +402,8 @@ void sort_adjacency_list(raft::handle_t const& handle, edge_t{0}); for (size_t i = 0; i < num_chunks; ++i) { size_t tmp_storage_bytes{0}; - auto offset_first = thrust::make_transform_iterator( - offsets.data() + h_vertex_offsets[i], shift_left_t{h_edge_offsets[i]}); + auto offset_first = thrust::make_transform_iterator(offsets.data() + h_vertex_offsets[i], + shift_left_t{h_edge_offsets[i]}); cub::DeviceSegmentedSort::SortPairs(static_cast(nullptr), tmp_storage_bytes, index_first + h_edge_offsets[i], diff --git a/cpp/tests/sampling/renumber_sampled_edgelist_test.cu b/cpp/tests/sampling/renumber_sampled_edgelist_test.cu index 49cb674214c..6d944314605 100644 --- a/cpp/tests/sampling/renumber_sampled_edgelist_test.cu +++ b/cpp/tests/sampling/renumber_sampled_edgelist_test.cu @@ -426,13 +426,13 @@ TEST_P(Tests_RenumberSampledEdgelist, CheckInt64) run_current_test(param); } -INSTANTIATE_TEST_SUITE_P(small_test, - Tests_RenumberSampledEdgelist, - ::testing::Values(RenumberSampledEdgelist_Usecase{1024, 4096, 1, 1, true}, - RenumberSampledEdgelist_Usecase{1024, 4096, 3, 1, true}, - RenumberSampledEdgelist_Usecase{ - 1024, 32768, 1, 256, true}, - RenumberSampledEdgelist_Usecase{1024, 32768, 3, 256, true})); +INSTANTIATE_TEST_SUITE_P( + small_test, + Tests_RenumberSampledEdgelist, + ::testing::Values(RenumberSampledEdgelist_Usecase{1024, 4096, 1, 1, true}, + RenumberSampledEdgelist_Usecase{1024, 4096, 3, 1, true}, + RenumberSampledEdgelist_Usecase{1024, 32768, 1, 256, true}, + RenumberSampledEdgelist_Usecase{1024, 32768, 3, 256, true})); INSTANTIATE_TEST_SUITE_P( benchmark_test, From cf7dff8593dc13577ef75f03f0e66d1443965328 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 20 Jul 2023 16:36:22 -0700 Subject: [PATCH 14/17] clang-format --- cpp/src/sampling/renumber_sampled_edgelist_impl.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh index ae9dd64348a..f35a89349b4 100644 --- a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh +++ b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh @@ -78,7 +78,7 @@ compute_renumber_map(raft::handle_t const& handle, rmm::device_uvector unique_label_src_pair_vertices( 0, handle.get_stream()); // sorted by (label, hop, src) std::optional> sorted_srcs{ - std::nullopt}; // sorted by (label, src), relevant only when edgelist_hops is valid + std::nullopt}; // sorted by (label, src), relevant only when edgelist_hops is valid { if (label_offsets) { rmm::device_uvector label_indices((*edgelist_label_indices).size(), From 948ad11ddea3f4bc85049ef755c4c385e2014080 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 20 Jul 2023 16:40:42 -0700 Subject: [PATCH 15/17] remove unnecessary template parameter --- cpp/src/sampling/renumber_sampled_edgelist_impl.cuh | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh index f35a89349b4..cde67f27896 100644 --- a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh +++ b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh @@ -45,7 +45,7 @@ namespace cugraph { namespace { -template +template std::tuple, std::optional>> compute_renumber_map(raft::handle_t const& handle, raft::device_span edgelist_srcs, @@ -183,8 +183,6 @@ compute_renumber_map(raft::handle_t const& handle, offset_first + 1, handle.get_stream()); } - d_tmp_storage.resize(0, handle.get_stream()); - d_tmp_storage.shrink_to_fit(handle.get_stream()); unique_label_src_pair_vertices = std::move(segment_sorted_srcs); } else { @@ -344,7 +342,6 @@ compute_renumber_map(raft::handle_t const& handle, offset_first + 1, handle.get_stream()); } - dsts.resize(0, handle.get_stream()); d_tmp_storage.resize(0, handle.get_stream()); dsts.shrink_to_fit(handle.get_stream()); @@ -515,7 +512,7 @@ renumber_sampled_edgelist( // 2. compute renumber_map auto [renumber_map, renumber_map_label_indices] = - compute_renumber_map( + compute_renumber_map( handle, raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), edgelist_hops, From b26523b0b42cc83a025eb5aac507139fb827d2c3 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 20 Jul 2023 16:55:47 -0700 Subject: [PATCH 16/17] clang-format --- .../sampling/renumber_sampled_edgelist_impl.cuh | 16 +++++++--------- 1 file changed, 7 insertions(+), 9 deletions(-) diff --git a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh index cde67f27896..dceb61b49d5 100644 --- a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh +++ b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh @@ -511,15 +511,13 @@ renumber_sampled_edgelist( // 2. compute renumber_map - auto [renumber_map, renumber_map_label_indices] = - compute_renumber_map( - handle, - raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), - edgelist_hops, - raft::device_span(edgelist_dsts.data(), edgelist_dsts.size()), - label_offsets - ? std::make_optional>(std::get<1>(*label_offsets)) - : std::nullopt); + auto [renumber_map, renumber_map_label_indices] = compute_renumber_map( + handle, + raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()), + edgelist_hops, + raft::device_span(edgelist_dsts.data(), edgelist_dsts.size()), + label_offsets ? std::make_optional>(std::get<1>(*label_offsets)) + : std::nullopt); // 3. compute renumber map offsets for each label From 1e4174e6dc421de885ccbe1ca2a61932e7b576e0 Mon Sep 17 00:00:00 2001 From: Charles Hastings Date: Fri, 21 Jul 2023 15:51:39 -0700 Subject: [PATCH 17/17] Testing with Seunghwa's branch merged in --- cpp/src/c_api/uniform_neighbor_sampling.cpp | 36 ++-- .../c_api/uniform_neighbor_sample_test.c | 177 +++++++++++++++--- 2 files changed, 168 insertions(+), 45 deletions(-) diff --git a/cpp/src/c_api/uniform_neighbor_sampling.cpp b/cpp/src/c_api/uniform_neighbor_sampling.cpp index fe29b6252bd..ff6a6c49437 100644 --- a/cpp/src/c_api/uniform_neighbor_sampling.cpp +++ b/cpp/src/c_api/uniform_neighbor_sampling.cpp @@ -232,18 +232,18 @@ struct uniform_neighbor_sampling_functor : public cugraph::c_api::abstract_funct std::optional> renumber_map{std::nullopt}; std::optional> renumber_map_offsets{std::nullopt}; -#if 0 - std::tie(src, dst, renumber_map, renumber_map_offsets) = - cugraph::renumber_sampled_edgelist(handle_, - std::move(src), - hop ? std::make_optional(raft::device_span{hop->data(), hop->size()}) : std::nullopt, - std::move(dst), - std::make_optional(std::make_tuple(raft::device_span{edge_label->data(), - edge_label->size()}, - raft::device_span{offsets->data(), - offsets->size()})) - do_expensive_check); -#endif + if (options_.renumber_results_) { + std::tie(src, dst, renumber_map, renumber_map_offsets) = cugraph::renumber_sampled_edgelist( + handle_, + std::move(src), + hop ? std::make_optional(raft::device_span{hop->data(), hop->size()}) + : std::nullopt, + std::move(dst), + std::make_optional(std::make_tuple( + raft::device_span{edge_label->data(), edge_label->size()}, + raft::device_span{offsets->data(), offsets->size()})), + do_expensive_check_); + } result_ = new cugraph::c_api::cugraph_sample_result_t{ new cugraph::c_api::cugraph_type_erased_device_array_t(src, graph_->vertex_type_), @@ -422,16 +422,20 @@ extern "C" cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_re const cugraph_sample_result_t* result) { auto internal_pointer = reinterpret_cast(result); - return reinterpret_cast( - internal_pointer->renumber_map_->view()); + return internal_pointer->renumber_map_ == nullptr + ? NULL + : reinterpret_cast( + internal_pointer->renumber_map_->view()); } extern "C" cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_renumber_map_offsets( const cugraph_sample_result_t* result) { auto internal_pointer = reinterpret_cast(result); - return reinterpret_cast( - internal_pointer->renumber_map_offsets_->view()); + return internal_pointer->renumber_map_ == nullptr + ? NULL + : reinterpret_cast( + internal_pointer->renumber_map_offsets_->view()); } extern "C" cugraph_error_code_t cugraph_test_uniform_neighborhood_sample_result_create( diff --git a/cpp/tests/c_api/uniform_neighbor_sample_test.c b/cpp/tests/c_api/uniform_neighbor_sample_test.c index c93c99ea7fc..a2c1e230485 100644 --- a/cpp/tests/c_api/uniform_neighbor_sample_test.c +++ b/cpp/tests/c_api/uniform_neighbor_sample_test.c @@ -21,6 +21,7 @@ #include #include +#include typedef int32_t vertex_t; typedef int32_t edge_t; @@ -32,6 +33,15 @@ data_type_id_t weight_tid = FLOAT32; data_type_id_t edge_id_tid = INT32; data_type_id_t edge_type_tid = INT32; +int vertex_id_compare_function(const void * a, const void * b) { + if (*((vertex_t *) a) < *((vertex_t *) b)) + return -1; + else if (*((vertex_t *) a) > *((vertex_t *) b)) + return 1; + else + return 0; +} + int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle, vertex_t *h_src, vertex_t *h_dst, @@ -48,7 +58,8 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle bool_t with_replacement, bool_t return_hops, cugraph_prior_sources_behavior_t prior_sources_behavior, - bool_t dedupe_sources) + bool_t dedupe_sources, + bool_t renumber_results) { // Create graph int test_ret_value = 0; @@ -119,6 +130,7 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle cugraph_sampling_set_return_hops(sampling_options, return_hops); cugraph_sampling_set_prior_sources_behavior(sampling_options, prior_sources_behavior); cugraph_sampling_set_dedupe_sources(sampling_options, dedupe_sources); + cugraph_sampling_set_renumber_results(sampling_options, renumber_results); ret_code = cugraph_uniform_neighbor_sample(handle, graph, @@ -150,19 +162,28 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle cugraph_type_erased_device_array_view_t* result_hops; cugraph_type_erased_device_array_view_t* result_offsets; cugraph_type_erased_device_array_view_t* result_labels; - - result_srcs = cugraph_sample_result_get_sources(result); - result_dsts = cugraph_sample_result_get_destinations(result); - result_edge_id = cugraph_sample_result_get_edge_id(result); - result_weights = cugraph_sample_result_get_edge_weight(result); - result_edge_types = cugraph_sample_result_get_edge_type(result); - result_hops = cugraph_sample_result_get_hop(result); - result_hops = cugraph_sample_result_get_hop(result); - result_offsets = cugraph_sample_result_get_offsets(result); - result_labels = cugraph_sample_result_get_start_labels(result); + cugraph_type_erased_device_array_view_t* result_renumber_map; + cugraph_type_erased_device_array_view_t* result_renumber_map_offsets; + + result_srcs = cugraph_sample_result_get_sources(result); + result_dsts = cugraph_sample_result_get_destinations(result); + result_edge_id = cugraph_sample_result_get_edge_id(result); + result_weights = cugraph_sample_result_get_edge_weight(result); + result_edge_types = cugraph_sample_result_get_edge_type(result); + result_hops = cugraph_sample_result_get_hop(result); + result_hops = cugraph_sample_result_get_hop(result); + result_offsets = cugraph_sample_result_get_offsets(result); + result_labels = cugraph_sample_result_get_start_labels(result); + result_renumber_map = cugraph_sample_result_get_renumber_map(result); + result_renumber_map_offsets = cugraph_sample_result_get_renumber_map_offsets(result); size_t result_size = cugraph_type_erased_device_array_view_size(result_srcs); size_t result_offsets_size = cugraph_type_erased_device_array_view_size(result_offsets); + size_t renumber_map_size = 0; + + if (renumber_results) { + renumber_map_size = cugraph_type_erased_device_array_view_size(result_renumber_map); + } vertex_t h_result_srcs[result_size]; vertex_t h_result_dsts[result_size]; @@ -172,6 +193,8 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle int32_t h_result_hops[result_size]; size_t h_result_offsets[result_offsets_size]; int h_result_labels[result_offsets_size-1]; + vertex_t h_renumber_map[renumber_map_size]; + size_t h_renumber_map_offsets[result_offsets_size]; ret_code = cugraph_type_erased_device_array_view_copy_to_host( handle, (byte_t*)h_result_srcs, result_srcs, &ret_error); @@ -205,6 +228,16 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle handle, (byte_t*)h_result_labels, result_labels, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + if (renumber_results) { + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_renumber_map, result_renumber_map, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_renumber_map_offsets, result_renumber_map_offsets, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + } + // First, check that all edges are actually part of the graph weight_t M_w[num_vertices][num_vertices]; edge_t M_edge_id[num_vertices][num_vertices]; @@ -223,16 +256,35 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle M_edge_type[h_src[i]][h_dst[i]] = h_edge_types[i]; } - for (int i = 0; (i < result_size) && (test_ret_value == 0); ++i) { - TEST_ASSERT(test_ret_value, - M_w[h_result_srcs[i]][h_result_dsts[i]] == h_result_weight[i], - "uniform_neighbor_sample got edge that doesn't exist"); - TEST_ASSERT(test_ret_value, - M_edge_id[h_result_srcs[i]][h_result_dsts[i]] == h_result_edge_id[i], - "uniform_neighbor_sample got edge that doesn't exist"); - TEST_ASSERT(test_ret_value, - M_edge_type[h_result_srcs[i]][h_result_dsts[i]] == h_result_edge_types[i], - "uniform_neighbor_sample got edge that doesn't exist"); + if (renumber_results) { + for (int label_id = 0 ; label_id < (result_offsets_size - 1) ; ++label_id) { + for (size_t i = h_result_offsets[label_id]; (i < h_result_offsets[label_id+1]) && (test_ret_value == 0) ; ++i) { + vertex_t src = h_renumber_map[h_renumber_map_offsets[label_id] + h_result_srcs[i]]; + vertex_t dst = h_renumber_map[h_renumber_map_offsets[label_id] + h_result_dsts[i]]; + + TEST_ASSERT(test_ret_value, + M_w[src][dst] == h_result_weight[i], + "uniform_neighbor_sample got edge that doesn't exist"); + TEST_ASSERT(test_ret_value, + M_edge_id[src][dst] == h_result_edge_id[i], + "uniform_neighbor_sample got edge that doesn't exist"); + TEST_ASSERT(test_ret_value, + M_edge_type[src][dst] == h_result_edge_types[i], + "uniform_neighbor_sample got edge that doesn't exist"); + } + } + } else { + for (int i = 0; (i < result_size) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + M_w[h_result_srcs[i]][h_result_dsts[i]] == h_result_weight[i], + "uniform_neighbor_sample got edge that doesn't exist"); + TEST_ASSERT(test_ret_value, + M_edge_id[h_result_srcs[i]][h_result_dsts[i]] == h_result_edge_id[i], + "uniform_neighbor_sample got edge that doesn't exist"); + TEST_ASSERT(test_ret_value, + M_edge_type[h_result_srcs[i]][h_result_dsts[i]] == h_result_edge_types[i], + "uniform_neighbor_sample got edge that doesn't exist"); + } } // @@ -264,6 +316,28 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle } } + if (renumber_results) { + size_t num_vertex_ids = 2 * (h_result_offsets[label_id+1] - h_result_offsets[label_id]); + vertex_t vertex_ids[num_vertex_ids]; + + for (size_t i = 0 ; (i < (h_result_offsets[label_id+1] - h_result_offsets[label_id])) && (test_ret_value == 0) ; ++i) { + vertex_ids[2*i] = h_result_srcs[h_result_offsets[label_id] + i]; + vertex_ids[2*i+1] = h_result_dsts[h_result_offsets[label_id] + i]; + } + + qsort(vertex_ids, num_vertex_ids, sizeof(vertex_t), vertex_id_compare_function); + + vertex_t current_v = 0; + for (size_t i = 0 ; (i < num_vertex_ids) && (test_ret_value == 0) ; ++i) { + if (vertex_ids[i] == current_v) + ++current_v; + else + TEST_ASSERT(test_ret_value, + vertex_ids[i] == (current_v - 1), + "vertices are not properly renumbered"); + } + } + for (int hop = 0 ; hop < fan_out_size ; ++hop) { if (prior_sources_behavior == CARRY_OVER) { destinations_size = sources_size; @@ -276,7 +350,8 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle if (h_result_hops[i] == hop) { bool found = false; for (size_t j = 0 ; (!found) && (j < sources_size) ; ++j) { - found = (h_result_srcs[i] == check_sources[j]); + found = renumber_results ? (h_renumber_map[h_renumber_map_offsets[label_id] + h_result_srcs[i]] == check_sources[j]) + : (h_result_srcs[i] == check_sources[j]); } TEST_ASSERT(test_ret_value, found, "encountered source vertex that was not part of previous frontier"); @@ -286,15 +361,16 @@ int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle // Make sure destination isn't already in the source list bool found = false; for (size_t j = 0 ; (!found) && (j < destinations_size) ; ++j) { - found = (h_result_dsts[i] == check_destinations[j]); + found = renumber_results ? (h_renumber_map[h_renumber_map_offsets[label_id] + h_result_dsts[i]] == check_destinations[j]) + : (h_result_dsts[i] == check_destinations[j]); } if (!found) { - check_destinations[destinations_size] = h_result_dsts[i]; + check_destinations[destinations_size] = renumber_results ? h_renumber_map[h_renumber_map_offsets[label_id] + h_result_dsts[i]] : h_result_dsts[i]; ++destinations_size; } } else { - check_destinations[destinations_size] = h_result_dsts[i]; + check_destinations[destinations_size] = renumber_results ? h_renumber_map[h_renumber_map_offsets[label_id] + h_result_dsts[i]] : h_result_dsts[i]; ++destinations_size; } } @@ -844,11 +920,12 @@ int test_uniform_neighbor_sample_clean(const cugraph_resource_handle_t* handle) bool_t return_hops = TRUE; cugraph_prior_sources_behavior_t prior_sources_behavior = DEFAULT; bool_t dedupe_sources = FALSE; + bool_t renumber_results = FALSE; return generic_uniform_neighbor_sample_test(handle, src, dst, weight, edge_ids, edge_types, num_vertices, num_edges, start, start_labels, num_starts, fan_out, fan_out_size, with_replacement, - return_hops, prior_sources_behavior, dedupe_sources); + return_hops, prior_sources_behavior, dedupe_sources, renumber_results); } int test_uniform_neighbor_sample_dedupe_sources(const cugraph_resource_handle_t* handle) @@ -881,11 +958,12 @@ int test_uniform_neighbor_sample_dedupe_sources(const cugraph_resource_handle_t* bool_t return_hops = TRUE; cugraph_prior_sources_behavior_t prior_sources_behavior = DEFAULT; bool_t dedupe_sources = TRUE; + bool_t renumber_results = FALSE; return generic_uniform_neighbor_sample_test(handle, src, dst, weight, edge_ids, edge_types, num_vertices, num_edges, start, start_labels, num_starts, fan_out, fan_out_size, with_replacement, - return_hops, prior_sources_behavior, dedupe_sources); + return_hops, prior_sources_behavior, dedupe_sources, renumber_results); } int test_uniform_neighbor_sample_unique_sources(const cugraph_resource_handle_t* handle) @@ -918,11 +996,12 @@ int test_uniform_neighbor_sample_unique_sources(const cugraph_resource_handle_t* bool_t return_hops = TRUE; cugraph_prior_sources_behavior_t prior_sources_behavior = EXCLUDE; bool_t dedupe_sources = FALSE; + bool_t renumber_results = FALSE; return generic_uniform_neighbor_sample_test(handle, src, dst, weight, edge_ids, edge_types, num_vertices, num_edges, start, start_labels, num_starts, fan_out, fan_out_size, with_replacement, - return_hops, prior_sources_behavior, dedupe_sources); + return_hops, prior_sources_behavior, dedupe_sources, renumber_results); } int test_uniform_neighbor_sample_carry_over_sources(const cugraph_resource_handle_t* handle) @@ -955,11 +1034,50 @@ int test_uniform_neighbor_sample_carry_over_sources(const cugraph_resource_handl bool_t return_hops = TRUE; cugraph_prior_sources_behavior_t prior_sources_behavior = CARRY_OVER; bool_t dedupe_sources = FALSE; + bool_t renumber_results = FALSE; + + return generic_uniform_neighbor_sample_test(handle, src, dst, weight, edge_ids, edge_types, num_vertices, num_edges, + start, start_labels, num_starts, + fan_out, fan_out_size, with_replacement, + return_hops, prior_sources_behavior, dedupe_sources, renumber_results); +} + +int test_uniform_neighbor_sample_renumber_results(const cugraph_resource_handle_t* handle) +{ + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + data_type_id_t edge_id_tid = INT32; + data_type_id_t edge_type_tid = INT32; + + size_t num_edges = 9; + size_t num_vertices = 6; + size_t fan_out_size = 3; + size_t num_starts = 2; + + vertex_t src[] = {0, 0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 2, 3, 4, 0, 1, 3, 5, 5}; + edge_t edge_ids[] = {0, 1, 2, 3, 4, 5, 6, 7, 8}; + weight_t weight[] = {0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, 0.9}; + int32_t edge_types[] = {8, 7, 6, 5, 4, 3, 2, 1, 0}; + vertex_t start[] = {2, 3}; + int start_labels[] = { 6, 12 }; + int fan_out[] = {-1, -1, -1}; + + int test_ret_value = 0; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + bool_t with_replacement = FALSE; + bool_t return_hops = TRUE; + cugraph_prior_sources_behavior_t prior_sources_behavior = DEFAULT; + bool_t dedupe_sources = FALSE; + bool_t renumber_results = TRUE; return generic_uniform_neighbor_sample_test(handle, src, dst, weight, edge_ids, edge_types, num_vertices, num_edges, start, start_labels, num_starts, fan_out, fan_out_size, with_replacement, - return_hops, prior_sources_behavior, dedupe_sources); + return_hops, prior_sources_behavior, dedupe_sources, renumber_results); } int main(int argc, char** argv) @@ -975,6 +1093,7 @@ int main(int argc, char** argv) result |= RUN_TEST_NEW(test_uniform_neighbor_sample_dedupe_sources, handle); result |= RUN_TEST_NEW(test_uniform_neighbor_sample_unique_sources, handle); result |= RUN_TEST_NEW(test_uniform_neighbor_sample_carry_over_sources, handle); + result |= RUN_TEST_NEW(test_uniform_neighbor_sample_renumber_results, handle); cugraph_free_resource_handle(handle);