diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 8d57c089d8f..800f0b46259 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -286,6 +286,12 @@ set(CUGRAPH_SOURCES src/community/k_truss_sg_v64_e64.cu src/community/k_truss_sg_v32_e32.cu src/community/k_truss_sg_v32_e64.cu + src/lookup/lookup_src_dst_mg_v32_e32.cu + src/lookup/lookup_src_dst_mg_v32_e64.cu + src/lookup/lookup_src_dst_mg_v64_e64.cu + src/lookup/lookup_src_dst_sg_v32_e32.cu + src/lookup/lookup_src_dst_sg_v32_e64.cu + src/lookup/lookup_src_dst_sg_v64_e64.cu src/sampling/random_walks_old_sg_v32_e32.cu src/sampling/random_walks_old_sg_v32_e64.cu src/sampling/random_walks_old_sg_v64_e64.cu diff --git a/cpp/include/cugraph/sampling_functions.hpp b/cpp/include/cugraph/sampling_functions.hpp index 6056fe72057..971a0197d6f 100644 --- a/cpp/include/cugraph/sampling_functions.hpp +++ b/cpp/include/cugraph/sampling_functions.hpp @@ -15,6 +15,8 @@ */ #pragma once +#include + #include #include @@ -449,5 +451,83 @@ sort_sampled_edgelist(raft::handle_t const& handle, size_t num_hops, bool src_is_major = true, bool do_expensive_check = false); +/* + * @brief Build map to lookup source and destination using edge id and type + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam edge_t Type of edge identifiers. Needs to be an integral type. + * @tparam edge_type_t Type of edge types. Needs to be an integral type. + * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param graph_view Graph view object. + * @param edge_id_view View object holding edge ids of the edges of the graph pointed @p graph_view + * @param edge_type_view View object holding edge types of the edges of the graph pointed @p + * graph_view + * @return An object of type cugraph::lookup_container_t that encapsulates edge id and type to + * source and destination lookup map. + */ +template +lookup_container_t build_edge_id_and_type_to_src_dst_lookup_map( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_id_view, + edge_property_view_t edge_type_view); + +/* + * @brief Lookup edge sources and destinations using edge ids and a single edge type. + * Use this function to lookup endpoints of edges belonging to the same edge type. + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam edge_t Type of edge identifiers. Needs to be an integral type. + * @tparam edge_type_t Type of edge types. Needs to be an integral type. + * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param lookup_container Object of type cugraph::lookup_container_t that encapsulates edge id and + * type to source and destination lookup map. + * @param edge_ids_to_lookup Device span of edge ids to lookup + * @param edge_type_to_lookup Type of the edges corresponding to edge ids in @p edge_ids_to_lookup + * @return A tuple of device vector containing edge sources and destinations for edge ids + * in @p edge_ids_to_lookup and edge type @. If an edge id in @p edge_ids_to_lookup or + * edge type @edge_type_to_lookup is not found, the corresponding entry in the device vectors of + * the returned tuple will contain cugraph::invalid_vertex_id. + */ +template +std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_single_type( + raft::handle_t const& handle, + lookup_container_t const& lookup_container, + raft::device_span edge_ids_to_lookup, + edge_type_t edge_type_to_lookup); + +/* + * @brief Lookup edge sources and destinations using edge ids and edge types. + * Use this function to lookup endpoints of edges belonging to different edge types. + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam edge_t Type of edge identifiers. Needs to be an integral type. + * @tparam edge_type_t Type of edge types. Needs to be an integral type. + * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param lookup_container Object of type cugraph::lookup_container_t that encapsulates edge id and + * type to source and destination lookup map. + * @param edge_ids_to_lookup Device span of edge ids to lookup + * @param edge_types_to_lookup Device span of edge types corresponding to the edge ids + * in @p edge_ids_to_lookup + * @return A tuple of device vector containing edge sources and destinations for the edge ids + * in @p edge_ids_to_lookup and the edge types in @p edge_types_to_lookup. If an edge id in + * @p edge_ids_to_lookup or edge type in @p edge_types_to_lookup is not found, the + * corresponding entry in the device vectors of the returned tuple will contain + * cugraph::invalid_vertex_id. + */ +template +std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_types( + raft::handle_t const& handle, + lookup_container_t const& lookup_container, + raft::device_span edge_ids_to_lookup, + raft::device_span edge_types_to_lookup); } // namespace cugraph diff --git a/cpp/include/cugraph/src_dst_lookup_container.hpp b/cpp/include/cugraph/src_dst_lookup_container.hpp new file mode 100644 index 00000000000..4b1509f0367 --- /dev/null +++ b/cpp/include/cugraph/src_dst_lookup_container.hpp @@ -0,0 +1,77 @@ +/* + * Copyright (c) 2024, 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 + +namespace cugraph { + +template > +class lookup_container_t { + template + struct lookup_container_impl; + std::unique_ptr> pimpl; + + public: + using edge_id_type = edge_id_t; + using edge_type_type = edge_type_t; + using value_type = value_t; + + static_assert(std::is_integral_v); + static_assert(std::is_integral_v); + static_assert(is_thrust_tuple_of_integral::value); + + ~lookup_container_t(); + lookup_container_t(); + lookup_container_t(raft::handle_t const& handle, + std::vector types, + std::vector type_counts); + lookup_container_t(const lookup_container_t&); + + void insert(raft::handle_t const& handle, + edge_type_t typ, + raft::device_span edge_ids_to_insert, + dataframe_buffer_type_t&& values_to_insert); + + dataframe_buffer_type_t lookup_from_edge_ids_and_single_type( + raft::handle_t const& handle, + raft::device_span edge_ids_to_lookup, + edge_type_t edge_type_to_lookup, + bool multi_gpu) const; + + dataframe_buffer_type_t lookup_from_edge_ids_and_types( + raft::handle_t const& handle, + raft::device_span edge_ids_to_lookup, + raft::device_span edge_types_to_lookup, + bool multi_gpu) const; +}; + +} // namespace cugraph diff --git a/cpp/include/cugraph/utilities/dataframe_buffer.hpp b/cpp/include/cugraph/utilities/dataframe_buffer.hpp index d52160abd19..450e816bd96 100644 --- a/cpp/include/cugraph/utilities/dataframe_buffer.hpp +++ b/cpp/include/cugraph/utilities/dataframe_buffer.hpp @@ -207,6 +207,14 @@ auto get_dataframe_buffer_end(BufferType& buffer) std::make_index_sequence::value>(), buffer); } +template +struct dataframe_buffer_type { + using type = decltype(allocate_dataframe_buffer(size_t{0}, rmm::cuda_stream_view{})); +}; + +template +using dataframe_buffer_type_t = typename dataframe_buffer_type::type; + template , rmm::device_uvector>::value>* = nullptr> diff --git a/cpp/include/cugraph/utilities/thrust_tuple_utils.hpp b/cpp/include/cugraph/utilities/thrust_tuple_utils.hpp index 304a5b94bd6..2c36ed33359 100644 --- a/cpp/include/cugraph/utilities/thrust_tuple_utils.hpp +++ b/cpp/include/cugraph/utilities/thrust_tuple_utils.hpp @@ -113,6 +113,19 @@ struct is_thrust_tuple_of_arithmetic> { static constexpr bool value = (... && is_valid); }; +template +struct is_thrust_tuple_of_integral : std::false_type {}; + +template +struct is_thrust_tuple_of_integral> { + private: + template + static constexpr bool is_valid = std::is_integral_v || std::is_same_v; + + public: + static constexpr bool value = (... && is_valid); +}; + template struct is_std_tuple : std::false_type {}; diff --git a/cpp/src/detail/graph_partition_utils.cuh b/cpp/src/detail/graph_partition_utils.cuh index 957436459cd..00931780266 100644 --- a/cpp/src/detail/graph_partition_utils.cuh +++ b/cpp/src/detail/graph_partition_utils.cuh @@ -50,6 +50,21 @@ struct compute_gpu_id_from_ext_vertex_t { } }; +template +struct compute_gpu_id_from_ext_edge_id_t { + int comm_size{0}; + int major_comm_size{0}; + int minor_comm_size{0}; + + __host__ __device__ int operator()(edge_t e) const + { + cuco::detail::MurmurHash3_32 hash_func{}; + auto vertex_partition_id = static_cast(hash_func(e) % comm_size); + return partition_manager::compute_global_comm_rank_from_vertex_partition_id( + major_comm_size, minor_comm_size, vertex_partition_id); + } +}; + template struct compute_gpu_id_from_int_vertex_t { raft::device_span vertex_partition_range_lasts{}; diff --git a/cpp/src/lookup/lookup_src_dst_impl.cuh b/cpp/src/lookup/lookup_src_dst_impl.cuh new file mode 100644 index 00000000000..4182cb5f65e --- /dev/null +++ b/cpp/src/lookup/lookup_src_dst_impl.cuh @@ -0,0 +1,841 @@ +/* + * Copyright (c) 2024, 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 "detail/graph_partition_utils.cuh" +#include "prims/extract_transform_e.cuh" +#include "prims/kv_store.cuh" +#include "utilities/collect_comm.cuh" + +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +namespace cugraph { + +template +template +struct lookup_container_t::lookup_container_impl { + static_assert(std::is_integral_v); + static_assert(std::is_integral_v); + + static_assert(std::is_same_v); + static_assert(std::is_same_v); + static_assert(std::is_same_v); + + ~lookup_container_impl() {} + lookup_container_impl() {} + lookup_container_impl(raft::handle_t const& handle, + std::vector types, + std::vector type_counts) + { + auto invalid_vertex_id = cugraph::invalid_vertex_id::value; + auto invalid_value = thrust::tuple(invalid_vertex_id, invalid_vertex_id); + + edge_type_to_kv_store = container_t{}; + edge_type_to_kv_store.reserve(types.size()); + + for (size_t idx = 0; idx < types.size(); idx++) { + auto typ = types[idx]; + assert(typ != empty_type_); + size_t store_capacity = type_counts[idx]; + + edge_type_to_kv_store.insert( + {typ, store_t(store_capacity, invalid_vertex_id, invalid_value, handle.get_stream())}); + } + + edge_type_to_kv_store.insert( + {empty_type_, store_t(0, invalid_vertex_id, invalid_value, handle.get_stream())}); + } + + void insert(raft::handle_t const& handle, + edge_type_t type, + raft::device_span edge_ids_to_insert, + dataframe_buffer_type_t&& values_to_insert) + { + auto itr = edge_type_to_kv_store.find(type); + + if (itr != edge_type_to_kv_store.end()) { + assert(itr->first == type); + + itr->second.insert(edge_ids_to_insert.begin(), + edge_ids_to_insert.end(), + cugraph::get_dataframe_buffer_begin(values_to_insert), + handle.get_stream()); + + } else { + assert(false); + } + } + + dataframe_buffer_type_t src_dst_from_edge_id_and_type( + raft::handle_t const& handle, + raft::device_span edge_ids_to_lookup, + edge_type_t edge_type_to_lookup, + bool multi_gpu) const + { + using store_t = typename container_t::mapped_type; + const store_t* kv_store_object{nullptr}; + + auto value_buffer = cugraph::allocate_dataframe_buffer(0, handle.get_stream()); + auto itr = edge_type_to_kv_store.find(edge_type_to_lookup); + + if (itr != edge_type_to_kv_store.end()) { + assert(edge_type_to_lookup == itr->first); + kv_store_object = &(itr->second); + + } else { + kv_store_object = &(edge_type_to_kv_store.find(empty_type_)->second); + } + + if (multi_gpu) { + auto& comm = handle.get_comms(); + auto const comm_size = comm.get_size(); + auto& major_comm = handle.get_subcomm(cugraph::partition_manager::major_comm_name()); + auto const major_comm_size = major_comm.get_size(); + auto& minor_comm = handle.get_subcomm(cugraph::partition_manager::minor_comm_name()); + auto const minor_comm_size = minor_comm.get_size(); + + value_buffer = cugraph::collect_values_for_keys( + handle, + kv_store_object->view(), + edge_ids_to_lookup.begin(), + edge_ids_to_lookup.end(), + cugraph::detail::compute_gpu_id_from_ext_edge_id_t{ + comm_size, major_comm_size, minor_comm_size}); + } else { + cugraph::resize_dataframe_buffer( + value_buffer, edge_ids_to_lookup.size(), handle.get_stream()); + + kv_store_object->view().find(edge_ids_to_lookup.begin(), + edge_ids_to_lookup.end(), + cugraph::get_dataframe_buffer_begin(value_buffer), + handle.get_stream()); + } + + return std::make_tuple(std::move(std::get<0>(value_buffer)), + std::move(std::get<1>(value_buffer))); + } + + dataframe_buffer_type_t src_dst_from_edge_id_and_type( + raft::handle_t const& handle, + raft::device_span edge_ids_to_lookup, + raft::device_span edge_types_to_lookup, + bool multi_gpu) const + { + static_assert(std::is_integral_v); + static_assert(std::is_integral_v); + + assert(edge_ids_to_lookup.size() == edge_types_to_lookup.size()); + + rmm::device_uvector tmp_edge_ids_to_lookup(edge_ids_to_lookup.size(), + handle.get_stream()); + + rmm::device_uvector tmp_edge_types_to_lookup(edge_types_to_lookup.size(), + handle.get_stream()); + + rmm::device_uvector original_idxs(edge_ids_to_lookup.size(), handle.get_stream()); + + thrust::sequence( + handle.get_thrust_policy(), original_idxs.begin(), original_idxs.end(), edge_id_t{0}); + + thrust::copy(handle.get_thrust_policy(), + edge_ids_to_lookup.begin(), + edge_ids_to_lookup.end(), + tmp_edge_ids_to_lookup.begin()); + + thrust::copy(handle.get_thrust_policy(), + edge_types_to_lookup.begin(), + edge_types_to_lookup.end(), + tmp_edge_types_to_lookup.begin()); + + thrust::sort_by_key(handle.get_thrust_policy(), + tmp_edge_types_to_lookup.begin(), + tmp_edge_types_to_lookup.end(), + thrust::make_zip_iterator(thrust::make_tuple(tmp_edge_ids_to_lookup.begin(), + original_idxs.begin()))); + + auto nr_uniqe_edge_types_to_lookup = thrust::count_if( + handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(tmp_edge_types_to_lookup.size()), + detail::is_first_in_run_t{tmp_edge_types_to_lookup.data()}); + + rmm::device_uvector unique_types(nr_uniqe_edge_types_to_lookup, + handle.get_stream()); + rmm::device_uvector type_offsets(nr_uniqe_edge_types_to_lookup + size_t{1}, + handle.get_stream()); + + thrust::copy_if(handle.get_thrust_policy(), + tmp_edge_types_to_lookup.begin(), + tmp_edge_types_to_lookup.end(), + thrust::make_counting_iterator(size_t{0}), + unique_types.begin(), + detail::is_first_in_run_t{tmp_edge_types_to_lookup.data()}); + + type_offsets.set_element_to_zero_async(0, handle.get_stream()); + thrust::upper_bound(handle.get_thrust_policy(), + tmp_edge_types_to_lookup.begin(), + tmp_edge_types_to_lookup.end(), + unique_types.begin(), + unique_types.end(), + type_offsets.begin() + 1); + + std::vector h_unique_types(unique_types.size()); + std::vector h_type_offsets(type_offsets.size()); + + raft::update_host( + h_unique_types.data(), unique_types.data(), unique_types.size(), handle.get_stream()); + + raft::update_host( + h_type_offsets.data(), type_offsets.data(), type_offsets.size(), handle.get_stream()); + + handle.sync_stream(); + + std::unordered_map typ_to_local_idx_map{}; + for (size_t idx = 0; idx < h_unique_types.size(); idx++) { + typ_to_local_idx_map[h_unique_types[idx]] = idx; + } + + auto output_value_buffer = + cugraph::allocate_dataframe_buffer(edge_ids_to_lookup.size(), handle.get_stream()); + if (multi_gpu) { + auto& comm = handle.get_comms(); + auto rx_counts = host_scalar_allgather(comm, unique_types.size(), handle.get_stream()); + std::vector rx_displacements(rx_counts.size()); + std::exclusive_scan(rx_counts.begin(), rx_counts.end(), rx_displacements.begin(), size_t{0}); + rmm::device_uvector rx_unique_types(rx_displacements.back() + rx_counts.back(), + handle.get_stream()); + + device_allgatherv(comm, + unique_types.begin(), + rx_unique_types.begin(), + rx_counts, + rx_displacements, + handle.get_stream()); + unique_types = std::move(rx_unique_types); + + thrust::sort(handle.get_thrust_policy(), unique_types.begin(), unique_types.end()); + + unique_types.resize( + thrust::distance( + unique_types.begin(), + thrust::unique(handle.get_thrust_policy(), unique_types.begin(), unique_types.end())), + handle.get_stream()); + } + + h_unique_types.resize(unique_types.size()); + raft::update_host( + h_unique_types.data(), unique_types.data(), unique_types.size(), handle.get_stream()); + + handle.sync_stream(); + + for (size_t idx = 0; idx < h_unique_types.size(); idx++) { + auto typ = h_unique_types[idx]; + + auto tmp_edge_ids_begin = tmp_edge_ids_to_lookup.begin(); + size_t span_size = 0; + + if (typ_to_local_idx_map.find(typ) != typ_to_local_idx_map.end()) { + auto local_idx = typ_to_local_idx_map[typ]; + tmp_edge_ids_begin = tmp_edge_ids_to_lookup.begin() + h_type_offsets[local_idx]; + span_size = h_type_offsets[local_idx + 1] - h_type_offsets[local_idx]; + } + + auto value_buffer_typ = src_dst_from_edge_id_and_type( + handle, raft::device_span{tmp_edge_ids_begin, span_size}, typ, multi_gpu); + + thrust::copy(handle.get_thrust_policy(), + cugraph::get_dataframe_buffer_begin(value_buffer_typ), + cugraph::get_dataframe_buffer_end(value_buffer_typ), + cugraph::get_dataframe_buffer_begin(output_value_buffer) + + h_type_offsets[typ_to_local_idx_map[typ]]); + } + + thrust::sort_by_key(handle.get_thrust_policy(), + original_idxs.begin(), + original_idxs.end(), + cugraph::get_dataframe_buffer_begin(output_value_buffer)); + + return std::make_tuple(std::move(std::get<0>(output_value_buffer)), + std::move(std::get<1>(output_value_buffer))); + } + + private: + using container_t = + std::unordered_map>; + using store_t = typename container_t::mapped_type; + container_t edge_type_to_kv_store; + edge_type_t empty_type_ = std::numeric_limits::max() - 1; +}; + +template +lookup_container_t::~lookup_container_t() +{ + pimpl.reset(); +} + +template +lookup_container_t::lookup_container_t() + : pimpl{std::make_unique>()} +{ +} + +template +lookup_container_t::lookup_container_t( + raft::handle_t const& handle, std::vector types, std::vector type_counts) + : pimpl{std::make_unique>( + handle, types, type_counts)} +{ +} + +template +lookup_container_t::lookup_container_t( + const lookup_container_t&) +{ +} + +template +void lookup_container_t::insert( + raft::handle_t const& handle, + edge_type_t type, + raft::device_span edge_ids_to_insert, + dataframe_buffer_type_t&& values_to_insert) +{ + pimpl->insert(handle, type, edge_ids_to_insert, std::move(values_to_insert)); +} + +template +dataframe_buffer_type_t +lookup_container_t::lookup_from_edge_ids_and_single_type( + raft::handle_t const& handle, + raft::device_span edge_ids_to_lookup, + edge_type_t edge_type_to_lookup, + bool multi_gpu) const +{ + return pimpl->src_dst_from_edge_id_and_type( + handle, edge_ids_to_lookup, edge_type_to_lookup, multi_gpu); +} + +template +dataframe_buffer_type_t +lookup_container_t::lookup_from_edge_ids_and_types( + raft::handle_t const& handle, + raft::device_span edge_ids_to_lookup, + raft::device_span edge_types_to_lookup, + bool multi_gpu) const +{ + return pimpl->src_dst_from_edge_id_and_type( + handle, edge_ids_to_lookup, edge_types_to_lookup, multi_gpu); +} + +namespace detail { + +template +EdgeTypeAndIdToSrcDstLookupContainerType build_edge_id_and_type_to_src_dst_lookup_map( + raft::handle_t const& handle, + GraphViewType const& graph_view, + EdgeIdInputWrapper edge_id_view, + EdgeTypeInputWrapper edge_type_view) +{ + static_assert(!std::is_same_v, + "Can not create edge id lookup table without edge ids"); + + using vertex_t = typename GraphViewType::vertex_type; + using edge_t = typename GraphViewType::edge_type; + using edge_type_t = typename EdgeTypeInputWrapper::value_type; + using edge_id_t = typename EdgeIdInputWrapper::value_type; + using value_t = typename EdgeTypeAndIdToSrcDstLookupContainerType::value_type; + + constexpr bool multi_gpu = GraphViewType::is_multi_gpu; + static_assert(std::is_integral_v); + static_assert(std::is_integral_v); + static_assert(std::is_same_v); + static_assert(std::is_same_v>); + + static_assert( + std::is_same_v, + "edge_type_t must match with EdgeTypeAndIdToSrcDstLookupContainerType::edge_type_type"); + + static_assert( + std::is_same_v, + "edge_id_t must match with typename EdgeTypeAndIdToSrcDstLookupContainerType::edge_id_type"); + + rmm::device_uvector unique_types(0, handle.get_stream()); + rmm::device_uvector unique_type_counts(0, handle.get_stream()); + + if constexpr (multi_gpu) { + auto& comm = handle.get_comms(); + auto const comm_size = comm.get_size(); + auto& major_comm = handle.get_subcomm(cugraph::partition_manager::major_comm_name()); + auto const major_comm_size = major_comm.get_size(); + auto& minor_comm = handle.get_subcomm(cugraph::partition_manager::minor_comm_name()); + auto const minor_comm_size = minor_comm.get_size(); + + auto [gpu_ids, edge_types] = + cugraph::extract_transform_e( + handle, + graph_view, + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + view_concat(edge_id_view, edge_type_view), + cuda::proclaim_return_type>>( + [key_func = + cugraph::detail::compute_gpu_id_from_ext_edge_id_t{ + comm_size, + major_comm_size, + minor_comm_size}] __device__(auto, + auto, + thrust::nullopt_t, + thrust::nullopt_t, + thrust::tuple id_and_type) { + return thrust::optional>{thrust::make_tuple( + key_func(thrust::get<0>(id_and_type)), thrust::get<1>(id_and_type))}; + })); + + auto type_and_gpu_id_pair_begin = + thrust::make_zip_iterator(thrust::make_tuple(edge_types.begin(), gpu_ids.begin())); + + thrust::sort(handle.get_thrust_policy(), + type_and_gpu_id_pair_begin, + type_and_gpu_id_pair_begin + edge_types.size()); + + auto nr_unique_paris = thrust::count_if( + handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(edge_types.size()), + detail::is_first_in_run_t{type_and_gpu_id_pair_begin}); + + auto unique_pairs_buffer = cugraph::allocate_dataframe_buffer< + typename thrust::iterator_traits::value_type>( + nr_unique_paris, handle.get_stream()); + + rmm::device_uvector unique_pair_counts(nr_unique_paris, handle.get_stream()); + + { + rmm::device_uvector unique_pair_end_offsets(nr_unique_paris, handle.get_stream()); + + thrust::copy_if(handle.get_thrust_policy(), + type_and_gpu_id_pair_begin, + type_and_gpu_id_pair_begin + edge_types.size(), + thrust::make_counting_iterator(size_t{0}), + cugraph::get_dataframe_buffer_begin(unique_pairs_buffer), + detail::is_first_in_run_t{ + type_and_gpu_id_pair_begin}); + + thrust::upper_bound(handle.get_thrust_policy(), + type_and_gpu_id_pair_begin, + type_and_gpu_id_pair_begin + edge_types.size(), + cugraph::get_dataframe_buffer_begin(unique_pairs_buffer), + cugraph::get_dataframe_buffer_end(unique_pairs_buffer), + unique_pair_end_offsets.begin()); + + thrust::adjacent_difference(handle.get_thrust_policy(), + unique_pair_end_offsets.begin(), + unique_pair_end_offsets.end(), + unique_pair_counts.begin()); + } + + edge_types.resize(0, handle.get_stream()); + gpu_ids.resize(0, handle.get_stream()); + edge_types.shrink_to_fit(handle.get_stream()); + gpu_ids.shrink_to_fit(handle.get_stream()); + + std::forward_as_tuple( + std::tie(std::get<0>(unique_pairs_buffer), std::ignore, unique_pair_counts), std::ignore) = + cugraph::groupby_gpu_id_and_shuffle_values( + handle.get_comms(), + thrust::make_zip_iterator(thrust::make_tuple(std::get<0>(unique_pairs_buffer).begin(), + std::get<1>(unique_pairs_buffer).begin(), + unique_pair_counts.begin())), + thrust::make_zip_iterator(thrust::make_tuple(std::get<0>(unique_pairs_buffer).end(), + std::get<1>(unique_pairs_buffer).end(), + unique_pair_counts.end())), + [] __device__(auto val) { return thrust::get<1>(val); }, + handle.get_stream()); + + // + // Count local #elments for all the types mapped to this GPU + // + + thrust::sort_by_key(handle.get_thrust_policy(), + std::get<0>(unique_pairs_buffer).begin(), + std::get<0>(unique_pairs_buffer).end(), + unique_pair_counts.begin()); + + auto nr_unique_types = thrust::count_if( + handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(std::get<0>(unique_pairs_buffer).size()), + detail::is_first_in_run_t{std::get<0>(unique_pairs_buffer).data()}); + + unique_types.resize(static_cast(nr_unique_types), handle.get_stream()); + unique_type_counts.resize(static_cast(nr_unique_types), handle.get_stream()); + + thrust::reduce_by_key(handle.get_thrust_policy(), + std::get<0>(unique_pairs_buffer).begin(), + std::get<0>(unique_pairs_buffer).end(), + unique_pair_counts.begin(), + unique_types.begin(), + unique_type_counts.begin()); + + } else { + auto edge_types = cugraph::extract_transform_e( + handle, + graph_view, + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + edge_type_view, + cuda::proclaim_return_type>( + [] __device__(auto, auto, thrust::nullopt_t, thrust::nullopt_t, edge_type_t et) { + return thrust::optional{et}; + })); + + thrust::sort(handle.get_thrust_policy(), edge_types.begin(), edge_types.end()); + + auto nr_unique_types = + thrust::count_if(handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(edge_types.size()), + detail::is_first_in_run_t{edge_types.data()}); + + unique_types.resize(static_cast(nr_unique_types), handle.get_stream()); + unique_type_counts.resize(static_cast(nr_unique_types), handle.get_stream()); + + { + rmm::device_uvector unique_type_end_offsets(nr_unique_types, handle.get_stream()); + thrust::copy_if(handle.get_thrust_policy(), + edge_types.begin(), + edge_types.end(), + thrust::make_counting_iterator(size_t{0}), + unique_types.begin(), + detail::is_first_in_run_t{edge_types.data()}); + + thrust::upper_bound(handle.get_thrust_policy(), + edge_types.begin(), + edge_types.end(), + unique_types.begin(), + unique_types.end(), + unique_type_end_offsets.begin()); + + thrust::adjacent_difference(handle.get_thrust_policy(), + unique_type_end_offsets.begin(), + unique_type_end_offsets.end(), + unique_type_counts.begin()); + } + } + + std::vector h_unique_types(unique_types.size()); + std::vector h_unique_type_counts(unique_types.size()); + + raft::update_host( + h_unique_types.data(), unique_types.data(), unique_types.size(), handle.get_stream()); + + raft::update_host(h_unique_type_counts.data(), + unique_type_counts.data(), + unique_type_counts.size(), + handle.get_stream()); + + handle.sync_stream(); + + auto search_container = + EdgeTypeAndIdToSrcDstLookupContainerType(handle, h_unique_types, h_unique_type_counts); + + // + // Populate the search container + // + + for (size_t local_ep_idx = 0; local_ep_idx < graph_view.number_of_local_edge_partitions(); + ++local_ep_idx) { + // + // decompress one edge_partition at a time + // + + auto edge_partition = edge_partition_device_view_t( + graph_view.local_edge_partition_view(local_ep_idx)); + + auto edge_partition_mask_view = + graph_view.has_edge_mask() + ? std::make_optional< + detail::edge_partition_edge_property_device_view_t>( + *(graph_view.edge_mask_view()), local_ep_idx) + : std::nullopt; + + auto number_of_local_edges = edge_partition.number_of_edges(); + if (graph_view.has_edge_mask()) { + number_of_local_edges = edge_partition.compute_number_of_edges_with_mask( + (*edge_partition_mask_view).value_first(), + thrust::make_counting_iterator(edge_partition.major_range_first()), + thrust::make_counting_iterator(edge_partition.major_range_last()), + handle.get_stream()); + } + + rmm::device_uvector edgelist_majors(number_of_local_edges, handle.get_stream()); + rmm::device_uvector edgelist_minors(edgelist_majors.size(), handle.get_stream()); + auto edgelist_ids = rmm::device_uvector(edgelist_majors.size(), handle.get_stream()); + auto edgelist_types = + rmm::device_uvector(edgelist_majors.size(), handle.get_stream()); + + detail::decompress_edge_partition_to_edgelist( + handle, + edge_partition, + std::nullopt, + std::make_optional>( + edge_id_view, local_ep_idx), + std::make_optional< + detail::edge_partition_edge_property_device_view_t>( + edge_type_view, local_ep_idx), + edge_partition_mask_view, + raft::device_span(edgelist_majors.data(), number_of_local_edges), + raft::device_span(edgelist_minors.data(), number_of_local_edges), + std::nullopt, + std::make_optional>(edgelist_ids.data(), number_of_local_edges), + std::make_optional>(edgelist_types.data(), + number_of_local_edges), + graph_view.local_edge_partition_segment_offsets(local_ep_idx)); + + // + // Shuffle to the right GPUs using edge ids as keys + // + + if constexpr (multi_gpu) { + auto const comm_size = handle.get_comms().get_size(); + auto& major_comm = handle.get_subcomm(cugraph::partition_manager::major_comm_name()); + auto const major_comm_size = major_comm.get_size(); + auto& minor_comm = handle.get_subcomm(cugraph::partition_manager::minor_comm_name()); + auto const minor_comm_size = minor_comm.get_size(); + + // Shuffle to the proper GPUs + std::forward_as_tuple( + std::tie(edgelist_majors, edgelist_minors, edgelist_ids, edgelist_types), std::ignore) = + cugraph::groupby_gpu_id_and_shuffle_values( + handle.get_comms(), + thrust::make_zip_iterator(thrust::make_tuple(edgelist_majors.begin(), + edgelist_minors.begin(), + edgelist_ids.begin(), + edgelist_types.begin())), + thrust::make_zip_iterator(thrust::make_tuple(edgelist_majors.end(), + edgelist_minors.end(), + edgelist_ids.end(), + edgelist_types.end())), + [key_func = + cugraph::detail::compute_gpu_id_from_ext_edge_id_t{ + comm_size, + major_comm_size, + minor_comm_size}] __device__(auto val) { return key_func(thrust::get<2>(val)); }, + handle.get_stream()); + } + + // + // Sort by edge types and insert to type specific kv_store_t object + // + + auto itr_to_triple = thrust::make_zip_iterator( + edgelist_majors.begin(), edgelist_minors.begin(), edgelist_ids.begin()); + + thrust::sort_by_key( + handle.get_thrust_policy(), edgelist_types.begin(), edgelist_types.end(), itr_to_triple); + + auto nr_uniqe_edge_types_partition = + thrust::count_if(handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(edgelist_types.size()), + detail::is_first_in_run_t{edgelist_types.data()}); + + rmm::device_uvector unique_types(nr_uniqe_edge_types_partition, + handle.get_stream()); + rmm::device_uvector type_offsets(nr_uniqe_edge_types_partition + 1, + handle.get_stream()); + + thrust::copy_if(handle.get_thrust_policy(), + edgelist_types.begin(), + edgelist_types.end(), + thrust::make_counting_iterator(size_t{0}), + unique_types.begin(), + detail::is_first_in_run_t{edgelist_types.data()}); + + type_offsets.set_element_to_zero_async(0, handle.get_stream()); + + thrust::upper_bound(handle.get_thrust_policy(), + edgelist_types.begin(), + edgelist_types.end(), + unique_types.begin(), + unique_types.end(), + type_offsets.begin() + 1); + + std::vector h_unique_types(unique_types.size()); + std::vector h_type_offsets(type_offsets.size()); + + raft::update_host( + h_unique_types.data(), unique_types.data(), unique_types.size(), handle.get_stream()); + + raft::update_host( + h_type_offsets.data(), type_offsets.data(), type_offsets.size(), handle.get_stream()); + handle.sync_stream(); + + for (size_t idx = 0; idx < h_unique_types.size(); idx++) { + auto typ = h_unique_types[idx]; + auto nr_elements_to_insert = (h_type_offsets[idx + 1] - h_type_offsets[idx]); + + auto values_to_insert = + cugraph::allocate_dataframe_buffer(nr_elements_to_insert, handle.get_stream()); + + auto zip_itr = thrust::make_zip_iterator( + thrust::make_tuple(edgelist_majors.begin(), edgelist_minors.begin())); + + thrust::copy(handle.get_thrust_policy(), + zip_itr + h_type_offsets[idx], + zip_itr + h_type_offsets[idx] + nr_elements_to_insert, + cugraph::get_dataframe_buffer_begin(values_to_insert)); + + static_assert(std::is_same_v< + typename thrust::iterator_traits::value_type, + value_t>); + + search_container.insert(handle, + typ, + raft::device_span(edgelist_ids.begin() + h_type_offsets[idx], + nr_elements_to_insert), + std::move(values_to_insert)); + } + } + + return search_container; +} + +template +std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_single_type( + raft::handle_t const& handle, + EdgeTypeAndIdToSrcDstLookupContainerType const& search_container, + raft::device_span edge_ids_to_lookup, + edge_type_t edge_type_to_lookup) +{ + using value_t = typename EdgeTypeAndIdToSrcDstLookupContainerType::value_type; + static_assert(std::is_integral_v); + static_assert(std::is_integral_v); + static_assert(std::is_same_v>); + + static_assert( + std::is_same_v, + "edge_id_t must match EdgeTypeAndIdToSrcDstLookupContainerType::edge_id_type"); + static_assert( + std::is_same_v, + "edge_type_t must match EdgeTypeAndIdToSrcDstLookupContainerType::edge_type_type "); + + auto value_buffer = search_container.lookup_from_edge_ids_and_single_type( + handle, edge_ids_to_lookup, edge_type_to_lookup, multi_gpu); + + return std::make_tuple(std::move(std::get<0>(value_buffer)), + std::move(std::get<1>(value_buffer))); +} + +template +std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_types( + raft::handle_t const& handle, + EdgeTypeAndIdToSrcDstLookupContainerType const& search_container, + raft::device_span edge_ids_to_lookup, + raft::device_span edge_types_to_lookup) +{ + using value_t = typename EdgeTypeAndIdToSrcDstLookupContainerType::value_type; + static_assert(std::is_integral_v); + static_assert(std::is_integral_v); + static_assert(std::is_same_v>); + + assert(edge_ids_to_lookup.size() == edge_types_to_lookup.size()); + + static_assert( + std::is_same_v, + "edge_id_t must match EdgeTypeAndIdToSrcDstLookupContainerType::edge_id_type"); + static_assert( + std::is_same_v, + "edge_type_t must match EdgeTypeAndIdToSrcDstLookupContainerType::edge_type_type "); + + auto value_buffer = search_container.lookup_from_edge_ids_and_types( + handle, edge_ids_to_lookup, edge_types_to_lookup, multi_gpu); + + return std::make_tuple(std::move(std::get<0>(value_buffer)), + std::move(std::get<1>(value_buffer))); +} +} // namespace detail + +template +std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_single_type( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + edge_type_t edge_type_to_lookup) +{ + using m_t = lookup_container_t; + return detail::lookup_endpoints_from_edge_ids_and_single_type( + handle, search_container, edge_ids_to_lookup, edge_type_to_lookup); +} + +template +std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_types( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + raft::device_span edge_types_to_lookup) +{ + using m_t = lookup_container_t; + return detail:: + lookup_endpoints_from_edge_ids_and_types( + handle, search_container, edge_ids_to_lookup, edge_types_to_lookup); +} + +template +lookup_container_t build_edge_id_and_type_to_src_dst_lookup_map( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_id_view, + edge_property_view_t edge_type_view) +{ + using graph_view_t = graph_view_t; + using return_t = lookup_container_t; + + return detail::build_edge_id_and_type_to_src_dst_lookup_map( + handle, graph_view, edge_id_view, edge_type_view); +} +} // namespace cugraph diff --git a/cpp/src/lookup/lookup_src_dst_mg_v32_e32.cu b/cpp/src/lookup/lookup_src_dst_mg_v32_e32.cu new file mode 100644 index 00000000000..5c644c5fe73 --- /dev/null +++ b/cpp/src/lookup/lookup_src_dst_mg_v32_e32.cu @@ -0,0 +1,43 @@ +/* + * Copyright (c) 2024, 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 "lookup/lookup_src_dst_impl.cuh" + +namespace cugraph { + +template class lookup_container_t; + +template lookup_container_t build_edge_id_and_type_to_src_dst_lookup_map( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_id_view, + edge_property_view_t edge_type_view); + +template std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_single_type( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + int32_t edge_type_to_lookup); + +template std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_types( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + raft::device_span edge_types_to_lookup); + +} // namespace cugraph diff --git a/cpp/src/lookup/lookup_src_dst_mg_v32_e64.cu b/cpp/src/lookup/lookup_src_dst_mg_v32_e64.cu new file mode 100644 index 00000000000..4e120f49f10 --- /dev/null +++ b/cpp/src/lookup/lookup_src_dst_mg_v32_e64.cu @@ -0,0 +1,43 @@ +/* + * Copyright (c) 2024, 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 "lookup/lookup_src_dst_impl.cuh" + +namespace cugraph { + +template class lookup_container_t; + +template lookup_container_t build_edge_id_and_type_to_src_dst_lookup_map( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_id_view, + edge_property_view_t edge_type_view); + +template std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_single_type( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + int32_t edge_type_to_lookup); + +template std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_types( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + raft::device_span edge_types_to_lookup); + +} // namespace cugraph diff --git a/cpp/src/lookup/lookup_src_dst_mg_v64_e64.cu b/cpp/src/lookup/lookup_src_dst_mg_v64_e64.cu new file mode 100644 index 00000000000..fc247eb2419 --- /dev/null +++ b/cpp/src/lookup/lookup_src_dst_mg_v64_e64.cu @@ -0,0 +1,43 @@ +/* + * Copyright (c) 2024, 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 "lookup/lookup_src_dst_impl.cuh" + +namespace cugraph { + +template class lookup_container_t; + +template lookup_container_t build_edge_id_and_type_to_src_dst_lookup_map( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_id_view, + edge_property_view_t edge_type_view); + +template std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_single_type( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + int32_t edge_type_to_lookup); + +template std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_types( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + raft::device_span edge_types_to_lookup); + +} // namespace cugraph diff --git a/cpp/src/lookup/lookup_src_dst_sg_v32_e32.cu b/cpp/src/lookup/lookup_src_dst_sg_v32_e32.cu new file mode 100644 index 00000000000..7aa330b5383 --- /dev/null +++ b/cpp/src/lookup/lookup_src_dst_sg_v32_e32.cu @@ -0,0 +1,43 @@ +/* + * Copyright (c) 2024, 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 "lookup/lookup_src_dst_impl.cuh" + +namespace cugraph { + +template class lookup_container_t; + +template lookup_container_t build_edge_id_and_type_to_src_dst_lookup_map( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_id_view, + edge_property_view_t edge_type_view); + +template std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_single_type( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + int32_t const edge_type_to_lookup); + +template std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_types( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + raft::device_span edge_types_to_lookup); + +} // namespace cugraph diff --git a/cpp/src/lookup/lookup_src_dst_sg_v32_e64.cu b/cpp/src/lookup/lookup_src_dst_sg_v32_e64.cu new file mode 100644 index 00000000000..46b62e05ed8 --- /dev/null +++ b/cpp/src/lookup/lookup_src_dst_sg_v32_e64.cu @@ -0,0 +1,43 @@ +/* + * Copyright (c) 2024, 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 "lookup/lookup_src_dst_impl.cuh" + +namespace cugraph { + +template class lookup_container_t; + +template lookup_container_t build_edge_id_and_type_to_src_dst_lookup_map( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_id_view, + edge_property_view_t edge_type_view); + +template std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_single_type( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + int32_t edge_type_to_lookup); + +template std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_types( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + raft::device_span edge_types_to_lookup); + +} // namespace cugraph diff --git a/cpp/src/lookup/lookup_src_dst_sg_v64_e64.cu b/cpp/src/lookup/lookup_src_dst_sg_v64_e64.cu new file mode 100644 index 00000000000..e449eb64de0 --- /dev/null +++ b/cpp/src/lookup/lookup_src_dst_sg_v64_e64.cu @@ -0,0 +1,42 @@ +/* + * Copyright (c) 2024, 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 "lookup/lookup_src_dst_impl.cuh" + +namespace cugraph { + +template class lookup_container_t; + +template lookup_container_t build_edge_id_and_type_to_src_dst_lookup_map( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_id_view, + edge_property_view_t edge_type_view); + +template std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_single_type( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + int32_t edge_type_to_lookup); +template std::tuple, rmm::device_uvector> +lookup_endpoints_from_edge_ids_and_types( + raft::handle_t const& handle, + lookup_container_t const& search_container, + raft::device_span edge_ids_to_lookup, + raft::device_span edge_types_to_lookup); + +} // namespace cugraph diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 2152de28ff9..fd356ff8b89 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -326,11 +326,11 @@ ConfigureTest(LOUVAIN_TEST community/louvain_test.cpp) ConfigureTest(LEIDEN_TEST community/leiden_test.cpp) ################################################################################################### -# - WEIGHTED MATCHING tests ---------------------------------------------------------------------------------- +# - WEIGHTED MATCHING tests ----------------------------------------------------------------------- ConfigureTest(WEIGHTED_MATCHING_TEST community/weighted_matching_test.cpp) ################################################################################################### -# - Legacy ECG tests ------------------------------------------------------------------------------------- +# - Legacy ECG tests ------------------------------------------------------------------------------ ConfigureTest(LEGACY_ECG_TEST community/legacy_ecg_test.cpp) ################################################################################################### @@ -455,13 +455,13 @@ ConfigureTest(EDGE_BETWEENNESS_CENTRALITY_TEST centrality/edge_betweenness_centr # - WEAKLY CONNECTED COMPONENTS tests ------------------------------------------------------------- ConfigureTest(WEAKLY_CONNECTED_COMPONENTS_TEST components/weakly_connected_components_test.cpp) -############################################################################################### -# - MIS tests ------------------------------------------------------------------------------ +################################################################################################### +# - MIS tests ------------------------------------------------------------------------------------- ConfigureTest(MIS_TEST components/mis_test.cu) target_include_directories(MIS_TEST PRIVATE "${CUGRAPH_SOURCE_DIR}/src") -############################################################################################### -# - VERTEX COLORING tests ------------------------------------------------------------------- +################################################################################################### +# - VERTEX COLORING tests ---------------------------------------------------------------------- ConfigureTest(VERTEX_COLORING_TEST components/vertex_coloring_test.cu) target_include_directories(VERTEX_COLORING_TEST PRIVATE "${CUGRAPH_SOURCE_DIR}/src") @@ -470,7 +470,7 @@ target_include_directories(VERTEX_COLORING_TEST PRIVATE "${CUGRAPH_SOURCE_DIR}/s ConfigureTest(SIMILARITY_TEST link_prediction/similarity_test.cu) ################################################################################################### -# - WEIGHTED_SIMILARITY tests ------------------------------------------------------------------------------ +# - WEIGHTED_SIMILARITY tests --------------------------------------------------------------------- ConfigureTest(WEIGHTED_SIMILARITY_TEST link_prediction/weighted_similarity_test.cpp) ################################################################################################### @@ -498,8 +498,8 @@ ConfigureTest(CORE_NUMBER_TEST cores/core_number_test.cpp) # - Core Number tests ----------------------------------------------------------------------------- ConfigureTest(K_CORE_TEST cores/k_core_test.cpp) -############################################################################################### -# - K-truss tests -------------------------------------------------------------------------- +################################################################################################### +# - K-truss tests --------------------------------------------------------------------------------- ConfigureTest(K_TRUSS_TEST community/k_truss_test.cpp) ################################################################################################### @@ -507,9 +507,12 @@ ConfigureTest(K_TRUSS_TEST community/k_truss_test.cpp) ConfigureTest(TRIANGLE_COUNT_TEST community/triangle_count_test.cpp) ################################################################################################### -# - Edge Triangle Count tests -------------------------------------------------------------------------- +# - Edge Triangle Count tests --------------------------------------------------------------------- ConfigureTest(EDGE_TRIANGLE_COUNT_TEST community/edge_triangle_count_test.cpp) +################################################################################################### +# - EDGE SOURCE DESTINATION LOOKUP tests ---------------------------------------------------------- +ConfigureTest(LOOKUP_SRC_DST_TEST lookup/lookup_src_dst_test.cpp) ################################################################################################### # - K-hop Neighbors tests ------------------------------------------------------------------------- @@ -596,7 +599,7 @@ if(BUILD_CUGRAPH_MG_TESTS) ConfigureTestMG(MG_LEIDEN_TEST community/mg_leiden_test.cpp) ############################################################################################### - # - MG WEIGHTED MATCHING tests -------------------------------------------------------------------------- + # - MG WEIGHTED MATCHING tests ---------------------------------------------------------------- ConfigureTestMG(MG_WEIGHTED_MATCHING_TEST community/mg_weighted_matching_test.cpp) ############################################################################################### @@ -612,7 +615,7 @@ if(BUILD_CUGRAPH_MG_TESTS) ConfigureTestMG(MG_EGONET_TEST community/mg_egonet_test.cu) ############################################################################################### - # - MG EDGE TRIANGLE COUNT tests -------------------------------------------------------------------------- + # - MG EDGE TRIANGLE COUNT tests -------------------------------------------------------------- ConfigureTestMG(MG_EDGE_TRIANGLE_COUNT_TEST community/mg_edge_triangle_count_test.cpp) ############################################################################################### @@ -620,6 +623,10 @@ if(BUILD_CUGRAPH_MG_TESTS) ConfigureTestMG(MG_WEAKLY_CONNECTED_COMPONENTS_TEST components/mg_weakly_connected_components_test.cpp) + ############################################################################################### + # - MG EDGE SOURCE DESTINATION LOOKUP tests --------------------------------------------------- + ConfigureTestMG(MG_LOOKUP_SRC_DST_TEST lookup/mg_lookup_src_dst_test.cpp) + ############################################################################################### # - MG MIS tests ------------------------------------------------------------------------------ ConfigureTestMG(MG_MIS_TEST components/mg_mis_test.cu) @@ -730,7 +737,7 @@ if(BUILD_CUGRAPH_MG_TESTS) ConfigureTestMG(MG_RANDOM_WALKS_TEST sampling/mg_random_walks_test.cpp) ############################################################################################### - # - MG WEIGHTED_SIMILARITY tests ----------------------------------------------------------------------- + # - MG WEIGHTED_SIMILARITY tests -------------------------------------------------------------- ConfigureTestMG(MG_WEIGHTED_SIMILARITY_TEST link_prediction/mg_weighted_similarity_test.cpp) ############################################################################################### diff --git a/cpp/tests/community/mg_weighted_matching_test.cpp b/cpp/tests/community/mg_weighted_matching_test.cpp index 4f36ee36902..8abd7646065 100644 --- a/cpp/tests/community/mg_weighted_matching_test.cpp +++ b/cpp/tests/community/mg_weighted_matching_test.cpp @@ -69,7 +69,7 @@ class Tests_MGWeightedMatching constexpr bool multi_gpu = true; - bool test_weighted = true; + bool test_weighted = false; bool renumber = true; bool drop_self_loops = false; bool drop_multi_edges = false; @@ -107,10 +107,23 @@ class Tests_MGWeightedMatching rmm::device_uvector mg_partners(0, handle_->get_stream()); weight_t mg_matching_weights; + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle_->get_comms().barrier(); + hr_timer.start("MG Approximate Weighted Matching"); + } + std::forward_as_tuple(mg_partners, mg_matching_weights) = cugraph::approximate_weighted_matching( *handle_, mg_graph_view, (*mg_edge_weights).view()); + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle_->get_comms().barrier(); + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + if (weighted_matching_usecase.check_correctness) { auto h_mg_partners = cugraph::test::to_host(*handle_, mg_partners); diff --git a/cpp/tests/lookup/lookup_src_dst_test.cpp b/cpp/tests/lookup/lookup_src_dst_test.cpp new file mode 100644 index 00000000000..b2e4355dff4 --- /dev/null +++ b/cpp/tests/lookup/lookup_src_dst_test.cpp @@ -0,0 +1,306 @@ + +/* + * Copyright (c) 2024, 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 governin_from_mtxg permissions and + * limitations under the License. + */ + +#include "utilities/base_fixture.hpp" +#include "utilities/conversion_utilities.hpp" +#include "utilities/property_generator_utilities.hpp" +#include "utilities/test_graphs.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include +#include + +struct EdgeSrcDstLookup_UseCase { + // FIXME: Test with edge mask once the graph generator is updated to generate edge ids and types + bool check_correctness{true}; +}; + +template +class Tests_SGLookupEdgeSrcDst + : public ::testing::TestWithParam> { + public: + Tests_SGLookupEdgeSrcDst() {} + + static void SetUpTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(std::tuple const& param) + { + auto [lookup_usecase, input_usecase] = param; + + raft::handle_t handle{}; + HighResTimer hr_timer{}; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + hr_timer.start("Construct graph"); + } + + constexpr bool multi_gpu = false; + + bool test_weighted = true; + bool renumber = true; + bool drop_self_loops = false; + bool drop_multi_edges = false; + + auto [sg_graph, sg_edge_weights, sg_renumber_map] = + cugraph::test::construct_graph( + handle, input_usecase, test_weighted, renumber, drop_self_loops, drop_multi_edges); + + std::tie(sg_graph, sg_edge_weights, sg_renumber_map) = cugraph::symmetrize_graph( + handle, std::move(sg_graph), std::move(sg_edge_weights), std::move(sg_renumber_map), false); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + auto sg_graph_view = sg_graph.view(); + auto sg_edge_weight_view = + sg_edge_weights ? std::make_optional((*sg_edge_weights).view()) : std::nullopt; + + std::optional> edge_mask{std::nullopt}; + + // + // FIXME: As the graph generator doesn't generate edge ids and types at the moment, generate + // edge ids and types for now and remove the code for generating edge ids and types from this + // file once the graph generator is updated to generate edge ids and types. + // + + int32_t number_of_edge_types = std::max( + 1 << 8, + static_cast(std::rand() % (1 + (sg_graph_view.number_of_vertices() / (1 << 16))))); + + std::optional> edge_types{ + std::nullopt}; + edge_types = cugraph::test::generate::edge_property( + handle, sg_graph_view, number_of_edge_types); + + std::optional> edge_ids{std::nullopt}; + + edge_ids = cugraph::test::generate::edge_property( + handle, sg_graph_view, 1); + + auto edge_counts = (*edge_ids).view().edge_counts(); + + std::vector type_freqs(number_of_edge_types, 0); + std::mutex mtx[number_of_edge_types]; + + for (size_t ep_idx = 0; ep_idx < edge_counts.size(); ep_idx++) { + auto ep_types = + cugraph::test::to_host(handle, + raft::device_span( + (*edge_types).view().value_firsts()[ep_idx], edge_counts[ep_idx])); + + std::for_each(std::execution::par, ep_types.begin(), ep_types.end(), [&](int32_t et) { + std::lock_guard guard(mtx[et]); + type_freqs[et]++; + }); + + auto ep_ids = + cugraph::test::to_host(handle, + raft::device_span( + (*edge_ids).view().value_firsts()[ep_idx], edge_counts[ep_idx])); + } + + assert(std::reduce(type_freqs.cbegin(), type_freqs.cend()) == + std::reduce(edge_counts.cbegin(), edge_counts.cend())); + + auto d_type_freqs = cugraph::test::to_device(handle, type_freqs); + + std::vector type_offsets(number_of_edge_types); + + std::copy(type_freqs.begin(), type_freqs.end(), type_offsets.begin()); + + assert(std::reduce(type_offsets.cbegin(), type_offsets.cend()) == + sg_graph_view.compute_number_of_edges(handle)); + + auto number_of_local_edges = std::reduce(edge_counts.cbegin(), edge_counts.cend()); + + for (size_t ep_idx = 0; ep_idx < edge_counts.size(); ep_idx++) { + auto ep_types = + cugraph::test::to_host(handle, + raft::device_span( + (*edge_types).view().value_firsts()[ep_idx], edge_counts[ep_idx])); + + auto ep_ids = + cugraph::test::to_host(handle, + raft::device_span( + (*edge_ids).view().value_firsts()[ep_idx], edge_counts[ep_idx])); + + std::transform(ep_types.cbegin(), ep_types.cend(), ep_ids.begin(), [&](int32_t et) { + edge_t val = type_offsets[et]; + type_offsets[et]++; + return val; + }); + + raft::update_device((*edge_ids).mutable_view().value_firsts()[ep_idx], + ep_ids.data(), + ep_ids.size(), + handle.get_stream()); + } + + auto search_container = + cugraph::build_edge_id_and_type_to_src_dst_lookup_map( + handle, sg_graph_view, (*edge_ids).view(), (*edge_types).view()); + + if (lookup_usecase.check_correctness) { + rmm::device_uvector d_mg_srcs(0, handle.get_stream()); + rmm::device_uvector d_mg_dsts(0, handle.get_stream()); + + std::optional> d_mg_edge_ids{std::nullopt}; + std::optional> d_mg_edge_types{std::nullopt}; + + std::tie(d_mg_srcs, d_mg_dsts, std::ignore, d_mg_edge_ids, d_mg_edge_types) = + cugraph::decompress_to_edgelist( + handle, + sg_graph_view, + std::optional>{std::nullopt}, + std::make_optional((*edge_ids).view()), + std::make_optional((*edge_types).view()), + std::optional>{std::nullopt}); + + auto number_of_edges = sg_graph_view.compute_number_of_edges(handle); + + auto h_mg_edge_ids = cugraph::test::to_host(handle, d_mg_edge_ids); + auto h_mg_edge_types = cugraph::test::to_host(handle, d_mg_edge_types); + + auto h_srcs_expected = cugraph::test::to_host(handle, d_mg_srcs); + auto h_dsts_expected = cugraph::test::to_host(handle, d_mg_dsts); + + if (number_of_local_edges > 0) { + int nr_wrong_ids_or_types = (std::rand() % number_of_local_edges); + + for (int k = 0; k < nr_wrong_ids_or_types; k++) { + auto id_or_type = std::rand() % 2; + auto random_idx = std::rand() % number_of_local_edges; + if (id_or_type) + (*h_mg_edge_ids)[random_idx] = std::numeric_limits::max(); + else + (*h_mg_edge_types)[random_idx] = std::numeric_limits::max() - 2; + + h_srcs_expected[random_idx] = cugraph::invalid_vertex_id::value; + h_dsts_expected[random_idx] = cugraph::invalid_vertex_id::value; + } + } + + d_mg_edge_ids = cugraph::test::to_device(handle, h_mg_edge_ids); + d_mg_edge_types = cugraph::test::to_device(handle, h_mg_edge_types); + + auto [srcs, dsts] = + cugraph::lookup_endpoints_from_edge_ids_and_types( + handle, + search_container, + raft::device_span((*d_mg_edge_ids).begin(), (*d_mg_edge_ids).size()), + raft::device_span((*d_mg_edge_types).begin(), (*d_mg_edge_types).size())); + + auto h_srcs_results = cugraph::test::to_host(handle, srcs); + auto h_dsts_results = cugraph::test::to_host(handle, dsts); + + EXPECT_EQ(h_srcs_expected.size(), h_srcs_results.size()); + ASSERT_TRUE( + std::equal(h_srcs_expected.begin(), h_srcs_expected.end(), h_srcs_results.begin())); + + EXPECT_EQ(h_dsts_expected.size(), h_dsts_results.size()); + ASSERT_TRUE( + std::equal(h_dsts_expected.begin(), h_dsts_expected.end(), h_dsts_results.begin())); + } + } +}; + +using Tests_SGLookupEdgeSrcDst_File = Tests_SGLookupEdgeSrcDst; +using Tests_SGLookupEdgeSrcDst_Rmat = Tests_SGLookupEdgeSrcDst; + +TEST_P(Tests_SGLookupEdgeSrcDst_File, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGLookupEdgeSrcDst_File, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGLookupEdgeSrcDst_File, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGLookupEdgeSrcDst_Rmat, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGLookupEdgeSrcDst_Rmat, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGLookupEdgeSrcDst_Rmat, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +INSTANTIATE_TEST_SUITE_P( + file_test, + Tests_SGLookupEdgeSrcDst_File, + ::testing::Combine(::testing::Values(EdgeSrcDstLookup_UseCase{}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +INSTANTIATE_TEST_SUITE_P(rmat_small_test, + Tests_SGLookupEdgeSrcDst_Rmat, + ::testing::Combine(::testing::Values(EdgeSrcDstLookup_UseCase{}), + ::testing::Values(cugraph::test::Rmat_Usecase( + 3, 3, 0.57, 0.19, 0.19, 0, true, false)))); + +INSTANTIATE_TEST_SUITE_P( + rmat_benchmark_test, /* note that scale & edge factor can be overridden in benchmarking (with + --gtest_filter to select only the rmat_benchmark_test with a specific + vertex & edge type combination) by command line arguments and do not + include more than one Rmat_Usecase that differ only in scale or edge + factor (to avoid running same benchmarks more than once) */ + Tests_SGLookupEdgeSrcDst_Rmat, + ::testing::Combine( + ::testing::Values(EdgeSrcDstLookup_UseCase{false}), + ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); + +CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/lookup/mg_lookup_src_dst_test.cpp b/cpp/tests/lookup/mg_lookup_src_dst_test.cpp new file mode 100644 index 00000000000..26119801b76 --- /dev/null +++ b/cpp/tests/lookup/mg_lookup_src_dst_test.cpp @@ -0,0 +1,349 @@ +/* + * Copyright (c) 2024, 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 governin_from_mtxg permissions and + * limitations under the License. + */ + +#include "utilities/base_fixture.hpp" +#include "utilities/conversion_utilities.hpp" +#include "utilities/device_comm_wrapper.hpp" +#include "utilities/mg_utilities.hpp" +#include "utilities/property_generator_utilities.hpp" +#include "utilities/test_graphs.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include +#include + +struct EdgeSrcDstLookup_UseCase { + // FIXME: Test with edge mask once the graph generator is updated to generate edge ids and types + bool check_correctness{true}; +}; + +template +class Tests_MGLookupEdgeSrcDst + : public ::testing::TestWithParam> { + public: + Tests_MGLookupEdgeSrcDst() {} + + static void SetUpTestCase() { handle_ = cugraph::test::initialize_mg_handle(); } + static void TearDownTestCase() { handle_.reset(); } + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(std::tuple const& param) + { + auto [lookup_usecase, input_usecase] = param; + + HighResTimer hr_timer{}; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + handle_->get_comms().barrier(); + hr_timer.start("MG Construct graph"); + } + + constexpr bool multi_gpu = true; + + bool test_weighted = false; + bool renumber = true; + bool drop_self_loops = false; + bool drop_multi_edges = false; + + auto [mg_graph, mg_edge_weights, mg_renumber_map] = + cugraph::test::construct_graph( + *handle_, input_usecase, test_weighted, renumber, drop_self_loops, drop_multi_edges); + + std::tie(mg_graph, mg_edge_weights, mg_renumber_map) = cugraph::symmetrize_graph( + *handle_, + std::move(mg_graph), + std::move(mg_edge_weights), + mg_renumber_map ? std::optional>(std::move(*mg_renumber_map)) + : std::nullopt, + false); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + handle_->get_comms().barrier(); + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + auto mg_graph_view = mg_graph.view(); + auto mg_edge_weight_view = + mg_edge_weights ? std::make_optional((*mg_edge_weights).view()) : std::nullopt; + + std::optional> edge_mask{std::nullopt}; + + // + // FIXME: As the graph generator doesn't generate edge ids and types at the moment, generate + // edge ids and types for now and remove the code for generating edge ids and types from this + // file once the graph generator is updated to generate edge ids and types. + // + + int number_of_edge_types = std::max( + 1 << 8, + static_cast(std::rand() % (1 + (mg_graph_view.number_of_vertices() / (1 << 20))))); + + std::optional> edge_types{ + std::nullopt}; + edge_types = cugraph::test::generate::edge_property( + *handle_, mg_graph_view, number_of_edge_types); + + std::optional> edge_ids{std::nullopt}; + + edge_ids = cugraph::test::generate::edge_property( + *handle_, mg_graph_view, 1); + + auto edge_counts = (*edge_ids).view().edge_counts(); + + auto const comm_rank = (*handle_).get_comms().get_rank(); + auto const comm_size = (*handle_).get_comms().get_size(); + + std::vector type_freqs(number_of_edge_types, 0); + std::mutex mtx[number_of_edge_types]; + + for (size_t ep_idx = 0; ep_idx < edge_counts.size(); ep_idx++) { + auto ep_types = + cugraph::test::to_host(*handle_, + raft::device_span( + (*edge_types).view().value_firsts()[ep_idx], edge_counts[ep_idx])); + + std::for_each(std::execution::par, ep_types.begin(), ep_types.end(), [&](int32_t et) { + std::lock_guard guard(mtx[et]); + type_freqs[et]++; + }); + + auto ep_ids = + cugraph::test::to_host(*handle_, + raft::device_span( + (*edge_ids).view().value_firsts()[ep_idx], edge_counts[ep_idx])); + } + + assert(std::reduce(type_freqs.cbegin(), type_freqs.cend()) == + std::reduce(edge_counts.cbegin(), edge_counts.cend())); + + auto d_type_freqs = cugraph::test::to_device(*handle_, type_freqs); + d_type_freqs = + cugraph::test::device_allgatherv(*handle_, d_type_freqs.data(), d_type_freqs.size()); + type_freqs = cugraph::test::to_host(*handle_, d_type_freqs); + + std::vector distributed_type_offsets(comm_size * number_of_edge_types); + + for (size_t i = 0; i < number_of_edge_types; i++) { + for (size_t j = 0; j < comm_size; j++) { + distributed_type_offsets[j + comm_size * i] = type_freqs[number_of_edge_types * j + i]; + } + } + + // prefix sum for each type + for (size_t i = 0; i < number_of_edge_types; i++) { + auto start = distributed_type_offsets.begin() + i * comm_size; + std::exclusive_scan(start, start + comm_size, start, 0); + } + + assert(std::reduce(distributed_type_offsets.cbegin(), distributed_type_offsets.cend()) == + mg_graph_view.compute_number_of_edges(*handle_)); + + auto number_of_local_edges = std::reduce(edge_counts.cbegin(), edge_counts.cend()); + + for (size_t ep_idx = 0; ep_idx < edge_counts.size(); ep_idx++) { + auto ep_types = + cugraph::test::to_host(*handle_, + raft::device_span( + (*edge_types).view().value_firsts()[ep_idx], edge_counts[ep_idx])); + + auto ep_ids = + cugraph::test::to_host(*handle_, + raft::device_span( + (*edge_ids).view().value_firsts()[ep_idx], edge_counts[ep_idx])); + + std::transform(ep_types.cbegin(), ep_types.cend(), ep_ids.begin(), [&](int32_t et) { + edge_t val = distributed_type_offsets[(comm_size * et + comm_rank)]; + distributed_type_offsets[(comm_size * et + comm_rank)]++; + return val; + }); + + raft::update_device((*edge_ids).mutable_view().value_firsts()[ep_idx], + ep_ids.data(), + ep_ids.size(), + handle_->get_stream()); + } + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle_->get_comms().barrier(); + hr_timer.start("MG Build Lookup Map"); + } + + auto search_container = + cugraph::build_edge_id_and_type_to_src_dst_lookup_map( + *handle_, mg_graph_view, (*edge_ids).view(), (*edge_types).view()); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement + handle_->get_comms().barrier(); + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + if (lookup_usecase.check_correctness) { + rmm::device_uvector d_mg_srcs(0, handle_->get_stream()); + rmm::device_uvector d_mg_dsts(0, handle_->get_stream()); + + std::optional> d_mg_edge_ids{std::nullopt}; + std::optional> d_mg_edge_types{std::nullopt}; + + std::tie(d_mg_srcs, d_mg_dsts, std::ignore, d_mg_edge_ids, d_mg_edge_types) = + cugraph::decompress_to_edgelist( + *handle_, + mg_graph_view, + std::optional>{std::nullopt}, + std::make_optional((*edge_ids).view()), + std::make_optional((*edge_types).view()), + std::optional>{std::nullopt}); + + auto number_of_edges = mg_graph_view.compute_number_of_edges(*handle_); + + auto h_mg_edge_ids = cugraph::test::to_host(*handle_, d_mg_edge_ids); + auto h_mg_edge_types = cugraph::test::to_host(*handle_, d_mg_edge_types); + + auto h_srcs_expected = cugraph::test::to_host(*handle_, d_mg_srcs); + auto h_dsts_expected = cugraph::test::to_host(*handle_, d_mg_dsts); + + if (number_of_local_edges > 0) { + int nr_wrong_ids_or_types = (std::rand() % number_of_local_edges); + + for (int k = 0; k < nr_wrong_ids_or_types; k++) { + auto id_or_type = std::rand() % 2; + auto random_idx = std::rand() % number_of_local_edges; + if (id_or_type) + (*h_mg_edge_ids)[random_idx] = std::numeric_limits::max(); + else + (*h_mg_edge_types)[random_idx] = std::numeric_limits::max() - 2; + + h_srcs_expected[random_idx] = cugraph::invalid_vertex_id::value; + h_dsts_expected[random_idx] = cugraph::invalid_vertex_id::value; + } + } + + d_mg_edge_ids = cugraph::test::to_device(*handle_, h_mg_edge_ids); + d_mg_edge_types = cugraph::test::to_device(*handle_, h_mg_edge_types); + + auto [srcs, dsts] = + cugraph::lookup_endpoints_from_edge_ids_and_types( + *handle_, + search_container, + raft::device_span((*d_mg_edge_ids).begin(), (*d_mg_edge_ids).size()), + raft::device_span((*d_mg_edge_types).begin(), (*d_mg_edge_types).size())); + + auto h_srcs_results = cugraph::test::to_host(*handle_, srcs); + auto h_dsts_results = cugraph::test::to_host(*handle_, dsts); + + EXPECT_EQ(h_srcs_expected.size(), h_srcs_results.size()); + ASSERT_TRUE( + std::equal(h_srcs_expected.begin(), h_srcs_expected.end(), h_srcs_results.begin())); + + EXPECT_EQ(h_dsts_expected.size(), h_dsts_results.size()); + ASSERT_TRUE( + std::equal(h_dsts_expected.begin(), h_dsts_expected.end(), h_dsts_results.begin())); + } + } + + private: + static std::unique_ptr handle_; +}; + +template +std::unique_ptr Tests_MGLookupEdgeSrcDst::handle_ = nullptr; + +using Tests_MGLookupEdgeSrcDst_File = Tests_MGLookupEdgeSrcDst; +using Tests_MGLookupEdgeSrcDst_Rmat = Tests_MGLookupEdgeSrcDst; + +TEST_P(Tests_MGLookupEdgeSrcDst_File, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGLookupEdgeSrcDst_File, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGLookupEdgeSrcDst_File, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGLookupEdgeSrcDst_Rmat, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGLookupEdgeSrcDst_Rmat, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGLookupEdgeSrcDst_Rmat, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +INSTANTIATE_TEST_SUITE_P( + file_test, + Tests_MGLookupEdgeSrcDst_File, + ::testing::Combine(::testing::Values(EdgeSrcDstLookup_UseCase{}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +INSTANTIATE_TEST_SUITE_P(rmat_small_test, + Tests_MGLookupEdgeSrcDst_Rmat, + ::testing::Combine(::testing::Values(EdgeSrcDstLookup_UseCase{}), + ::testing::Values(cugraph::test::Rmat_Usecase( + 3, 2, 0.57, 0.19, 0.19, 0, true, false)))); + +INSTANTIATE_TEST_SUITE_P( + rmat_benchmark_test, /* note that scale & edge factor can be overridden in benchmarking (with + --gtest_filter to select only the rmat_benchmark_test with a specific + vertex & edge type combination) by command line arguments and do not + include more than one Rmat_Usecase that differ only in scale or edge + factor (to avoid running same benchmarks more than once) */ + Tests_MGLookupEdgeSrcDst_Rmat, + ::testing::Combine( + ::testing::Values(EdgeSrcDstLookup_UseCase{false}), + ::testing::Values(cugraph::test::Rmat_Usecase(5, 32, 0.57, 0.19, 0.19, 0, true, false)))); + +CUGRAPH_MG_TEST_PROGRAM_MAIN()