From 80a34592cfc43946e6bd921fd7649a6418064fcc Mon Sep 17 00:00:00 2001 From: Dante Gama Dessavre Date: Wed, 25 Aug 2021 14:25:26 -0500 Subject: [PATCH] Update thrust/RMM deprecated calls (#1789) PR essentially backports the fixes in #1707 for the deprecated RMM exec policy, and unpins the RMM nighlty to unblock local builds and CI and local development. Conflicts with the bigger PR should be minimal, and mainly around a single change: ``` rmm::exec_policy(handle.get_stream()) ``` to ``` handle.get_thrust_policy() ``` Authors: - Dante Gama Dessavre (https://github.com/dantegd) Approvers: - AJ Schmidt (https://github.com/ajschmidt8) - Chuck Hastings (https://github.com/ChuckHastings) - Seunghwa Kang (https://github.com/seunghwak) URL: https://github.com/rapidsai/cugraph/pull/1789 --- ci/gpu/build.sh | 5 +- conda/environments/cugraph_dev_cuda11.0.yml | 4 +- conda/environments/cugraph_dev_cuda11.2.yml | 4 +- conda/environments/cugraph_dev_cuda11.4.yml | 4 +- conda/recipes/libcugraph/meta.yaml | 3 +- cpp/cmake/thirdparty/get_cuhornet.cmake | 2 +- cpp/cmake/thirdparty/get_rmm.cmake | 6 +- cpp/include/cugraph/compute_partition.cuh | 2 +- cpp/include/cugraph/detail/graph_utils.cuh | 10 +-- .../prims/copy_to_adj_matrix_row_col.cuh | 22 ++--- .../copy_v_transform_reduce_in_out_nbr.cuh | 8 +- ...ransform_reduce_key_aggregated_out_nbr.cuh | 83 +++++++++---------- cpp/include/cugraph/prims/count_if_v.cuh | 8 +- cpp/include/cugraph/prims/reduce_v.cuh | 4 +- ...orm_reduce_by_adj_matrix_row_col_key_e.cuh | 10 +-- .../cugraph/prims/transform_reduce_e.cuh | 6 +- .../cugraph/prims/transform_reduce_v.cuh | 4 +- .../update_frontier_v_push_if_out_nbr.cuh | 51 ++++++------ cpp/include/cugraph/prims/vertex_frontier.cuh | 70 +++++++--------- cpp/src/centrality/betweenness_centrality.cu | 1 + cpp/src/centrality/betweenness_centrality.cuh | 3 +- cpp/src/centrality/katz_centrality.cu | 10 +-- cpp/src/community/flatten_dendrogram.cuh | 4 +- cpp/src/community/legacy/egonet.cu | 2 +- .../legacy/extract_subgraph_by_vertex.cu | 9 +- cpp/src/community/legacy/leiden.cu | 2 +- .../community/legacy/spectral_clustering.cu | 11 +-- .../community/legacy/triangles_counting.cu | 3 +- cpp/src/community/louvain.cu | 8 +- cpp/src/components/weak_cc.cuh | 3 +- cpp/src/converters/permute_graph.cuh | 15 ++-- cpp/src/cores/core_number.cu | 12 +-- cpp/src/link_analysis/pagerank.cu | 12 +-- cpp/src/link_prediction/overlap.cu | 3 +- cpp/src/serialization/serializer.cu | 2 +- cpp/src/structure/coarsen_graph.cu | 78 ++++++++--------- .../structure/create_graph_from_edgelist.cpp | 2 +- cpp/src/structure/renumber_edgelist.cu | 2 +- cpp/src/traversal/bfs.cu | 4 +- cpp/src/traversal/legacy/bfs.cuh | 3 +- cpp/src/traversal/legacy/mg/bfs.cuh | 5 +- cpp/src/traversal/legacy/mg/common_utils.cuh | 30 +++---- .../traversal/legacy/mg/frontier_expand.cuh | 1 + .../traversal/legacy/mg/vertex_binning.cuh | 6 +- .../legacy/mg/vertex_binning_kernels.cuh | 3 +- cpp/src/traversal/legacy/sssp.cu | 1 + cpp/src/traversal/legacy/sssp.cuh | 3 +- cpp/src/traversal/legacy/traversal_common.cuh | 2 +- cpp/src/traversal/sssp.cu | 6 +- cpp/src/traversal/tsp.hpp | 2 +- cpp/src/traversal/two_hop_neighbors.cu | 18 ++-- cpp/src/tree/mst.cu | 2 +- cpp/src/utilities/cython.cu | 14 ++-- cpp/src/utilities/path_retrieval.cu | 7 +- cpp/src/utilities/spmv_1D.cuh | 3 +- .../legacy/betweenness_centrality_test.cu | 2 + .../edge_betweenness_centrality_test.cu | 2 + .../centrality/legacy/katz_centrality_test.cu | 13 ++- cpp/tests/community/balanced_edge_test.cpp | 3 +- cpp/tests/community/ecg_test.cpp | 3 +- cpp/tests/community/egonet_test.cu | 2 +- cpp/tests/community/leiden_test.cpp | 2 +- cpp/tests/community/mg_louvain_helper.cu | 40 ++++----- cpp/tests/community/triangle_test.cu | 3 +- cpp/tests/components/con_comp_test.cu | 2 + cpp/tests/components/scc_test.cu | 2 + cpp/tests/generators/erdos_renyi_test.cpp | 1 + cpp/tests/layout/force_atlas2_test.cu | 2 +- cpp/tests/linear_assignment/hungarian_test.cu | 2 +- 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 | 4 +- cpp/tests/sampling/random_walks_test.cu | 4 +- cpp/tests/sampling/random_walks_utils.cuh | 2 +- cpp/tests/sampling/rw_biased_seg_sort.cu | 2 +- cpp/tests/sampling/rw_low_level_test.cu | 2 +- cpp/tests/serialization/un_serialize_test.cpp | 2 +- cpp/tests/traversal/legacy/bfs_test.cu | 3 +- cpp/tests/traversal/legacy/sssp_test.cu | 2 + cpp/tests/traversal/ms_bfs_test.cpp | 2 +- cpp/tests/utilities/base_fixture.hpp | 2 +- .../utilities/matrix_market_file_utilities.cu | 26 +++--- 83 files changed, 357 insertions(+), 381 deletions(-) diff --git a/ci/gpu/build.sh b/ci/gpu/build.sh index 31232d41e83..7d7e315fb73 100755 --- a/ci/gpu/build.sh +++ b/ci/gpu/build.sh @@ -53,13 +53,12 @@ gpuci_logger "Activate conda env" . /opt/conda/etc/profile.d/conda.sh conda activate rapids -# FIXME: return librmm and RMM to ${MINOR_VERSION} gpuci_logger "Install dependencies" gpuci_mamba_retry install -y \ "libcudf=${MINOR_VERSION}" \ "cudf=${MINOR_VERSION}" \ - "librmm=21.10.00a210812" \ - "rmm=21.10.00a210813" \ + "librmm=${MINOR_VERSION}" \ + "rmm=${MINOR_VERSION}" \ "cudatoolkit=$CUDA_REL" \ "dask-cudf=${MINOR_VERSION}" \ "dask-cuda=${MINOR_VERSION}" \ diff --git a/conda/environments/cugraph_dev_cuda11.0.yml b/conda/environments/cugraph_dev_cuda11.0.yml index 78e1efecafb..2b72b4b8570 100644 --- a/conda/environments/cugraph_dev_cuda11.0.yml +++ b/conda/environments/cugraph_dev_cuda11.0.yml @@ -8,8 +8,8 @@ dependencies: - cudatoolkit=11.0 - cudf=21.10.* - libcudf=21.10.* -- rmm=21.10.00a210813 -- librmm=21.10.00a210812 +- rmm=21.10.* +- librmm=21.10.* - dask>=2021.6.0 - distributed>=2021.6.0 - dask-cuda=21.10.* diff --git a/conda/environments/cugraph_dev_cuda11.2.yml b/conda/environments/cugraph_dev_cuda11.2.yml index 1941aa22b76..f97dede33ae 100644 --- a/conda/environments/cugraph_dev_cuda11.2.yml +++ b/conda/environments/cugraph_dev_cuda11.2.yml @@ -8,8 +8,8 @@ dependencies: - cudatoolkit=11.2 - cudf=21.10.* - libcudf=21.10.* -- rmm=21.10.00a210813 -- librmm=21.10.00a210812 +- rmm=21.10.* +- librmm=21.10.* - dask>=2021.6.0 - distributed>=2021.6.0 - dask-cuda=21.10.* diff --git a/conda/environments/cugraph_dev_cuda11.4.yml b/conda/environments/cugraph_dev_cuda11.4.yml index defd2e633a4..c6fe1364ce9 100644 --- a/conda/environments/cugraph_dev_cuda11.4.yml +++ b/conda/environments/cugraph_dev_cuda11.4.yml @@ -8,8 +8,8 @@ dependencies: - cudatoolkit=11.4 - cudf=21.10.* - libcudf=21.10.* -- rmm=21.10.00a210813 -- librmm=21.10.00a210812 +- rmm=21.10.* +- librmm=21.10.* - dask>=2021.6.0 - distributed>=2021.6.0 - dask-cuda=21.10.* diff --git a/conda/recipes/libcugraph/meta.yaml b/conda/recipes/libcugraph/meta.yaml index c4d7f30afc1..570a0ec09b2 100644 --- a/conda/recipes/libcugraph/meta.yaml +++ b/conda/recipes/libcugraph/meta.yaml @@ -29,12 +29,11 @@ build: - CMAKE_CXX_COMPILER_LAUNCHER - CMAKE_CUDA_COMPILER_LAUNCHER -# FIXME: return librmm to {{ minor_version }}.* requirements: build: - cmake>=3.20.1 - cudatoolkit {{ cuda_version }}.* - - librmm=21.10.00a210812 + - librmm {{ minor_version }}.* - boost-cpp>=1.66 - nccl>=2.9.9 - ucx-proc=*=gpu diff --git a/cpp/cmake/thirdparty/get_cuhornet.cmake b/cpp/cmake/thirdparty/get_cuhornet.cmake index 28c83161ff4..df5448f479e 100644 --- a/cpp/cmake/thirdparty/get_cuhornet.cmake +++ b/cpp/cmake/thirdparty/get_cuhornet.cmake @@ -21,7 +21,7 @@ function(find_and_configure_cuhornet) FetchContent_Declare( cuhornet GIT_REPOSITORY https://github.com/rapidsai/cuhornet.git - GIT_TAG 261399356e62bd76fa7628880f1a847aee713eed + GIT_TAG 4a1daa18405c0242370e16ce302dfa7eb5d9e857 SOURCE_SUBDIR hornet ) FetchContent_GetProperties(cuhornet) diff --git a/cpp/cmake/thirdparty/get_rmm.cmake b/cpp/cmake/thirdparty/get_rmm.cmake index 80fd8f329ad..aecb6489f92 100644 --- a/cpp/cmake/thirdparty/get_rmm.cmake +++ b/cpp/cmake/thirdparty/get_rmm.cmake @@ -26,16 +26,14 @@ function(find_and_configure_rmm VERSION) return() endif() - # FIXME: turn GIT_SHALLOW back to TRUE when changing GIT_TAG back - # to branch-${MAJOR_AND_MINOR} rapids_cpm_find(rmm ${VERSION} GLOBAL_TARGETS rmm::rmm BUILD_EXPORT_SET cugraph-exports INSTALL_EXPORT_SET cugraph-exports CPM_ARGS GIT_REPOSITORY https://github.com/rapidsai/rmm.git - GIT_TAG 23bbe745af1d988224b5498f7b8e3fe3720532d4 - GIT_SHALLOW FALSE + GIT_TAG branch-${MAJOR_AND_MINOR} + GIT_SHALLOW TRUE OPTIONS "BUILD_TESTS OFF" "BUILD_BENCHMARKS OFF" "CUDA_STATIC_RUNTIME ${CUDA_STATIC_RUNTIME}" diff --git a/cpp/include/cugraph/compute_partition.cuh b/cpp/include/cugraph/compute_partition.cuh index 3fd66c95d27..ca311452a73 100644 --- a/cpp/include/cugraph/compute_partition.cuh +++ b/cpp/include/cugraph/compute_partition.cuh @@ -19,7 +19,7 @@ #include -#include +#include namespace cugraph { namespace detail { diff --git a/cpp/include/cugraph/detail/graph_utils.cuh b/cpp/include/cugraph/detail/graph_utils.cuh index f7acc6dcefa..7f22699b62c 100644 --- a/cpp/include/cugraph/detail/graph_utils.cuh +++ b/cpp/include/cugraph/detail/graph_utils.cuh @@ -20,9 +20,9 @@ #include #include -#include #include #include +#include #include #include @@ -78,7 +78,7 @@ 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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), thrust::make_counting_iterator(vertex_t{0}), thrust::make_counting_iterator(major_hypersparse_first - major_first), local_degrees.begin(), @@ -86,11 +86,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())->on(handle.get_stream()), + thrust::fill(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::for_each(rmm::exec_policy(handle.get_stream()), thrust::make_counting_iterator(vertex_t{0}), thrust::make_counting_iterator(dcs_nzd_vertex_count), [p_offsets, @@ -123,7 +123,7 @@ 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())->on(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]; }); 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 96aefa016fa..435340f84dc 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 @@ -26,8 +26,8 @@ #include #include -#include #include +#include #include #include @@ -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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::gather(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::scatter(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::gather(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::scatter(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::scatter(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::scatter(rmm::exec_policy(handle.get_stream()), 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 857ed1b7da2..335b34828e5 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 @@ -25,8 +25,8 @@ #include #include -#include #include +#include #include #include @@ -439,12 +439,12 @@ void copy_v_transform_reduce_nbr(raft::handle_t const& handle, } if (GraphViewType::is_multi_gpu) { - thrust::fill(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + thrust::fill(rmm::exec_policy(handle.get_stream()), minor_buffer_first, minor_buffer_first + minor_tmp_buffer_size, minor_init); } else { - thrust::fill(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + thrust::fill(rmm::exec_policy(handle.get_stream()), vertex_value_output_first, vertex_value_output_first + graph_view.get_number_of_local_vertices(), minor_init); @@ -546,7 +546,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())->on(handle.get_stream()), + thrust::fill(rmm::exec_policy(handle.get_stream()), 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 52b936090aa..5ae32a6f56a 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 @@ -153,7 +153,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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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 +167,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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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 +183,7 @@ void decompress_matrix_partition_to_fill_edgelist_majors( } } else { thrust::for_each( - rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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 +340,12 @@ 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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), map_key_first, map_key_last, map_keys.begin() + map_displacements[row_comm_rank]); thrust::copy( - rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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 +420,12 @@ 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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), *(matrix_partition.get_weights()), *(matrix_partition.get_weights()) + matrix_partition.get_number_of_edges(), tmp_key_aggregated_edge_weights.begin()); @@ -448,32 +448,30 @@ 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())->on(handle.get_stream()), + thrust::sort_by_key(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), - 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()))); + reduced_size = thrust::distance( + output_key_first, + thrust::get<0>(thrust::reduce_by_key(rmm::exec_policy(handle.get_stream()), + 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())->on(handle.get_stream()), + thrust::sort(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), - input_key_first, - input_key_first + tmp_major_vertices.size(), - thrust::make_constant_iterator(weight_t{1.0}), - output_key_first, - reduced_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()), + input_key_first, + input_key_first + tmp_major_vertices.size(), + thrust::make_constant_iterator(weight_t{1.0}), + output_key_first, + reduced_key_aggregated_edge_weights.begin()))); } tmp_major_vertices = std::move(reduced_major_vertices); tmp_minor_keys = std::move(reduced_minor_keys); @@ -517,21 +515,20 @@ 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())->on(handle.get_stream()), + thrust::sort_by_key(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), - pair_first, - pair_first + rx_major_vertices.size(), - rx_key_aggregated_edge_weights.begin(), - thrust::make_zip_iterator(thrust::make_tuple( - tmp_major_vertices.begin(), tmp_minor_keys.begin())), - tmp_key_aggregated_edge_weights.begin()); + auto pair_it = thrust::reduce_by_key(rmm::exec_policy(handle.get_stream()), + pair_first, + pair_first + rx_major_vertices.size(), + rx_key_aggregated_edge_weights.begin(), + thrust::make_zip_iterator(thrust::make_tuple( + tmp_major_vertices.begin(), tmp_minor_keys.begin())), + tmp_key_aggregated_edge_weights.begin()); tmp_major_vertices.resize( thrust::distance(tmp_key_aggregated_edge_weights.begin(), thrust::get<1>(pair_it)), handle.get_stream()); @@ -549,7 +546,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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), triplet_first, triplet_first + tmp_major_vertices.size(), tmp_e_op_result_buffer_first, @@ -635,17 +632,17 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( #endif } - thrust::fill(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + thrust::fill(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::sort_by_key(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), thrust::make_counting_iterator(size_t{0}), thrust::make_counting_iterator(major_vertices.size()), [major_vertices = major_vertices.data()] __device__(auto i) { @@ -661,13 +658,13 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( : invalid_vertex_id::value; }); thrust::copy_if( - rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), major_vertices.begin(), major_vertices.end(), get_dataframe_buffer_begin(e_op_result_buffer), @@ -683,7 +680,7 @@ void copy_v_transform_reduce_key_aggregated_out_nbr( thrust::equal_to{}, reduce_op); - thrust::transform(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), 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 d1296c73b93..5a7684d19aa 100644 --- a/cpp/include/cugraph/prims/count_if_v.cuh +++ b/cpp/include/cugraph/prims/count_if_v.cuh @@ -19,8 +19,8 @@ #include #include -#include #include +#include #include #include @@ -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())->on(handle.get_stream()), + thrust::count_if(rmm::exec_policy(handle.get_stream()), vertex_value_input_first, vertex_value_input_first + graph_view.get_number_of_local_vertices(), v_op); @@ -92,8 +92,8 @@ 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())->on(handle.get_stream()), input_first, input_last, v_op); + auto count = + thrust::count_if(rmm::exec_policy(handle.get_stream()), 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 21a732e7750..ef737a153df 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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 f6735946b52..f8583d71f5c 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 @@ -319,12 +319,12 @@ template std::tuple, BufferType> reduce_to_unique_kv_pairs( rmm::device_uvector&& keys, BufferType&& value_buffer, cudaStream_t stream) { - thrust::sort_by_key(rmm::exec_policy(stream)->on(stream), + thrust::sort_by_key(rmm::exec_policy(stream), keys.begin(), keys.end(), get_dataframe_buffer_begin(value_buffer)); auto num_uniques = - thrust::count_if(rmm::exec_policy(stream)->on(stream), + thrust::count_if(rmm::exec_policy(stream), thrust::make_counting_iterator(size_t{0}), thrust::make_counting_iterator(keys.size()), [keys = keys.data()] __device__(auto i) { @@ -333,7 +333,7 @@ std::tuple, BufferType> reduce_to_unique_kv_pairs( rmm::device_uvector unique_keys(num_uniques, stream); auto value_for_unique_key_buffer = allocate_dataframe_buffer(unique_keys.size(), stream); - thrust::reduce_by_key(rmm::exec_policy(stream)->on(stream), + thrust::reduce_by_key(rmm::exec_policy(stream), keys.begin(), keys.end(), get_dataframe_buffer_begin(value_buffer), @@ -530,11 +530,11 @@ 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())->on(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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), 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 f3381cae37f..f46a00d37e4 100644 --- a/cpp/include/cugraph/prims/transform_reduce_e.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_e.cuh @@ -21,8 +21,8 @@ #include #include -#include #include +#include #include @@ -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())->on(handle.get_stream()), + thrust::fill(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + auto result = thrust::reduce(rmm::exec_policy(handle.get_stream()), 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 60a2470d592..696d004e89b 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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 3fb0c600263..1d04dd7fa87 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 @@ -30,9 +30,9 @@ #include #include -#include #include #include +#include #include #include @@ -528,11 +528,11 @@ size_t sort_and_reduce_buffer_elements(raft::handle_t const& handle, typename optional_payload_buffer_value_type_t::value; if constexpr (std::is_same_v) { - thrust::sort(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + thrust::sort(rmm::exec_policy(handle.get_stream()), buffer_key_output_first, buffer_key_output_first + num_buffer_elements); } else { - thrust::sort_by_key(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + thrust::sort_by_key(rmm::exec_policy(handle.get_stream()), buffer_key_output_first, buffer_key_output_first + num_buffer_elements, buffer_payload_output_first); @@ -540,7 +540,7 @@ 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())->on(handle.get_stream()), + auto it = thrust::unique(rmm::exec_policy(handle.get_stream()), buffer_key_output_first, buffer_key_output_first + num_buffer_elements); num_reduced_buffer_elements = @@ -548,7 +548,7 @@ size_t sort_and_reduce_buffer_elements(raft::handle_t const& handle, } 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())->on(handle.get_stream()), + auto it = thrust::unique_by_key(rmm::exec_policy(handle.get_stream()), buffer_key_output_first, buffer_key_output_first + num_buffer_elements, buffer_payload_output_first); @@ -567,7 +567,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())->on(handle.get_stream()), + auto it = thrust::reduce_by_key(rmm::exec_policy(handle.get_stream()), buffer_key_output_first, buffer_key_output_first + num_buffer_elements, buffer_payload_output_first, @@ -578,11 +578,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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), keys.begin(), keys.begin() + num_reduced_buffer_elements, buffer_key_output_first); - thrust::copy(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), get_dataframe_buffer_begin(value_buffer), get_dataframe_buffer_begin(value_buffer) + num_reduced_buffer_elements, buffer_payload_output_first); @@ -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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), frontier_key_first, frontier_key_last, get_dataframe_buffer_begin(matrix_partition_frontier_key_buffer)); @@ -938,7 +938,7 @@ void update_frontier_v_push_if_out_nbr( auto max_pushes = use_dcs ? thrust::transform_reduce( - rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), matrix_partition_frontier_row_first, matrix_partition_frontier_row_last, [matrix_partition, @@ -963,7 +963,7 @@ void update_frontier_v_push_if_out_nbr( edge_t{0}, thrust::plus()) : thrust::transform_reduce( - rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), matrix_partition_frontier_row_first, matrix_partition_frontier_row_last, [matrix_partition] __device__(auto row) { @@ -1007,7 +1007,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())->on(handle.get_stream()), + thrust::lower_bound(rmm::exec_policy(handle.get_stream()), matrix_partition_frontier_row_first, matrix_partition_frontier_row_last, d_thresholds.begin(), @@ -1170,7 +1170,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())->on(handle.get_stream()), + thrust::lower_bound(rmm::exec_policy(handle.get_stream()), row_first, row_first + num_buffer_elements, d_vertex_lasts.begin(), @@ -1234,7 +1234,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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), key_payload_pair_first, key_payload_pair_first + num_buffer_elements, bucket_indices.begin(), @@ -1266,7 +1266,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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), get_dataframe_buffer_begin(key_buffer), get_dataframe_buffer_begin(key_buffer) + num_buffer_elements, bucket_indices.begin(), @@ -1285,12 +1285,11 @@ void update_frontier_v_push_if_out_nbr( auto bucket_key_pair_first = thrust::make_zip_iterator( 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())->on(handle.get_stream()), - bucket_key_pair_first, - bucket_key_pair_first + num_buffer_elements, - detail::check_invalid_bucket_idx_t())), + thrust::distance(bucket_key_pair_first, + thrust::remove_if(rmm::exec_policy(handle.get_stream()), + bucket_key_pair_first, + bucket_key_pair_first + num_buffer_elements, + detail::check_invalid_bucket_idx_t())), handle.get_stream()); resize_dataframe_buffer(key_buffer, bucket_indices.size(), handle.get_stream()); bucket_indices.shrink_to_fit(handle.get_stream()); diff --git a/cpp/include/cugraph/prims/vertex_frontier.cuh b/cpp/include/cugraph/prims/vertex_frontier.cuh index 8d34131cb45..c66444e4a77 100644 --- a/cpp/include/cugraph/prims/vertex_frontier.cuh +++ b/cpp/include/cugraph/prims/vertex_frontier.cuh @@ -96,10 +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())->on(handle_ptr_->get_stream()), - pair_first, - pair_first + 1, - key); + thrust::fill(rmm::exec_policy(handle_ptr_->get_stream()), pair_first, pair_first + 1, key); } } @@ -122,27 +119,24 @@ 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())->on(handle_ptr_->get_stream()), + thrust::merge(rmm::exec_policy(handle_ptr_->get_stream()), 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())->on(handle_ptr_->get_stream()), - merged_vertices.begin(), - merged_vertices.end())), + 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.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())->on(handle_ptr_->get_stream()), - vertex_first, - vertex_last, - vertices_.begin()); + thrust::copy( + rmm::exec_policy(handle_ptr_->get_stream()), vertex_first, vertex_last, vertices_.begin()); } } @@ -170,18 +164,17 @@ 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())->on(handle_ptr_->get_stream()), + thrust::merge(rmm::exec_policy(handle_ptr_->get_stream()), old_pair_first, old_pair_first + vertices_.size(), key_first, key_last, merged_pair_first); merged_vertices.resize( - thrust::distance( - merged_pair_first, - thrust::unique(rmm::exec_policy(handle_ptr_->get_stream())->on(handle_ptr_->get_stream()), - merged_pair_first, - merged_pair_first + merged_vertices.size())), + thrust::distance(merged_pair_first, + thrust::unique(rmm::exec_policy(handle_ptr_->get_stream()), + merged_pair_first, + merged_pair_first + merged_vertices.size())), handle_ptr_->get_stream()); merged_tags.resize(merged_vertices.size(), handle_ptr_->get_stream()); merged_vertices.shrink_to_fit(handle_ptr_->get_stream()); @@ -191,7 +184,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())->on(handle_ptr_->get_stream()), + thrust::copy(rmm::exec_policy(handle_ptr_->get_stream()), key_first, key_last, thrust::make_zip_iterator(thrust::make_tuple(vertices_.begin(), tags_.begin()))); @@ -332,7 +325,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())->on(handle_ptr_->get_stream()), + rmm::exec_policy(handle_ptr_->get_stream()), this_bucket.begin(), this_bucket.end(), bucket_indices.begin(), @@ -347,13 +340,13 @@ 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())->on(handle_ptr_->get_stream()), - pair_first, - pair_first + bucket_indices.size(), - [] __device__(auto pair) { - return thrust::get<0>(pair) == static_cast(kInvalidBucketIdx); - })), + thrust::remove_if(rmm::exec_policy(handle_ptr_->get_stream()), + pair_first, + pair_first + bucket_indices.size(), + [] __device__(auto pair) { + return thrust::get<0>(pair) == + static_cast(kInvalidBucketIdx); + })), handle_ptr_->get_stream()); this_bucket.resize(bucket_indices.size()); bucket_indices.shrink_to_fit(handle_ptr_->get_stream()); @@ -366,7 +359,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())->on(handle_ptr_->get_stream()), + rmm::exec_policy(handle_ptr_->get_stream()), pair_first, pair_first + bucket_indices.size(), [this_bucket_idx = static_cast(this_bucket_idx)] __device__(auto pair) { @@ -406,7 +399,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())->on(handle_ptr_->get_stream()), + rmm::exec_policy(handle_ptr_->get_stream()), pair_first, pair_last, [next_bucket_idx = static_cast(to_bucket_indices[0])] __device__(auto pair) { @@ -419,19 +412,18 @@ 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())->on(handle_ptr_->get_stream()), + rmm::exec_policy(handle_ptr_->get_stream()), 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())->on(handle_ptr_->get_stream()), - bucket_idx_first, - bucket_idx_last, - thrust::make_constant_iterator(size_t{1}), - d_indices.begin(), - d_counts.begin()); + auto it = thrust::reduce_by_key(rmm::exec_policy(handle_ptr_->get_stream()), + bucket_idx_first, + bucket_idx_last, + thrust::make_constant_iterator(size_t{1}), + d_indices.begin(), + d_counts.begin()); d_indices.resize(thrust::distance(d_indices.begin(), thrust::get<0>(it)), handle_ptr_->get_stream()); d_counts.resize(d_indices.size(), handle_ptr_->get_stream()); diff --git a/cpp/src/centrality/betweenness_centrality.cu b/cpp/src/centrality/betweenness_centrality.cu index f60152ed525..70b1c87fbe3 100644 --- a/cpp/src/centrality/betweenness_centrality.cu +++ b/cpp/src/centrality/betweenness_centrality.cu @@ -25,6 +25,7 @@ #include #include +#include #include #include diff --git a/cpp/src/centrality/betweenness_centrality.cuh b/cpp/src/centrality/betweenness_centrality.cuh index 706b8bfebac..fe8093367cb 100644 --- a/cpp/src/centrality/betweenness_centrality.cuh +++ b/cpp/src/centrality/betweenness_centrality.cuh @@ -17,7 +17,8 @@ // Author: Xavier Cadet xcadet@nvidia.com #pragma once -#include +#include +#include namespace cugraph { namespace detail { diff --git a/cpp/src/centrality/katz_centrality.cu b/cpp/src/centrality/katz_centrality.cu index f574d2cbf11..a638694153b 100644 --- a/cpp/src/centrality/katz_centrality.cu +++ b/cpp/src/centrality/katz_centrality.cu @@ -22,8 +22,8 @@ #include #include -#include #include +#include #include #include @@ -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())->on(handle.get_stream()), + thrust::fill(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), 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 f2b8cb7a6b7..37c3c7278d7 100644 --- a/cpp/src/community/flatten_dendrogram.cuh +++ b/cpp/src/community/flatten_dendrogram.cuh @@ -18,8 +18,8 @@ #include #include -#include #include +#include namespace cugraph { @@ -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())->on(handle.get_stream()), + thrust::sequence(rmm::exec_policy(handle.get_stream()), 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/egonet.cu b/cpp/src/community/legacy/egonet.cu index 11df5a98262..a84bb16f3dc 100644 --- a/cpp/src/community/legacy/egonet.cu +++ b/cpp/src/community/legacy/egonet.cu @@ -21,8 +21,8 @@ #include #include -#include #include +#include #include #include diff --git a/cpp/src/community/legacy/extract_subgraph_by_vertex.cu b/cpp/src/community/legacy/extract_subgraph_by_vertex.cu index 224a3417caf..69443b977ea 100644 --- a/cpp/src/community/legacy/extract_subgraph_by_vertex.cu +++ b/cpp/src/community/legacy/extract_subgraph_by_vertex.cu @@ -18,8 +18,9 @@ #include #include -#include #include +#include +#include namespace { @@ -39,7 +40,7 @@ std::unique_ptr> extract_s int64_t* d_error_count = error_count_v.data().get(); thrust::for_each( - rmm::exec_policy(stream)->on(stream), + rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(num_vertices), [vertices, d_vertex_used, d_error_count, graph_num_verts] __device__(vertex_t idx) { @@ -60,7 +61,7 @@ std::unique_ptr> extract_s // iterate over the edges and count how many make it into the output int64_t count = thrust::count_if( - rmm::exec_policy(stream)->on(stream), + rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(graph.number_of_edges), [graph_src, graph_dst, d_vertex_used, num_vertices] __device__(edge_t e) { @@ -78,7 +79,7 @@ std::unique_ptr> extract_s weight_t* d_new_weight = result->edge_data(); // reusing error_count as a vertex counter... - thrust::for_each(rmm::exec_policy(stream)->on(stream), + thrust::for_each(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(graph.number_of_edges), [graph_src, diff --git a/cpp/src/community/legacy/leiden.cu b/cpp/src/community/legacy/leiden.cu index e1ab917e1e9..7044004d8ed 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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), 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/spectral_clustering.cu b/cpp/src/community/legacy/spectral_clustering.cu index 4dd27a56b70..c56b8eb641b 100644 --- a/cpp/src/community/legacy/spectral_clustering.cu +++ b/cpp/src/community/legacy/spectral_clustering.cu @@ -22,9 +22,10 @@ #include -#include #include #include +#include +#include #include #include @@ -71,7 +72,7 @@ void balancedCutClustering_impl(legacy::GraphCSRView raft::handle_t handle; auto stream = handle.get_stream(); auto exec = rmm::exec_policy(stream); - auto t_exe_p = exec->on(stream); + auto t_exe_p = exec; int evs_max_it{4000}; int kmean_max_it{200}; @@ -142,7 +143,7 @@ void spectralModularityMaximization_impl( raft::handle_t handle; auto stream = handle.get_stream(); auto exec = rmm::exec_policy(stream); - auto t_exe_p = exec->on(stream); + auto t_exe_p = exec; int evs_max_it{4000}; int kmean_max_it{200}; @@ -195,7 +196,7 @@ void analyzeModularityClustering_impl(legacy::GraphCSRViewon(stream); + auto t_exe_p = exec; using index_type = vertex_t; using value_type = weight_t; @@ -217,7 +218,7 @@ void analyzeBalancedCut_impl(legacy::GraphCSRView co raft::handle_t handle; auto stream = handle.get_stream(); auto exec = rmm::exec_policy(stream); - auto t_exe_p = exec->on(stream); + auto t_exe_p = exec; RAFT_EXPECTS(n_clusters <= graph.number_of_vertices, "API error: number of clusters must be <= number of vertices"); diff --git a/cpp/src/community/legacy/triangles_counting.cu b/cpp/src/community/legacy/triangles_counting.cu index 97543d28c62..8922f92336d 100644 --- a/cpp/src/community/legacy/triangles_counting.cu +++ b/cpp/src/community/legacy/triangles_counting.cu @@ -22,8 +22,9 @@ #include -#include #include +#include +#include #include diff --git a/cpp/src/community/louvain.cu b/cpp/src/community/louvain.cu index d221d5c5d53..d6bd224fedf 100644 --- a/cpp/src/community/louvain.cu +++ b/cpp/src/community/louvain.cu @@ -81,10 +81,8 @@ 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())->on(handle.get_stream()), - vertex_ids_v.begin(), - vertex_ids_v.end(), - vertex_t{0}); + thrust::sequence( + rmm::exec_policy(handle.get_stream()), vertex_ids_v.begin(), vertex_ids_v.end(), vertex_t{0}); partition_at_level( handle, dendrogram, vertex_ids_v.data(), clustering, dendrogram.num_levels()); @@ -100,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())->on(handle.get_stream()), + thrust::sequence(rmm::exec_policy(handle.get_stream()), vertex_ids_v.begin(), vertex_ids_v.end(), graph_view.get_local_vertex_first()); diff --git a/cpp/src/components/weak_cc.cuh b/cpp/src/components/weak_cc.cuh index e0da23c2ae8..31beda96342 100644 --- a/cpp/src/components/weak_cc.cuh +++ b/cpp/src/components/weak_cc.cuh @@ -28,7 +28,8 @@ #include #include -#include +#include +#include #include "utils.h" namespace MLCommon { diff --git a/cpp/src/converters/permute_graph.cuh b/cpp/src/converters/permute_graph.cuh index 024dfc2f3a7..cbff8ad3f7c 100644 --- a/cpp/src/converters/permute_graph.cuh +++ b/cpp/src/converters/permute_graph.cuh @@ -13,9 +13,10 @@ * See the License for the specific language governing permissions and * limitations under the License. */ -#include #include #include +#include +#include #include #include "converters/COOtoCSR.cuh" @@ -59,22 +60,18 @@ void permute_graph(legacy::GraphCSRView const& graph graph.get_source_indices(d_src); if (graph.has_data()) - thrust::copy(rmm::exec_policy(stream)->on(stream), + thrust::copy(rmm::exec_policy(stream), graph.edge_data, graph.edge_data + graph.number_of_edges, d_weights); // Permute the src_indices permutation_functor pf(permutation); - thrust::transform( - rmm::exec_policy(stream)->on(stream), d_src, d_src + graph.number_of_edges, d_src, pf); + thrust::transform(rmm::exec_policy(stream), d_src, d_src + graph.number_of_edges, d_src, pf); // Permute the destination indices - thrust::transform(rmm::exec_policy(stream)->on(stream), - graph.indices, - graph.indices + graph.number_of_edges, - d_dst, - pf); + thrust::transform( + rmm::exec_policy(stream), graph.indices, graph.indices + graph.number_of_edges, d_dst, pf); legacy::GraphCOOView graph_coo; diff --git a/cpp/src/cores/core_number.cu b/cpp/src/cores/core_number.cu index b23e7a25405..509d408e229 100644 --- a/cpp/src/cores/core_number.cu +++ b/cpp/src/cores/core_number.cu @@ -14,11 +14,12 @@ * limitations under the License. */ -#include #include #include #include #include +#include +#include //#include namespace cugraph { @@ -66,7 +67,7 @@ void extract_edges(legacy::GraphCOOView const& i_graph, thrust::make_tuple(i_graph.src_indices, i_graph.dst_indices, i_graph.edge_data)); auto outEdge = thrust::make_zip_iterator( thrust::make_tuple(o_graph.src_indices, o_graph.dst_indices, o_graph.edge_data)); - auto ptr = thrust::copy_if(rmm::exec_policy(stream)->on(stream), + auto ptr = thrust::copy_if(rmm::exec_policy(stream), inEdge, inEdge + i_graph.number_of_edges, outEdge, @@ -79,7 +80,7 @@ void extract_edges(legacy::GraphCOOView const& i_graph, thrust::make_zip_iterator(thrust::make_tuple(i_graph.src_indices, i_graph.dst_indices)); auto outEdge = thrust::make_zip_iterator(thrust::make_tuple(o_graph.src_indices, o_graph.dst_indices)); - auto ptr = thrust::copy_if(rmm::exec_policy(stream)->on(stream), + auto ptr = thrust::copy_if(rmm::exec_policy(stream), inEdge, inEdge + i_graph.number_of_edges, outEdge, @@ -109,8 +110,7 @@ std::unique_ptr> extract_subgraph( rmm::device_vector sorted_core_num(in_graph.number_of_vertices); - thrust::scatter( - rmm::exec_policy(stream)->on(stream), core_num, core_num + len, vid, sorted_core_num.begin()); + thrust::scatter(rmm::exec_policy(stream), core_num, core_num + len, vid, sorted_core_num.begin()); VT* d_sorted_core_num = sorted_core_num.data().get(); @@ -121,7 +121,7 @@ std::unique_ptr> extract_subgraph( auto out_graph = std::make_unique>( in_graph.number_of_vertices, - thrust::count_if(rmm::exec_policy(stream)->on(stream), + thrust::count_if(rmm::exec_policy(stream), edge, edge + in_graph.number_of_edges, detail::FilterEdges(k, d_sorted_core_num)), diff --git a/cpp/src/link_analysis/pagerank.cu b/cpp/src/link_analysis/pagerank.cu index 69d5927f629..9a569fafae6 100644 --- a/cpp/src/link_analysis/pagerank.cu +++ b/cpp/src/link_analysis/pagerank.cu @@ -24,8 +24,8 @@ #include #include -#include #include +#include #include #include @@ -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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::fill(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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 7b7470da7fc..446d0c8cfdb 100644 --- a/cpp/src/link_prediction/overlap.cu +++ b/cpp/src/link_prediction/overlap.cu @@ -19,9 +19,10 @@ * @file jaccard.cu * ---------------------------------------------------------------------------**/ -#include #include #include +#include +#include #include namespace cugraph { diff --git a/cpp/src/serialization/serializer.cu b/cpp/src/serialization/serializer.cu index 28529c9f3ed..2f4c8268a67 100644 --- a/cpp/src/serialization/serializer.cu +++ b/cpp/src/serialization/serializer.cu @@ -23,7 +23,7 @@ #include -#include +#include #include diff --git a/cpp/src/structure/coarsen_graph.cu b/cpp/src/structure/coarsen_graph.cu index a7abb4846bd..c66cc24932a 100644 --- a/cpp/src/structure/coarsen_graph.cu +++ b/cpp/src/structure/coarsen_graph.cu @@ -24,9 +24,9 @@ #include #include -#include #include #include +#include #include #include @@ -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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), *(matrix_partition.get_weights()), *(matrix_partition.get_weights()) + number_of_edges, (*edgelist_weights).data()); @@ -89,17 +89,15 @@ edge_t groupby_e_and_coarsen_edgelist(vertex_t* edgelist_major_vertices /* [INOU thrust::make_zip_iterator(thrust::make_tuple(edgelist_major_vertices, edgelist_minor_vertices)); if (edgelist_weights) { - thrust::sort_by_key(rmm::exec_policy(stream)->on(stream), - pair_first, - pair_first + number_of_edges, - *edgelist_weights); + thrust::sort_by_key( + rmm::exec_policy(stream), pair_first, pair_first + number_of_edges, *edgelist_weights); rmm::device_uvector tmp_edgelist_major_vertices(number_of_edges, stream); rmm::device_uvector tmp_edgelist_minor_vertices(tmp_edgelist_major_vertices.size(), stream); rmm::device_uvector tmp_edgelist_weights(tmp_edgelist_major_vertices.size(), stream); auto it = thrust::reduce_by_key( - rmm::exec_policy(stream)->on(stream), + rmm::exec_policy(stream), pair_first, pair_first + number_of_edges, (*edgelist_weights), @@ -113,7 +111,7 @@ edge_t groupby_e_and_coarsen_edgelist(vertex_t* edgelist_major_vertices /* [INOU thrust::make_zip_iterator(thrust::make_tuple(tmp_edgelist_major_vertices.begin(), tmp_edgelist_minor_vertices.begin(), tmp_edgelist_weights.begin())); - thrust::copy(rmm::exec_policy(stream)->on(stream), + thrust::copy(rmm::exec_policy(stream), edge_first, edge_first + ret, thrust::make_zip_iterator(thrust::make_tuple( @@ -121,11 +119,10 @@ edge_t groupby_e_and_coarsen_edgelist(vertex_t* edgelist_major_vertices /* [INOU return ret; } else { - thrust::sort(rmm::exec_policy(stream)->on(stream), pair_first, pair_first + number_of_edges); + thrust::sort(rmm::exec_policy(stream), pair_first, pair_first + number_of_edges); return static_cast(thrust::distance( pair_first, - thrust::unique( - rmm::exec_policy(stream)->on(stream), pair_first, pair_first + number_of_edges))); + thrust::unique(rmm::exec_policy(stream), pair_first, pair_first + number_of_edges))); } } @@ -148,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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), pair_first, pair_first + edgelist_major_vertices.size(), pair_first, @@ -263,7 +260,7 @@ 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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), labels, labels + major_labels.size(), major_labels.begin()); @@ -349,7 +346,7 @@ coarsen_graph( coarsened_edgelist_minor_vertices[j].begin(), (*coarsened_edgelist_weights)[j].begin())) + cur_size; - thrust::copy(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), src_edge_first, src_edge_first + number_of_partition_edges, dst_edge_first); @@ -361,7 +358,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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), src_edge_first, src_edge_first + number_of_partition_edges, dst_edge_first); @@ -391,31 +388,25 @@ 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())->on(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())->on(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())->on(handle.get_stream()), - unique_labels.begin(), - unique_labels.end())), - handle.get_stream()); + 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()); unique_labels = cugraph::detail::shuffle_vertices_by_gpu_id(handle, std::move(unique_labels)); - thrust::sort(rmm::exec_policy(handle.get_stream())->on(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())->on(handle.get_stream()), - unique_labels.begin(), - unique_labels.end())), - handle.get_stream()); + 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()); // 4. renumber @@ -510,19 +501,16 @@ coarsen_graph( rmm::device_uvector unique_labels(graph_view.get_number_of_vertices(), handle.get_stream()); - thrust::copy(rmm::exec_policy(handle.get_stream())->on(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())->on(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())->on(handle.get_stream()), - unique_labels.begin(), - unique_labels.end())), - handle.get_stream()); + 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()); 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 151ba4e307b..8d91206671c 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 diff --git a/cpp/src/structure/renumber_edgelist.cu b/cpp/src/structure/renumber_edgelist.cu index 9e593e3f169..3cf9954926b 100644 --- a/cpp/src/structure/renumber_edgelist.cu +++ b/cpp/src/structure/renumber_edgelist.cu @@ -356,7 +356,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())->on(handle.get_stream()), + thrust::upper_bound(rmm::exec_policy(handle.get_stream()), counts.begin(), counts.end(), d_thresholds.begin(), diff --git a/cpp/src/traversal/bfs.cu b/cpp/src/traversal/bfs.cu index 9c5449bcbd7..fa653b7ddb3 100644 --- a/cpp/src/traversal/bfs.cu +++ b/cpp/src/traversal/bfs.cu @@ -22,8 +22,8 @@ #include #include -#include #include +#include #include #include @@ -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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), 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 6bf8e0d0197..6fec3bde68d 100644 --- a/cpp/src/traversal/legacy/bfs.cuh +++ b/cpp/src/traversal/legacy/bfs.cuh @@ -11,8 +11,9 @@ #pragma once -#include #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 e6c8c3bf700..1059a8fa1de 100644 --- a/cpp/src/traversal/legacy/mg/bfs.cuh +++ b/cpp/src/traversal/legacy/mg/bfs.cuh @@ -17,6 +17,7 @@ #pragma once #include +#include #include "../traversal_common.cuh" #include "common_utils.cuh" #include "frontier_expand.cuh" @@ -83,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)->on(stream), + thrust::fill(rmm::exec_policy(stream), output_frontier_bmap.begin(), output_frontier_bmap.end(), static_cast(0)); @@ -129,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)->on(stream), + thrust::fill(rmm::exec_policy(stream), 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 9a95aba7901..6c1a4514619 100644 --- a/cpp/src/traversal/legacy/mg/common_utils.cuh +++ b/cpp/src/traversal/legacy/mg/common_utils.cuh @@ -18,9 +18,10 @@ #include "../traversal_common.cuh" -#include #include #include +#include +#include #include #include @@ -159,7 +160,7 @@ vertex_t populate_isolated_vertices( vertex_begin_ = 0; vertex_end_ = graph.number_of_vertices; } - auto count = thrust::copy_if(rmm::exec_policy(stream)->on(stream), + auto count = thrust::copy_if(rmm::exec_policy(stream), thrust::make_counting_iterator(vertex_begin_), thrust::make_counting_iterator(vertex_end_), thrust::make_counting_iterator(0), @@ -212,10 +213,8 @@ void add_to_bitmap(raft::handle_t const& handle, return_t count) { cudaStream_t stream = handle.get_stream(); - thrust::for_each(rmm::exec_policy(stream)->on(stream), - id.begin(), - id.begin() + count, - set_nth_bit(bmap.data().get())); + thrust::for_each( + rmm::exec_policy(stream), id.begin(), id.begin() + count, set_nth_bit(bmap.data().get())); CHECK_CUDA(stream); } @@ -247,10 +246,9 @@ 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)->on(stream), data.begin(), data.begin() + data_len); + thrust::sort(rmm::exec_policy(stream), data.begin(), data.begin() + data_len); auto unique_count = - thrust::unique(rmm::exec_policy(stream)->on(stream), data.begin(), data.begin() + data_len) - - data.begin(); + thrust::unique(rmm::exec_policy(stream), data.begin(), data.begin() + data_len) - data.begin(); return static_cast(unique_count); } @@ -372,8 +370,7 @@ return_t remove_duplicates(raft::handle_t const& handle, rmm::device_vector unique_count(1, 0); - thrust::fill( - rmm::exec_policy(stream)->on(stream), bmap.begin(), bmap.end(), static_cast(0)); + thrust::fill(rmm::exec_policy(stream), 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(), @@ -404,8 +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)->on(stream), bmap.begin(), bmap.end(), static_cast(0)); + thrust::fill(rmm::exec_policy(stream), 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(), @@ -436,8 +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)->on(stream), bmap.begin(), bmap.end(), static_cast(0)); + thrust::fill(rmm::exec_policy(stream), 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(), @@ -483,9 +478,8 @@ vertex_t get_global_vertex_count( cugraph::legacy::GraphCSRView const& graph) { rmm::device_vector id(1); - id[0] = *thrust::max_element(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), - graph.indices, - graph.indices + graph.number_of_edges); + id[0] = *thrust::max_element( + rmm::exec_policy(handle.get_stream()), 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/frontier_expand.cuh b/cpp/src/traversal/legacy/mg/frontier_expand.cuh index 078ab085724..8390dbaf5e3 100644 --- a/cpp/src/traversal/legacy/mg/frontier_expand.cuh +++ b/cpp/src/traversal/legacy/mg/frontier_expand.cuh @@ -17,6 +17,7 @@ #pragma once #include +#include #include "frontier_expand_kernels.cuh" #include "vertex_binning.cuh" diff --git a/cpp/src/traversal/legacy/mg/vertex_binning.cuh b/cpp/src/traversal/legacy/mg/vertex_binning.cuh index b4ed881a06e..22973f6d1a9 100644 --- a/cpp/src/traversal/legacy/mg/vertex_binning.cuh +++ b/cpp/src/traversal/legacy/mg/vertex_binning.cuh @@ -16,6 +16,7 @@ #pragma once +#include #include "common_utils.cuh" #include "vertex_binning_kernels.cuh" @@ -95,9 +96,8 @@ template LogDistribution VertexBinner::run( rmm::device_vector& reorganized_vertices, cudaStream_t stream) { - thrust::fill( - rmm::exec_policy(stream)->on(stream), bin_offsets_.begin(), bin_offsets_.end(), edge_t{0}); - thrust::fill(rmm::exec_policy(stream)->on(stream), tempBins_.begin(), tempBins_.end(), edge_t{0}); + thrust::fill(rmm::exec_policy(stream), bin_offsets_.begin(), bin_offsets_.end(), edge_t{0}); + thrust::fill(rmm::exec_policy(stream), tempBins_.begin(), tempBins_.end(), edge_t{0}); bin_vertices(reorganized_vertices, bin_offsets_, tempBins_, diff --git a/cpp/src/traversal/legacy/mg/vertex_binning_kernels.cuh b/cpp/src/traversal/legacy/mg/vertex_binning_kernels.cuh index 57574965a3a..012c65785a7 100644 --- a/cpp/src/traversal/legacy/mg/vertex_binning_kernels.cuh +++ b/cpp/src/traversal/legacy/mg/vertex_binning_kernels.cuh @@ -16,7 +16,8 @@ #pragma once -#include +#include +#include #include "../traversal_common.cuh" namespace cugraph { diff --git a/cpp/src/traversal/legacy/sssp.cu b/cpp/src/traversal/legacy/sssp.cu index 61225dd7fd6..394a22ad37a 100644 --- a/cpp/src/traversal/legacy/sssp.cu +++ b/cpp/src/traversal/legacy/sssp.cu @@ -18,6 +18,7 @@ #include #include +#include #include diff --git a/cpp/src/traversal/legacy/sssp.cuh b/cpp/src/traversal/legacy/sssp.cuh index df3f3621956..26388136eb4 100644 --- a/cpp/src/traversal/legacy/sssp.cuh +++ b/cpp/src/traversal/legacy/sssp.cuh @@ -17,7 +17,8 @@ // Author: Prasun Gera pgera@nvidia.com #pragma once -#include +#include +#include namespace cugraph { namespace detail { diff --git a/cpp/src/traversal/legacy/traversal_common.cuh b/cpp/src/traversal/legacy/traversal_common.cuh index ea77173870e..6e27df9bbb8 100644 --- a/cpp/src/traversal/legacy/traversal_common.cuh +++ b/cpp/src/traversal/legacy/traversal_common.cuh @@ -426,7 +426,7 @@ template void exclusive_sum(IndexType* d_in, IndexType* d_out, IndexType num_items, cudaStream_t m_stream) { if (num_items <= 1) return; // DeviceScan fails if n==1 - thrust::exclusive_scan(rmm::exec_policy(m_stream)->on(m_stream), d_in, d_in + num_items, d_out); + thrust::exclusive_scan(rmm::exec_policy(m_stream), d_in, d_in + num_items, d_out); } // diff --git a/cpp/src/traversal/sssp.cu b/cpp/src/traversal/sssp.cu index 8402a74181b..4301bcec431 100644 --- a/cpp/src/traversal/sssp.cu +++ b/cpp/src/traversal/sssp.cu @@ -26,7 +26,7 @@ #include #include -#include +#include #include #include @@ -93,7 +93,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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), 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 +143,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())->on(handle.get_stream()), + thrust::fill(rmm::exec_policy(handle.get_stream()), adj_matrix_row_distances.begin(), adj_matrix_row_distances.end(), std::numeric_limits::max()); diff --git a/cpp/src/traversal/tsp.hpp b/cpp/src/traversal/tsp.hpp index 6073f46ab28..4195b07fb45 100644 --- a/cpp/src/traversal/tsp.hpp +++ b/cpp/src/traversal/tsp.hpp @@ -21,9 +21,9 @@ #include #include -#include #include #include +#include namespace cugraph { namespace detail { diff --git a/cpp/src/traversal/two_hop_neighbors.cu b/cpp/src/traversal/two_hop_neighbors.cu index e1fce911130..ab6f5bcfaff 100644 --- a/cpp/src/traversal/two_hop_neighbors.cu +++ b/cpp/src/traversal/two_hop_neighbors.cu @@ -19,10 +19,11 @@ * @file two_hop_neighbors.cu * ---------------------------------------------------------------------------**/ -#include #include #include #include +#include +#include #include "two_hop_neighbors.cuh" #include @@ -44,14 +45,14 @@ std::unique_ptr> get_two_hop_neighbors( degree_iterator deg_it(graph.offsets); deref_functor, ET> deref(deg_it); exsum_degree[0] = ET{0}; - thrust::transform(rmm::exec_policy(stream)->on(stream), + thrust::transform(rmm::exec_policy(stream), graph.indices, graph.indices + graph.number_of_edges, d_exsum_degree + 1, deref); // Take the inclusive sum of the degrees - thrust::inclusive_scan(rmm::exec_policy(stream)->on(stream), + thrust::inclusive_scan(rmm::exec_policy(stream), d_exsum_degree + 1, d_exsum_degree + graph.number_of_edges + 1, d_exsum_degree + 1); @@ -98,13 +99,10 @@ std::unique_ptr> get_two_hop_neighbors( // Remove duplicates and self pairings auto tuple_start = thrust::make_zip_iterator(thrust::make_tuple(d_first_pair, d_second_pair)); auto tuple_end = tuple_start + output_size; - thrust::sort(rmm::exec_policy(stream)->on(stream), tuple_start, tuple_end); - tuple_end = thrust::copy_if(rmm::exec_policy(stream)->on(stream), - tuple_start, - tuple_end, - tuple_start, - self_loop_flagger()); - tuple_end = thrust::unique(rmm::exec_policy(stream)->on(stream), tuple_start, tuple_end); + thrust::sort(rmm::exec_policy(stream), tuple_start, tuple_end); + tuple_end = thrust::copy_if( + rmm::exec_policy(stream), tuple_start, tuple_end, tuple_start, self_loop_flagger()); + tuple_end = thrust::unique(rmm::exec_policy(stream), tuple_start, tuple_end); // Get things ready to return ET outputSize = tuple_end - tuple_start; diff --git a/cpp/src/tree/mst.cu b/cpp/src/tree/mst.cu index e6caa629cd1..e1c22fb87da 100644 --- a/cpp/src/tree/mst.cu +++ b/cpp/src/tree/mst.cu @@ -24,9 +24,9 @@ #include #include -#include #include #include +#include #include #include diff --git a/cpp/src/utilities/cython.cu b/cpp/src/utilities/cython.cu index 5574685f581..12948373192 100644 --- a/cpp/src/utilities/cython.cu +++ b/cpp/src/utilities/cython.cu @@ -30,8 +30,8 @@ #include -#include #include +#include #include #include @@ -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())->on(handle.get_stream()), + auto it = thrust::reduce_by_key(rmm::exec_policy(handle.get_stream()), key_first, key_first + graph_container.num_local_edges, thrust::make_constant_iterator(edge_t{1}), @@ -101,11 +101,9 @@ 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())->on(handle.get_stream()), - d_counts.begin(), - d_counts.end(), - edge_t{0}); - thrust::scatter(rmm::exec_policy(handle.get_stream())->on(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()), d_edge_counts.begin(), thrust::get<1>(it), d_local_partition_ids.begin(), @@ -510,7 +508,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())->on(handle.get_stream()), + thrust::copy(rmm::exec_policy(handle.get_stream()), 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 765cccc1916..b51593b117f 100644 --- a/cpp/src/utilities/path_retrieval.cu +++ b/cpp/src/utilities/path_retrieval.cu @@ -14,8 +14,8 @@ * limitations under the License. */ -#include #include +#include #include @@ -74,10 +74,9 @@ 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)->on(stream), vtx_map, vtx_map + num_vertices); + thrust::sequence(rmm::exec_policy(stream), vtx_map, vtx_map + num_vertices); - thrust::stable_sort_by_key( - rmm::exec_policy(stream)->on(stream), vtx_keys, vtx_keys + num_vertices, vtx_map); + thrust::stable_sort_by_key(rmm::exec_policy(stream), 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 31af0c75585..b45011d7f26 100644 --- a/cpp/src/utilities/spmv_1D.cuh +++ b/cpp/src/utilities/spmv_1D.cuh @@ -15,9 +15,10 @@ */ #pragma once -#include #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 10ac06c1536..203bf506811 100644 --- a/cpp/tests/centrality/legacy/betweenness_centrality_test.cu +++ b/cpp/tests/centrality/legacy/betweenness_centrality_test.cu @@ -24,6 +24,8 @@ #include #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 60b74497455..c5fd7af2bf6 100644 --- a/cpp/tests/centrality/legacy/edge_betweenness_centrality_test.cu +++ b/cpp/tests/centrality/legacy/edge_betweenness_centrality_test.cu @@ -21,6 +21,8 @@ #include #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 ee2df5347fc..34097cb244d 100644 --- a/cpp/tests/centrality/legacy/katz_centrality_test.cu +++ b/cpp/tests/centrality/legacy/katz_centrality_test.cu @@ -20,6 +20,8 @@ #include +#include + #include #include @@ -46,12 +48,9 @@ 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)->on(stream), id.begin(), id.end()); - thrust::sort_by_key(rmm::exec_policy(stream)->on(stream), - p_katz, - p_katz + count, - id.begin(), - thrust::greater()); + 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()); std::vector topK(k); thrust::copy(id.begin(), id.begin() + k, topK.begin()); return topK; @@ -65,7 +64,7 @@ int getMaxDegree(cugraph::legacy::GraphCSRView const& g) 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)->on(stream), + ET max_out_degree = thrust::reduce(rmm::exec_policy(stream), p_degree, p_degree + g.number_of_vertices, static_cast(-1), diff --git a/cpp/tests/community/balanced_edge_test.cpp b/cpp/tests/community/balanced_edge_test.cpp index d4c5edf3f35..38c64257b0a 100644 --- a/cpp/tests/community/balanced_edge_test.cpp +++ b/cpp/tests/community/balanced_edge_test.cpp @@ -12,7 +12,8 @@ #include -#include +#include +#include TEST(balanced_edge, success) { diff --git a/cpp/tests/community/ecg_test.cpp b/cpp/tests/community/ecg_test.cpp index f174d882937..7906ca19a9a 100644 --- a/cpp/tests/community/ecg_test.cpp +++ b/cpp/tests/community/ecg_test.cpp @@ -13,7 +13,8 @@ #include #include -#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/egonet_test.cu b/cpp/tests/community/egonet_test.cu index bd71266b549..f5e08d53340 100644 --- a/cpp/tests/community/egonet_test.cu +++ b/cpp/tests/community/egonet_test.cu @@ -30,8 +30,8 @@ #include #include -#include #include +#include #include #include diff --git a/cpp/tests/community/leiden_test.cpp b/cpp/tests/community/leiden_test.cpp index 13e139666f6..64b861720ab 100644 --- a/cpp/tests/community/leiden_test.cpp +++ b/cpp/tests/community/leiden_test.cpp @@ -15,7 +15,7 @@ #include -#include +#include TEST(leiden_karate, success) { diff --git a/cpp/tests/community/mg_louvain_helper.cu b/cpp/tests/community/mg_louvain_helper.cu index a160e480027..f5814af9820 100644 --- a/cpp/tests/community/mg_louvain_helper.cu +++ b/cpp/tests/community/mg_louvain_helper.cu @@ -22,7 +22,7 @@ #include #include -#include +#include #include #include @@ -40,19 +40,19 @@ 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()); thrust::for_each( - rmm::exec_policy(handle.get_stream())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), edgelist_cols_v.begin(), edgelist_cols_v.end(), edgelist_cols_v.begin(), @@ -84,7 +84,7 @@ 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)->on(stream), + thrust::for_each(rmm::exec_policy(stream), thrust::make_counting_iterator(major_first), thrust::make_counting_iterator(major_last), [compressed_sparse_offsets, @@ -94,12 +94,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)->on(stream), + thrust::copy(rmm::exec_policy(stream), compressed_sparse_indices, compressed_sparse_indices + number_of_edges, edgelist_minor_vertices.begin()); if (compressed_sparse_weights) { - thrust::copy(rmm::exec_policy(stream)->on(stream), + thrust::copy(rmm::exec_policy(stream), (*compressed_sparse_weights), (*compressed_sparse_weights) + number_of_edges, (*edgelist_weights).data()); @@ -122,7 +122,7 @@ void sort_and_coarsen_edgelist( size_t number_of_edges{0}; if (edgelist_weights) { - thrust::sort_by_key(rmm::exec_policy(stream)->on(stream), + thrust::sort_by_key(rmm::exec_policy(stream), pair_first, pair_first + edgelist_major_vertices.size(), (*edgelist_weights).begin()); @@ -133,7 +133,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)->on(stream), + rmm::exec_policy(stream), pair_first, pair_first + edgelist_major_vertices.size(), (*edgelist_weights).begin(), @@ -146,12 +146,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)->on(stream), - pair_first, - pair_first + edgelist_major_vertices.size()); - auto it = thrust::unique(rmm::exec_policy(stream)->on(stream), - pair_first, - pair_first + edgelist_major_vertices.size()); + 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()); number_of_edges = thrust::distance(pair_first, it); } @@ -195,7 +192,7 @@ compressed_sparse_to_relabeled_and_sorted_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(stream)->on(stream), + rmm::exec_policy(stream), pair_first, pair_first + edgelist_major_vertices.size(), pair_first, @@ -246,12 +243,11 @@ 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())->on(handle.get_stream()), - labels, - labels + graph_view.get_number_of_vertices(), - vertex_t{0}, - thrust::maximum()); + vertex_t new_number_of_vertices = 1 + thrust::reduce(rmm::exec_policy(handle.get_stream()), + labels, + labels + graph_view.get_number_of_vertices(), + vertex_t{0}, + thrust::maximum()); return std::make_unique>( handle, diff --git a/cpp/tests/community/triangle_test.cu b/cpp/tests/community/triangle_test.cu index 4c51e15b111..f607ed6bf87 100644 --- a/cpp/tests/community/triangle_test.cu +++ b/cpp/tests/community/triangle_test.cu @@ -13,7 +13,8 @@ #include #include -#include +#include +#include TEST(triangle, dolphin) { diff --git a/cpp/tests/components/con_comp_test.cu b/cpp/tests/components/con_comp_test.cu index d58ebc03a11..331ba53b3a7 100644 --- a/cpp/tests/components/con_comp_test.cu +++ b/cpp/tests/components/con_comp_test.cu @@ -18,6 +18,8 @@ #include +#include + #include #include #include diff --git a/cpp/tests/components/scc_test.cu b/cpp/tests/components/scc_test.cu index e458b5c4be7..d97bb62201b 100644 --- a/cpp/tests/components/scc_test.cu +++ b/cpp/tests/components/scc_test.cu @@ -16,6 +16,8 @@ #include #include +#include + #include #include #include diff --git a/cpp/tests/generators/erdos_renyi_test.cpp b/cpp/tests/generators/erdos_renyi_test.cpp index b99b9a21194..3fdf8c1eda3 100644 --- a/cpp/tests/generators/erdos_renyi_test.cpp +++ b/cpp/tests/generators/erdos_renyi_test.cpp @@ -20,6 +20,7 @@ #include #include +#include #include #include diff --git a/cpp/tests/layout/force_atlas2_test.cu b/cpp/tests/layout/force_atlas2_test.cu index 251b2a85c87..6f02a5dace9 100644 --- a/cpp/tests/layout/force_atlas2_test.cu +++ b/cpp/tests/layout/force_atlas2_test.cu @@ -20,8 +20,8 @@ #include #include -#include #include +#include #include diff --git a/cpp/tests/linear_assignment/hungarian_test.cu b/cpp/tests/linear_assignment/hungarian_test.cu index f806a217a8f..81ffec90221 100644 --- a/cpp/tests/linear_assignment/hungarian_test.cu +++ b/cpp/tests/linear_assignment/hungarian_test.cu @@ -359,7 +359,7 @@ void random_test(int32_t num_rows, int32_t num_cols, int32_t upper_bound, int re //int64_t seed{85}; int64_t seed{time(nullptr)}; - thrust::for_each(rmm::exec_policy(stream)->on(stream), + thrust::for_each(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(num_rows * num_cols), [d_data, seed, upper_bound] __device__ (int32_t e) { diff --git a/cpp/tests/prims/mg_count_if_v.cu b/cpp/tests/prims/mg_count_if_v.cu index 2684c3121d9..f4d4c24dd71 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())->on(handle.get_stream()), + thrust::count_if(rmm::exec_policy(handle.get_stream()), 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 37c0cf22ef4..1c5c1c261b1 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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), 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())->on(handle.get_stream()), + thrust::reduce(rmm::exec_policy(handle.get_stream()), 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 b896da31b0b..4f7f3a5a724 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())->on(handle.get_stream()), + rmm::exec_policy(handle.get_stream()), 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 819e6a10ddb..ca105a482b9 100644 --- a/cpp/tests/sampling/random_walks_profiling.cu +++ b/cpp/tests/sampling/random_walks_profiling.cu @@ -24,7 +24,7 @@ #include #include -#include +#include #include #include @@ -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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), 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 c291810e363..f3603549eb5 100644 --- a/cpp/tests/sampling/random_walks_test.cu +++ b/cpp/tests/sampling/random_walks_test.cu @@ -20,8 +20,8 @@ #include #include -#include #include +#include #include #include @@ -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())->on(handle.get_stream()), + thrust::transform(rmm::exec_policy(handle.get_stream()), thrust::make_counting_iterator(0), thrust::make_counting_iterator(num_paths), diff --git a/cpp/tests/sampling/random_walks_utils.cuh b/cpp/tests/sampling/random_walks_utils.cuh index 5b8130aab30..df3bd08ce9b 100644 --- a/cpp/tests/sampling/random_walks_utils.cuh +++ b/cpp/tests/sampling/random_walks_utils.cuh @@ -15,7 +15,7 @@ */ #pragma once -#include +#include #include #include diff --git a/cpp/tests/sampling/rw_biased_seg_sort.cu b/cpp/tests/sampling/rw_biased_seg_sort.cu index 620f0aa6ff5..ab0436a2fab 100644 --- a/cpp/tests/sampling/rw_biased_seg_sort.cu +++ b/cpp/tests/sampling/rw_biased_seg_sort.cu @@ -20,7 +20,7 @@ #include #include -#include +#include #include diff --git a/cpp/tests/sampling/rw_low_level_test.cu b/cpp/tests/sampling/rw_low_level_test.cu index e3feb3adcf6..d1e444c0513 100644 --- a/cpp/tests/sampling/rw_low_level_test.cu +++ b/cpp/tests/sampling/rw_low_level_test.cu @@ -21,8 +21,8 @@ #include #include -#include #include +#include #include #include diff --git a/cpp/tests/serialization/un_serialize_test.cpp b/cpp/tests/serialization/un_serialize_test.cpp index e65d37fd77a..d1a08818978 100644 --- a/cpp/tests/serialization/un_serialize_test.cpp +++ b/cpp/tests/serialization/un_serialize_test.cpp @@ -19,8 +19,8 @@ #include #include -#include #include +#include #include diff --git a/cpp/tests/traversal/legacy/bfs_test.cu b/cpp/tests/traversal/legacy/bfs_test.cu index b0da605a0a0..fbc89439662 100644 --- a/cpp/tests/traversal/legacy/bfs_test.cu +++ b/cpp/tests/traversal/legacy/bfs_test.cu @@ -21,7 +21,8 @@ #include -#include +#include +#include #include diff --git a/cpp/tests/traversal/legacy/sssp_test.cu b/cpp/tests/traversal/legacy/sssp_test.cu index 5a19457543b..74257256dca 100644 --- a/cpp/tests/traversal/legacy/sssp_test.cu +++ b/cpp/tests/traversal/legacy/sssp_test.cu @@ -13,6 +13,8 @@ #include #include +#include + #include #include #include diff --git a/cpp/tests/traversal/ms_bfs_test.cpp b/cpp/tests/traversal/ms_bfs_test.cpp index ae81c8f1444..50ddced1bd5 100644 --- a/cpp/tests/traversal/ms_bfs_test.cpp +++ b/cpp/tests/traversal/ms_bfs_test.cpp @@ -31,9 +31,9 @@ #include #include -#include #include #include +#include #include #include diff --git a/cpp/tests/utilities/base_fixture.hpp b/cpp/tests/utilities/base_fixture.hpp index c2964fb9805..d58d71d9bd2 100644 --- a/cpp/tests/utilities/base_fixture.hpp +++ b/cpp/tests/utilities/base_fixture.hpp @@ -20,7 +20,7 @@ #include #include -#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 42f047db811..22498a124aa 100644 --- a/cpp/tests/utilities/matrix_market_file_utilities.cu +++ b/cpp/tests/utilities/matrix_market_file_utilities.cu @@ -23,7 +23,7 @@ #include #include -#include +#include #include @@ -335,10 +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())->on(handle.get_stream()), - d_vertices.begin(), - d_vertices.end(), - vertex_t{0}); + thrust::sequence( + rmm::exec_policy(handle.get_stream()), d_vertices.begin(), d_vertices.end(), vertex_t{0}); handle.get_stream_view().synchronize(); if (multi_gpu) { @@ -352,14 +350,12 @@ 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())->on(handle.get_stream()), - d_vertices.begin(), - d_vertices.end(), - [comm_rank, key_func = vertex_key_func] __device__(auto val) { - return key_func(val) != comm_rank; - })), + thrust::distance(d_vertices.begin(), + thrust::remove_if(rmm::exec_policy(handle.get_stream()), + d_vertices.begin(), + d_vertices.end(), + [comm_rank, key_func = vertex_key_func] __device__( + auto val) { return key_func(val) != comm_rank; })), handle.get_stream()); d_vertices.shrink_to_fit(handle.get_stream()); @@ -371,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())->on(handle.get_stream()), + thrust::remove_if(rmm::exec_policy(handle.get_stream()), edge_first, edge_first + d_edgelist_rows.size(), [comm_rank, key_func = edge_key_func] __device__(auto e) { @@ -384,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())->on(handle.get_stream()), + thrust::remove_if(rmm::exec_policy(handle.get_stream()), edge_first, edge_first + d_edgelist_rows.size(), [comm_rank, key_func = edge_key_func] __device__(auto e) {