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) {