diff --git a/cpp/include/cugraph/graph_functions.hpp b/cpp/include/cugraph/graph_functions.hpp index be79f0eec84..4de12eac16e 100644 --- a/cpp/include/cugraph/graph_functions.hpp +++ b/cpp/include/cugraph/graph_functions.hpp @@ -19,6 +19,7 @@ #include #include +#include #include #include @@ -471,8 +472,8 @@ std::tuple, extract_induced_subgraphs( raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets /* size == num_subgraphs + 1 */, - vertex_t const* subgraph_vertices /* size == subgraph_offsets[num_subgraphs] */, + raft::device_span subgraph_offsets /* size == num_subgraphs + 1 */, + raft::device_span subgraph_vertices /* size == subgraph_offsets[num_subgraphs] */, size_t num_subgraphs, bool do_expensive_check = false); diff --git a/cpp/src/community/legacy/egonet.cu b/cpp/src/community/legacy/egonet.cu index 273c272ef07..2d80b37f8b8 100644 --- a/cpp/src/community/legacy/egonet.cu +++ b/cpp/src/community/legacy/egonet.cu @@ -169,7 +169,12 @@ extract(raft::handle_t const& handle, // extract return cugraph::extract_induced_subgraphs( - handle, csr_view, neighbors_offsets.data().get(), neighbors.data().get(), n_subgraphs); + handle, + csr_view, + raft::device_span(neighbors_offsets.data().get(), neighbors_offsets.size()), + raft::device_span(neighbors.data().get(), neighbors.size()), + n_subgraphs, + false); } } // namespace diff --git a/cpp/src/structure/induced_subgraph_impl.cuh b/cpp/src/structure/induced_subgraph_impl.cuh index b8ad7fc5c2b..33b3fa546fd 100644 --- a/cpp/src/structure/induced_subgraph_impl.cuh +++ b/cpp/src/structure/induced_subgraph_impl.cuh @@ -14,7 +14,6 @@ * limitations under the License. */ #pragma once - #include #include #include @@ -48,8 +47,8 @@ std::tuple, extract_induced_subgraphs( raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets /* size == num_subgraphs + 1 */, - vertex_t const* subgraph_vertices /* size == subgraph_offsets[num_subgraphs] */, + raft::device_span subgraph_offsets /* size == num_subgraphs + 1 */, + raft::device_span subgraph_vertices /* size == subgraph_offsets[num_subgraphs] */, size_t num_subgraphs, bool do_expensive_check) { @@ -69,24 +68,25 @@ extract_induced_subgraphs( if (do_expensive_check) { size_t should_be_zero{std::numeric_limits::max()}; size_t num_aggregate_subgraph_vertices{}; - raft::update_host(&should_be_zero, subgraph_offsets, 1, handle.get_stream()); - raft::update_host( - &num_aggregate_subgraph_vertices, subgraph_offsets + num_subgraphs, 1, handle.get_stream()); + raft::update_host(&should_be_zero, subgraph_offsets.data(), 1, handle.get_stream()); + raft::update_host(&num_aggregate_subgraph_vertices, + subgraph_offsets.data() + num_subgraphs, + 1, + handle.get_stream()); handle.sync_stream(); CUGRAPH_EXPECTS(should_be_zero == 0, "Invalid input argument: subgraph_offsets[0] should be 0."); - CUGRAPH_EXPECTS( - thrust::is_sorted( - handle.get_thrust_policy(), subgraph_offsets, subgraph_offsets + (num_subgraphs + 1)), - "Invalid input argument: subgraph_offsets is not sorted."); + CUGRAPH_EXPECTS(thrust::is_sorted( + handle.get_thrust_policy(), subgraph_offsets.begin(), subgraph_offsets.end()), + "Invalid input argument: subgraph_offsets is not sorted."); auto vertex_partition = vertex_partition_device_view_t(graph_view.local_vertex_partition_view()); CUGRAPH_EXPECTS( thrust::count_if(handle.get_thrust_policy(), - subgraph_vertices, - subgraph_vertices + num_aggregate_subgraph_vertices, + subgraph_vertices.begin(), + subgraph_vertices.end(), [vertex_partition] __device__(auto v) { return !vertex_partition.is_valid_vertex(v) || !vertex_partition.in_local_vertex_partition_range_nocheck(v); @@ -101,8 +101,8 @@ extract_induced_subgraphs( [subgraph_offsets, subgraph_vertices] __device__(auto i) { // vertices are sorted and unique return !thrust::is_sorted(thrust::seq, - subgraph_vertices + subgraph_offsets[i], - subgraph_vertices + subgraph_offsets[i + 1]) || + subgraph_vertices.begin() + subgraph_offsets[i], + subgraph_vertices.begin() + subgraph_offsets[i + 1]) || (thrust::count_if( thrust::seq, thrust::make_counting_iterator(subgraph_offsets[i]), @@ -127,8 +127,10 @@ extract_induced_subgraphs( // 2-1. Phase 1: calculate memory requirements size_t num_aggregate_subgraph_vertices{}; - raft::update_host( - &num_aggregate_subgraph_vertices, subgraph_offsets + num_subgraphs, 1, handle.get_stream()); + raft::update_host(&num_aggregate_subgraph_vertices, + subgraph_offsets.data() + num_subgraphs, + 1, + handle.get_stream()); handle.sync_stream(); rmm::device_uvector subgraph_vertex_output_offsets( @@ -145,9 +147,10 @@ extract_induced_subgraphs( thrust::make_counting_iterator(num_aggregate_subgraph_vertices), subgraph_vertex_output_offsets.begin(), [subgraph_offsets, subgraph_vertices, num_subgraphs, edge_partition] __device__(auto i) { - auto subgraph_idx = thrust::distance( - subgraph_offsets + 1, - thrust::upper_bound(thrust::seq, subgraph_offsets, subgraph_offsets + num_subgraphs, i)); + auto subgraph_idx = + thrust::distance(subgraph_offsets.begin() + 1, + thrust::upper_bound( + thrust::seq, subgraph_offsets.begin(), subgraph_offsets.end() - 1, i)); vertex_t const* indices{nullptr}; thrust::optional weights{thrust::nullopt}; edge_t local_degree{}; @@ -158,9 +161,9 @@ extract_induced_subgraphs( thrust::seq, indices, indices + local_degree, - [vertex_first = subgraph_vertices + subgraph_offsets[subgraph_idx], + [vertex_first = subgraph_vertices.begin() + subgraph_offsets[subgraph_idx], vertex_last = - subgraph_vertices + subgraph_offsets[subgraph_idx + 1]] __device__(auto nbr) { + subgraph_vertices.begin() + subgraph_offsets[subgraph_idx + 1]] __device__(auto nbr) { return thrust::binary_search(thrust::seq, vertex_first, vertex_last, nbr); }); }); @@ -201,9 +204,9 @@ extract_induced_subgraphs( edge_weights = edge_weights ? thrust::optional{(*edge_weights).data()} : thrust::nullopt] __device__(auto i) { auto subgraph_idx = thrust::distance( - subgraph_offsets + 1, + subgraph_offsets.begin() + 1, thrust::upper_bound( - thrust::seq, subgraph_offsets, subgraph_offsets + num_subgraphs, size_t{i})); + thrust::seq, subgraph_offsets.begin(), subgraph_offsets.end() - 1, size_t{i})); vertex_t const* indices{nullptr}; thrust::optional weights{thrust::nullopt}; edge_t local_degree{}; @@ -219,9 +222,9 @@ extract_induced_subgraphs( triplet_first + local_degree, thrust::make_zip_iterator(thrust::make_tuple(edge_majors, edge_minors, *edge_weights)) + subgraph_vertex_output_offsets[i], - [vertex_first = subgraph_vertices + subgraph_offsets[subgraph_idx], + [vertex_first = subgraph_vertices.begin() + subgraph_offsets[subgraph_idx], vertex_last = - subgraph_vertices + subgraph_offsets[subgraph_idx + 1]] __device__(auto t) { + subgraph_vertices.begin() + subgraph_offsets[subgraph_idx + 1]] __device__(auto t) { return thrust::binary_search( thrust::seq, vertex_first, vertex_last, thrust::get<1>(t)); }); @@ -229,24 +232,25 @@ extract_induced_subgraphs( auto pair_first = thrust::make_zip_iterator( thrust::make_tuple(thrust::make_constant_iterator(subgraph_vertices[i]), indices)); // FIXME: this is inefficient for high local degree vertices - thrust::copy_if(thrust::seq, - pair_first, - pair_first + local_degree, - thrust::make_zip_iterator(thrust::make_tuple(edge_majors, edge_minors)) + - subgraph_vertex_output_offsets[i], - [vertex_first = subgraph_vertices + subgraph_offsets[subgraph_idx], - vertex_last = subgraph_vertices + - subgraph_offsets[subgraph_idx + 1]] __device__(auto t) { - return thrust::binary_search( - thrust::seq, vertex_first, vertex_last, thrust::get<1>(t)); - }); + thrust::copy_if( + thrust::seq, + pair_first, + pair_first + local_degree, + thrust::make_zip_iterator(thrust::make_tuple(edge_majors, edge_minors)) + + subgraph_vertex_output_offsets[i], + [vertex_first = subgraph_vertices.begin() + subgraph_offsets[subgraph_idx], + vertex_last = + subgraph_vertices.begin() + subgraph_offsets[subgraph_idx + 1]] __device__(auto t) { + return thrust::binary_search( + thrust::seq, vertex_first, vertex_last, thrust::get<1>(t)); + }); } }); rmm::device_uvector subgraph_edge_offsets(num_subgraphs + 1, handle.get_stream()); thrust::gather(handle.get_thrust_policy(), - subgraph_offsets, - subgraph_offsets + (num_subgraphs + 1), + subgraph_offsets.begin(), + subgraph_offsets.end(), subgraph_vertex_output_offsets.begin(), subgraph_edge_offsets.begin()); #ifdef TIMING diff --git a/cpp/src/structure/induced_subgraph_mg.cu b/cpp/src/structure/induced_subgraph_mg.cu index 8f7b242f367..8ddcc78e3bc 100644 --- a/cpp/src/structure/induced_subgraph_mg.cu +++ b/cpp/src/structure/induced_subgraph_mg.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -25,8 +25,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -36,8 +36,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -47,8 +47,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -58,8 +58,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -69,8 +69,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -80,8 +80,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -91,8 +91,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -102,8 +102,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -113,8 +113,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int64_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -124,8 +124,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int64_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -135,8 +135,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int64_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -146,8 +146,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int64_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); diff --git a/cpp/src/structure/induced_subgraph_sg.cu b/cpp/src/structure/induced_subgraph_sg.cu index a98a7cc1994..484b628f2f5 100644 --- a/cpp/src/structure/induced_subgraph_sg.cu +++ b/cpp/src/structure/induced_subgraph_sg.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -25,8 +25,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -36,8 +36,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -47,8 +47,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -58,8 +58,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -69,8 +69,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -80,8 +80,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -91,8 +91,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -102,8 +102,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int32_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -113,8 +113,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int64_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -124,8 +124,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int64_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -135,8 +135,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int64_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); @@ -146,8 +146,8 @@ template std::tuple, rmm::device_uvector> extract_induced_subgraphs(raft::handle_t const& handle, graph_view_t const& graph_view, - size_t const* subgraph_offsets, - int64_t const* subgraph_vertices, + raft::device_span subgraph_offsets, + raft::device_span subgraph_vertices, size_t num_subgraphs, bool do_expensive_check); diff --git a/cpp/tests/community/induced_subgraph_test.cpp b/cpp/tests/community/induced_subgraph_test.cpp index 3fe636aa028..ed2a43e0317 100644 --- a/cpp/tests/community/induced_subgraph_test.cpp +++ b/cpp/tests/community/induced_subgraph_test.cpp @@ -23,6 +23,7 @@ #include #include +#include #include #include @@ -199,12 +200,13 @@ class Tests_InducedSubgraph : public ::testing::TestWithParam(d_subgraph_offsets.data(), d_subgraph_offsets.size()), + raft::device_span(d_subgraph_vertices.data(), d_subgraph_vertices.size()), + configuration.subgraph_sizes.size(), + true); RAFT_CUDA_TRY(cudaDeviceSynchronize()); // for consistent performance measurement diff --git a/cpp/tests/sampling/sg_uniform_neighbor_sampling.cu b/cpp/tests/sampling/sg_uniform_neighbor_sampling.cu index 20197fe629b..8624c56e828 100644 --- a/cpp/tests/sampling/sg_uniform_neighbor_sampling.cu +++ b/cpp/tests/sampling/sg_uniform_neighbor_sampling.cu @@ -104,7 +104,12 @@ class Tests_Uniform_Neighbor_Sampling handle.get_stream()); auto [d_src_in, d_dst_in, d_indices_in, d_ignore] = extract_induced_subgraphs( - handle, graph_view, d_subgraph_offsets.data(), d_vertices.data(), 1, true); + handle, + graph_view, + raft::device_span(d_subgraph_offsets.data(), d_subgraph_offsets.size()), + raft::device_span(d_vertices.data(), d_vertices.size()), + 1, + true); cugraph::test::validate_extracted_graph_is_subgraph( handle, d_src_in, d_dst_in, *d_indices_in, d_src_out, d_dst_out, d_indices);