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