From 65ca8767b9e5357fb7bcf6b17c5755cfb2e359ef Mon Sep 17 00:00:00 2001 From: Victor Lafargue Date: Fri, 27 Aug 2021 21:15:34 +0200 Subject: [PATCH] Apply modifications to account for RAFT changes (#1707) This PR apply modifications to the cuGraph codebase to account for changes in RAFT and RMM : - https://github.com/rapidsai/raft/pull/283 - https://github.com/rapidsai/raft/pull/285 - https://github.com/rapidsai/raft/pull/286 - https://github.com/rapidsai/rmm/pull/816 This PR requires some changes in the cuHornet dependency : https://github.com/rapidsai/cuhornet/pull/52 Authors: - Victor Lafargue (https://github.com/viclafargue) Approvers: - Brad Rees (https://github.com/BradReesWork) - AJ Schmidt (https://github.com/ajschmidt8) - Seunghwa Kang (https://github.com/seunghwak) - Chuck Hastings (https://github.com/ChuckHastings) URL: https://github.com/rapidsai/cugraph/pull/1707 --- cpp/include/cugraph/detail/graph_utils.cuh | 15 +-- .../prims/copy_to_adj_matrix_row_col.cuh | 20 ++-- .../copy_v_transform_reduce_in_out_nbr.cuh | 7 +- ...ransform_reduce_key_aggregated_out_nbr.cuh | 48 +++++----- cpp/include/cugraph/prims/count_if_v.cuh | 5 +- cpp/include/cugraph/prims/reduce_v.cuh | 4 +- ...orm_reduce_by_adj_matrix_row_col_key_e.cuh | 8 +- .../cugraph/prims/transform_reduce_e.cuh | 4 +- .../cugraph/prims/transform_reduce_v.cuh | 4 +- .../update_frontier_v_push_if_out_nbr.cuh | 49 +++++----- cpp/include/cugraph/prims/vertex_frontier.cuh | 36 ++++--- .../cugraph/utilities/host_barrier.hpp | 1 - cpp/src/centrality/betweenness_centrality.cu | 23 ++--- cpp/src/centrality/betweenness_centrality.cuh | 1 - cpp/src/centrality/katz_centrality.cu | 8 +- cpp/src/community/flatten_dendrogram.cuh | 2 +- cpp/src/community/legacy/ecg.cu | 4 +- cpp/src/community/legacy/leiden.cu | 2 +- cpp/src/community/legacy/leiden.cuh | 13 ++- cpp/src/community/legacy/louvain.cuh | 79 +++++++-------- .../community/legacy/spectral_clustering.cu | 21 +--- .../community/legacy/triangles_counting.cu | 4 - cpp/src/community/louvain.cu | 4 +- cpp/src/community/louvain.cuh | 21 ++-- cpp/src/components/weak_cc.cuh | 1 - .../components/weakly_connected_components.cu | 54 +++++------ cpp/src/converters/COOtoCSR.cuh | 1 - cpp/src/generators/erdos_renyi_generator.cu | 7 +- cpp/src/generators/generate_rmat_edgelist.cu | 3 +- cpp/src/generators/generator_tools.cu | 22 ++--- cpp/src/generators/simple_generators.cu | 25 ++--- cpp/src/layout/barnes_hut.cuh | 18 ++-- cpp/src/layout/exact_fa2.cuh | 18 ++-- cpp/src/linear_assignment/hungarian.cu | 32 +++---- cpp/src/link_analysis/pagerank.cu | 10 +- cpp/src/link_prediction/overlap.cu | 1 - cpp/src/sampling/random_walks.cuh | 53 +++++----- cpp/src/sampling/rw_traversals.hpp | 3 +- cpp/src/serialization/serializer.cu | 8 +- cpp/src/structure/coarsen_graph.cu | 65 ++++++------- .../structure/create_graph_from_edgelist.cpp | 7 +- cpp/src/structure/graph_view.cu | 28 +++--- cpp/src/structure/induced_subgraph.cu | 21 ++-- cpp/src/structure/relabel.cu | 31 +++--- cpp/src/structure/renumber_edgelist.cu | 96 ++++++++----------- cpp/src/structure/renumber_utils.cu | 31 +++--- cpp/src/topology/topology.cuh | 18 ++-- cpp/src/traversal/bfs.cu | 2 +- cpp/src/traversal/legacy/bfs.cuh | 1 - cpp/src/traversal/legacy/mg/bfs.cuh | 4 +- cpp/src/traversal/legacy/mg/common_utils.cuh | 20 ++-- .../legacy/mg/vertex_binning_kernels.cuh | 1 - cpp/src/traversal/legacy/sssp.cuh | 1 - cpp/src/traversal/sssp.cu | 6 +- cpp/src/traversal/two_hop_neighbors.cu | 1 - cpp/src/utilities/cython.cu | 9 +- cpp/src/utilities/path_retrieval.cu | 5 +- cpp/src/utilities/spmv_1D.cuh | 1 - .../legacy/betweenness_centrality_test.cu | 1 + .../edge_betweenness_centrality_test.cu | 1 + .../centrality/legacy/katz_centrality_test.cu | 15 +-- cpp/tests/community/ecg_test.cpp | 1 - cpp/tests/community/mg_louvain_helper.cu | 28 +++--- cpp/tests/components/con_comp_test.cu | 2 + cpp/tests/components/scc_test.cu | 1 + cpp/tests/components/wcc_graphs.cu | 10 +- cpp/tests/prims/mg_count_if_v.cu | 2 +- cpp/tests/prims/mg_reduce_v.cu | 6 +- cpp/tests/prims/mg_transform_reduce_v.cu | 2 +- cpp/tests/sampling/random_walks_profiling.cu | 2 +- cpp/tests/sampling/random_walks_test.cu | 2 +- cpp/tests/sampling/rw_low_level_test.cu | 2 +- cpp/tests/traversal/legacy/sssp_test.cu | 1 + .../utilities/matrix_market_file_utilities.cu | 10 +- cpp/tests/utilities/thrust_wrapper.cu | 24 ++--- 75 files changed, 499 insertions(+), 598 deletions(-) diff --git a/cpp/include/cugraph/detail/graph_utils.cuh b/cpp/include/cugraph/detail/graph_utils.cuh index 7f22699b62c..98ebce63b1c 100644 --- a/cpp/include/cugraph/detail/graph_utils.cuh +++ b/cpp/include/cugraph/detail/graph_utils.cuh @@ -78,7 +78,8 @@ rmm::device_uvector compute_major_degrees( [(detail::num_sparse_segments_per_vertex_partition + 2) * i + detail::num_sparse_segments_per_vertex_partition] : major_last; - thrust::transform(rmm::exec_policy(handle.get_stream()), + auto execution_policy = handle.get_thrust_policy(); + thrust::transform(execution_policy, thrust::make_counting_iterator(vertex_t{0}), thrust::make_counting_iterator(major_hypersparse_first - major_first), local_degrees.begin(), @@ -86,11 +87,11 @@ rmm::device_uvector compute_major_degrees( if (use_dcs) { auto p_dcs_nzd_vertices = (*adj_matrix_partition_dcs_nzd_vertices)[i]; auto dcs_nzd_vertex_count = (*adj_matrix_partition_dcs_nzd_vertex_counts)[i]; - thrust::fill(rmm::exec_policy(handle.get_stream()), + thrust::fill(execution_policy, local_degrees.begin() + (major_hypersparse_first - major_first), local_degrees.begin() + (major_last - major_first), edge_t{0}); - thrust::for_each(rmm::exec_policy(handle.get_stream()), + thrust::for_each(execution_policy, thrust::make_counting_iterator(vertex_t{0}), thrust::make_counting_iterator(dcs_nzd_vertex_count), [p_offsets, @@ -123,10 +124,10 @@ rmm::device_uvector compute_major_degrees(raft::handle_t const& handle, vertex_t number_of_vertices) { rmm::device_uvector degrees(number_of_vertices, handle.get_stream()); - thrust::tabulate(rmm::exec_policy(handle.get_stream()), - degrees.begin(), - degrees.end(), - [offsets] __device__(auto i) { return offsets[i + 1] - offsets[i]; }); + thrust::tabulate( + handle.get_thrust_policy(), degrees.begin(), degrees.end(), [offsets] __device__(auto i) { + return offsets[i + 1] - offsets[i]; + }); return degrees; } diff --git a/cpp/include/cugraph/prims/copy_to_adj_matrix_row_col.cuh b/cpp/include/cugraph/prims/copy_to_adj_matrix_row_col.cuh index 435340f84dc..af5081a33d1 100644 --- a/cpp/include/cugraph/prims/copy_to_adj_matrix_row_col.cuh +++ b/cpp/include/cugraph/prims/copy_to_adj_matrix_row_col.cuh @@ -98,7 +98,7 @@ void copy_to_matrix_major(raft::handle_t const& handle, assert(graph_view.get_number_of_local_vertices() == GraphViewType::is_adj_matrix_transposed ? graph_view.get_number_of_local_adj_matrix_partition_cols() : graph_view.get_number_of_local_adj_matrix_partition_rows()); - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(handle.get_thrust_policy(), vertex_value_input_first, vertex_value_input_first + graph_view.get_number_of_local_vertices(), matrix_major_value_output_first); @@ -169,7 +169,7 @@ void copy_to_matrix_major(raft::handle_t const& handle, }); // FIXME: this gather (and temporary buffer) is unnecessary if NCCL directly takes a // permutation iterator (and directly gathers to the internal buffer) - thrust::gather(rmm::exec_policy(handle.get_stream()), + thrust::gather(handle.get_thrust_policy(), map_first, map_first + thrust::distance(vertex_first, vertex_last), vertex_value_input_first, @@ -190,7 +190,7 @@ void copy_to_matrix_major(raft::handle_t const& handle, // FIXME: this scatter is unnecessary if NCCL directly takes a permutation iterator (and // directly scatters from the internal buffer) thrust::scatter( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), rx_value_first, rx_value_first + rx_counts[i], map_first, @@ -203,7 +203,7 @@ void copy_to_matrix_major(raft::handle_t const& handle, // FIXME: this scatter is unnecessary if NCCL directly takes a permutation iterator (and // directly scatters from the internal buffer) thrust::scatter( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), rx_value_first, rx_value_first + rx_counts[i], map_first, @@ -226,7 +226,7 @@ void copy_to_matrix_major(raft::handle_t const& handle, ? graph_view.get_number_of_local_adj_matrix_partition_cols() : graph_view.get_number_of_local_adj_matrix_partition_rows()); auto val_first = thrust::make_permutation_iterator(vertex_value_input_first, vertex_first); - thrust::scatter(rmm::exec_policy(handle.get_stream()), + thrust::scatter(handle.get_thrust_policy(), val_first, val_first + thrust::distance(vertex_first, vertex_last), vertex_first, @@ -290,7 +290,7 @@ void copy_to_matrix_minor(raft::handle_t const& handle, assert(graph_view.get_number_of_local_vertices() == GraphViewType::is_adj_matrix_transposed ? graph_view.get_number_of_local_adj_matrix_partition_rows() : graph_view.get_number_of_local_adj_matrix_partition_cols()); - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(handle.get_thrust_policy(), vertex_value_input_first, vertex_value_input_first + graph_view.get_number_of_local_vertices(), matrix_minor_value_output_first); @@ -360,7 +360,7 @@ void copy_to_matrix_minor(raft::handle_t const& handle, }); // FIXME: this gather (and temporary buffer) is unnecessary if NCCL directly takes a // permutation iterator (and directly gathers to the internal buffer) - thrust::gather(rmm::exec_policy(handle.get_stream()), + thrust::gather(handle.get_thrust_policy(), map_first, map_first + thrust::distance(vertex_first, vertex_last), vertex_value_input_first, @@ -380,7 +380,7 @@ void copy_to_matrix_minor(raft::handle_t const& handle, }); // FIXME: this scatter is unnecessary if NCCL directly takes a permutation iterator (and // directly scatters from the internal buffer) - thrust::scatter(rmm::exec_policy(handle.get_stream()), + thrust::scatter(handle.get_thrust_policy(), rx_value_first, rx_value_first + rx_counts[i], map_first, @@ -392,7 +392,7 @@ void copy_to_matrix_minor(raft::handle_t const& handle, }); // FIXME: this scatter is unnecessary if NCCL directly takes a permutation iterator (and // directly scatters from the internal buffer) - thrust::scatter(rmm::exec_policy(handle.get_stream()), + thrust::scatter(handle.get_thrust_policy(), rx_value_first, rx_value_first + rx_counts[i], map_first, @@ -414,7 +414,7 @@ void copy_to_matrix_minor(raft::handle_t const& handle, assert(graph_view.get_number_of_local_vertices() == graph_view.get_number_of_local_adj_matrix_partition_rows()); auto val_first = thrust::make_permutation_iterator(vertex_value_input_first, vertex_first); - thrust::scatter(rmm::exec_policy(handle.get_stream()), + thrust::scatter(handle.get_thrust_policy(), val_first, val_first + thrust::distance(vertex_first, vertex_last), vertex_first, diff --git a/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh b/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh index 335b34828e5..117e7525c25 100644 --- a/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh +++ b/cpp/include/cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh @@ -438,13 +438,14 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, minor_init = (row_comm_rank == 0) ? init : T{}; } + auto execution_policy = handle.get_thrust_policy(); if (GraphViewType::is_multi_gpu) { - thrust::fill(rmm::exec_policy(handle.get_stream()), + thrust::fill(execution_policy, minor_buffer_first, minor_buffer_first + minor_tmp_buffer_size, minor_init); } else { - thrust::fill(rmm::exec_policy(handle.get_stream()), + thrust::fill(execution_policy, vertex_value_output_first, vertex_value_output_first + graph_view.get_number_of_local_vertices(), minor_init); @@ -546,7 +547,7 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, if constexpr (update_major) { // this is necessary as we don't visit every vertex in the // hypersparse segment in // for_all_major_for_all_nbr_hypersparse - thrust::fill(rmm::exec_policy(handle.get_stream()), + thrust::fill(handle.get_thrust_policy(), output_buffer_first + (*segment_offsets)[3], output_buffer_first + (*segment_offsets)[4], major_init); diff --git a/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh b/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh index 5ae32a6f56a..f7f9dae9dd7 100644 --- a/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh +++ b/cpp/include/cugraph/prims/copy_v_transform_reduce_key_aggregated_out_nbr.cuh @@ -116,6 +116,7 @@ void decompress_matrix_partition_to_fill_edgelist_majors( vertex_t* majors, std::optional> const& segment_offsets) { + auto execution_policy = handle.get_thrust_policy(); if (segment_offsets) { // FIXME: we may further improve performance by 1) concurrently running kernels on different // segments; 2) individually tuning block sizes for different segments; and 3) adding one more @@ -153,7 +154,7 @@ void decompress_matrix_partition_to_fill_edgelist_majors( } if ((*segment_offsets)[3] - (*segment_offsets)[2] > 0) { thrust::for_each( - rmm::exec_policy(handle.get_stream()), + execution_policy, thrust::make_counting_iterator(matrix_partition.get_major_first()) + (*segment_offsets)[2], thrust::make_counting_iterator(matrix_partition.get_major_first()) + (*segment_offsets)[3], [matrix_partition, majors] __device__(auto major) { @@ -167,7 +168,7 @@ void decompress_matrix_partition_to_fill_edgelist_majors( if (matrix_partition.get_dcs_nzd_vertex_count() && (*(matrix_partition.get_dcs_nzd_vertex_count()) > 0)) { thrust::for_each( - rmm::exec_policy(handle.get_stream()), + execution_policy, thrust::make_counting_iterator(vertex_t{0}), thrust::make_counting_iterator(*(matrix_partition.get_dcs_nzd_vertex_count())), [matrix_partition, major_start_offset = (*segment_offsets)[3], majors] __device__( @@ -183,7 +184,7 @@ void decompress_matrix_partition_to_fill_edgelist_majors( } } else { thrust::for_each( - rmm::exec_policy(handle.get_stream()), + execution_policy, thrust::make_counting_iterator(matrix_partition.get_major_first()), thrust::make_counting_iterator(matrix_partition.get_major_first()) + matrix_partition.get_major_size(), @@ -340,12 +341,13 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( } // FIXME: these copies are unnecessary, better fix RAFT comm's bcast to take separate input & // output pointers - thrust::copy(rmm::exec_policy(handle.get_stream()), + auto execution_policy = handle.get_thrust_policy(); + thrust::copy(execution_policy, map_key_first, map_key_last, map_keys.begin() + map_displacements[row_comm_rank]); thrust::copy( - rmm::exec_policy(handle.get_stream()), + execution_policy, map_value_first, map_value_first + thrust::distance(map_key_first, map_key_last), get_dataframe_buffer_begin(map_value_buffer) + map_displacements[row_comm_rank]); @@ -420,12 +422,13 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( matrix_partition.get_indices(), detail::minor_to_key_t{adj_matrix_col_key_first, matrix_partition.get_minor_first()}); - thrust::copy(rmm::exec_policy(handle.get_stream()), + auto execution_policy = handle.get_thrust_policy(); + thrust::copy(execution_policy, minor_key_first, minor_key_first + matrix_partition.get_number_of_edges(), tmp_minor_keys.begin()); if (graph_view.is_weighted()) { - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(execution_policy, *(matrix_partition.get_weights()), *(matrix_partition.get_weights()) + matrix_partition.get_number_of_edges(), tmp_key_aggregated_edge_weights.begin()); @@ -448,25 +451,24 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( auto output_key_first = thrust::make_zip_iterator( thrust::make_tuple(reduced_major_vertices.begin(), reduced_minor_keys.begin())); if (graph_view.is_weighted()) { - thrust::sort_by_key(rmm::exec_policy(handle.get_stream()), + thrust::sort_by_key(execution_policy, input_key_first, input_key_first + tmp_major_vertices.size(), tmp_key_aggregated_edge_weights.begin()); reduced_size = thrust::distance( output_key_first, - thrust::get<0>(thrust::reduce_by_key(rmm::exec_policy(handle.get_stream()), + thrust::get<0>(thrust::reduce_by_key(execution_policy, input_key_first, input_key_first + tmp_major_vertices.size(), tmp_key_aggregated_edge_weights.begin(), output_key_first, reduced_key_aggregated_edge_weights.begin()))); } else { - thrust::sort(rmm::exec_policy(handle.get_stream()), - input_key_first, - input_key_first + tmp_major_vertices.size()); + thrust::sort( + execution_policy, input_key_first, input_key_first + tmp_major_vertices.size()); reduced_size = thrust::distance( output_key_first, - thrust::get<0>(thrust::reduce_by_key(rmm::exec_policy(handle.get_stream()), + thrust::get<0>(thrust::reduce_by_key(execution_policy, input_key_first, input_key_first + tmp_major_vertices.size(), thrust::make_constant_iterator(weight_t{1.0}), @@ -515,14 +517,15 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( auto pair_first = thrust::make_zip_iterator( thrust::make_tuple(rx_major_vertices.begin(), rx_minor_keys.begin())); - thrust::sort_by_key(rmm::exec_policy(handle.get_stream()), + auto execution_policy = handle.get_thrust_policy(); + thrust::sort_by_key(execution_policy, pair_first, pair_first + rx_major_vertices.size(), rx_key_aggregated_edge_weights.begin()); tmp_major_vertices.resize(rx_major_vertices.size(), handle.get_stream()); tmp_minor_keys.resize(tmp_major_vertices.size(), handle.get_stream()); tmp_key_aggregated_edge_weights.resize(tmp_major_vertices.size(), handle.get_stream()); - auto pair_it = thrust::reduce_by_key(rmm::exec_policy(handle.get_stream()), + auto pair_it = thrust::reduce_by_key(execution_policy, pair_first, pair_first + rx_major_vertices.size(), rx_key_aggregated_edge_weights.begin(), @@ -546,7 +549,7 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( auto triplet_first = thrust::make_zip_iterator(thrust::make_tuple( tmp_major_vertices.begin(), tmp_minor_keys.begin(), tmp_key_aggregated_edge_weights.begin())); thrust::transform( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), triplet_first, triplet_first + tmp_major_vertices.size(), tmp_e_op_result_buffer_first, @@ -632,17 +635,18 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( #endif } - thrust::fill(rmm::exec_policy(handle.get_stream()), + auto execution_policy = handle.get_thrust_policy(); + thrust::fill(execution_policy, vertex_value_output_first, vertex_value_output_first + graph_view.get_number_of_local_vertices(), T{}); - thrust::sort_by_key(rmm::exec_policy(handle.get_stream()), + thrust::sort_by_key(execution_policy, major_vertices.begin(), major_vertices.end(), get_dataframe_buffer_begin(e_op_result_buffer)); auto num_uniques = thrust::count_if( - rmm::exec_policy(handle.get_stream()), + execution_policy, thrust::make_counting_iterator(size_t{0}), thrust::make_counting_iterator(major_vertices.size()), [major_vertices = major_vertices.data()] __device__(auto i) { @@ -658,13 +662,13 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( : invalid_vertex_id::value; }); thrust::copy_if( - rmm::exec_policy(handle.get_stream()), + execution_policy, major_vertex_first, major_vertex_first + major_vertices.size(), unique_major_vertices.begin(), [] __device__(auto major) { return major != invalid_vertex_id::value; }); thrust::reduce_by_key( - rmm::exec_policy(handle.get_stream()), + execution_policy, major_vertices.begin(), major_vertices.end(), get_dataframe_buffer_begin(e_op_result_buffer), @@ -680,7 +684,7 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( thrust::equal_to{}, reduce_op); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(execution_policy, vertex_value_output_first, vertex_value_output_first + graph_view.get_number_of_local_vertices(), vertex_value_output_first, diff --git a/cpp/include/cugraph/prims/count_if_v.cuh b/cpp/include/cugraph/prims/count_if_v.cuh index 5a7684d19aa..b2d4283d859 100644 --- a/cpp/include/cugraph/prims/count_if_v.cuh +++ b/cpp/include/cugraph/prims/count_if_v.cuh @@ -54,7 +54,7 @@ typename GraphViewType::vertex_type count_if_v(raft::handle_t const& handle, VertexOp v_op) { auto count = - thrust::count_if(rmm::exec_policy(handle.get_stream()), + thrust::count_if(handle.get_thrust_policy(), vertex_value_input_first, vertex_value_input_first + graph_view.get_number_of_local_vertices(), v_op); @@ -92,8 +92,7 @@ typename GraphViewType::vertex_type count_if_v(raft::handle_t const& handle, InputIterator input_last, VertexOp v_op) { - auto count = - thrust::count_if(rmm::exec_policy(handle.get_stream()), input_first, input_last, v_op); + auto count = thrust::count_if(handle.get_thrust_policy(), input_first, input_last, v_op); if (GraphViewType::is_multi_gpu) { count = host_scalar_allreduce(handle.get_comms(), count, handle.get_stream()); } diff --git a/cpp/include/cugraph/prims/reduce_v.cuh b/cpp/include/cugraph/prims/reduce_v.cuh index ef737a153df..f41774675fb 100644 --- a/cpp/include/cugraph/prims/reduce_v.cuh +++ b/cpp/include/cugraph/prims/reduce_v.cuh @@ -52,7 +52,7 @@ T reduce_v(raft::handle_t const& handle, T init) { auto ret = thrust::reduce( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), vertex_value_input_first, vertex_value_input_first + graph_view.get_number_of_local_vertices(), ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() == 0)) ? init : T{}, @@ -89,7 +89,7 @@ T reduce_v(raft::handle_t const& handle, T init) { auto ret = thrust::reduce( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), input_first, input_last, ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() == 0)) ? init : T{}, diff --git a/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh b/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh index f8583d71f5c..70a9afa32c0 100644 --- a/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_by_adj_matrix_row_col_key_e.cuh @@ -530,11 +530,9 @@ transform_reduce_by_adj_matrix_row_col_key_e( keys.resize(cur_size + tmp_keys.size(), handle.get_stream()); resize_dataframe_buffer(value_buffer, keys.size(), handle.get_stream()); - thrust::copy(rmm::exec_policy(handle.get_stream()), - tmp_keys.begin(), - tmp_keys.end(), - keys.begin() + cur_size); - thrust::copy(rmm::exec_policy(handle.get_stream()), + auto execution_policy = handle.get_thrust_policy(); + thrust::copy(execution_policy, tmp_keys.begin(), tmp_keys.end(), keys.begin() + cur_size); + thrust::copy(execution_policy, get_dataframe_buffer_begin(tmp_value_buffer), get_dataframe_buffer_begin(tmp_value_buffer) + tmp_keys.size(), get_dataframe_buffer_begin(value_buffer) + cur_size); diff --git a/cpp/include/cugraph/prims/transform_reduce_e.cuh b/cpp/include/cugraph/prims/transform_reduce_e.cuh index f46a00d37e4..000800a9862 100644 --- a/cpp/include/cugraph/prims/transform_reduce_e.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_e.cuh @@ -406,7 +406,7 @@ T transform_reduce_e(raft::handle_t const& handle, property_add edge_property_add{}; auto result_buffer = allocate_dataframe_buffer(1, handle.get_stream()); - thrust::fill(rmm::exec_policy(handle.get_stream()), + thrust::fill(handle.get_thrust_policy(), get_dataframe_buffer_begin(result_buffer), get_dataframe_buffer_begin(result_buffer) + 1, T{}); @@ -503,7 +503,7 @@ T transform_reduce_e(raft::handle_t const& handle, } } - auto result = thrust::reduce(rmm::exec_policy(handle.get_stream()), + auto result = thrust::reduce(handle.get_thrust_policy(), get_dataframe_buffer_begin(result_buffer), get_dataframe_buffer_begin(result_buffer) + 1, T{}, diff --git a/cpp/include/cugraph/prims/transform_reduce_v.cuh b/cpp/include/cugraph/prims/transform_reduce_v.cuh index 696d004e89b..118db15b38a 100644 --- a/cpp/include/cugraph/prims/transform_reduce_v.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_v.cuh @@ -56,7 +56,7 @@ T transform_reduce_v(raft::handle_t const& handle, T init) { auto ret = thrust::transform_reduce( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), vertex_value_input_first, vertex_value_input_first + graph_view.get_number_of_local_vertices(), v_op, @@ -99,7 +99,7 @@ T transform_reduce_v(raft::handle_t const& handle, T init) { auto ret = thrust::transform_reduce( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), input_first, input_last, v_op, diff --git a/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh b/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh index 1d04dd7fa87..ffa15663376 100644 --- a/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh +++ b/cpp/include/cugraph/prims/update_frontier_v_push_if_out_nbr.cuh @@ -527,12 +527,12 @@ size_t sort_and_reduce_buffer_elements(raft::handle_t const& handle, using payload_t = typename optional_payload_buffer_value_type_t::value; + auto execution_policy = handle.get_thrust_policy(); if constexpr (std::is_same_v) { - thrust::sort(rmm::exec_policy(handle.get_stream()), - buffer_key_output_first, - buffer_key_output_first + num_buffer_elements); + thrust::sort( + execution_policy, buffer_key_output_first, buffer_key_output_first + num_buffer_elements); } else { - thrust::sort_by_key(rmm::exec_policy(handle.get_stream()), + thrust::sort_by_key(execution_policy, buffer_key_output_first, buffer_key_output_first + num_buffer_elements, buffer_payload_output_first); @@ -540,15 +540,14 @@ size_t sort_and_reduce_buffer_elements(raft::handle_t const& handle, size_t num_reduced_buffer_elements{}; if constexpr (std::is_same_v) { - auto it = thrust::unique(rmm::exec_policy(handle.get_stream()), - buffer_key_output_first, - buffer_key_output_first + num_buffer_elements); + auto it = thrust::unique( + execution_policy, buffer_key_output_first, buffer_key_output_first + num_buffer_elements); num_reduced_buffer_elements = static_cast(thrust::distance(buffer_key_output_first, it)); } else if constexpr (std::is_same>::value) { // FIXME: if ReducOp is any, we may have a cheaper alternative than sort & uique (i.e. discard // non-first elements) - auto it = thrust::unique_by_key(rmm::exec_policy(handle.get_stream()), + auto it = thrust::unique_by_key(execution_policy, buffer_key_output_first, buffer_key_output_first + num_buffer_elements, buffer_payload_output_first); @@ -567,7 +566,7 @@ size_t sort_and_reduce_buffer_elements(raft::handle_t const& handle, rmm::device_uvector keys(num_buffer_elements, handle.get_stream()); auto value_buffer = allocate_dataframe_buffer(num_buffer_elements, handle.get_stream()); - auto it = thrust::reduce_by_key(rmm::exec_policy(handle.get_stream()), + auto it = thrust::reduce_by_key(execution_policy, buffer_key_output_first, buffer_key_output_first + num_buffer_elements, buffer_payload_output_first, @@ -578,11 +577,11 @@ size_t sort_and_reduce_buffer_elements(raft::handle_t const& handle, num_reduced_buffer_elements = static_cast(thrust::distance(keys.begin(), thrust::get<0>(it))); // FIXME: this copy can be replaced by move - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(execution_policy, keys.begin(), keys.begin() + num_reduced_buffer_elements, buffer_key_output_first); - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(execution_policy, get_dataframe_buffer_begin(value_buffer), get_dataframe_buffer_begin(value_buffer) + num_reduced_buffer_elements, buffer_payload_output_first); @@ -648,6 +647,7 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( matrix_partition_device_view_t( graph_view.get_matrix_partition_view(i)); + auto execution_policy = handle.get_thrust_policy(); if (GraphViewType::is_multi_gpu) { auto& col_comm = handle.get_subcomm(cugraph::partition_2d::key_naming_t().col_name()); auto const col_comm_rank = col_comm.get_rank(); @@ -657,7 +657,7 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( // FIXME: this copy is unnecessary, better fix RAFT comm's bcast to take const iterators for // input if (col_comm_rank == static_cast(i)) { - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(execution_policy, local_frontier_vertex_first, local_frontier_vertex_last, frontier_vertices.begin()); @@ -678,7 +678,7 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( ret += use_dcs ? thrust::transform_reduce( - rmm::exec_policy(handle.get_stream()), + execution_policy, frontier_vertices.begin(), frontier_vertices.end(), [matrix_partition, @@ -703,7 +703,7 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( edge_t{0}, thrust::plus()) : thrust::transform_reduce( - rmm::exec_policy(handle.get_stream()), + execution_policy, frontier_vertices.begin(), frontier_vertices.end(), [matrix_partition] __device__(auto major) { @@ -715,7 +715,7 @@ typename GraphViewType::edge_type compute_num_out_nbrs_from_frontier( } else { assert(i == 0); ret += thrust::transform_reduce( - rmm::exec_policy(handle.get_stream()), + execution_policy, local_frontier_vertex_first, local_frontier_vertex_last, [matrix_partition] __device__(auto major) { @@ -894,7 +894,7 @@ void update_frontier_v_push_if_out_nbr( matrix_partition_frontier_key_buffer, matrix_partition_frontier_size, handle.get_stream()); if (static_cast(col_comm_rank) == i) { - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(handle.get_thrust_policy(), frontier_key_first, frontier_key_last, get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer)); @@ -909,7 +909,7 @@ void update_frontier_v_push_if_out_nbr( } else { resize_dataframe_buffer( matrix_partition_frontier_key_buffer, matrix_partition_frontier_size, handle.get_stream()); - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(handle.get_thrust_policy(), frontier_key_first, frontier_key_last, get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer)); @@ -936,9 +936,10 @@ void update_frontier_v_push_if_out_nbr( ? ((*segment_offsets).size() > (detail::num_sparse_segments_per_vertex_partition + 1)) : false; + auto execution_policy = handle.get_thrust_policy(); auto max_pushes = use_dcs ? thrust::transform_reduce( - rmm::exec_policy(handle.get_stream()), + execution_policy, matrix_partition_frontier_row_first, matrix_partition_frontier_row_last, [matrix_partition, @@ -963,7 +964,7 @@ void update_frontier_v_push_if_out_nbr( edge_t{0}, thrust::plus()) : thrust::transform_reduce( - rmm::exec_policy(handle.get_stream()), + execution_policy, matrix_partition_frontier_row_first, matrix_partition_frontier_row_last, [matrix_partition] __device__(auto row) { @@ -1007,7 +1008,7 @@ void update_frontier_v_push_if_out_nbr( raft::update_device( d_thresholds.data(), h_thresholds.data(), h_thresholds.size(), handle.get_stream()); rmm::device_uvector d_offsets(d_thresholds.size(), handle.get_stream()); - thrust::lower_bound(rmm::exec_policy(handle.get_stream()), + thrust::lower_bound(handle.get_thrust_policy(), matrix_partition_frontier_row_first, matrix_partition_frontier_row_last, d_thresholds.begin(), @@ -1170,7 +1171,7 @@ void update_frontier_v_push_if_out_nbr( row_first = thrust::get<0>(get_dataframe_buffer_begin(key_buffer).get_iterator_tuple()); } - thrust::lower_bound(rmm::exec_policy(handle.get_stream()), + thrust::lower_bound(handle.get_thrust_policy(), row_first, row_first + num_buffer_elements, d_vertex_lasts.begin(), @@ -1234,7 +1235,7 @@ void update_frontier_v_push_if_out_nbr( thrust::make_tuple(get_dataframe_buffer_begin(key_buffer), detail::get_optional_payload_buffer_begin(payload_buffer))); thrust::transform( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), key_payload_pair_first, key_payload_pair_first + num_buffer_elements, bucket_indices.begin(), @@ -1266,7 +1267,7 @@ void update_frontier_v_push_if_out_nbr( shrink_to_fit_dataframe_buffer(payload_buffer, handle.get_stream()); } else { thrust::transform( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), get_dataframe_buffer_begin(key_buffer), get_dataframe_buffer_begin(key_buffer) + num_buffer_elements, bucket_indices.begin(), @@ -1286,7 +1287,7 @@ void update_frontier_v_push_if_out_nbr( thrust::make_tuple(bucket_indices.begin(), get_dataframe_buffer_begin(key_buffer))); bucket_indices.resize( thrust::distance(bucket_key_pair_first, - thrust::remove_if(rmm::exec_policy(handle.get_stream()), + thrust::remove_if(handle.get_thrust_policy(), bucket_key_pair_first, bucket_key_pair_first + num_buffer_elements, detail::check_invalid_bucket_idx_t())), diff --git a/cpp/include/cugraph/prims/vertex_frontier.cuh b/cpp/include/cugraph/prims/vertex_frontier.cuh index c66444e4a77..5f5a3225bdc 100644 --- a/cpp/include/cugraph/prims/vertex_frontier.cuh +++ b/cpp/include/cugraph/prims/vertex_frontier.cuh @@ -96,7 +96,7 @@ class SortedUniqueKeyBucket { tags_.resize(1, handle_ptr_->get_stream()); auto pair_first = thrust::make_tuple(thrust::make_zip_iterator(vertices_.begin(), tags_.begin())); - thrust::fill(rmm::exec_policy(handle_ptr_->get_stream()), pair_first, pair_first + 1, key); + thrust::fill(handle_ptr_->get_thrust_policy(), pair_first, pair_first + 1, key); } } @@ -119,24 +119,22 @@ class SortedUniqueKeyBucket { if (vertices_.size() > 0) { rmm::device_uvector merged_vertices( vertices_.size() + thrust::distance(vertex_first, vertex_last), handle_ptr_->get_stream()); - thrust::merge(rmm::exec_policy(handle_ptr_->get_stream()), + thrust::merge(handle_ptr_->get_thrust_policy(), vertices_.begin(), vertices_.end(), vertex_first, vertex_last, merged_vertices.begin()); - merged_vertices.resize( - thrust::distance(merged_vertices.begin(), - thrust::unique(rmm::exec_policy(handle_ptr_->get_stream()), - merged_vertices.begin(), - merged_vertices.end())), - handle_ptr_->get_stream()); + merged_vertices.resize(thrust::distance(merged_vertices.begin(), + thrust::unique(handle_ptr_->get_thrust_policy(), + merged_vertices.begin(), + merged_vertices.end())), + handle_ptr_->get_stream()); merged_vertices.shrink_to_fit(handle_ptr_->get_stream()); vertices_ = std::move(merged_vertices); } else { vertices_.resize(thrust::distance(vertex_first, vertex_last), handle_ptr_->get_stream()); - thrust::copy( - rmm::exec_policy(handle_ptr_->get_stream()), vertex_first, vertex_last, vertices_.begin()); + thrust::copy(handle_ptr_->get_thrust_policy(), vertex_first, vertex_last, vertices_.begin()); } } @@ -164,7 +162,7 @@ class SortedUniqueKeyBucket { thrust::make_zip_iterator(thrust::make_tuple(vertices_.begin(), tags_.begin())); auto merged_pair_first = thrust::make_zip_iterator(thrust::make_tuple(merged_vertices.begin(), merged_tags.begin())); - thrust::merge(rmm::exec_policy(handle_ptr_->get_stream()), + thrust::merge(handle_ptr_->get_thrust_policy(), old_pair_first, old_pair_first + vertices_.size(), key_first, @@ -172,7 +170,7 @@ class SortedUniqueKeyBucket { merged_pair_first); merged_vertices.resize( thrust::distance(merged_pair_first, - thrust::unique(rmm::exec_policy(handle_ptr_->get_stream()), + thrust::unique(handle_ptr_->get_thrust_policy(), merged_pair_first, merged_pair_first + merged_vertices.size())), handle_ptr_->get_stream()); @@ -184,7 +182,7 @@ class SortedUniqueKeyBucket { } else { vertices_.resize(thrust::distance(key_first, key_last), handle_ptr_->get_stream()); tags_.resize(thrust::distance(key_first, key_last), handle_ptr_->get_stream()); - thrust::copy(rmm::exec_policy(handle_ptr_->get_stream()), + thrust::copy(handle_ptr_->get_thrust_policy(), key_first, key_last, thrust::make_zip_iterator(thrust::make_tuple(vertices_.begin(), tags_.begin()))); @@ -325,7 +323,7 @@ class VertexFrontier { static_assert(kNumBuckets <= std::numeric_limits::max()); rmm::device_uvector bucket_indices(this_bucket.size(), handle_ptr_->get_stream()); thrust::transform( - rmm::exec_policy(handle_ptr_->get_stream()), + handle_ptr_->get_thrust_policy(), this_bucket.begin(), this_bucket.end(), bucket_indices.begin(), @@ -340,7 +338,7 @@ class VertexFrontier { thrust::make_zip_iterator(thrust::make_tuple(bucket_indices.begin(), this_bucket.begin())); bucket_indices.resize( thrust::distance(pair_first, - thrust::remove_if(rmm::exec_policy(handle_ptr_->get_stream()), + thrust::remove_if(handle_ptr_->get_thrust_policy(), pair_first, pair_first + bucket_indices.size(), [] __device__(auto pair) { @@ -359,7 +357,7 @@ class VertexFrontier { auto new_this_bucket_size = static_cast(thrust::distance( pair_first, thrust::stable_partition( // stalbe_partition to maintain sorted order within each bucket - rmm::exec_policy(handle_ptr_->get_stream()), + handle_ptr_->get_thrust_policy(), pair_first, pair_first + bucket_indices.size(), [this_bucket_idx = static_cast(this_bucket_idx)] __device__(auto pair) { @@ -399,7 +397,7 @@ class VertexFrontier { auto next_bucket_size = static_cast(thrust::distance( pair_first, thrust::stable_partition( // stalbe_partition to maintain sorted order within each bucket - rmm::exec_policy(handle_ptr_->get_stream()), + handle_ptr_->get_thrust_policy(), pair_first, pair_last, [next_bucket_idx = static_cast(to_bucket_indices[0])] __device__(auto pair) { @@ -412,13 +410,13 @@ class VertexFrontier { static_cast(thrust::distance(pair_first + next_bucket_size, pair_last))}; } else { thrust::stable_sort( // stalbe_sort to maintain sorted order within each bucket - rmm::exec_policy(handle_ptr_->get_stream()), + handle_ptr_->get_thrust_policy(), pair_first, pair_last, [] __device__(auto lhs, auto rhs) { return thrust::get<0>(lhs) < thrust::get<0>(rhs); }); rmm::device_uvector d_indices(to_bucket_indices.size(), handle_ptr_->get_stream()); rmm::device_uvector d_counts(d_indices.size(), handle_ptr_->get_stream()); - auto it = thrust::reduce_by_key(rmm::exec_policy(handle_ptr_->get_stream()), + auto it = thrust::reduce_by_key(handle_ptr_->get_thrust_policy(), bucket_idx_first, bucket_idx_last, thrust::make_constant_iterator(size_t{1}), diff --git a/cpp/include/cugraph/utilities/host_barrier.hpp b/cpp/include/cugraph/utilities/host_barrier.hpp index aeafa5b28db..6825814eb93 100644 --- a/cpp/include/cugraph/utilities/host_barrier.hpp +++ b/cpp/include/cugraph/utilities/host_barrier.hpp @@ -16,7 +16,6 @@ #pragma once #include -#include namespace cugraph { diff --git a/cpp/src/centrality/betweenness_centrality.cu b/cpp/src/centrality/betweenness_centrality.cu index 70b1c87fbe3..927850769cb 100644 --- a/cpp/src/centrality/betweenness_centrality.cu +++ b/cpp/src/centrality/betweenness_centrality.cu @@ -26,7 +26,6 @@ #include #include -#include #include #include "betweenness_centrality.cuh" @@ -230,13 +229,13 @@ void BC::compute_single_source(vertex_t so // the traversal, this value is avalaible within the bfs implementation and // there could be a way to access it directly and avoid both replace and the // max - thrust::replace(rmm::exec_policy(handle_.get_stream_view()), + thrust::replace(handle_.get_thrust_policy(), distances_, distances_ + number_of_vertices_, std::numeric_limits::max(), static_cast(-1)); - auto current_max_depth = thrust::max_element( - rmm::exec_policy(handle_.get_stream_view()), distances_, distances_ + number_of_vertices_); + auto current_max_depth = + thrust::max_element(handle_.get_thrust_policy(), distances_, distances_ + number_of_vertices_); vertex_t max_depth = 0; CUDA_TRY(cudaMemcpy(&max_depth, current_max_depth, sizeof(vertex_t), cudaMemcpyDeviceToHost)); // Step 2) Dependency accumulation @@ -266,10 +265,8 @@ void BC::accumulate(vertex_t source_vertex template void BC::initialize_dependencies() { - thrust::fill(rmm::exec_policy(handle_.get_stream_view()), - deltas_, - deltas_ + number_of_vertices_, - static_cast(0)); + thrust::fill( + handle_.get_thrust_policy(), deltas_, deltas_ + number_of_vertices_, static_cast(0)); } template void BC::accumulate_edges(vertex_t max_depth, @@ -316,13 +313,13 @@ template ::add_reached_endpoints_to_source_betweenness( vertex_t source_vertex) { - vertex_t number_of_unvisited_vertices = thrust::count( - rmm::exec_policy(handle_.get_stream_view()), distances_, distances_ + number_of_vertices_, -1); + vertex_t number_of_unvisited_vertices = + thrust::count(handle_.get_thrust_policy(), distances_, distances_ + number_of_vertices_, -1); vertex_t number_of_visited_vertices_except_source = number_of_vertices_ - number_of_unvisited_vertices - 1; rmm::device_vector buffer(1); buffer[0] = number_of_visited_vertices_except_source; - thrust::transform(rmm::exec_policy(handle_.get_stream_view()), + thrust::transform(handle_.get_thrust_policy(), buffer.begin(), buffer.end(), betweenness_ + source_vertex, @@ -333,7 +330,7 @@ void BC::add_reached_endpoints_to_source_b template void BC::add_vertices_dependencies_to_betweenness() { - thrust::transform(rmm::exec_policy(handle_.get_stream_view()), + thrust::transform(handle_.get_thrust_policy(), deltas_, deltas_ + number_of_vertices_, betweenness_, @@ -418,7 +415,7 @@ void BC::apply_rescale_factor_to_betweenne { size_t result_size = number_of_vertices_; if (is_edge_betweenness_) result_size = number_of_edges_; - thrust::transform(rmm::exec_policy(handle_.get_stream_view()), + thrust::transform(handle_.get_thrust_policy(), betweenness_, betweenness_ + result_size, thrust::make_constant_iterator(rescale_factor), diff --git a/cpp/src/centrality/betweenness_centrality.cuh b/cpp/src/centrality/betweenness_centrality.cuh index fe8093367cb..e808e9450b4 100644 --- a/cpp/src/centrality/betweenness_centrality.cuh +++ b/cpp/src/centrality/betweenness_centrality.cuh @@ -18,7 +18,6 @@ #pragma once #include -#include namespace cugraph { namespace detail { diff --git a/cpp/src/centrality/katz_centrality.cu b/cpp/src/centrality/katz_centrality.cu index a638694153b..0dc50c08374 100644 --- a/cpp/src/centrality/katz_centrality.cu +++ b/cpp/src/centrality/katz_centrality.cu @@ -80,7 +80,7 @@ void katz_centrality(raft::handle_t const& handle, // 2. initialize katz centrality values if (!has_initial_guess) { - thrust::fill(rmm::exec_policy(handle.get_stream()), + thrust::fill(handle.get_thrust_policy(), katz_centralities, katz_centralities + pull_graph_view.get_number_of_local_vertices(), result_t{0.0}); @@ -115,7 +115,7 @@ void katz_centrality(raft::handle_t const& handle, if (betas != nullptr) { auto val_first = thrust::make_zip_iterator(thrust::make_tuple(new_katz_centralities, betas)); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), val_first, val_first + pull_graph_view.get_number_of_local_vertices(), new_katz_centralities, @@ -143,7 +143,7 @@ void katz_centrality(raft::handle_t const& handle, } if (new_katz_centralities != katz_centralities) { - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(handle.get_thrust_policy(), new_katz_centralities, new_katz_centralities + pull_graph_view.get_number_of_local_vertices(), katz_centralities); @@ -159,7 +159,7 @@ void katz_centrality(raft::handle_t const& handle, l2_norm = std::sqrt(l2_norm); CUGRAPH_EXPECTS(l2_norm > 0.0, "L2 norm of the computed Katz Centrality values should be positive."); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), katz_centralities, katz_centralities + pull_graph_view.get_number_of_local_vertices(), katz_centralities, diff --git a/cpp/src/community/flatten_dendrogram.cuh b/cpp/src/community/flatten_dendrogram.cuh index 37c3c7278d7..8f3656d6533 100644 --- a/cpp/src/community/flatten_dendrogram.cuh +++ b/cpp/src/community/flatten_dendrogram.cuh @@ -40,7 +40,7 @@ void partition_at_level(raft::handle_t const& handle, thrust::make_counting_iterator(level), [&handle, &dendrogram, &local_vertex_ids_v, d_vertex_ids, &d_partition, local_num_verts]( size_t l) { - thrust::sequence(rmm::exec_policy(handle.get_stream()), + thrust::sequence(handle.get_thrust_policy(), local_vertex_ids_v.begin(), local_vertex_ids_v.begin() + dendrogram.get_level_size_nocheck(l), dendrogram.get_level_first_index_nocheck(l)); diff --git a/cpp/src/community/legacy/ecg.cu b/cpp/src/community/legacy/ecg.cu index bdaa02871b1..30af37ec2e5 100644 --- a/cpp/src/community/legacy/ecg.cu +++ b/cpp/src/community/legacy/ecg.cu @@ -149,7 +149,7 @@ void ecg(raft::handle_t const& handle, rmm::device_uvector ecg_weights_v(graph.number_of_edges, handle.get_stream_view()); - thrust::copy(rmm::exec_policy(handle.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), graph.edge_data, graph.edge_data + graph.number_of_edges, ecg_weights_v.data()); @@ -182,7 +182,7 @@ void ecg(raft::handle_t const& handle, // Set weights = min_weight + (1 - min-weight)*sum/ensemble_size update_functor uf(min_weight, ensemble_size); - thrust::transform(rmm::exec_policy(handle.get_stream_view()), + thrust::transform(handle.get_thrust_policy(), ecg_weights_v.begin(), ecg_weights_v.end(), ecg_weights_v.begin(), diff --git a/cpp/src/community/legacy/leiden.cu b/cpp/src/community/legacy/leiden.cu index 7044004d8ed..0cc7a991734 100644 --- a/cpp/src/community/legacy/leiden.cu +++ b/cpp/src/community/legacy/leiden.cu @@ -39,7 +39,7 @@ std::pair leiden(raft::handle_t const& handle, rmm::device_uvector vertex_ids_v(graph.number_of_vertices, handle.get_stream()); - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(handle.get_thrust_policy(), thrust::make_counting_iterator(0), // MNMG - base vertex id thrust::make_counting_iterator( graph.number_of_vertices), // MNMG - base vertex id + number_of_vertices diff --git a/cpp/src/community/legacy/leiden.cuh b/cpp/src/community/legacy/leiden.cuh index c4edfa12193..36778d9ab37 100644 --- a/cpp/src/community/legacy/leiden.cuh +++ b/cpp/src/community/legacy/leiden.cuh @@ -59,7 +59,7 @@ class Leiden : public Louvain { weight_t* d_delta_Q = delta_Q_v.data(); vertex_t* d_constraint = constraint_v_.data(); - thrust::copy(rmm::exec_policy(this->handle_.get_stream_view()), + thrust::copy(this->handle_.get_thrust_policy(), this->dendrogram_->current_level_begin(), this->dendrogram_->current_level_end(), next_cluster_v.data()); @@ -82,7 +82,7 @@ class Leiden : public Louvain { // Filter out positive delta_Q values for nodes not in the same constraint group thrust::for_each( - rmm::exec_policy(this->handle_.get_stream_view()), + this->handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(graph.number_of_edges), [d_src_indices, d_dst_indices, d_constraint, d_delta_Q] __device__(vertex_t i) { @@ -98,7 +98,7 @@ class Leiden : public Louvain { new_Q = this->modularity(total_edge_weight, resolution, graph, next_cluster_v.data()); if (new_Q > cur_Q) { - thrust::copy(rmm::exec_policy(this->handle_.get_stream_view()), + thrust::copy(this->handle_.get_thrust_policy(), next_cluster_v.begin(), next_cluster_v.end(), this->dendrogram_->current_level_begin()); @@ -113,9 +113,8 @@ class Leiden : public Louvain { { size_t num_level{0}; - weight_t total_edge_weight = thrust::reduce(rmm::exec_policy(this->handle_.get_stream_view()), - this->weights_v_.begin(), - this->weights_v_.end()); + weight_t total_edge_weight = thrust::reduce( + this->handle_.get_thrust_policy(), this->weights_v_.begin(), this->weights_v_.end()); weight_t best_modularity = weight_t{-1}; @@ -138,7 +137,7 @@ class Leiden : public Louvain { this->dendrogram_->add_level( 0, current_graph.number_of_vertices, this->handle_.get_stream_view()); - thrust::sequence(rmm::exec_policy(this->handle_.get_stream_view()), + thrust::sequence(this->handle_.get_thrust_policy(), this->dendrogram_->current_level_begin(), this->dendrogram_->current_level_end()); diff --git a/cpp/src/community/legacy/louvain.cuh b/cpp/src/community/legacy/louvain.cuh index 0c14552aecc..c7292c2590a 100644 --- a/cpp/src/community/legacy/louvain.cuh +++ b/cpp/src/community/legacy/louvain.cuh @@ -65,17 +65,17 @@ class Louvain { number_of_vertices_(graph.number_of_vertices), number_of_edges_(graph.number_of_edges) { - thrust::copy(rmm::exec_policy(handle_.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), graph.offsets, graph.offsets + graph.number_of_vertices + 1, offsets_v_.begin()); - thrust::copy(rmm::exec_policy(handle_.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), graph.indices, graph.indices + graph.number_of_edges, indices_v_.begin()); - thrust::copy(rmm::exec_policy(handle_.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), graph.edge_data, graph.edge_data + graph.number_of_edges, weights_v_.begin()); @@ -93,16 +93,14 @@ class Louvain { rmm::device_uvector inc(n_verts, handle_.get_stream_view()); rmm::device_uvector deg(n_verts, handle_.get_stream_view()); - thrust::fill( - rmm::exec_policy(handle_.get_stream_view()), inc.begin(), inc.end(), weight_t{0.0}); - thrust::fill( - rmm::exec_policy(handle_.get_stream_view()), deg.begin(), deg.end(), weight_t{0.0}); + thrust::fill(handle_.get_thrust_policy(), inc.begin(), inc.end(), weight_t{0.0}); + thrust::fill(handle_.get_thrust_policy(), deg.begin(), deg.end(), weight_t{0.0}); // FIXME: Already have weighted degree computed in main loop, // could pass that in rather than computing d_deg... which // would save an atomicAdd (synchronization) // - thrust::for_each(rmm::exec_policy(handle_.get_stream_view()), + thrust::for_each(handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(graph.number_of_vertices), [d_inc = inc.data(), @@ -126,7 +124,7 @@ class Louvain { }); weight_t Q = thrust::transform_reduce( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(graph.number_of_vertices), [d_deg = deg.data(), d_inc = inc.data(), total_edge_weight, resolution] __device__( @@ -149,8 +147,8 @@ class Louvain { virtual weight_t operator()(size_t max_level, weight_t resolution) { - weight_t total_edge_weight = thrust::reduce( - rmm::exec_policy(handle_.get_stream_view()), weights_v_.begin(), weights_v_.end()); + weight_t total_edge_weight = + thrust::reduce(handle_.get_thrust_policy(), weights_v_.begin(), weights_v_.end()); weight_t best_modularity = weight_t{-1}; @@ -215,7 +213,7 @@ class Louvain { { dendrogram_->add_level(0, num_vertices, handle_.get_stream_view()); - thrust::sequence(rmm::exec_policy(handle_.get_stream_view()), + thrust::sequence(handle_.get_thrust_policy(), dendrogram_->current_level_begin(), dendrogram_->current_level_end()); } @@ -235,7 +233,7 @@ class Louvain { // MNMG: copy_v_transform_reduce_out_nbr, then copy // thrust::for_each( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(graph.number_of_vertices), [d_offsets, d_indices, d_weights, d_vertex_weights, d_cluster_weights] __device__( @@ -268,7 +266,7 @@ class Louvain { weight_t* d_cluster_weights = cluster_weights_v_.data(); weight_t* d_delta_Q = delta_Q_v.data(); - thrust::copy(rmm::exec_policy(handle_.get_stream_view()), + thrust::copy(handle_.get_thrust_policy(), dendrogram_->current_level_begin(), dendrogram_->current_level_end(), next_cluster_v.data()); @@ -296,7 +294,7 @@ class Louvain { new_Q = modularity(total_edge_weight, resolution, graph, next_cluster_v.data()); if (new_Q > cur_Q) { - thrust::copy(rmm::exec_policy(handle_.get_stream_view()), + thrust::copy(handle_.get_thrust_policy(), next_cluster_v.begin(), next_cluster_v.end(), dendrogram_->current_level_begin()); @@ -325,20 +323,15 @@ class Louvain { weight_t* d_old_cluster_sum = old_cluster_sum_v.data(); weight_t* d_new_cluster_sum = d_delta_Q; - thrust::fill(rmm::exec_policy(handle_.get_stream_view()), - cluster_hash_v.begin(), - cluster_hash_v.end(), - vertex_t{-1}); - thrust::fill(rmm::exec_policy(handle_.get_stream_view()), - delta_Q_v.begin(), - delta_Q_v.end(), - weight_t{0.0}); - thrust::fill(rmm::exec_policy(handle_.get_stream_view()), + thrust::fill( + handle_.get_thrust_policy(), cluster_hash_v.begin(), cluster_hash_v.end(), vertex_t{-1}); + thrust::fill(handle_.get_thrust_policy(), delta_Q_v.begin(), delta_Q_v.end(), weight_t{0.0}); + thrust::fill(handle_.get_thrust_policy(), old_cluster_sum_v.begin(), old_cluster_sum_v.end(), weight_t{0.0}); - thrust::for_each(rmm::exec_policy(handle_.get_stream_view()), + thrust::for_each(handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(graph.number_of_edges), [d_src_indices = src_indices_v_.data(), @@ -377,7 +370,7 @@ class Louvain { }); thrust::for_each( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(graph.number_of_edges), [total_edge_weight, @@ -423,15 +416,11 @@ class Louvain { rmm::device_uvector temp_delta_Q_v(graph.number_of_vertices, handle_.get_stream_view()); - thrust::fill(rmm::exec_policy(handle_.get_stream_view()), - temp_cluster_v.begin(), - temp_cluster_v.end(), - vertex_t{-1}); + thrust::fill( + handle_.get_thrust_policy(), temp_cluster_v.begin(), temp_cluster_v.end(), vertex_t{-1}); - thrust::fill(rmm::exec_policy(handle_.get_stream_view()), - temp_delta_Q_v.begin(), - temp_delta_Q_v.end(), - weight_t{0}); + thrust::fill( + handle_.get_thrust_policy(), temp_delta_Q_v.begin(), temp_delta_Q_v.end(), weight_t{0}); auto cluster_reduce_iterator = thrust::make_zip_iterator(thrust::make_tuple(cluster_hash_v.begin(), delta_Q_v.begin())); @@ -440,7 +429,7 @@ class Louvain { thrust::make_zip_iterator(thrust::make_tuple(temp_cluster_v.begin(), temp_delta_Q_v.begin())); auto cluster_reduce_end = - thrust::reduce_by_key(rmm::exec_policy(handle_.get_stream_view()), + thrust::reduce_by_key(handle_.get_thrust_policy(), src_indices_v_.begin(), src_indices_v_.end(), cluster_reduce_iterator, @@ -459,7 +448,7 @@ class Louvain { vertex_t final_size = thrust::distance(temp_vertices_v.data(), cluster_reduce_end.first); - thrust::for_each(rmm::exec_policy(handle_.get_stream_view()), + thrust::for_each(handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(final_size), [up_down, @@ -509,7 +498,7 @@ class Louvain { // // New technique. Initialize cluster_inverse_v_ to 0 // - thrust::fill(rmm::exec_policy(handle_.get_stream_view()), + thrust::fill(handle_.get_thrust_policy(), cluster_inverse_v_.begin(), cluster_inverse_v_.end(), vertex_t{0}); @@ -520,7 +509,7 @@ class Louvain { auto first_1 = thrust::make_constant_iterator(1); auto last_1 = first_1 + old_num_clusters; - thrust::scatter(rmm::exec_policy(handle_.get_stream_view()), + thrust::scatter(handle_.get_thrust_policy(), first_1, last_1, dendrogram_->current_level_begin(), @@ -530,7 +519,7 @@ class Louvain { // Now we'll copy all of the clusters that have a value of 1 into a temporary array // auto copy_end = thrust::copy_if( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(old_num_clusters), tmp_arr_v_.begin(), @@ -542,14 +531,14 @@ class Louvain { // // Now we can set each value in cluster_inverse of a cluster to its index // - thrust::for_each(rmm::exec_policy(handle_.get_stream_view()), + thrust::for_each(handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(new_num_clusters), [d_cluster_inverse, d_tmp_array] __device__(const vertex_t idx) { d_cluster_inverse[d_tmp_array[idx]] = idx; }); - thrust::for_each(rmm::exec_policy(handle_.get_stream_view()), + thrust::for_each(handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(old_num_clusters), [d_cluster, d_cluster_inverse] __device__(vertex_t i) { @@ -570,7 +559,7 @@ class Louvain { // // Renumber the COO // - thrust::for_each(rmm::exec_policy(handle_.get_stream_view()), + thrust::for_each(handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(graph.number_of_edges), [d_old_src = src_indices_v_.data(), @@ -586,12 +575,12 @@ class Louvain { }); thrust::stable_sort_by_key( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), new_dst_v.begin(), new_dst_v.end(), thrust::make_zip_iterator(thrust::make_tuple(new_src_v.begin(), new_weight_v.begin()))); thrust::stable_sort_by_key( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), new_src_v.begin(), new_src_v.end(), thrust::make_zip_iterator(thrust::make_tuple(new_dst_v.begin(), new_weight_v.begin()))); @@ -604,7 +593,7 @@ class Louvain { thrust::make_zip_iterator(thrust::make_tuple(new_src_v.begin(), new_dst_v.begin())); auto new_start = thrust::make_zip_iterator(thrust::make_tuple(src_indices_v_.data(), graph.indices)); - auto new_end = thrust::reduce_by_key(rmm::exec_policy(handle_.get_stream_view()), + auto new_end = thrust::reduce_by_key(handle_.get_thrust_policy(), start, start + graph.number_of_edges, new_weight_v.begin(), diff --git a/cpp/src/community/legacy/spectral_clustering.cu b/cpp/src/community/legacy/spectral_clustering.cu index c56b8eb641b..8534665530b 100644 --- a/cpp/src/community/legacy/spectral_clustering.cu +++ b/cpp/src/community/legacy/spectral_clustering.cu @@ -70,9 +70,6 @@ void balancedCutClustering_impl(legacy::GraphCSRView RAFT_EXPECTS(eig_vects != nullptr, "API error, must specify valid eigenvectors"); raft::handle_t handle; - auto stream = handle.get_stream(); - auto exec = rmm::exec_policy(stream); - auto t_exe_p = exec; int evs_max_it{4000}; int kmean_max_it{200}; @@ -106,7 +103,7 @@ void balancedCutClustering_impl(legacy::GraphCSRView raft::kmeans_solver_t cluster_solver{clust_cfg}; raft::spectral::partition( - handle, t_exe_p, r_csr_m, eig_solver, cluster_solver, clustering, eig_vals, eig_vects); + handle, r_csr_m, eig_solver, cluster_solver, clustering, eig_vals, eig_vects); } template @@ -141,9 +138,6 @@ void spectralModularityMaximization_impl( RAFT_EXPECTS(eig_vects != nullptr, "API error, must specify valid eigenvectors"); raft::handle_t handle; - auto stream = handle.get_stream(); - auto exec = rmm::exec_policy(stream); - auto t_exe_p = exec; int evs_max_it{4000}; int kmean_max_it{200}; @@ -179,7 +173,7 @@ void spectralModularityMaximization_impl( // not returned... // auto result = raft::spectral::modularity_maximization( - handle, t_exe_p, r_csr_m, eig_solver, cluster_solver, clustering, eig_vals, eig_vects); + handle, r_csr_m, eig_solver, cluster_solver, clustering, eig_vals, eig_vects); // not returned... // int iters_lanczos, iters_kmeans; @@ -194,9 +188,6 @@ void analyzeModularityClustering_impl(legacy::GraphCSRView const r_csr_m{handle, graph}; weight_t mod; - raft::spectral::analyzeModularity(handle, t_exe_p, r_csr_m, n_clusters, clustering, mod); + raft::spectral::analyzeModularity(handle, r_csr_m, n_clusters, clustering, mod); *modularity = mod; } @@ -216,9 +207,6 @@ void analyzeBalancedCut_impl(legacy::GraphCSRView co weight_t* ratioCut) { raft::handle_t handle; - auto stream = handle.get_stream(); - auto exec = rmm::exec_policy(stream); - auto t_exe_p = exec; RAFT_EXPECTS(n_clusters <= graph.number_of_vertices, "API error: number of clusters must be <= number of vertices"); @@ -232,8 +220,7 @@ void analyzeBalancedCut_impl(legacy::GraphCSRView co raft::matrix::sparse_matrix_t const r_csr_m{handle, graph}; - raft::spectral::analyzePartition( - handle, t_exe_p, r_csr_m, n_clusters, clustering, edge_cut, cost); + raft::spectral::analyzePartition(handle, r_csr_m, n_clusters, clustering, edge_cut, cost); *edgeCut = edge_cut; *ratioCut = cost; diff --git a/cpp/src/community/legacy/triangles_counting.cu b/cpp/src/community/legacy/triangles_counting.cu index 8922f92336d..e4fe6f09a1b 100644 --- a/cpp/src/community/legacy/triangles_counting.cu +++ b/cpp/src/community/legacy/triangles_counting.cu @@ -19,16 +19,12 @@ #include #include #include - #include - #include #include -#include #include -#include #include "cub/cub.cuh" #define TH_CENT_K_LOCLEN (34) diff --git a/cpp/src/community/louvain.cu b/cpp/src/community/louvain.cu index d6bd224fedf..c65bfe4faf6 100644 --- a/cpp/src/community/louvain.cu +++ b/cpp/src/community/louvain.cu @@ -82,7 +82,7 @@ void flatten_dendrogram(raft::handle_t const& handle, rmm::device_uvector vertex_ids_v(graph_view.number_of_vertices, handle.get_stream()); thrust::sequence( - rmm::exec_policy(handle.get_stream()), vertex_ids_v.begin(), vertex_ids_v.end(), vertex_t{0}); + handle.get_thrust_policy(), vertex_ids_v.begin(), vertex_ids_v.end(), vertex_t{0}); partition_at_level( handle, dendrogram, vertex_ids_v.data(), clustering, dendrogram.num_levels()); @@ -98,7 +98,7 @@ void flatten_dendrogram( rmm::device_uvector vertex_ids_v(graph_view.get_number_of_vertices(), handle.get_stream()); - thrust::sequence(rmm::exec_policy(handle.get_stream()), + thrust::sequence(handle.get_thrust_policy(), vertex_ids_v.begin(), vertex_ids_v.end(), graph_view.get_local_vertex_first()); diff --git a/cpp/src/community/louvain.cuh b/cpp/src/community/louvain.cuh index 09189c95e38..a7ecd2802e8 100644 --- a/cpp/src/community/louvain.cuh +++ b/cpp/src/community/louvain.cuh @@ -28,6 +28,9 @@ #include #include +#include +#include + #include #include @@ -153,7 +156,7 @@ class Louvain { dendrogram_->add_level( current_graph_view_.get_local_vertex_first(), num_vertices, handle_.get_stream_view()); - thrust::sequence(rmm::exec_policy(handle_.get_stream_view()), + thrust::sequence(handle_.get_thrust_policy(), dendrogram_->current_level_begin(), dendrogram_->current_level_end(), current_graph_view_.get_local_vertex_first()); @@ -163,7 +166,7 @@ class Louvain { weight_t modularity(weight_t total_edge_weight, weight_t resolution) { weight_t sum_degree_squared = thrust::transform_reduce( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), cluster_weights_v_.begin(), cluster_weights_v_.end(), [] __device__(weight_t p) { return p * p; }, @@ -203,7 +206,7 @@ class Louvain { cluster_keys_v_.resize(vertex_weights_v_.size(), handle_.get_stream_view()); cluster_weights_v_.resize(vertex_weights_v_.size(), handle_.get_stream_view()); - thrust::sequence(rmm::exec_policy(handle_.get_stream_view()), + thrust::sequence(handle_.get_thrust_policy(), cluster_keys_v_.begin(), cluster_keys_v_.end(), current_graph_view_.get_local_vertex_first()); @@ -337,7 +340,7 @@ class Louvain { cugraph::get_dataframe_buffer_begin>(output_buffer)); thrust::transform( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), cugraph::get_dataframe_buffer_begin>(output_buffer), cugraph::get_dataframe_buffer_begin>(output_buffer) + current_graph_view_.get_number_of_local_vertices(), @@ -345,7 +348,7 @@ class Louvain { [] __device__(auto p) { return thrust::get<1>(p); }); thrust::transform( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), cugraph::get_dataframe_buffer_begin>(output_buffer), cugraph::get_dataframe_buffer_begin>(output_buffer) + current_graph_view_.get_number_of_local_vertices(), @@ -392,12 +395,12 @@ class Louvain { map_key_last = cluster_keys_v_.end(); map_value_first = cluster_weights_v_.begin(); } else { - thrust::sort_by_key(rmm::exec_policy(handle_.get_stream_view()), + thrust::sort_by_key(handle_.get_thrust_policy(), cluster_keys_v_.begin(), cluster_keys_v_.end(), cluster_weights_v_.begin()); - thrust::transform(rmm::exec_policy(handle_.get_stream_view()), + thrust::transform(handle_.get_thrust_policy(), next_cluster_v.begin(), next_cluster_v.end(), src_cluster_weights_v.begin(), @@ -464,7 +467,7 @@ class Louvain { cugraph::get_dataframe_buffer_begin>(output_buffer)); thrust::transform( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), next_cluster_v.begin(), next_cluster_v.end(), cugraph::get_dataframe_buffer_begin>(output_buffer), @@ -504,7 +507,7 @@ class Louvain { current_graph_view_ = current_graph_->view(); rmm::device_uvector numbering_indices(numbering_map.size(), handle_.get_stream()); - thrust::sequence(rmm::exec_policy(handle_.get_stream_view()), + thrust::sequence(handle_.get_thrust_policy(), numbering_indices.begin(), numbering_indices.end(), current_graph_view_.get_local_vertex_first()); diff --git a/cpp/src/components/weak_cc.cuh b/cpp/src/components/weak_cc.cuh index 31beda96342..6ae1e8d096c 100644 --- a/cpp/src/components/weak_cc.cuh +++ b/cpp/src/components/weak_cc.cuh @@ -29,7 +29,6 @@ #include #include -#include #include "utils.h" namespace MLCommon { diff --git a/cpp/src/components/weakly_connected_components.cu b/cpp/src/components/weakly_connected_components.cu index f20356a6d58..192bef6c432 100644 --- a/cpp/src/components/weakly_connected_components.cu +++ b/cpp/src/components/weakly_connected_components.cu @@ -27,7 +27,6 @@ #include #include -#include #include #include @@ -96,7 +95,7 @@ accumulate_new_roots(raft::handle_t const& handle, static_cast(thrust::distance( output_pair_first, thrust::copy_if( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), input_pair_first, input_pair_first + scan_size, output_pair_first, @@ -112,18 +111,18 @@ accumulate_new_roots(raft::handle_t const& handle, rmm::device_uvector tmp_cumulative_degrees(tmp_new_roots.size(), handle.get_stream_view()); thrust::transform( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), tmp_new_roots.begin(), tmp_new_roots.end(), tmp_cumulative_degrees.begin(), [vertex_partition, degrees] __device__(auto v) { return degrees[vertex_partition.get_local_vertex_offset_from_vertex_nocheck(v)]; }); - thrust::inclusive_scan(rmm::exec_policy(handle.get_stream_view()), + thrust::inclusive_scan(handle.get_thrust_policy(), tmp_cumulative_degrees.begin(), tmp_cumulative_degrees.end(), tmp_cumulative_degrees.begin()); - auto last = thrust::lower_bound(rmm::exec_policy(handle.get_stream_view()), + auto last = thrust::lower_bound(handle.get_thrust_policy(), tmp_cumulative_degrees.begin(), tmp_cumulative_degrees.end(), degree_sum_threshold - degree_sum); @@ -132,7 +131,7 @@ accumulate_new_roots(raft::handle_t const& handle, std::min(static_cast(thrust::distance(tmp_cumulative_degrees.begin(), last)), max_new_roots - num_new_roots); - thrust::copy(rmm::exec_policy(handle.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), tmp_new_roots.begin(), tmp_new_roots.begin() + tmp_num_new_roots, new_roots.begin() + num_new_roots); @@ -284,7 +283,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, auto pair_first = thrust::make_zip_iterator(thrust::make_tuple( thrust::make_counting_iterator(level_graph_view.get_local_vertex_first()), degrees.begin())); - thrust::transform(rmm::exec_policy(handle.get_stream_view()), + thrust::transform(handle.get_thrust_policy(), pair_first, pair_first + level_graph_view.get_number_of_local_vertices(), level_components, @@ -312,7 +311,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, thrust::distance( new_root_candidates.begin(), thrust::copy_if( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), thrust::make_counting_iterator(level_graph_view.get_local_vertex_first()), thrust::make_counting_iterator(level_graph_view.get_local_vertex_last()), new_root_candidates.begin(), @@ -322,7 +321,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, })), handle.get_stream_view()); auto high_degree_partition_last = thrust::stable_partition( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), new_root_candidates.begin(), new_root_candidates.end(), [vertex_partition, @@ -332,7 +331,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, return degrees[vertex_partition.get_local_vertex_offset_from_vertex_nocheck(v)] >= threshold; }); - thrust::shuffle(rmm::exec_policy(handle.get_stream_view()), + thrust::shuffle(handle.get_thrust_policy(), high_degree_partition_last, new_root_candidates.end(), thrust::default_random_engine()); @@ -350,7 +349,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, auto const comm_size = comm.get_size(); auto first_candidate_degree = thrust::transform_reduce( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), new_root_candidates.begin(), new_root_candidates.begin() + (new_root_candidates.size() > 0 ? 1 : 0), [vertex_partition, degrees = degrees.data()] __device__(auto v) { @@ -463,7 +462,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, : vertex_t{0}, handle.get_stream_view()); if (GraphViewType::is_multi_gpu) { - thrust::fill(rmm::exec_policy(handle.get_stream_view()), + thrust::fill(handle.get_thrust_policy(), col_components.begin(), col_components.end(), invalid_component_id::value); @@ -487,11 +486,10 @@ void weakly_connected_components_impl(raft::handle_t const& handle, next_candidate_offset += num_scanned; edge_count += degree_sum; - thrust::sort( - rmm::exec_policy(handle.get_stream_view()), new_roots.begin(), new_roots.end()); + thrust::sort(handle.get_thrust_policy(), new_roots.begin(), new_roots.end()); thrust::for_each( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), new_roots.begin(), new_roots.end(), [vertex_partition, components = level_components] __device__(auto c) { @@ -585,7 +583,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, resize_dataframe_buffer>( edge_buffer, cur_num_edge_inserts + conflict_bucket.size(), handle.get_stream()); thrust::for_each( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), conflict_bucket.begin(), conflict_bucket.end(), [vertex_partition, @@ -613,7 +611,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, if (new_num_edge_inserts > old_num_edge_inserts) { auto edge_first = get_dataframe_buffer_begin>(edge_buffer); - thrust::sort(rmm::exec_policy(handle.get_stream_view()), + thrust::sort(handle.get_thrust_policy(), edge_first + old_num_edge_inserts, edge_first + new_num_edge_inserts); if (old_num_edge_inserts > 0) { @@ -621,7 +619,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, new_num_edge_inserts, handle.get_stream()); auto tmp_edge_first = get_dataframe_buffer_begin>(tmp_edge_buffer); - thrust::merge(rmm::exec_policy(handle.get_stream_view()), + thrust::merge(handle.get_thrust_policy(), edge_first, edge_first + old_num_edge_inserts, edge_first + old_num_edge_inserts, @@ -630,9 +628,8 @@ void weakly_connected_components_impl(raft::handle_t const& handle, edge_buffer = std::move(tmp_edge_buffer); } edge_first = get_dataframe_buffer_begin>(edge_buffer); - auto unique_edge_last = thrust::unique(rmm::exec_policy(handle.get_stream_view()), - edge_first, - edge_first + new_num_edge_inserts); + auto unique_edge_last = + thrust::unique(handle.get_thrust_policy(), edge_first, edge_first + new_num_edge_inserts); auto num_unique_edges = static_cast(thrust::distance(edge_first, unique_edge_last)); num_edge_inserts.set_value_async(num_unique_edges, handle.get_stream_view()); } @@ -642,7 +639,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, vertex_frontier.swap_buckets(static_cast(Bucket::cur), static_cast(Bucket::next)); edge_count = thrust::transform_reduce( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), thrust::get<0>(vertex_frontier.get_bucket(static_cast(Bucket::cur)) .begin() .get_iterator_tuple()), @@ -674,10 +671,8 @@ void weakly_connected_components_impl(raft::handle_t const& handle, thrust::make_tuple(thrust::get<1>(input_first.get_iterator_tuple()), thrust::get<0>(input_first.get_iterator_tuple()))) + num_inserts; - thrust::copy(rmm::exec_policy(handle.get_stream_view()), - input_first, - input_first + num_inserts, - output_first); + thrust::copy( + handle.get_thrust_policy(), input_first, input_first + num_inserts, output_first); if (GraphViewType::is_multi_gpu) { auto& comm = handle.get_comms(); @@ -700,9 +695,8 @@ void weakly_connected_components_impl(raft::handle_t const& handle, auto edge_first = get_dataframe_buffer_begin>(edge_buffer); auto edge_last = get_dataframe_buffer_end>(edge_buffer); - thrust::sort(rmm::exec_policy(handle.get_stream_view()), edge_first, edge_last); - auto unique_edge_last = - thrust::unique(rmm::exec_policy(handle.get_stream_view()), edge_first, edge_last); + thrust::sort(handle.get_thrust_policy(), edge_first, edge_last); + auto unique_edge_last = thrust::unique(handle.get_thrust_policy(), edge_first, edge_last); resize_dataframe_buffer>( edge_buffer, static_cast(thrust::distance(edge_first, unique_edge_last)), @@ -738,7 +732,7 @@ void weakly_connected_components_impl(raft::handle_t const& handle, rmm::device_uvector next_local_vertices(level_renumber_map_vectors[next_level].size(), handle.get_stream_view()); - thrust::sequence(rmm::exec_policy(handle.get_stream_view()), + thrust::sequence(handle.get_thrust_policy(), next_local_vertices.begin(), next_local_vertices.end(), level_local_vertex_first_vectors[next_level]); diff --git a/cpp/src/converters/COOtoCSR.cuh b/cpp/src/converters/COOtoCSR.cuh index 641b037efdd..a790782be11 100644 --- a/cpp/src/converters/COOtoCSR.cuh +++ b/cpp/src/converters/COOtoCSR.cuh @@ -22,7 +22,6 @@ #pragma once -#include #include #include #include diff --git a/cpp/src/generators/erdos_renyi_generator.cu b/cpp/src/generators/erdos_renyi_generator.cu index 8452a613174..3f8f558e4fe 100644 --- a/cpp/src/generators/erdos_renyi_generator.cu +++ b/cpp/src/generators/erdos_renyi_generator.cu @@ -18,7 +18,6 @@ #include #include -#include #include #include @@ -47,14 +46,14 @@ generate_erdos_renyi_graph_edgelist_gnp(raft::handle_t const& handle, return dist(rng); }); - size_t count = thrust::count_if(rmm::exec_policy(handle.get_stream()), + size_t count = thrust::count_if(handle.get_thrust_policy(), random_iterator, random_iterator + num_vertices * num_vertices, [p] __device__(float prob) { return prob < p; }); rmm::device_uvector indices_v(count, handle.get_stream()); - thrust::copy_if(rmm::exec_policy(handle.get_stream()), + thrust::copy_if(handle.get_thrust_policy(), random_iterator, random_iterator + num_vertices * num_vertices, indices_v.begin(), @@ -63,7 +62,7 @@ generate_erdos_renyi_graph_edgelist_gnp(raft::handle_t const& handle, rmm::device_uvector src_v(count, handle.get_stream()); rmm::device_uvector dst_v(count, handle.get_stream()); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), indices_v.begin(), indices_v.end(), thrust::make_zip_iterator(thrust::make_tuple(src_v.begin(), src_v.end())), diff --git a/cpp/src/generators/generate_rmat_edgelist.cu b/cpp/src/generators/generate_rmat_edgelist.cu index c7d8a5682bc..e43bce51872 100644 --- a/cpp/src/generators/generate_rmat_edgelist.cu +++ b/cpp/src/generators/generate_rmat_edgelist.cu @@ -20,7 +20,6 @@ #include #include -#include #include #include @@ -69,7 +68,7 @@ std::tuple, rmm::device_uvector> generat seed += num_edges_to_generate * 2 * scale; thrust::transform( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), thrust::make_counting_iterator(size_t{0}), thrust::make_counting_iterator(num_edges_to_generate), pair_first, diff --git a/cpp/src/generators/generator_tools.cu b/cpp/src/generators/generator_tools.cu index 800194ce2b9..f07c59e18a6 100644 --- a/cpp/src/generators/generator_tools.cu +++ b/cpp/src/generators/generator_tools.cu @@ -21,9 +21,7 @@ #include #include -#include -#include #include #include @@ -71,7 +69,7 @@ void scramble_vertex_ids(raft::handle_t const& handle, vertex_t scale = 1 + raft::log2(d_src_v.size()); auto pair_first = thrust::make_zip_iterator(thrust::make_tuple(d_src_v.begin(), d_dst_v.begin())); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), pair_first, pair_first + d_src_v.size(), pair_first, @@ -138,21 +136,19 @@ combine_edgelists(raft::handle_t const& handle, if (optional_d_weights) { thrust::sort( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), thrust::make_zip_iterator( thrust::make_tuple(srcs_v.begin(), dsts_v.begin(), weights_v.begin())), thrust::make_zip_iterator(thrust::make_tuple(srcs_v.end(), dsts_v.end(), weights_v.end()))); auto pair_first = thrust::make_zip_iterator(thrust::make_tuple(srcs_v.begin(), dsts_v.begin())); - auto end_iter = thrust::unique_by_key(rmm::exec_policy(handle.get_stream()), - pair_first, - pair_first + srcs_v.size(), - weights_v.begin()); + auto end_iter = thrust::unique_by_key( + handle.get_thrust_policy(), pair_first, pair_first + srcs_v.size(), weights_v.begin()); number_of_edges = thrust::distance(pair_first, thrust::get<0>(end_iter)); } else { - thrust::sort(rmm::exec_policy(handle.get_stream()), + thrust::sort(handle.get_thrust_policy(), thrust::make_zip_iterator(thrust::make_tuple(srcs_v.begin(), dsts_v.begin())), thrust::make_zip_iterator(thrust::make_tuple(srcs_v.end(), dsts_v.end()))); @@ -160,7 +156,7 @@ combine_edgelists(raft::handle_t const& handle, thrust::make_zip_iterator(thrust::make_tuple(srcs_v.begin(), dsts_v.begin())); auto end_iter = thrust::unique( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), thrust::make_zip_iterator(thrust::make_tuple(srcs_v.begin(), dsts_v.begin())), thrust::make_zip_iterator(thrust::make_tuple(srcs_v.end(), dsts_v.end()))); @@ -199,17 +195,17 @@ symmetrize_edgelist(raft::handle_t const& handle, d_src_v.resize(offset * 2, handle.get_stream_view()); d_dst_v.resize(offset * 2, handle.get_stream_view()); - thrust::copy(rmm::exec_policy(handle.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), d_dst_v.begin(), d_dst_v.begin() + offset, d_src_v.begin() + offset); - thrust::copy(rmm::exec_policy(handle.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), d_src_v.begin(), d_src_v.begin() + offset, d_dst_v.begin() + offset); if (optional_d_weights_v) { optional_d_weights_v->resize(d_src_v.size(), handle.get_stream_view()); - thrust::copy(rmm::exec_policy(handle.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), optional_d_weights_v->begin(), optional_d_weights_v->begin() + offset, optional_d_weights_v->begin() + offset); diff --git a/cpp/src/generators/simple_generators.cu b/cpp/src/generators/simple_generators.cu index 413e08962e7..5f003a04219 100644 --- a/cpp/src/generators/simple_generators.cu +++ b/cpp/src/generators/simple_generators.cu @@ -18,7 +18,6 @@ #include #include -#include #include @@ -68,15 +67,11 @@ generate_path_graph_edgelist(raft::handle_t const& handle, if (edge_off_end) ++num_edges; - thrust::sequence(rmm::exec_policy(handle.get_stream()), - src_iterator, - src_iterator + num_edges, - base_vertex_id); + thrust::sequence( + handle.get_thrust_policy(), src_iterator, src_iterator + num_edges, base_vertex_id); - thrust::sequence(rmm::exec_policy(handle.get_stream()), - dst_iterator, - dst_iterator + num_edges, - base_vertex_id + 1); + thrust::sequence( + handle.get_thrust_policy(), dst_iterator, dst_iterator + num_edges, base_vertex_id + 1); src_iterator += num_edges; dst_iterator += num_edges; @@ -122,7 +117,7 @@ generate_2d_mesh_graph_edgelist( thrust::make_tuple(thrust::make_counting_iterator(base_vertex_id), thrust::make_counting_iterator(base_vertex_id + 1))); - output_iterator = thrust::copy_if(rmm::exec_policy(handle.get_stream()), + output_iterator = thrust::copy_if(handle.get_thrust_policy(), x_iterator, x_iterator + num_vertices - 1, output_iterator, @@ -136,7 +131,7 @@ generate_2d_mesh_graph_edgelist( thrust::make_tuple(thrust::make_counting_iterator(base_vertex_id), thrust::make_counting_iterator(base_vertex_id + x))); - output_iterator = thrust::copy_if(rmm::exec_policy(handle.get_stream()), + output_iterator = thrust::copy_if(handle.get_thrust_policy(), y_iterator, y_iterator + num_vertices - x, output_iterator, @@ -188,7 +183,7 @@ generate_3d_mesh_graph_edgelist( thrust::make_tuple(thrust::make_counting_iterator(base_vertex_id), thrust::make_counting_iterator(base_vertex_id + 1))); - output_iterator = thrust::copy_if(rmm::exec_policy(handle.get_stream()), + output_iterator = thrust::copy_if(handle.get_thrust_policy(), x_iterator, x_iterator + num_vertices - 1, output_iterator, @@ -202,7 +197,7 @@ generate_3d_mesh_graph_edgelist( thrust::make_tuple(thrust::make_counting_iterator(base_vertex_id), thrust::make_counting_iterator(base_vertex_id + x))); - output_iterator = thrust::copy_if(rmm::exec_policy(handle.get_stream()), + output_iterator = thrust::copy_if(handle.get_thrust_policy(), y_iterator, y_iterator + num_vertices - x, output_iterator, @@ -216,7 +211,7 @@ generate_3d_mesh_graph_edgelist( thrust::make_tuple(thrust::make_counting_iterator(base_vertex_id), thrust::make_counting_iterator(base_vertex_id + x * y))); - output_iterator = thrust::copy_if(rmm::exec_policy(handle.get_stream()), + output_iterator = thrust::copy_if(handle.get_thrust_policy(), z_iterator, z_iterator + num_vertices - x * y, output_iterator, @@ -287,7 +282,7 @@ generate_complete_graph_edgelist( return thrust::make_tuple(src, dst); }); - output_iterator = thrust::copy_if(rmm::exec_policy(handle.get_stream()), + output_iterator = thrust::copy_if(handle.get_thrust_policy(), transform_iter, transform_iter + num_vertices * num_vertices, output_iterator, diff --git a/cpp/src/layout/barnes_hut.cuh b/cpp/src/layout/barnes_hut.cuh index d05c6051d8b..61e47b03b5c 100644 --- a/cpp/src/layout/barnes_hut.cuh +++ b/cpp/src/layout/barnes_hut.cuh @@ -91,7 +91,7 @@ void barnes_hut(raft::handle_t const& handle, rmm::device_uvector d_childl((nnodes + 1) * 4, stream_view); // FA2 requires degree + 1 rmm::device_uvector d_massl(nnodes + 1, stream_view); - thrust::fill(rmm::exec_policy(stream_view), d_massl.begin(), d_massl.end(), 1); + thrust::fill(handle.get_thrust_policy(), d_massl.begin(), d_massl.end(), 1); rmm::device_uvector d_maxxl(blocks * FACTOR1, stream_view); rmm::device_uvector d_maxyl(blocks * FACTOR1, stream_view); @@ -154,7 +154,7 @@ void barnes_hut(raft::handle_t const& handle, swinging = d_swinging.data(); traction = d_traction.data(); - thrust::fill(rmm::exec_policy(stream_view), d_old_forces.begin(), d_old_forces.end(), 0.f); + thrust::fill(handle.get_thrust_policy(), d_old_forces.begin(), d_old_forces.end(), 0.f); // Sort COO for coalesced memory access. sort(graph, stream_view.value()); @@ -175,7 +175,7 @@ void barnes_hut(raft::handle_t const& handle, // If outboundAttractionDistribution active, compensate. if (outbound_attraction_distribution) { - int sum = thrust::reduce(rmm::exec_policy(stream_view), d_massl.begin(), d_massl.begin() + n); + int sum = thrust::reduce(handle.get_thrust_policy(), d_massl.begin(), d_massl.begin() + n); outbound_att_compensation = sum / (float)n; } @@ -198,10 +198,10 @@ void barnes_hut(raft::handle_t const& handle, for (int iter = 0; iter < max_iter; ++iter) { // Reset force values - thrust::fill(rmm::exec_policy(stream_view), d_rep_forces.begin(), d_rep_forces.end(), 0.f); - thrust::fill(rmm::exec_policy(stream_view), d_attract.begin(), d_attract.end(), 0.f); - thrust::fill(rmm::exec_policy(stream_view), d_swinging.begin(), d_swinging.end(), 0.f); - thrust::fill(rmm::exec_policy(stream_view), d_traction.begin(), d_traction.end(), 0.f); + thrust::fill(handle.get_thrust_policy(), d_rep_forces.begin(), d_rep_forces.end(), 0.f); + thrust::fill(handle.get_thrust_policy(), d_attract.begin(), d_attract.end(), 0.f); + thrust::fill(handle.get_thrust_policy(), d_swinging.begin(), d_swinging.end(), 0.f); + thrust::fill(handle.get_thrust_policy(), d_traction.begin(), d_traction.end(), 0.f); ResetKernel<<<1, 1, 0, stream_view.value()>>>(radiusd_squared, bottomd, NNODES, radiusd); CHECK_CUDA(stream_view.value()); @@ -304,10 +304,10 @@ void barnes_hut(raft::handle_t const& handle, // Compute global swinging and traction values const float s = - thrust::reduce(rmm::exec_policy(stream_view), d_swinging.begin(), d_swinging.end()); + thrust::reduce(handle.get_thrust_policy(), d_swinging.begin(), d_swinging.end()); const float t = - thrust::reduce(rmm::exec_policy(stream_view), d_traction.begin(), d_traction.end()); + thrust::reduce(handle.get_thrust_policy(), d_traction.begin(), d_traction.end()); // Compute global speed based on gloab and local swinging and traction. adapt_speed(jitter_tolerance, &jt, &speed, &speed_efficiency, s, t, n); diff --git a/cpp/src/layout/exact_fa2.cuh b/cpp/src/layout/exact_fa2.cuh index 5b5c3f5e82e..db84594c8b8 100644 --- a/cpp/src/layout/exact_fa2.cuh +++ b/cpp/src/layout/exact_fa2.cuh @@ -65,10 +65,10 @@ void exact_fa2(raft::handle_t const& handle, rmm::device_uvector repel(n * 2, stream_view); rmm::device_uvector attract(n * 2, stream_view); rmm::device_uvector old_forces(n * 2, stream_view); - thrust::fill(rmm::exec_policy(stream_view), old_forces.begin(), old_forces.end(), 0.f); + thrust::fill(handle.get_thrust_policy(), old_forces.begin(), old_forces.end(), 0.f); // FA2 requires degree + 1. rmm::device_uvector mass(n, stream_view); - thrust::fill(rmm::exec_policy(stream_view), mass.begin(), mass.end(), 1); + thrust::fill(handle.get_thrust_policy(), mass.begin(), mass.end(), 1); rmm::device_uvector swinging(n, stream_view); rmm::device_uvector traction(n, stream_view); @@ -103,7 +103,7 @@ void exact_fa2(raft::handle_t const& handle, float jt = 0.f; if (outbound_attraction_distribution) { - int sum = thrust::reduce(rmm::exec_policy(stream_view), mass.begin(), mass.end()); + int sum = thrust::reduce(handle.get_thrust_policy(), mass.begin(), mass.end()); outbound_att_compensation = sum / (float)n; } @@ -114,10 +114,10 @@ void exact_fa2(raft::handle_t const& handle, for (int iter = 0; iter < max_iter; ++iter) { // Reset force arrays - thrust::fill(rmm::exec_policy(stream_view), repel.begin(), repel.end(), 0.f); - thrust::fill(rmm::exec_policy(stream_view), attract.begin(), attract.end(), 0.f); - thrust::fill(rmm::exec_policy(stream_view), swinging.begin(), swinging.end(), 0.f); - thrust::fill(rmm::exec_policy(stream_view), traction.begin(), traction.end(), 0.f); + thrust::fill(handle.get_thrust_policy(), repel.begin(), repel.end(), 0.f); + thrust::fill(handle.get_thrust_policy(), attract.begin(), attract.end(), 0.f); + thrust::fill(handle.get_thrust_policy(), swinging.begin(), swinging.end(), 0.f); + thrust::fill(handle.get_thrust_policy(), traction.begin(), traction.end(), 0.f); // Exact repulsion apply_repulsion( @@ -162,8 +162,8 @@ void exact_fa2(raft::handle_t const& handle, stream_view.value()); // Compute global swinging and traction values. - const float s = thrust::reduce(rmm::exec_policy(stream_view), swinging.begin(), swinging.end()); - const float t = thrust::reduce(rmm::exec_policy(stream_view), traction.begin(), traction.end()); + const float s = thrust::reduce(handle.get_thrust_policy(), swinging.begin(), swinging.end()); + const float t = thrust::reduce(handle.get_thrust_policy(), traction.begin(), traction.end()); adapt_speed(jitter_tolerance, &jt, &speed, &speed_efficiency, s, t, n); diff --git a/cpp/src/linear_assignment/hungarian.cu b/cpp/src/linear_assignment/hungarian.cu index 368e119e93c..7af829da2b3 100644 --- a/cpp/src/linear_assignment/hungarian.cu +++ b/cpp/src/linear_assignment/hungarian.cu @@ -20,7 +20,6 @@ #include #include -#include #include #include @@ -79,7 +78,7 @@ weight_t hungarian(raft::handle_t const& handle, // Fill the extra rows/columns with max(d_original_cost) // index_t n = std::max(num_rows, num_cols); - weight_t max_cost = thrust::reduce(rmm::exec_policy(handle.get_stream_view()), + weight_t max_cost = thrust::reduce(handle.get_thrust_policy(), d_original_cost, d_original_cost + (num_rows * num_cols), weight_t{0}, @@ -89,7 +88,7 @@ weight_t hungarian(raft::handle_t const& handle, rmm::device_uvector tmp_row_assignment_v(n, handle.get_stream_view()); rmm::device_uvector tmp_col_assignment_v(n, handle.get_stream_view()); - thrust::transform(rmm::exec_policy(handle.get_stream_view()), + thrust::transform(handle.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(n * n), tmp_cost_v.begin(), @@ -160,15 +159,14 @@ weight_t hungarian_sparse(raft::handle_t const& handle, // Renumber vertices internally. Workers will become // rows, tasks will become columns // - thrust::sequence( - rmm::exec_policy(handle.get_stream_view()), temp_tasks_v.begin(), temp_tasks_v.end()); + thrust::sequence(handle.get_thrust_policy(), temp_tasks_v.begin(), temp_tasks_v.end()); - thrust::for_each(rmm::exec_policy(handle.get_stream_view()), + thrust::for_each(handle.get_thrust_policy(), workers, workers + num_workers, [d_temp_tasks] __device__(vertex_t v) { d_temp_tasks[v] = -1; }); - auto temp_end = thrust::copy_if(rmm::exec_policy(handle.get_stream_view()), + auto temp_end = thrust::copy_if(handle.get_thrust_policy(), temp_tasks_v.begin(), temp_tasks_v.end(), d_tasks, @@ -180,30 +178,24 @@ weight_t hungarian_sparse(raft::handle_t const& handle, // // Now we'll assign costs into the dense array // - thrust::fill(rmm::exec_policy(handle.get_stream_view()), - temp_workers_v.begin(), - temp_workers_v.end(), - vertex_t{-1}); - thrust::fill(rmm::exec_policy(handle.get_stream_view()), - temp_tasks_v.begin(), - temp_tasks_v.end(), - vertex_t{-1}); thrust::fill( - rmm::exec_policy(handle.get_stream_view()), cost_v.begin(), cost_v.end(), weight_t{0}); + handle.get_thrust_policy(), temp_workers_v.begin(), temp_workers_v.end(), vertex_t{-1}); + thrust::fill(handle.get_thrust_policy(), temp_tasks_v.begin(), temp_tasks_v.end(), vertex_t{-1}); + thrust::fill(handle.get_thrust_policy(), cost_v.begin(), cost_v.end(), weight_t{0}); thrust::for_each( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(num_rows), [d_temp_workers, workers] __device__(vertex_t v) { d_temp_workers[workers[v]] = v; }); thrust::for_each( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(num_cols), [d_temp_tasks, d_tasks] __device__(vertex_t v) { d_temp_tasks[d_tasks[v]] = v; }); - thrust::for_each(rmm::exec_policy(handle.get_stream_view()), + thrust::for_each(handle.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(graph.number_of_edges), [d_temp_workers, @@ -246,7 +238,7 @@ weight_t hungarian_sparse(raft::handle_t const& handle, // // Translate the assignment back to the original vertex ids // - thrust::for_each(rmm::exec_policy(handle.get_stream_view()), + thrust::for_each(handle.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(num_rows), [d_tasks, d_temp_assignment, assignment] __device__(vertex_t id) { diff --git a/cpp/src/link_analysis/pagerank.cu b/cpp/src/link_analysis/pagerank.cu index 9a569fafae6..0ba1cda6704 100644 --- a/cpp/src/link_analysis/pagerank.cu +++ b/cpp/src/link_analysis/pagerank.cu @@ -157,13 +157,13 @@ void pagerank( CUGRAPH_EXPECTS(sum > 0.0, "Invalid input argument: sum of the PageRank initial " "guess values should be positive."); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), pageranks, pageranks + pull_graph_view.get_number_of_local_vertices(), pageranks, [sum] __device__(auto val) { return val / sum; }); } else { - thrust::fill(rmm::exec_policy(handle.get_stream()), + thrust::fill(handle.get_thrust_policy(), pageranks, pageranks + pull_graph_view.get_number_of_local_vertices(), result_t{1.0} / static_cast(num_vertices)); @@ -192,7 +192,7 @@ void pagerank( pull_graph_view.get_number_of_local_adj_matrix_partition_rows(), handle.get_stream()); size_t iter{0}; while (true) { - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(handle.get_thrust_policy(), pageranks, pageranks + pull_graph_view.get_number_of_local_vertices(), old_pageranks.data()); @@ -211,7 +211,7 @@ void pagerank( }, result_t{0.0}); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), vertex_val_first, vertex_val_first + pull_graph_view.get_number_of_local_vertices(), pageranks, @@ -247,7 +247,7 @@ void pagerank( auto val_first = thrust::make_zip_iterator( thrust::make_tuple(*personalization_vertices, *personalization_values)); thrust::for_each( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), val_first, val_first + *personalization_vector_size, [vertex_partition, pageranks, dangling_sum, personalization_sum, alpha] __device__( diff --git a/cpp/src/link_prediction/overlap.cu b/cpp/src/link_prediction/overlap.cu index 446d0c8cfdb..1cfae153719 100644 --- a/cpp/src/link_prediction/overlap.cu +++ b/cpp/src/link_prediction/overlap.cu @@ -22,7 +22,6 @@ #include #include #include -#include #include namespace cugraph { diff --git a/cpp/src/sampling/random_walks.cuh b/cpp/src/sampling/random_walks.cuh index ef4d1739463..159c1b455ea 100644 --- a/cpp/src/sampling/random_walks.cuh +++ b/cpp/src/sampling/random_walks.cuh @@ -28,7 +28,6 @@ #include #include -#include #include #include @@ -119,7 +118,7 @@ struct rrandom_gen_t { { auto const* d_ptr_out_degs = d_crt_out_deg.data(); thrust::transform_if( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), d_ptr_random_, d_ptr_random_ + num_paths_, // input1 d_ptr_out_degs, // input2 @@ -229,7 +228,7 @@ struct col_indx_extract_t(0), thrust::make_counting_iterator(num_paths_), // input1 d_v_col_indx.begin(), // input2 @@ -263,7 +262,7 @@ struct col_indx_extract_t& d_coalesced_w, // out: set of coalesced weights real_t tag) // otherwise. ambiguity with the other operator() { - thrust::for_each(rmm::exec_policy(handle_.get_stream_view()), + thrust::for_each(handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(num_paths_), // input1 [max_depth = max_depth_, @@ -386,7 +385,7 @@ struct random_walker_t { // intialize path sizes to 1, as they contain at least one vertex each: // the initial set: d_src_init_v; // - thrust::copy_n(rmm::exec_policy(handle_.get_stream_view()), + thrust::copy_n(handle_.get_thrust_policy(), thrust::make_constant_iterator(1), num_paths_, d_sizes.begin()); @@ -400,7 +399,7 @@ struct random_walker_t { auto map_it_begin = thrust::make_transform_iterator(thrust::make_counting_iterator(0), dlambda); - thrust::scatter(rmm::exec_policy(handle_.get_stream_view()), + thrust::scatter(handle_.get_thrust_policy(), d_src_init_v.begin(), d_src_init_v.end(), map_it_begin, @@ -478,7 +477,7 @@ struct random_walker_t { bool all_paths_stopped(device_vec_t const& d_crt_out_degs) const { auto how_many_stopped = - thrust::count_if(rmm::exec_policy(handle_.get_stream_view()), + thrust::count_if(handle_.get_thrust_policy(), d_crt_out_degs.begin(), d_crt_out_degs.end(), [] __device__(auto crt_out_deg) { return crt_out_deg == 0; }); @@ -510,13 +509,13 @@ struct random_walker_t { return (col_indx >= ptr_d_sizes[row_indx] - 1); }; - auto new_end_v = thrust::remove_if(rmm::exec_policy(handle_.get_stream_view()), + auto new_end_v = thrust::remove_if(handle_.get_thrust_policy(), d_coalesced_v.begin(), d_coalesced_v.end(), thrust::make_counting_iterator(0), predicate_v); - auto new_end_w = thrust::remove_if(rmm::exec_policy(handle_.get_stream_view()), + auto new_end_w = thrust::remove_if(handle_.get_thrust_policy(), d_coalesced_w.begin(), d_coalesced_w.end(), thrust::make_counting_iterator(0), @@ -556,7 +555,7 @@ struct random_walker_t { auto map_it_begin = thrust::make_transform_iterator(thrust::make_counting_iterator(0), dlambda); - thrust::gather(rmm::exec_policy(handle_.get_stream_view()), + thrust::gather(handle_.get_thrust_policy(), map_it_begin, map_it_begin + nelems, d_src.begin(), @@ -603,7 +602,7 @@ struct random_walker_t { auto map_it_begin = thrust::make_transform_iterator(thrust::make_counting_iterator(0), dlambda); - thrust::scatter_if(rmm::exec_policy(handle_.get_stream_view()), + thrust::scatter_if(handle_.get_thrust_policy(), d_src.begin(), d_src.end(), map_it_begin, @@ -642,7 +641,7 @@ struct random_walker_t { device_vec_t& d_sizes) const { thrust::transform_if( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), d_sizes.begin(), d_sizes.end(), // input d_crt_out_degs.begin(), // stencil @@ -663,12 +662,12 @@ struct random_walker_t { void init_padding(device_vec_t& d_coalesced_v, device_vec_t& d_coalesced_w) const { - thrust::fill(rmm::exec_policy(handle_.get_stream_view()), + thrust::fill(handle_.get_thrust_policy(), d_coalesced_v.begin(), d_coalesced_v.end(), vertex_padding_value_); - thrust::fill(rmm::exec_policy(handle_.get_stream_view()), + thrust::fill(handle_.get_thrust_policy(), d_coalesced_w.begin(), d_coalesced_w.end(), weight_padding_value_); @@ -745,7 +744,7 @@ random_walks_impl(raft::handle_t const& handle, vertex_t num_vertices = graph.get_number_of_vertices(); - auto how_many_valid = thrust::count_if(rmm::exec_policy(handle.get_stream_view()), + auto how_many_valid = thrust::count_if(handle.get_thrust_policy(), d_v_start.begin(), d_v_start.end(), [num_vertices] __device__(auto crt_vertex) { @@ -924,7 +923,7 @@ struct coo_convertor_t { // and edge_paths_sz == 0 don't contribute // anything): // - auto new_end_it = thrust::copy_if(rmm::exec_policy(handle_.get_stream_view()), + auto new_end_it = thrust::copy_if(handle_.get_thrust_policy(), d_sizes.begin(), d_sizes.end(), d_sz_w_scan.begin(), @@ -940,7 +939,7 @@ struct coo_convertor_t { // edge_path_sz = (vertex_path_sz-1): // thrust::transform_exclusive_scan( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), d_sz_w_scan.begin(), d_sz_w_scan.end(), d_sz_w_scan.begin(), @@ -956,7 +955,7 @@ struct coo_convertor_t { { device_vec_t d_scan(num_paths_, handle_.get_stream()); thrust::inclusive_scan( - rmm::exec_policy(handle_.get_stream_view()), d_sizes.begin(), d_sizes.end(), d_scan.begin()); + handle_.get_thrust_policy(), d_sizes.begin(), d_sizes.end(), d_scan.begin()); index_t total_sz{0}; CUDA_TRY(cudaMemcpy( @@ -966,7 +965,7 @@ struct coo_convertor_t { // initialize stencil to all 1's: // - thrust::copy_n(rmm::exec_policy(handle_.get_stream_view()), + thrust::copy_n(handle_.get_thrust_policy(), thrust::make_constant_iterator(1), d_stencil.size(), d_stencil.begin()); @@ -976,7 +975,7 @@ struct coo_convertor_t { // and the next one starts, hence there cannot be an edge // between a path ending vertex and next path starting vertex; // - thrust::scatter(rmm::exec_policy(handle_.get_stream_view()), + thrust::scatter(handle_.get_thrust_policy(), thrust::make_constant_iterator(0), thrust::make_constant_iterator(0) + num_paths_, d_scan.begin(), @@ -999,7 +998,7 @@ struct coo_convertor_t { // in stencil is not 0; (if it is, there's no "next" // or dst index, because the path has ended); // - thrust::copy_if(rmm::exec_policy(handle_.get_stream_view()), + thrust::copy_if(handle_.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(total_sz_v - 1), valid_src_indx.begin(), @@ -1018,7 +1017,7 @@ struct coo_convertor_t { // generated at the previous step; // thrust::transform( - rmm::exec_policy(handle_.get_stream_view()), + handle_.get_thrust_policy(), valid_src_indx.begin(), valid_src_indx.end(), thrust::make_zip_iterator(thrust::make_tuple(d_src_v.begin(), d_dst_v.begin())), // start_zip @@ -1220,12 +1219,10 @@ query_rw_sizes_offsets(raft::handle_t const& handle, index_t num_paths, index_t rmm::device_uvector d_weight_sizes(num_paths, handle.get_stream()); rmm::device_uvector d_weight_offsets(num_paths, handle.get_stream()); - thrust::exclusive_scan(rmm::exec_policy(handle.get_stream_view()), - ptr_d_sizes, - ptr_d_sizes + num_paths, - d_vertex_offsets.begin()); + thrust::exclusive_scan( + handle.get_thrust_policy(), ptr_d_sizes, ptr_d_sizes + num_paths, d_vertex_offsets.begin()); - thrust::transform(rmm::exec_policy(handle.get_stream_view()), + thrust::transform(handle.get_thrust_policy(), ptr_d_sizes, ptr_d_sizes + num_paths, d_weight_sizes.begin(), @@ -1233,7 +1230,7 @@ query_rw_sizes_offsets(raft::handle_t const& handle, index_t num_paths, index_t handle.get_stream_view().synchronize(); - thrust::exclusive_scan(rmm::exec_policy(handle.get_stream_view()), + thrust::exclusive_scan(handle.get_thrust_policy(), d_weight_sizes.begin(), d_weight_sizes.end(), d_weight_offsets.begin()); diff --git a/cpp/src/sampling/rw_traversals.hpp b/cpp/src/sampling/rw_traversals.hpp index f1c5083a98a..3d3ffc4e161 100644 --- a/cpp/src/sampling/rw_traversals.hpp +++ b/cpp/src/sampling/rw_traversals.hpp @@ -31,7 +31,6 @@ #include #include -#include #include #include @@ -459,7 +458,7 @@ struct horizontal_traversal_t { // start from 1, as 0-th was initialized above: // - thrust::for_each(rmm::exec_policy(handle.get_stream_view()), + thrust::for_each(handle.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(num_paths_), [max_depth = max_depth_, diff --git a/cpp/src/serialization/serializer.cu b/cpp/src/serialization/serializer.cu index 2f4c8268a67..5849c50ea47 100644 --- a/cpp/src/serialization/serializer.cu +++ b/cpp/src/serialization/serializer.cu @@ -23,10 +23,6 @@ #include -#include - -#include - #include #include @@ -65,7 +61,7 @@ void serializer_t::serialize(value_t const* p_d_src, size_t size) auto it_end = begin_ + byte_buff_sz; byte_t const* byte_buff = reinterpret_cast(p_d_src); - thrust::copy_n(rmm::exec_policy(handle_.get_stream_view()), byte_buff, byte_buff_sz, begin_); + thrust::copy_n(handle_.get_thrust_policy(), byte_buff, byte_buff_sz, begin_); begin_ = it_end; } @@ -77,7 +73,7 @@ rmm::device_uvector serializer_t::unserialize(size_t size) rmm::device_uvector d_dest(size, handle_.get_stream()); byte_t* byte_buff = reinterpret_cast(d_dest.data()); - thrust::copy_n(rmm::exec_policy(handle_.get_stream_view()), cbegin_, byte_buff_sz, byte_buff); + thrust::copy_n(handle_.get_thrust_policy(), cbegin_, byte_buff_sz, byte_buff); cbegin_ += byte_buff_sz; return d_dest; diff --git a/cpp/src/structure/coarsen_graph.cu b/cpp/src/structure/coarsen_graph.cu index c66cc24932a..727681daa73 100644 --- a/cpp/src/structure/coarsen_graph.cu +++ b/cpp/src/structure/coarsen_graph.cu @@ -62,12 +62,12 @@ decompress_matrix_partition_to_edgelist( decompress_matrix_partition_to_fill_edgelist_majors( handle, matrix_partition, edgelist_major_vertices.data(), segment_offsets); - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(handle.get_thrust_policy(), matrix_partition.get_indices(), matrix_partition.get_indices() + number_of_edges, edgelist_minor_vertices.begin()); if (edgelist_weights) { - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(handle.get_thrust_policy(), *(matrix_partition.get_weights()), *(matrix_partition.get_weights()) + number_of_edges, (*edgelist_weights).data()); @@ -145,7 +145,7 @@ decompress_matrix_partition_to_relabeled_and_grouped_and_coarsened_edgelist( auto pair_first = thrust::make_zip_iterator( thrust::make_tuple(edgelist_major_vertices.begin(), edgelist_minor_vertices.begin())); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), pair_first, pair_first + edgelist_major_vertices.size(), pair_first, @@ -260,10 +260,8 @@ coarsen_graph( if (col_comm_rank == static_cast(i)) { // FIXME: this copy is unnecessary, beter fix RAFT comm's bcast to take const iterators for // input - thrust::copy(rmm::exec_policy(handle.get_stream()), - labels, - labels + major_labels.size(), - major_labels.begin()); + thrust::copy( + handle.get_thrust_policy(), labels, labels + major_labels.size(), major_labels.begin()); } device_bcast(col_comm, major_labels.data(), @@ -332,6 +330,7 @@ coarsen_graph( handle.get_stream()); coarsened_edgelist_minor_vertices[j].resize(coarsened_edgelist_major_vertices[j].size(), handle.get_stream()); + if (coarsened_edgelist_weights) { (*coarsened_edgelist_weights)[j].resize(coarsened_edgelist_major_vertices[j].size(), handle.get_stream()); @@ -346,7 +345,7 @@ coarsen_graph( coarsened_edgelist_minor_vertices[j].begin(), (*coarsened_edgelist_weights)[j].begin())) + cur_size; - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(handle.get_thrust_policy(), src_edge_first, src_edge_first + number_of_partition_edges, dst_edge_first); @@ -358,7 +357,7 @@ coarsen_graph( thrust::make_tuple(coarsened_edgelist_major_vertices[j].begin(), coarsened_edgelist_minor_vertices[j].begin())) + cur_size; - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(handle.get_thrust_policy(), src_edge_first, src_edge_first + number_of_partition_edges, dst_edge_first); @@ -388,25 +387,23 @@ coarsen_graph( rmm::device_uvector unique_labels(graph_view.get_number_of_local_vertices(), handle.get_stream()); - thrust::copy(rmm::exec_policy(handle.get_stream()), - labels, - labels + unique_labels.size(), - unique_labels.begin()); - thrust::sort(rmm::exec_policy(handle.get_stream()), unique_labels.begin(), unique_labels.end()); - unique_labels.resize(thrust::distance(unique_labels.begin(), - thrust::unique(rmm::exec_policy(handle.get_stream()), - unique_labels.begin(), - unique_labels.end())), - handle.get_stream()); + thrust::copy( + handle.get_thrust_policy(), labels, labels + unique_labels.size(), unique_labels.begin()); + thrust::sort(handle.get_thrust_policy(), unique_labels.begin(), unique_labels.end()); + unique_labels.resize( + thrust::distance( + unique_labels.begin(), + thrust::unique(handle.get_thrust_policy(), unique_labels.begin(), unique_labels.end())), + handle.get_stream()); unique_labels = cugraph::detail::shuffle_vertices_by_gpu_id(handle, std::move(unique_labels)); - thrust::sort(rmm::exec_policy(handle.get_stream()), unique_labels.begin(), unique_labels.end()); - unique_labels.resize(thrust::distance(unique_labels.begin(), - thrust::unique(rmm::exec_policy(handle.get_stream()), - unique_labels.begin(), - unique_labels.end())), - handle.get_stream()); + thrust::sort(handle.get_thrust_policy(), unique_labels.begin(), unique_labels.end()); + unique_labels.resize( + thrust::distance( + unique_labels.begin(), + thrust::unique(handle.get_thrust_policy(), unique_labels.begin(), unique_labels.end())), + handle.get_stream()); // 4. renumber @@ -501,16 +498,14 @@ coarsen_graph( rmm::device_uvector unique_labels(graph_view.get_number_of_vertices(), handle.get_stream()); - thrust::copy(rmm::exec_policy(handle.get_stream()), - labels, - labels + unique_labels.size(), - unique_labels.begin()); - thrust::sort(rmm::exec_policy(handle.get_stream()), unique_labels.begin(), unique_labels.end()); - unique_labels.resize(thrust::distance(unique_labels.begin(), - thrust::unique(rmm::exec_policy(handle.get_stream()), - unique_labels.begin(), - unique_labels.end())), - handle.get_stream()); + thrust::copy( + handle.get_thrust_policy(), labels, labels + unique_labels.size(), unique_labels.begin()); + thrust::sort(handle.get_thrust_policy(), unique_labels.begin(), unique_labels.end()); + unique_labels.resize( + thrust::distance( + unique_labels.begin(), + thrust::unique(handle.get_thrust_policy(), unique_labels.begin(), unique_labels.end())), + handle.get_stream()); auto [renumber_map_labels, segment_offsets] = renumber_edgelist( handle, diff --git a/cpp/src/structure/create_graph_from_edgelist.cpp b/cpp/src/structure/create_graph_from_edgelist.cpp index 8d91206671c..d3a385b05bf 100644 --- a/cpp/src/structure/create_graph_from_edgelist.cpp +++ b/cpp/src/structure/create_graph_from_edgelist.cpp @@ -19,7 +19,7 @@ #include #include -#include +#include #include #include @@ -172,8 +172,9 @@ create_graph_from_edgelist_impl(raft::handle_t const& handle, std::tie(*renumber_map_labels, *segment_offsets) = cugraph::renumber_edgelist( handle, - std::optional>{ - std::make_tuple((*vertex_span).data(), static_cast((*vertex_span).size()))}, + vertex_span ? std::optional>{std::make_tuple( + (*vertex_span).data(), static_cast((*vertex_span).size()))} + : std::nullopt, store_transposed ? edgelist_cols.data() : edgelist_rows.data(), store_transposed ? edgelist_rows.data() : edgelist_cols.data(), static_cast(edgelist_rows.size())); diff --git a/cpp/src/structure/graph_view.cu b/cpp/src/structure/graph_view.cu index 088ed214a74..725357e3ec3 100644 --- a/cpp/src/structure/graph_view.cu +++ b/cpp/src/structure/graph_view.cu @@ -562,8 +562,7 @@ graph_view_t ret(edge_t{0}, handle.get_stream()); device_allreduce(handle.get_comms(), it != in_degrees.end() ? it : ret.data(), @@ -588,8 +587,7 @@ edge_t graph_view_t ret(edge_t{0}, handle.get_stream()); device_allreduce(handle.get_comms(), it != out_degrees.end() ? it : ret.data(), @@ -632,8 +629,7 @@ edge_t graph_view_t ret(weight_t{0.0}, handle.get_stream()); device_allreduce(handle.get_comms(), it != in_weight_sums.end() ? it : ret.data(), @@ -676,8 +672,8 @@ weight_t graph_view_t ret(weight_t{0.0}, handle.get_stream()); device_allreduce(handle.get_comms(), it != out_weight_sums.end() ? it : ret.data(), @@ -720,8 +716,8 @@ weight_t graph_view_t< std::enable_if_t>::compute_max_out_weight_sum(raft::handle_t const& handle) const { auto out_weight_sums = compute_out_weight_sums(handle); - auto it = thrust::max_element( - rmm::exec_policy(handle.get_stream_view()), out_weight_sums.begin(), out_weight_sums.end()); + auto it = + thrust::max_element(handle.get_thrust_policy(), out_weight_sums.begin(), out_weight_sums.end()); weight_t ret{0.0}; if (it != out_weight_sums.end()) { raft::update_host(&ret, it, 1, handle.get_stream()); } handle.get_stream_view().synchronize(); diff --git a/cpp/src/structure/induced_subgraph.cu b/cpp/src/structure/induced_subgraph.cu index d56a46af4e8..952ffddec3a 100644 --- a/cpp/src/structure/induced_subgraph.cu +++ b/cpp/src/structure/induced_subgraph.cu @@ -22,7 +22,6 @@ #include #include -#include #include #include @@ -76,13 +75,13 @@ extract_induced_subgraphs( CUGRAPH_EXPECTS(should_be_zero == 0, "Invalid input argument: subgraph_offsets[0] should be 0."); - CUGRAPH_EXPECTS(thrust::is_sorted(rmm::exec_policy(handle.get_stream_view()), - 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, subgraph_offsets + (num_subgraphs + 1)), + "Invalid input argument: subgraph_offsets is not sorted."); auto vertex_partition = vertex_partition_device_view_t(graph_view.get_vertex_partition_view()); - CUGRAPH_EXPECTS(thrust::count_if(rmm::exec_policy(handle.get_stream_view()), + CUGRAPH_EXPECTS(thrust::count_if(handle.get_thrust_policy(), subgraph_vertices, subgraph_vertices + num_aggregate_subgraph_vertices, [vertex_partition] __device__(auto v) { @@ -93,7 +92,7 @@ extract_induced_subgraphs( CUGRAPH_EXPECTS( thrust::count_if( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), thrust::make_counting_iterator(size_t{0}), thrust::make_counting_iterator(num_subgraphs), [subgraph_offsets, subgraph_vertices] __device__(auto i) { @@ -138,7 +137,7 @@ extract_induced_subgraphs( // count the numbers of the induced subgraph edges for each vertex in the aggregate subgraph // vertex list. thrust::transform( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), thrust::make_counting_iterator(size_t{0}), thrust::make_counting_iterator(num_aggregate_subgraph_vertices), subgraph_vertex_output_offsets.begin(), @@ -164,7 +163,7 @@ extract_induced_subgraphs( return thrust::binary_search(thrust::seq, vertex_first, vertex_last, nbr); }); }); - thrust::exclusive_scan(rmm::exec_policy(handle.get_stream_view()), + thrust::exclusive_scan(handle.get_thrust_policy(), subgraph_vertex_output_offsets.begin(), subgraph_vertex_output_offsets.end(), subgraph_vertex_output_offsets.begin()); @@ -188,7 +187,7 @@ extract_induced_subgraphs( // fill the edge list buffer (to be returned) for each vetex in the aggregate subgraph vertex // list (use the offsets computed in the Phase 1) thrust::for_each( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), thrust::make_counting_iterator(size_t{0}), thrust::make_counting_iterator(num_aggregate_subgraph_vertices), [subgraph_offsets, @@ -246,7 +245,7 @@ extract_induced_subgraphs( }); rmm::device_uvector subgraph_edge_offsets(num_subgraphs + 1, handle.get_stream_view()); - thrust::gather(rmm::exec_policy(handle.get_stream_view()), + thrust::gather(handle.get_thrust_policy(), subgraph_offsets, subgraph_offsets + (num_subgraphs + 1), subgraph_vertex_output_offsets.begin(), diff --git a/cpp/src/structure/relabel.cu b/cpp/src/structure/relabel.cu index d01143a922e..230fc691628 100644 --- a/cpp/src/structure/relabel.cu +++ b/cpp/src/structure/relabel.cu @@ -25,7 +25,6 @@ #include #include #include -#include #include #include @@ -63,19 +62,13 @@ void relabel(raft::handle_t const& handle, // find unique old labels (to be relabeled) rmm::device_uvector unique_old_labels(num_labels, handle.get_stream_view()); - thrust::copy(rmm::exec_policy(handle.get_stream_view()), - labels, - labels + num_labels, - unique_old_labels.data()); - thrust::sort(rmm::exec_policy(handle.get_stream_view()), - unique_old_labels.begin(), - unique_old_labels.end()); - unique_old_labels.resize( - thrust::distance(unique_old_labels.begin(), - thrust::unique(rmm::exec_policy(handle.get_stream_view()), - unique_old_labels.begin(), - unique_old_labels.end())), - handle.get_stream_view()); + thrust::copy(handle.get_thrust_policy(), labels, labels + num_labels, unique_old_labels.data()); + thrust::sort(handle.get_thrust_policy(), unique_old_labels.begin(), unique_old_labels.end()); + unique_old_labels.resize(thrust::distance(unique_old_labels.begin(), + thrust::unique(handle.get_thrust_policy(), + unique_old_labels.begin(), + unique_old_labels.end())), + handle.get_stream_view()); unique_old_labels.shrink_to_fit(handle.get_stream_view()); // collect new labels for the unique old labels @@ -92,11 +85,11 @@ void relabel(raft::handle_t const& handle, handle.get_stream_view()); rmm::device_uvector label_pair_new_labels(num_label_pairs, handle.get_stream_view()); - thrust::copy(rmm::exec_policy(handle.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), std::get<0>(old_new_label_pairs), std::get<0>(old_new_label_pairs) + num_label_pairs, label_pair_old_labels.begin()); - thrust::copy(rmm::exec_policy(handle.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), std::get<1>(old_new_label_pairs), std::get<1>(old_new_label_pairs) + num_label_pairs, label_pair_new_labels.begin()); @@ -153,7 +146,7 @@ void relabel(raft::handle_t const& handle, handle.get_stream_view().synchronize(); // cuco::static_map currently does not take stream if (skip_missing_labels) { - thrust::transform(rmm::exec_policy(handle.get_stream_view()), + thrust::transform(handle.get_thrust_policy(), rx_unique_old_labels.begin(), rx_unique_old_labels.end(), rx_unique_old_labels.begin(), @@ -212,7 +205,7 @@ void relabel(raft::handle_t const& handle, thrust::make_tuple(std::get<0>(old_new_label_pairs), std::get<1>(old_new_label_pairs))); relabel_map.insert(pair_first, pair_first + num_label_pairs); if (skip_missing_labels) { - thrust::transform(rmm::exec_policy(handle.get_stream_view()), + thrust::transform(handle.get_thrust_policy(), labels, labels + num_labels, labels, @@ -229,7 +222,7 @@ void relabel(raft::handle_t const& handle, if (do_expensive_check && !skip_missing_labels) { CUGRAPH_EXPECTS( - thrust::count(rmm::exec_policy(handle.get_stream_view()), + thrust::count(handle.get_thrust_policy(), labels, labels + num_labels, invalid_vertex_id::value) == 0, diff --git a/cpp/src/structure/renumber_edgelist.cu b/cpp/src/structure/renumber_edgelist.cu index 3cf9954926b..4123bb5f218 100644 --- a/cpp/src/structure/renumber_edgelist.cu +++ b/cpp/src/structure/renumber_edgelist.cu @@ -26,7 +26,6 @@ #include #include #include -#include #include #include @@ -82,16 +81,15 @@ std::tuple, std::vector> compute_renumbe { rmm::device_uvector sorted_major_labels(edgelist_edge_counts[i], handle.get_stream()); - thrust::copy(rmm::exec_policy(handle.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), edgelist_major_vertices[i], edgelist_major_vertices[i] + edgelist_edge_counts[i], sorted_major_labels.begin()); // FIXME: better refactor this sort-count_if-reduce_by_key routine for reuse - thrust::sort(rmm::exec_policy(handle.get_stream_view()), - sorted_major_labels.begin(), - sorted_major_labels.end()); + thrust::sort( + handle.get_thrust_policy(), sorted_major_labels.begin(), sorted_major_labels.end()); auto num_unique_labels = - thrust::count_if(rmm::exec_policy(handle.get_stream_view()), + thrust::count_if(handle.get_thrust_policy(), thrust::make_counting_iterator(size_t{0}), thrust::make_counting_iterator(sorted_major_labels.size()), [labels = sorted_major_labels.data()] __device__(auto i) { @@ -99,7 +97,7 @@ std::tuple, std::vector> compute_renumbe }); tmp_major_labels.resize(num_unique_labels, handle.get_stream()); tmp_major_counts.resize(tmp_major_labels.size(), handle.get_stream()); - thrust::reduce_by_key(rmm::exec_policy(handle.get_stream_view()), + thrust::reduce_by_key(handle.get_thrust_policy(), sorted_major_labels.begin(), sorted_major_labels.end(), thrust::make_constant_iterator(edge_t{1}), @@ -145,11 +143,9 @@ std::tuple, std::vector> compute_renumbe } if (multi_gpu) { // FIXME: better refactor this sort-count_if-reduce_by_key routine for reuse - thrust::sort_by_key(rmm::exec_policy(handle.get_stream_view()), - major_labels.begin(), - major_labels.end(), - major_counts.begin()); - auto num_unique_labels = thrust::count_if(rmm::exec_policy(handle.get_stream_view()), + thrust::sort_by_key( + handle.get_thrust_policy(), major_labels.begin(), major_labels.end(), major_counts.begin()); + auto num_unique_labels = thrust::count_if(handle.get_thrust_policy(), thrust::make_counting_iterator(size_t{0}), thrust::make_counting_iterator(major_labels.size()), [labels = major_labels.data()] __device__(auto i) { @@ -157,7 +153,7 @@ std::tuple, std::vector> compute_renumbe }); rmm::device_uvector tmp_major_labels(num_unique_labels, handle.get_stream()); rmm::device_uvector tmp_major_counts(tmp_major_labels.size(), handle.get_stream()); - thrust::reduce_by_key(rmm::exec_policy(handle.get_stream_view()), + thrust::reduce_by_key(handle.get_thrust_policy(), major_labels.begin(), major_labels.end(), major_counts.begin(), @@ -175,18 +171,17 @@ std::tuple, std::vector> compute_renumbe rmm::device_uvector minor_labels(minor_displs.back() + edgelist_edge_counts.back(), handle.get_stream()); for (size_t i = 0; i < edgelist_minor_vertices.size(); ++i) { - thrust::copy(rmm::exec_policy(handle.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), edgelist_minor_vertices[i], edgelist_minor_vertices[i] + edgelist_edge_counts[i], minor_labels.begin() + minor_displs[i]); } - thrust::sort( - rmm::exec_policy(handle.get_stream_view()), minor_labels.begin(), minor_labels.end()); - minor_labels.resize(thrust::distance(minor_labels.begin(), - thrust::unique(rmm::exec_policy(handle.get_stream_view()), - minor_labels.begin(), - minor_labels.end())), - handle.get_stream()); + thrust::sort(handle.get_thrust_policy(), minor_labels.begin(), minor_labels.end()); + minor_labels.resize( + thrust::distance( + minor_labels.begin(), + thrust::unique(handle.get_thrust_policy(), minor_labels.begin(), minor_labels.end())), + handle.get_stream()); if (multi_gpu) { auto& comm = handle.get_comms(); auto& row_comm = handle.get_subcomm(cugraph::partition_2d::key_naming_t().row_name()); @@ -212,14 +207,12 @@ std::tuple, std::vector> compute_renumbe [key_func = detail::compute_gpu_id_from_vertex_t{row_comm_size}] __device__( auto val) { return key_func(val); }, handle.get_stream()); - thrust::sort( - rmm::exec_policy(handle.get_stream_view()), rx_minor_labels.begin(), rx_minor_labels.end()); - rx_minor_labels.resize( - thrust::distance(rx_minor_labels.begin(), - thrust::unique(rmm::exec_policy(handle.get_stream_view()), - rx_minor_labels.begin(), - rx_minor_labels.end())), - handle.get_stream()); + thrust::sort(handle.get_thrust_policy(), rx_minor_labels.begin(), rx_minor_labels.end()); + rx_minor_labels.resize(thrust::distance(rx_minor_labels.begin(), + thrust::unique(handle.get_thrust_policy(), + rx_minor_labels.begin(), + rx_minor_labels.end())), + handle.get_stream()); minor_labels = std::move(rx_minor_labels); } @@ -242,7 +235,7 @@ std::tuple, std::vector> compute_renumbe rmm::device_uvector merged_labels(major_labels.size() + minor_labels.size(), handle.get_stream_view()); rmm::device_uvector merged_counts(merged_labels.size(), handle.get_stream_view()); - thrust::merge_by_key(rmm::exec_policy(handle.get_stream_view()), + thrust::merge_by_key(handle.get_thrust_policy(), major_labels.begin(), major_labels.end(), minor_labels.begin(), @@ -261,7 +254,7 @@ std::tuple, std::vector> compute_renumbe rmm::device_uvector labels(merged_labels.size(), handle.get_stream()); rmm::device_uvector counts(labels.size(), handle.get_stream()); - auto pair_it = thrust::reduce_by_key(rmm::exec_policy(handle.get_stream_view()), + auto pair_it = thrust::reduce_by_key(handle.get_thrust_policy(), merged_labels.begin(), merged_labels.end(), merged_counts.begin(), @@ -285,14 +278,14 @@ std::tuple, std::vector> compute_renumbe auto [vertices, num_vertices] = *vertex_span; auto num_isolated_vertices = thrust::count_if( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), vertices, vertices + num_vertices, [label_first = labels.begin(), label_last = labels.end()] __device__(auto v) { return !thrust::binary_search(thrust::seq, label_first, label_last, v); }); isolated_vertices.resize(num_isolated_vertices, handle.get_stream()); - thrust::copy_if(rmm::exec_policy(handle.get_stream_view()), + thrust::copy_if(handle.get_thrust_policy(), vertices, vertices + num_vertices, isolated_vertices.begin(), @@ -302,7 +295,7 @@ std::tuple, std::vector> compute_renumbe if (isolated_vertices.size() > 0) { labels.resize(labels.size() + isolated_vertices.size(), handle.get_stream()); - thrust::copy(rmm::exec_policy(handle.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), isolated_vertices.begin(), isolated_vertices.end(), labels.end() - isolated_vertices.size()); @@ -311,7 +304,7 @@ std::tuple, std::vector> compute_renumbe // 5. sort non-isolated vertices by degree - thrust::sort_by_key(rmm::exec_policy(handle.get_stream_view()), + thrust::sort_by_key(handle.get_thrust_policy(), counts.begin(), counts.begin() + num_non_isolated_vertices, labels.begin(), @@ -356,7 +349,7 @@ std::tuple, std::vector> compute_renumbe d_segment_offsets.set_element_async( num_segments_per_vertex_partition, vertex_count, handle.get_stream()); - thrust::upper_bound(rmm::exec_policy(handle.get_stream()), + thrust::upper_bound(handle.get_thrust_policy(), counts.begin(), counts.end(), d_thresholds.begin(), @@ -387,16 +380,13 @@ void expensive_check_edgelist( if (vertex_span) { auto [vertices, num_vertices] = *vertex_span; sorted_local_vertices.resize(num_vertices, handle.get_stream()); - thrust::copy(rmm::exec_policy(handle.get_stream_view()), - vertices, - vertices + num_vertices, - sorted_local_vertices.begin()); - thrust::sort(rmm::exec_policy(handle.get_stream_view()), - sorted_local_vertices.begin(), - sorted_local_vertices.end()); + thrust::copy( + handle.get_thrust_policy(), vertices, vertices + num_vertices, sorted_local_vertices.begin()); + thrust::sort( + handle.get_thrust_policy(), sorted_local_vertices.begin(), sorted_local_vertices.end()); CUGRAPH_EXPECTS(static_cast(thrust::distance( sorted_local_vertices.begin(), - thrust::unique(rmm::exec_policy(handle.get_stream_view()), + thrust::unique(handle.get_thrust_policy(), sorted_local_vertices.begin(), sorted_local_vertices.end()))) == sorted_local_vertices.size(), "Invalid input argument: local_vertices should not have duplicates."); @@ -421,7 +411,7 @@ void expensive_check_edgelist( auto [local_vertices, num_local_vertices] = *vertex_span; CUGRAPH_EXPECTS( thrust::count_if( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), local_vertices, local_vertices + num_local_vertices, [comm_rank, @@ -436,7 +426,7 @@ void expensive_check_edgelist( thrust::make_tuple(edgelist_major_vertices[i], edgelist_minor_vertices[i])); CUGRAPH_EXPECTS( thrust::count_if( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), edge_first, edge_first + edgelist_edge_counts[i], [comm_size, @@ -488,9 +478,8 @@ void expensive_check_edgelist( recvcounts, displacements, handle.get_stream()); - thrust::sort(rmm::exec_policy(handle.get_stream_view()), - sorted_major_vertices.begin(), - sorted_major_vertices.end()); + thrust::sort( + handle.get_thrust_policy(), sorted_major_vertices.begin(), sorted_major_vertices.end()); } // barrier is necessary here to avoid potential overlap (which can leads to deadlock) @@ -518,9 +507,8 @@ void expensive_check_edgelist( recvcounts, displacements, handle.get_stream()); - thrust::sort(rmm::exec_policy(handle.get_stream_view()), - sorted_minor_vertices.begin(), - sorted_minor_vertices.end()); + thrust::sort( + handle.get_thrust_policy(), sorted_minor_vertices.begin(), sorted_minor_vertices.end()); } // barrier is necessary here to avoid potential overlap (which can leads to deadlock) @@ -538,7 +526,7 @@ void expensive_check_edgelist( thrust::make_tuple(edgelist_major_vertices[i], edgelist_minor_vertices[i])); CUGRAPH_EXPECTS( thrust::count_if( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), edge_first, edge_first + edgelist_edge_counts[i], [num_major_vertices = static_cast(sorted_major_vertices.size()), @@ -587,7 +575,7 @@ void expensive_check_edgelist( thrust::make_tuple(edgelist_major_vertices[0], edgelist_minor_vertices[0])); CUGRAPH_EXPECTS( thrust::count_if( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), edge_first, edge_first + edgelist_edge_counts[0], [sorted_local_vertices = sorted_local_vertices.data(), diff --git a/cpp/src/structure/renumber_utils.cu b/cpp/src/structure/renumber_utils.cu index 90d9c7e7a43..eb9e535adeb 100644 --- a/cpp/src/structure/renumber_utils.cu +++ b/cpp/src/structure/renumber_utils.cu @@ -50,14 +50,13 @@ void renumber_ext_vertices(raft::handle_t const& handle, if (do_expensive_check) { rmm::device_uvector labels(local_int_vertex_last - local_int_vertex_first, handle.get_stream_view()); - thrust::copy(rmm::exec_policy(handle.get_stream_view()), + thrust::copy(handle.get_thrust_policy(), renumber_map_labels, renumber_map_labels + labels.size(), labels.begin()); - thrust::sort(rmm::exec_policy(handle.get_stream_view()), labels.begin(), labels.end()); + thrust::sort(handle.get_thrust_policy(), labels.begin(), labels.end()); CUGRAPH_EXPECTS( - thrust::unique(rmm::exec_policy(handle.get_stream_view()), labels.begin(), labels.end()) == - labels.end(), + thrust::unique(handle.get_thrust_policy(), labels.begin(), labels.end()) == labels.end(), "Invalid input arguments: renumber_map_labels have duplicate elements."); } @@ -78,18 +77,18 @@ void renumber_ext_vertices(raft::handle_t const& handle, sorted_unique_ext_vertices.resize( thrust::distance( sorted_unique_ext_vertices.begin(), - thrust::copy_if(rmm::exec_policy(handle.get_stream_view()), + thrust::copy_if(handle.get_thrust_policy(), vertices, vertices + num_vertices, sorted_unique_ext_vertices.begin(), [] __device__(auto v) { return v != invalid_vertex_id::value; })), handle.get_stream_view()); - thrust::sort(rmm::exec_policy(handle.get_stream_view()), + thrust::sort(handle.get_thrust_policy(), sorted_unique_ext_vertices.begin(), sorted_unique_ext_vertices.end()); sorted_unique_ext_vertices.resize( thrust::distance(sorted_unique_ext_vertices.begin(), - thrust::unique(rmm::exec_policy(handle.get_stream_view()), + thrust::unique(handle.get_thrust_policy(), sorted_unique_ext_vertices.begin(), sorted_unique_ext_vertices.end())), handle.get_stream_view()); @@ -146,7 +145,7 @@ void renumber_ext_vertices(raft::handle_t const& handle, rmm::device_uvector contains(num_vertices, handle.get_stream_view()); renumber_map_ptr->contains(vertices, vertices + num_vertices, contains.begin()); auto vc_pair_first = thrust::make_zip_iterator(thrust::make_tuple(vertices, contains.begin())); - CUGRAPH_EXPECTS(thrust::count_if(rmm::exec_policy(handle.get_stream_view()), + CUGRAPH_EXPECTS(thrust::count_if(handle.get_thrust_policy(), vc_pair_first, vc_pair_first + num_vertices, [] __device__(auto pair) { @@ -175,7 +174,7 @@ void unrenumber_local_int_vertices( { if (do_expensive_check) { CUGRAPH_EXPECTS( - thrust::count_if(rmm::exec_policy(handle.get_stream_view()), + thrust::count_if(handle.get_thrust_policy(), vertices, vertices + num_vertices, [local_int_vertex_first, local_int_vertex_last] __device__(auto v) { @@ -186,7 +185,7 @@ void unrenumber_local_int_vertices( "+ num_vertices)."); } - thrust::transform(rmm::exec_policy(handle.get_stream_view()), + thrust::transform(handle.get_thrust_policy(), vertices, vertices + num_vertices, vertices, @@ -211,7 +210,7 @@ void unrenumber_int_vertices(raft::handle_t const& handle, if (do_expensive_check) { CUGRAPH_EXPECTS( - thrust::count_if(rmm::exec_policy(handle.get_stream_view()), + thrust::count_if(handle.get_thrust_policy(), vertices, vertices + num_vertices, [int_vertex_last = vertex_partition_lasts.back()] __device__(auto v) { @@ -231,18 +230,18 @@ void unrenumber_int_vertices(raft::handle_t const& handle, sorted_unique_int_vertices.resize( thrust::distance( sorted_unique_int_vertices.begin(), - thrust::copy_if(rmm::exec_policy(handle.get_stream_view()), + thrust::copy_if(handle.get_thrust_policy(), vertices, vertices + num_vertices, sorted_unique_int_vertices.begin(), [] __device__(auto v) { return v != invalid_vertex_id::value; })), handle.get_stream_view()); - thrust::sort(rmm::exec_policy(handle.get_stream_view()), + thrust::sort(handle.get_thrust_policy(), sorted_unique_int_vertices.begin(), sorted_unique_int_vertices.end()); sorted_unique_int_vertices.resize( thrust::distance(sorted_unique_int_vertices.begin(), - thrust::unique(rmm::exec_policy(handle.get_stream_view()), + thrust::unique(handle.get_thrust_policy(), sorted_unique_int_vertices.begin(), sorted_unique_int_vertices.end())), handle.get_stream_view()); @@ -255,7 +254,7 @@ void unrenumber_int_vertices(raft::handle_t const& handle, handle.get_stream()); rmm::device_uvector d_tx_int_vertex_offsets(d_vertex_partition_lasts.size(), handle.get_stream_view()); - thrust::lower_bound(rmm::exec_policy(handle.get_stream_view()), + thrust::lower_bound(handle.get_thrust_policy(), sorted_unique_int_vertices.begin(), sorted_unique_int_vertices.end(), d_vertex_partition_lasts.begin(), @@ -276,7 +275,7 @@ void unrenumber_int_vertices(raft::handle_t const& handle, comm, sorted_unique_int_vertices.begin(), h_tx_int_vertex_counts, handle.get_stream_view()); auto tx_ext_vertices = std::move(rx_int_vertices); - thrust::transform(rmm::exec_policy(handle.get_stream_view()), + thrust::transform(handle.get_thrust_policy(), tx_ext_vertices.begin(), tx_ext_vertices.end(), tx_ext_vertices.begin(), diff --git a/cpp/src/topology/topology.cuh b/cpp/src/topology/topology.cuh index c3b6c8bae5c..a06a325680c 100644 --- a/cpp/src/topology/topology.cuh +++ b/cpp/src/topology/topology.cuh @@ -73,13 +73,13 @@ bool check_symmetry(raft::handle_t const& handle, { using BoolT = bool; rmm::device_uvector d_flags(nrows, handle.get_stream()); - thrust::fill_n(rmm::exec_policy(handle.get_stream_view()), d_flags.begin(), nrows, true); + thrust::fill_n(handle.get_thrust_policy(), d_flags.begin(), nrows, true); BoolT* start_flags = d_flags.data(); // d_flags.begin(); BoolT* end_flags = start_flags + nrows; BoolT init{1}; return thrust::transform_reduce( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), start_flags, end_flags, [ptr_r_o, ptr_c_i, start_flags, nnz] __device__(BoolT & crt_flag) { @@ -145,13 +145,12 @@ struct thrust_segment_sorter_by_weights_t { // cannot use counting iterator, because d_keys gets passed to sort-by-key() // - thrust::sequence( - rmm::exec_policy(handle_.get_stream_view()), d_keys.begin(), d_keys.end(), edge_t{0}); + thrust::sequence(handle.get_thrust_policy(), d_keys.begin(), d_keys.end(), edge_t{0}); // d_segs = map each key(i.e., edge index), to corresponding // segment (i.e., partition = out-going set) index // - thrust::upper_bound(rmm::exec_policy(handle_.get_stream_view()), + thrust::upper_bound(handle.get_thrust_policy(), ptr_d_offsets_, ptr_d_offsets_ + num_vertices_ + 1, d_keys.begin(), @@ -159,7 +158,7 @@ struct thrust_segment_sorter_by_weights_t { d_segs.begin()); thrust::sort_by_key( - rmm::exec_policy(handle_.get_stream_view()), + handle.get_thrust_policy(), d_keys.begin(), d_keys.end(), thrust::make_zip_iterator(thrust::make_tuple(ptr_d_indices_, ptr_d_weights_)), @@ -309,7 +308,7 @@ bool check_segmented_sort(raft::handle_t const& handle, // that are _not_ ordered increasingly: // auto it = thrust::find_if( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), thrust::make_counting_iterator(0), end, [ptr_d_segs, ptr_d_weights] __device__(auto indx) { @@ -346,13 +345,12 @@ bool check_segmented_sort(raft::handle_t const& handle, // cannot use counting iterator, because d_keys gets passed to sort-by-key() // - thrust::sequence( - rmm::exec_policy(handle.get_stream_view()), d_keys.begin(), d_keys.end(), edge_t{0}); + thrust::sequence(handle.get_thrust_policy(), d_keys.begin(), d_keys.end(), edge_t{0}); // d_segs = map each key(i.e., edge index), to corresponding // segment (i.e., partition = out-going set) index // - thrust::upper_bound(rmm::exec_policy(handle.get_stream_view()), + thrust::upper_bound(handle.get_thrust_policy(), ptr_d_offsets, ptr_d_offsets + num_vertices + 1, d_keys.begin(), diff --git a/cpp/src/traversal/bfs.cu b/cpp/src/traversal/bfs.cu index fa653b7ddb3..1781cc10738 100644 --- a/cpp/src/traversal/bfs.cu +++ b/cpp/src/traversal/bfs.cu @@ -78,7 +78,7 @@ void bfs(raft::handle_t const& handle, auto constexpr invalid_vertex = invalid_vertex_id::value; auto val_first = thrust::make_zip_iterator(thrust::make_tuple(distances, predecessor_first)); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), thrust::make_counting_iterator(push_graph_view.get_local_vertex_first()), thrust::make_counting_iterator(push_graph_view.get_local_vertex_last()), val_first, diff --git a/cpp/src/traversal/legacy/bfs.cuh b/cpp/src/traversal/legacy/bfs.cuh index 6fec3bde68d..dd636a2c97c 100644 --- a/cpp/src/traversal/legacy/bfs.cuh +++ b/cpp/src/traversal/legacy/bfs.cuh @@ -13,7 +13,6 @@ #include #include -#include #define TRAVERSAL_DEFAULT_ALPHA 15 diff --git a/cpp/src/traversal/legacy/mg/bfs.cuh b/cpp/src/traversal/legacy/mg/bfs.cuh index 1059a8fa1de..5e53ea78e04 100644 --- a/cpp/src/traversal/legacy/mg/bfs.cuh +++ b/cpp/src/traversal/legacy/mg/bfs.cuh @@ -84,7 +84,7 @@ void bfs_traverse(raft::handle_t const& handle, input_frontier.swap(output_frontier); // Clear output frontier bitmap - thrust::fill(rmm::exec_policy(stream), + thrust::fill(handle.get_thrust_policy(), output_frontier_bmap.begin(), output_frontier_bmap.end(), static_cast(0)); @@ -130,7 +130,7 @@ void bfs(raft::handle_t const& handle, cudaStream_t stream = handle.get_stream(); // Set all predecessors to be invalid vertex ids - thrust::fill(rmm::exec_policy(stream), + thrust::fill(handle.get_thrust_policy(), predecessors, predecessors + global_number_of_vertices, cugraph::legacy::invalid_idx::value); diff --git a/cpp/src/traversal/legacy/mg/common_utils.cuh b/cpp/src/traversal/legacy/mg/common_utils.cuh index 6c1a4514619..ad3c6b71659 100644 --- a/cpp/src/traversal/legacy/mg/common_utils.cuh +++ b/cpp/src/traversal/legacy/mg/common_utils.cuh @@ -147,9 +147,8 @@ vertex_t populate_isolated_vertices( cugraph::legacy::GraphCSRView const& graph, rmm::device_vector& isolated_vertex_ids) { - bool is_mg = (handle.comms_initialized() && (graph.local_vertices != nullptr) && + bool is_mg = (handle.comms_initialized() && (graph.local_vertices != nullptr) && (graph.local_offsets != nullptr)); - cudaStream_t stream = handle.get_stream(); edge_t vertex_begin_, vertex_end_; if (is_mg) { @@ -160,7 +159,7 @@ vertex_t populate_isolated_vertices( vertex_begin_ = 0; vertex_end_ = graph.number_of_vertices; } - auto count = thrust::copy_if(rmm::exec_policy(stream), + auto count = thrust::copy_if(handle.get_thrust_policy(), thrust::make_counting_iterator(vertex_begin_), thrust::make_counting_iterator(vertex_end_), thrust::make_counting_iterator(0), @@ -214,7 +213,7 @@ void add_to_bitmap(raft::handle_t const& handle, { cudaStream_t stream = handle.get_stream(); thrust::for_each( - rmm::exec_policy(stream), id.begin(), id.begin() + count, set_nth_bit(bmap.data().get())); + handle.get_thrust_policy(), id.begin(), id.begin() + count, set_nth_bit(bmap.data().get())); CHECK_CUDA(stream); } @@ -246,9 +245,10 @@ return_t remove_duplicates(raft::handle_t const& handle, return_t data_len) { cudaStream_t stream = handle.get_stream(); - thrust::sort(rmm::exec_policy(stream), data.begin(), data.begin() + data_len); + thrust::sort(handle.get_thrust_policy(), data.begin(), data.begin() + data_len); auto unique_count = - thrust::unique(rmm::exec_policy(stream), data.begin(), data.begin() + data_len) - data.begin(); + thrust::unique(handle.get_thrust_policy(), data.begin(), data.begin() + data_len) - + data.begin(); return static_cast(unique_count); } @@ -370,7 +370,7 @@ return_t remove_duplicates(raft::handle_t const& handle, rmm::device_vector unique_count(1, 0); - thrust::fill(rmm::exec_policy(stream), bmap.begin(), bmap.end(), static_cast(0)); + thrust::fill(handle.get_thrust_policy(), bmap.begin(), bmap.end(), static_cast(0)); constexpr return_t threads = 256; return_t blocks = raft::div_rounding_up_safe(data_len, threads); remove_duplicates_kernel<<>>(bmap.data().get(), @@ -401,7 +401,7 @@ vertex_t preprocess_input_frontier( graph.local_vertices[handle.get_comms().get_rank()]; rmm::device_vector unique_count(1, 0); - thrust::fill(rmm::exec_policy(stream), bmap.begin(), bmap.end(), static_cast(0)); + thrust::fill(handle.get_thrust_policy(), bmap.begin(), bmap.end(), static_cast(0)); constexpr vertex_t threads = 256; vertex_t blocks = raft::div_rounding_up_safe(input_frontier_len, threads); remove_duplicates_kernel<<>>(bmap.data().get(), @@ -432,7 +432,7 @@ vertex_t preprocess_input_frontier( graph.local_vertices[handle.get_comms().get_rank()]; rmm::device_vector unique_count(1, 0); - thrust::fill(rmm::exec_policy(stream), bmap.begin(), bmap.end(), static_cast(0)); + thrust::fill(handle.get_thrust_policy(), bmap.begin(), bmap.end(), static_cast(0)); constexpr vertex_t threads = 256; vertex_t blocks = raft::div_rounding_up_safe(input_frontier_len, threads); remove_duplicates_kernel<<>>(bmap.data().get(), @@ -479,7 +479,7 @@ vertex_t get_global_vertex_count( { rmm::device_vector id(1); id[0] = *thrust::max_element( - rmm::exec_policy(handle.get_stream()), graph.indices, graph.indices + graph.number_of_edges); + handle.get_thrust_policy(), graph.indices, graph.indices + graph.number_of_edges); handle.get_comms().allreduce( id.data().get(), id.data().get(), 1, raft::comms::op_t::MAX, handle.get_stream()); vertex_t max_vertex_id = id[0]; diff --git a/cpp/src/traversal/legacy/mg/vertex_binning_kernels.cuh b/cpp/src/traversal/legacy/mg/vertex_binning_kernels.cuh index 012c65785a7..6af2df61f14 100644 --- a/cpp/src/traversal/legacy/mg/vertex_binning_kernels.cuh +++ b/cpp/src/traversal/legacy/mg/vertex_binning_kernels.cuh @@ -17,7 +17,6 @@ #pragma once #include -#include #include "../traversal_common.cuh" namespace cugraph { diff --git a/cpp/src/traversal/legacy/sssp.cuh b/cpp/src/traversal/legacy/sssp.cuh index 26388136eb4..c14f1f33708 100644 --- a/cpp/src/traversal/legacy/sssp.cuh +++ b/cpp/src/traversal/legacy/sssp.cuh @@ -18,7 +18,6 @@ #pragma once #include -#include namespace cugraph { namespace detail { diff --git a/cpp/src/traversal/sssp.cu b/cpp/src/traversal/sssp.cu index 4301bcec431..742218b5214 100644 --- a/cpp/src/traversal/sssp.cu +++ b/cpp/src/traversal/sssp.cu @@ -26,7 +26,6 @@ #include #include -#include #include #include @@ -93,7 +92,7 @@ void sssp(raft::handle_t const& handle, auto constexpr invalid_vertex = invalid_vertex_id::value; auto val_first = thrust::make_zip_iterator(thrust::make_tuple(distances, predecessor_first)); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), thrust::make_counting_iterator(push_graph_view.get_local_vertex_first()), thrust::make_counting_iterator(push_graph_view.get_local_vertex_last()), val_first, @@ -143,7 +142,7 @@ void sssp(raft::handle_t const& handle, if (!vertex_and_adj_matrix_row_ranges_coincide) { adj_matrix_row_distances.resize(push_graph_view.get_number_of_local_adj_matrix_partition_rows(), handle.get_stream()); - thrust::fill(rmm::exec_policy(handle.get_stream()), + thrust::fill(handle.get_thrust_policy(), adj_matrix_row_distances.begin(), adj_matrix_row_distances.end(), std::numeric_limits::max()); @@ -379,5 +378,4 @@ template void sssp(raft::handle_t const& handle, int64_t source_vertex, double cutoff, bool do_expensive_check); - } // namespace cugraph diff --git a/cpp/src/traversal/two_hop_neighbors.cu b/cpp/src/traversal/two_hop_neighbors.cu index ab6f5bcfaff..c6ef012b1d3 100644 --- a/cpp/src/traversal/two_hop_neighbors.cu +++ b/cpp/src/traversal/two_hop_neighbors.cu @@ -26,7 +26,6 @@ #include #include "two_hop_neighbors.cuh" -#include #include #include diff --git a/cpp/src/utilities/cython.cu b/cpp/src/utilities/cython.cu index 12948373192..90ca5090fe8 100644 --- a/cpp/src/utilities/cython.cu +++ b/cpp/src/utilities/cython.cu @@ -92,7 +92,7 @@ std::vector compute_edge_counts(raft::handle_t const& handle, major_vertices, compute_local_partition_id_t{d_lasts.data(), num_local_partitions}); rmm::device_uvector d_local_partition_ids(num_local_partitions, handle.get_stream()); rmm::device_uvector d_edge_counts(d_local_partition_ids.size(), handle.get_stream()); - auto it = thrust::reduce_by_key(rmm::exec_policy(handle.get_stream()), + auto it = thrust::reduce_by_key(handle.get_thrust_policy(), key_first, key_first + graph_container.num_local_edges, thrust::make_constant_iterator(edge_t{1}), @@ -101,9 +101,8 @@ std::vector compute_edge_counts(raft::handle_t const& handle, if (static_cast(thrust::distance(d_local_partition_ids.begin(), thrust::get<0>(it))) < num_local_partitions) { rmm::device_uvector d_counts(num_local_partitions, handle.get_stream()); - thrust::fill( - rmm::exec_policy(handle.get_stream()), d_counts.begin(), d_counts.end(), edge_t{0}); - thrust::scatter(rmm::exec_policy(handle.get_stream()), + thrust::fill(handle.get_thrust_policy(), d_counts.begin(), d_counts.end(), edge_t{0}); + thrust::scatter(handle.get_thrust_policy(), d_edge_counts.begin(), thrust::get<1>(it), d_local_partition_ids.begin(), @@ -508,7 +507,7 @@ class louvain_functor { std::pair operator()(raft::handle_t const& handle, graph_view_t const& graph_view) { - thrust::copy(rmm::exec_policy(handle.get_stream()), + thrust::copy(handle.get_thrust_policy(), thrust::make_counting_iterator(graph_view.get_local_vertex_first()), thrust::make_counting_iterator(graph_view.get_local_vertex_last()), reinterpret_cast(identifiers_)); diff --git a/cpp/src/utilities/path_retrieval.cu b/cpp/src/utilities/path_retrieval.cu index b51593b117f..102c0f33560 100644 --- a/cpp/src/utilities/path_retrieval.cu +++ b/cpp/src/utilities/path_retrieval.cu @@ -74,9 +74,10 @@ void get_traversed_cost_impl(raft::handle_t const& handle, vertex_t* vtx_keys = vtx_keys_v.data(); raft::copy(vtx_keys, vertices, num_vertices, stream); - thrust::sequence(rmm::exec_policy(stream), vtx_map, vtx_map + num_vertices); + thrust::sequence(handle.get_thrust_policy(), vtx_map, vtx_map + num_vertices); - thrust::stable_sort_by_key(rmm::exec_policy(stream), vtx_keys, vtx_keys + num_vertices, vtx_map); + thrust::stable_sort_by_key( + handle.get_thrust_policy(), vtx_keys, vtx_keys + num_vertices, vtx_map); get_traversed_cost_kernel<<>>( vertices, preds, vtx_map, info_weights, out, stop_vertex, num_vertices); diff --git a/cpp/src/utilities/spmv_1D.cuh b/cpp/src/utilities/spmv_1D.cuh index b45011d7f26..54221aafb24 100644 --- a/cpp/src/utilities/spmv_1D.cuh +++ b/cpp/src/utilities/spmv_1D.cuh @@ -18,7 +18,6 @@ #include #include #include -#include namespace cugraph { namespace mg { diff --git a/cpp/tests/centrality/legacy/betweenness_centrality_test.cu b/cpp/tests/centrality/legacy/betweenness_centrality_test.cu index 203bf506811..53d55b52ed4 100644 --- a/cpp/tests/centrality/legacy/betweenness_centrality_test.cu +++ b/cpp/tests/centrality/legacy/betweenness_centrality_test.cu @@ -23,6 +23,7 @@ #include #include +#include #include diff --git a/cpp/tests/centrality/legacy/edge_betweenness_centrality_test.cu b/cpp/tests/centrality/legacy/edge_betweenness_centrality_test.cu index c5fd7af2bf6..46a95695fbc 100644 --- a/cpp/tests/centrality/legacy/edge_betweenness_centrality_test.cu +++ b/cpp/tests/centrality/legacy/edge_betweenness_centrality_test.cu @@ -20,6 +20,7 @@ #include #include +#include #include diff --git a/cpp/tests/centrality/legacy/katz_centrality_test.cu b/cpp/tests/centrality/legacy/katz_centrality_test.cu index 34097cb244d..e9d3f8d7a69 100644 --- a/cpp/tests/centrality/legacy/katz_centrality_test.cu +++ b/cpp/tests/centrality/legacy/katz_centrality_test.cu @@ -24,6 +24,7 @@ #include #include +#include #include #include @@ -46,11 +47,13 @@ std::vector getGoldenTopKIds(std::ifstream& fs_result, int k = 10) std::vector getTopKIds(double* p_katz, int count, int k = 10) { - cudaStream_t stream = nullptr; rmm::device_vector id(count); - thrust::sequence(rmm::exec_policy(stream), id.begin(), id.end()); - thrust::sort_by_key( - rmm::exec_policy(stream), p_katz, p_katz + count, id.begin(), thrust::greater()); + thrust::sequence(rmm::exec_policy(rmm::cuda_stream_default), id.begin(), id.end()); + thrust::sort_by_key(rmm::exec_policy(rmm::cuda_stream_default), + p_katz, + p_katz + count, + id.begin(), + thrust::greater()); std::vector topK(k); thrust::copy(id.begin(), id.begin() + k, topK.begin()); return topK; @@ -59,12 +62,10 @@ std::vector getTopKIds(double* p_katz, int count, int k = 10) template int getMaxDegree(cugraph::legacy::GraphCSRView const& g) { - cudaStream_t stream{nullptr}; - rmm::device_vector degree_vector(g.number_of_vertices); ET* p_degree = degree_vector.data().get(); g.degree(p_degree, cugraph::legacy::DegreeDirection::OUT); - ET max_out_degree = thrust::reduce(rmm::exec_policy(stream), + ET max_out_degree = thrust::reduce(rmm::exec_policy(rmm::cuda_stream_default), p_degree, p_degree + g.number_of_vertices, static_cast(-1), diff --git a/cpp/tests/community/ecg_test.cpp b/cpp/tests/community/ecg_test.cpp index 7906ca19a9a..f8d2ebf13f0 100644 --- a/cpp/tests/community/ecg_test.cpp +++ b/cpp/tests/community/ecg_test.cpp @@ -14,7 +14,6 @@ #include #include -#include // FIXME: Temporarily disable this test. Something is wrong with // ECG, or the expectation of this test. If I run ensemble size diff --git a/cpp/tests/community/mg_louvain_helper.cu b/cpp/tests/community/mg_louvain_helper.cu index f5814af9820..5909ab177cd 100644 --- a/cpp/tests/community/mg_louvain_helper.cu +++ b/cpp/tests/community/mg_louvain_helper.cu @@ -39,20 +39,21 @@ void single_gpu_renumber_edgelist_given_number_map(raft::handle_t const& handle, { rmm::device_uvector index_v(renumber_map_gathered_v.size(), handle.get_stream()); + auto execution_policy = handle.get_thrust_policy(); thrust::for_each( - rmm::exec_policy(handle.get_stream()), + execution_policy, thrust::make_counting_iterator(0), thrust::make_counting_iterator(renumber_map_gathered_v.size()), [d_renumber_map_gathered = renumber_map_gathered_v.data(), d_index = index_v.data()] __device__( auto idx) { d_index[d_renumber_map_gathered[idx]] = idx; }); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(execution_policy, edgelist_rows_v.begin(), edgelist_rows_v.end(), edgelist_rows_v.begin(), [d_index = index_v.data()] __device__(auto v) { return d_index[v]; }); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(execution_policy, edgelist_cols_v.begin(), edgelist_cols_v.end(), edgelist_cols_v.begin(), @@ -84,7 +85,8 @@ compressed_sparse_to_edgelist(edge_t const* compressed_sparse_offsets, // FIXME: this is highly inefficient for very high-degree vertices, for better performance, we can // fill high-degree vertices using one CUDA block per vertex, mid-degree vertices using one CUDA // warp per vertex, and low-degree vertices using one CUDA thread per block - thrust::for_each(rmm::exec_policy(stream), + auto execution_policy = handle.get_thrust_policy(); + thrust::for_each(execution_policy, thrust::make_counting_iterator(major_first), thrust::make_counting_iterator(major_last), [compressed_sparse_offsets, @@ -94,12 +96,12 @@ compressed_sparse_to_edgelist(edge_t const* compressed_sparse_offsets, auto last = compressed_sparse_offsets[v - major_first + 1]; thrust::fill(thrust::seq, p_majors + first, p_majors + last, v); }); - thrust::copy(rmm::exec_policy(stream), + thrust::copy(execution_policy, compressed_sparse_indices, compressed_sparse_indices + number_of_edges, edgelist_minor_vertices.begin()); if (compressed_sparse_weights) { - thrust::copy(rmm::exec_policy(stream), + thrust::copy(execution_policy, (*compressed_sparse_weights), (*compressed_sparse_weights) + number_of_edges, (*edgelist_weights).data()); @@ -121,8 +123,10 @@ void sort_and_coarsen_edgelist( thrust::make_tuple(edgelist_major_vertices.begin(), edgelist_minor_vertices.begin())); size_t number_of_edges{0}; + + auto execution_policy = handle.get_thrust_policy(); if (edgelist_weights) { - thrust::sort_by_key(rmm::exec_policy(stream), + thrust::sort_by_key(execution_policy, pair_first, pair_first + edgelist_major_vertices.size(), (*edgelist_weights).begin()); @@ -133,7 +137,7 @@ void sort_and_coarsen_edgelist( stream); rmm::device_uvector tmp_edgelist_weights(tmp_edgelist_major_vertices.size(), stream); auto it = thrust::reduce_by_key( - rmm::exec_policy(stream), + execution_policy, pair_first, pair_first + edgelist_major_vertices.size(), (*edgelist_weights).begin(), @@ -146,9 +150,9 @@ void sort_and_coarsen_edgelist( edgelist_minor_vertices = std::move(tmp_edgelist_minor_vertices); (*edgelist_weights) = std::move(tmp_edgelist_weights); } else { - thrust::sort(rmm::exec_policy(stream), pair_first, pair_first + edgelist_major_vertices.size()); - auto it = thrust::unique( - rmm::exec_policy(stream), pair_first, pair_first + edgelist_major_vertices.size()); + thrust::sort(execution_policy, pair_first, pair_first + edgelist_major_vertices.size()); + auto it = + thrust::unique(execution_policy, pair_first, pair_first + edgelist_major_vertices.size()); number_of_edges = thrust::distance(pair_first, it); } @@ -243,7 +247,7 @@ coarsen_graph( : std::nullopt; edgelist.number_of_edges = static_cast(coarsened_edgelist_major_vertices.size()); - vertex_t new_number_of_vertices = 1 + thrust::reduce(rmm::exec_policy(handle.get_stream()), + vertex_t new_number_of_vertices = 1 + thrust::reduce(handle.get_thrust_policy(), labels, labels + graph_view.get_number_of_vertices(), vertex_t{0}, diff --git a/cpp/tests/components/con_comp_test.cu b/cpp/tests/components/con_comp_test.cu index 331ba53b3a7..97758058adc 100644 --- a/cpp/tests/components/con_comp_test.cu +++ b/cpp/tests/components/con_comp_test.cu @@ -24,6 +24,8 @@ #include #include +#include + #include #include diff --git a/cpp/tests/components/scc_test.cu b/cpp/tests/components/scc_test.cu index d97bb62201b..1a0d22cdb26 100644 --- a/cpp/tests/components/scc_test.cu +++ b/cpp/tests/components/scc_test.cu @@ -22,6 +22,7 @@ #include #include #include +#include #include #include diff --git a/cpp/tests/components/wcc_graphs.cu b/cpp/tests/components/wcc_graphs.cu index 8254eaa1b1a..3429ad3cf21 100644 --- a/cpp/tests/components/wcc_graphs.cu +++ b/cpp/tests/components/wcc_graphs.cu @@ -42,14 +42,13 @@ LineGraph_Usecase::construct_graph(raft::handle_t const& handle, rmm::device_uvector dst_v(num_edges, handle.get_stream()); rmm::device_uvector order_v(num_vertices_, handle.get_stream()); - thrust::sequence( - rmm::exec_policy(handle.get_stream()), vertices_v.begin(), vertices_v.end(), vertex_t{0}); + auto execution_policy = handle.get_thrust_policy(); + thrust::sequence(execution_policy, vertices_v.begin(), vertices_v.end(), vertex_t{0}); cugraph::detail::uniform_random_fill( handle.get_stream_view(), order_v.data(), num_vertices_, double{0.0}, double{1.0}, seed); - thrust::sort_by_key( - rmm::exec_policy(handle.get_stream()), order_v.begin(), order_v.end(), vertices_v.begin()); + thrust::sort_by_key(execution_policy, order_v.begin(), order_v.end(), vertices_v.begin()); raft::copy(src_v.begin(), vertices_v.begin(), (num_vertices_ - 1), handle.get_stream()); raft::copy(dst_v.begin(), vertices_v.begin() + 1, (num_vertices_ - 1), handle.get_stream()); @@ -63,8 +62,7 @@ LineGraph_Usecase::construct_graph(raft::handle_t const& handle, (num_vertices_ - 1), handle.get_stream()); - thrust::sequence( - rmm::exec_policy(handle.get_stream()), vertices_v.begin(), vertices_v.end(), vertex_t{0}); + thrust::sequence(execution_policy, vertices_v.begin(), vertices_v.end(), vertex_t{0}); handle.get_stream_view().synchronize(); diff --git a/cpp/tests/prims/mg_count_if_v.cu b/cpp/tests/prims/mg_count_if_v.cu index f4d4c24dd71..888832b2efe 100644 --- a/cpp/tests/prims/mg_count_if_v.cu +++ b/cpp/tests/prims/mg_count_if_v.cu @@ -138,7 +138,7 @@ class Tests_MG_CountIfV handle, true, false); auto sg_graph_view = sg_graph.view(); auto expected_vertex_count = - thrust::count_if(rmm::exec_policy(handle.get_stream()), + thrust::count_if(handle.get_thrust_policy(), thrust::make_counting_iterator(sg_graph_view.get_local_vertex_first()), thrust::make_counting_iterator(sg_graph_view.get_local_vertex_last()), test_predicate(hash_bin_count)); diff --git a/cpp/tests/prims/mg_reduce_v.cu b/cpp/tests/prims/mg_reduce_v.cu index 1c5c1c261b1..8f44d00100d 100644 --- a/cpp/tests/prims/mg_reduce_v.cu +++ b/cpp/tests/prims/mg_reduce_v.cu @@ -92,7 +92,7 @@ struct generate_impl { { auto data = std::make_tuple(rmm::device_uvector(labels.size(), handle.get_stream())...); auto zip = get_zip_iterator(data); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), labels.begin(), labels.end(), zip, @@ -108,7 +108,7 @@ struct generate_impl { auto length = thrust::distance(begin, end); auto data = std::make_tuple(rmm::device_uvector(length, handle.get_stream())...); auto zip = get_zip_iterator(data); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), begin, end, zip, @@ -272,7 +272,7 @@ class Tests_MG_ReduceV using property_t = decltype(property_initial_value); auto expected_result = - thrust::reduce(rmm::exec_policy(handle.get_stream()), + thrust::reduce(handle.get_thrust_policy(), sg_property_iter, sg_property_iter + sg_graph_view.get_number_of_local_vertices(), property_initial_value, diff --git a/cpp/tests/prims/mg_transform_reduce_v.cu b/cpp/tests/prims/mg_transform_reduce_v.cu index 4f7f3a5a724..5e68b52dd02 100644 --- a/cpp/tests/prims/mg_transform_reduce_v.cu +++ b/cpp/tests/prims/mg_transform_reduce_v.cu @@ -211,7 +211,7 @@ class Tests_MG_TransformReduceV using property_t = decltype(property_initial_value); auto expected_result = thrust::transform_reduce( - rmm::exec_policy(handle.get_stream()), + handle.get_thrust_policy(), thrust::make_counting_iterator(sg_graph_view.get_local_vertex_first()), thrust::make_counting_iterator(sg_graph_view.get_local_vertex_last()), prop, diff --git a/cpp/tests/sampling/random_walks_profiling.cu b/cpp/tests/sampling/random_walks_profiling.cu index ca105a482b9..b5aa787ec28 100644 --- a/cpp/tests/sampling/random_walks_profiling.cu +++ b/cpp/tests/sampling/random_walks_profiling.cu @@ -47,7 +47,7 @@ void fill_start(raft::handle_t const& handle, { index_t num_paths = d_start.size(); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(num_paths), diff --git a/cpp/tests/sampling/random_walks_test.cu b/cpp/tests/sampling/random_walks_test.cu index f3603549eb5..7c35440a9b5 100644 --- a/cpp/tests/sampling/random_walks_test.cu +++ b/cpp/tests/sampling/random_walks_test.cu @@ -47,7 +47,7 @@ void fill_start(raft::handle_t const& handle, { index_t num_paths = d_start.size(); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(handle.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(num_paths), diff --git a/cpp/tests/sampling/rw_low_level_test.cu b/cpp/tests/sampling/rw_low_level_test.cu index d1e444c0513..0977d1031bf 100644 --- a/cpp/tests/sampling/rw_low_level_test.cu +++ b/cpp/tests/sampling/rw_low_level_test.cu @@ -53,7 +53,7 @@ bool check_col_indices(raft::handle_t const& handle, index_t num_paths) { bool all_indices_within_degs = thrust::all_of( - rmm::exec_policy(handle.get_stream_view()), + handle.get_thrust_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(num_paths), [p_d_col_indx = cugraph::detail::raw_const_ptr(d_col_indx), diff --git a/cpp/tests/traversal/legacy/sssp_test.cu b/cpp/tests/traversal/legacy/sssp_test.cu index 74257256dca..ffa04f6d649 100644 --- a/cpp/tests/traversal/legacy/sssp_test.cu +++ b/cpp/tests/traversal/legacy/sssp_test.cu @@ -18,6 +18,7 @@ #include #include #include +#include #include diff --git a/cpp/tests/utilities/matrix_market_file_utilities.cu b/cpp/tests/utilities/matrix_market_file_utilities.cu index 22498a124aa..80ce4509ea4 100644 --- a/cpp/tests/utilities/matrix_market_file_utilities.cu +++ b/cpp/tests/utilities/matrix_market_file_utilities.cu @@ -335,8 +335,8 @@ read_graph_from_matrix_market_file(raft::handle_t const& handle, handle, graph_file_full_path, test_weighted); rmm::device_uvector d_vertices(number_of_vertices, handle.get_stream()); - thrust::sequence( - rmm::exec_policy(handle.get_stream()), d_vertices.begin(), d_vertices.end(), vertex_t{0}); + auto execution_policy = handle.get_thrust_policy(); + thrust::sequence(execution_policy, d_vertices.begin(), d_vertices.end(), vertex_t{0}); handle.get_stream_view().synchronize(); if (multi_gpu) { @@ -351,7 +351,7 @@ read_graph_from_matrix_market_file(raft::handle_t const& handle, auto vertex_key_func = cugraph::detail::compute_gpu_id_from_vertex_t{comm_size}; d_vertices.resize( thrust::distance(d_vertices.begin(), - thrust::remove_if(rmm::exec_policy(handle.get_stream()), + thrust::remove_if(execution_policy, d_vertices.begin(), d_vertices.end(), [comm_rank, key_func = vertex_key_func] __device__( @@ -367,7 +367,7 @@ read_graph_from_matrix_market_file(raft::handle_t const& handle, d_edgelist_rows.begin(), d_edgelist_cols.begin(), (*d_edgelist_weights).begin())); number_of_local_edges = thrust::distance( edge_first, - thrust::remove_if(rmm::exec_policy(handle.get_stream()), + thrust::remove_if(execution_policy, edge_first, edge_first + d_edgelist_rows.size(), [comm_rank, key_func = edge_key_func] __device__(auto e) { @@ -380,7 +380,7 @@ read_graph_from_matrix_market_file(raft::handle_t const& handle, thrust::make_tuple(d_edgelist_rows.begin(), d_edgelist_cols.begin())); number_of_local_edges = thrust::distance( edge_first, - thrust::remove_if(rmm::exec_policy(handle.get_stream()), + thrust::remove_if(execution_policy, edge_first, edge_first + d_edgelist_rows.size(), [comm_rank, key_func = edge_key_func] __device__(auto e) { diff --git a/cpp/tests/utilities/thrust_wrapper.cu b/cpp/tests/utilities/thrust_wrapper.cu index ae36582d18d..82048955abd 100644 --- a/cpp/tests/utilities/thrust_wrapper.cu +++ b/cpp/tests/utilities/thrust_wrapper.cu @@ -32,15 +32,12 @@ std::tuple, rmm::device_uvector> sort_by_ rmm::device_uvector sorted_keys(num_pairs, handle.get_stream_view()); rmm::device_uvector sorted_values(num_pairs, handle.get_stream_view()); - thrust::copy( - rmm::exec_policy(handle.get_stream_view()), keys, keys + num_pairs, sorted_keys.begin()); - thrust::copy( - rmm::exec_policy(handle.get_stream_view()), values, values + num_pairs, sorted_values.begin()); + auto execution_policy = handle.get_thrust_policy(); + thrust::copy(execution_policy, keys, keys + num_pairs, sorted_keys.begin()); + thrust::copy(execution_policy, values, values + num_pairs, sorted_values.begin()); - thrust::sort_by_key(rmm::exec_policy(handle.get_stream_view()), - sorted_keys.begin(), - sorted_keys.end(), - sorted_values.begin()); + thrust::sort_by_key( + execution_policy, sorted_keys.begin(), sorted_keys.end(), sorted_values.begin()); return std::make_tuple(std::move(sorted_keys), std::move(sorted_values)); } @@ -87,13 +84,14 @@ void translate_vertex_ids(raft::handle_t const& handle, rmm::device_uvector& d_dst_v, vertex_t vertex_id_offset) { - thrust::transform(rmm::exec_policy(handle.get_stream()), + auto execution_policy = handle.get_thrust_policy(); + thrust::transform(execution_policy, d_src_v.begin(), d_src_v.end(), d_src_v.begin(), [offset = vertex_id_offset] __device__(vertex_t v) { return offset + v; }); - thrust::transform(rmm::exec_policy(handle.get_stream()), + thrust::transform(execution_policy, d_dst_v.begin(), d_dst_v.end(), d_dst_v.begin(), @@ -105,10 +103,8 @@ void populate_vertex_ids(raft::handle_t const& handle, rmm::device_uvector& d_vertices_v, vertex_t vertex_id_offset) { - thrust::sequence(rmm::exec_policy(handle.get_stream()), - d_vertices_v.begin(), - d_vertices_v.end(), - vertex_id_offset); + thrust::sequence( + handle.get_thrust_policy(), d_vertices_v.begin(), d_vertices_v.end(), vertex_id_offset); } template void translate_vertex_ids(raft::handle_t const& handle,