diff --git a/cpp/cmake/thirdparty/get_cuco.cmake b/cpp/cmake/thirdparty/get_cuco.cmake index 63453071ece..ac1eb993d35 100644 --- a/cpp/cmake/thirdparty/get_cuco.cmake +++ b/cpp/cmake/thirdparty/get_cuco.cmake @@ -22,7 +22,7 @@ function(find_and_configure_cuco VERSION) CPM_ARGS EXCLUDE_FROM_ALL TRUE GIT_REPOSITORY https://github.com/NVIDIA/cuCollections.git - GIT_TAG 0ca860b824f5dc22cf8a41f09912e62e11f07d82 + GIT_TAG 55029034c3f82bca36148c9be29941b37492394d OPTIONS "BUILD_TESTS OFF" "BUILD_BENCHMARKS OFF" "BUILD_EXAMPLES OFF" diff --git a/cpp/include/cugraph/prims/detail/nbr_intersection.cuh b/cpp/include/cugraph/prims/detail/nbr_intersection.cuh index c0826d2bc12..76011217548 100644 --- a/cpp/include/cugraph/prims/detail/nbr_intersection.cuh +++ b/cpp/include/cugraph/prims/detail/nbr_intersection.cuh @@ -834,8 +834,8 @@ nbr_intersection(raft::handle_t const& handle, // cuco::static_map requires at least one empty slot std::max(static_cast(static_cast(unique_majors.size()) / load_factor), static_cast(unique_majors.size()) + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()); auto pair_first = thrust::make_zip_iterator(unique_majors.begin(), diff --git a/cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh b/cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh index ddf579f93ed..5c2e3947b04 100644 --- a/cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh +++ b/cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh @@ -297,8 +297,8 @@ void per_v_transform_reduce_dst_key_aggregated_outgoing_e( static_cast(thrust::distance(map_unique_key_first, map_unique_key_last)) / load_factor), static_cast(thrust::distance(map_unique_key_first, map_unique_key_last)) + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()); @@ -570,8 +570,8 @@ void per_v_transform_reduce_dst_key_aggregated_outgoing_e( auto multi_gpu_kv_map_ptr = std::make_unique< cuco::static_map>( size_t{0}, - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()); // relevant only when GraphViewType::is_multi_gpu is true if constexpr (GraphViewType::is_multi_gpu) { @@ -608,8 +608,8 @@ void per_v_transform_reduce_dst_key_aggregated_outgoing_e( // cuco::static_map requires at least one empty slot std::max(static_cast(static_cast(unique_minor_keys.size()) / load_factor), static_cast(unique_minor_keys.size()) + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()); diff --git a/cpp/include/cugraph/utilities/collect_comm.cuh b/cpp/include/cugraph/utilities/collect_comm.cuh index d41383a5188..84debe0b6a4 100644 --- a/cpp/include/cugraph/utilities/collect_comm.cuh +++ b/cpp/include/cugraph/utilities/collect_comm.cuh @@ -73,8 +73,8 @@ collect_values_for_keys(raft::comms::comms_t const& comm, std::max(static_cast( static_cast(thrust::distance(map_key_first, map_key_last)) / load_factor), static_cast(thrust::distance(map_key_first, map_key_last)) + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, stream_view); { @@ -136,8 +136,8 @@ collect_values_for_keys(raft::comms::comms_t const& comm, // cuco::static_map requires at least one empty slot std::max(static_cast(static_cast(unique_keys.size()) / load_factor), unique_keys.size() + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, stream_view); { @@ -233,8 +233,8 @@ collect_values_for_unique_keys( std::max(static_cast( static_cast(thrust::distance(map_key_first, map_key_last)) / load_factor), static_cast(thrust::distance(map_key_first, map_key_last)) + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, stream_view); { diff --git a/cpp/src/structure/relabel_impl.cuh b/cpp/src/structure/relabel_impl.cuh index 0c1c9475e2a..5e95348ddb5 100644 --- a/cpp/src/structure/relabel_impl.cuh +++ b/cpp/src/structure/relabel_impl.cuh @@ -113,8 +113,8 @@ void relabel(raft::handle_t const& handle, std::max(static_cast( static_cast(rx_label_pair_old_labels.size()) / load_factor), rx_label_pair_old_labels.size() + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()}; @@ -178,8 +178,8 @@ void relabel(raft::handle_t const& handle, // cuco::static_map requires at least one empty slot std::max(static_cast(static_cast(unique_old_labels.size()) / load_factor), unique_old_labels.size() + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()}; @@ -205,8 +205,8 @@ void relabel(raft::handle_t const& handle, // cuco::static_map requires at least one empty slot std::max(static_cast(static_cast(num_label_pairs) / load_factor), static_cast(num_label_pairs) + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()); diff --git a/cpp/src/structure/renumber_edgelist_impl.cuh b/cpp/src/structure/renumber_edgelist_impl.cuh index 92b2db13b7e..5dc787c1244 100644 --- a/cpp/src/structure/renumber_edgelist_impl.cuh +++ b/cpp/src/structure/renumber_edgelist_impl.cuh @@ -747,8 +747,8 @@ renumber_edgelist( static_cast(partition.local_edge_partition_major_range_size(i)) / load_factor), static_cast(partition.local_edge_partition_major_range_size(i)) + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()}; auto pair_first = thrust::make_zip_iterator(thrust::make_tuple( @@ -798,8 +798,8 @@ renumber_edgelist( renumber_map{// cuco::static_map requires at least one empty slot std::max(static_cast(static_cast(segment_size) / load_factor), static_cast(segment_size) + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()}; auto pair_first = thrust::make_zip_iterator(thrust::make_tuple( @@ -844,8 +844,8 @@ renumber_edgelist( std::max(static_cast( static_cast(renumber_map_minor_labels.size()) / load_factor), renumber_map_minor_labels.size() + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()}; auto pair_first = thrust::make_zip_iterator(thrust::make_tuple( @@ -917,8 +917,8 @@ renumber_edgelist(raft::handle_t const& handle, // cuco::static_map requires at least one empty slot std::max(static_cast(static_cast(renumber_map_labels.size()) / load_factor), renumber_map_labels.size() + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()}; auto pair_first = thrust::make_zip_iterator( diff --git a/cpp/src/structure/renumber_utils_impl.cuh b/cpp/src/structure/renumber_utils_impl.cuh index c61f8148cfb..a934335acdf 100644 --- a/cpp/src/structure/renumber_utils_impl.cuh +++ b/cpp/src/structure/renumber_utils_impl.cuh @@ -199,8 +199,8 @@ void unrenumber_local_int_edges( std::max(static_cast( static_cast(edge_partition_major_range_size) / load_factor), static_cast(edge_partition_major_range_size) + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()}; auto pair_first = thrust::make_zip_iterator( @@ -262,8 +262,8 @@ void unrenumber_local_int_edges( renumber_map{// cuco::static_map requires at least one empty slot std::max(static_cast(static_cast(segment_size) / load_factor), static_cast(segment_size) + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()}; auto pair_first = thrust::make_zip_iterator( @@ -315,8 +315,8 @@ void unrenumber_local_int_edges( std::max(static_cast( static_cast(renumber_map_minor_labels.size()) / load_factor), renumber_map_minor_labels.size() + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()}; auto pair_first = thrust::make_zip_iterator( @@ -369,8 +369,8 @@ void renumber_ext_vertices(raft::handle_t const& handle, auto renumber_map_ptr = std::make_unique< cuco::static_map>( size_t{0}, - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()); if (multi_gpu) { @@ -417,8 +417,8 @@ void renumber_ext_vertices(raft::handle_t const& handle, std::max( static_cast(static_cast(sorted_unique_ext_vertices.size()) / load_factor), sorted_unique_ext_vertices.size() + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()); @@ -438,8 +438,8 @@ void renumber_ext_vertices(raft::handle_t const& handle, std::max(static_cast( static_cast(local_int_vertex_last - local_int_vertex_first) / load_factor), static_cast(local_int_vertex_last - local_int_vertex_first) + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()); @@ -504,8 +504,8 @@ void renumber_local_ext_vertices(raft::handle_t const& handle, std::max(static_cast( static_cast(local_int_vertex_last - local_int_vertex_first) / load_factor), static_cast(local_int_vertex_last - local_int_vertex_first) + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()); @@ -673,8 +673,8 @@ void unrenumber_int_vertices(raft::handle_t const& handle, std::max(static_cast( static_cast(sorted_unique_int_vertices.size()) / load_factor), sorted_unique_int_vertices.size() + 1), - invalid_vertex_id::value, - invalid_vertex_id::value, + cuco::sentinel::empty_key{invalid_vertex_id::value}, + cuco::sentinel::empty_value{0}, stream_adapter, handle.get_stream()};