Support raft updating to new version of cuco (#2360)

RAFT is updating to the newest version of `cuco` in this PR: rapidsai/raft#714

This PR makes the necessary changes to cugraph to support the latest `cuco`.  It should be merged after the RAFT PR.
 * update version of cuco
 * update constructor calls to cuco::static_map

Authors:
  - Chuck Hastings (https://github.com/ChuckHastings)

Approvers:
  - Brad Rees (https://github.com/BradReesWork)

URL: #2360

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"

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<size_t>(static_cast<double>(unique_majors.size()) / load_factor),
                  static_cast<size_t>(unique_majors.size()) + 1),
-        invalid_vertex_id<vertex_t>::value,
-        invalid_vertex_id<vertex_t>::value,
+        cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+        cuco::sentinel::empty_value<vertex_t>{0},
         stream_adapter,
         handle.get_stream());
       auto pair_first = thrust::make_zip_iterator(unique_majors.begin(),

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<double>(thrust::distance(map_unique_key_first, map_unique_key_last)) /
           load_factor),
         static_cast<size_t>(thrust::distance(map_unique_key_first, map_unique_key_last)) + 1),
-      invalid_vertex_id<vertex_t>::value,
-      invalid_vertex_id<vertex_t>::value,
+      cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+      cuco::sentinel::empty_value<value_t>{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<vertex_t, value_t, cuda::thread_scope_device, decltype(stream_adapter)>>(
       size_t{0},
-      invalid_vertex_id<vertex_t>::value,
-      invalid_vertex_id<vertex_t>::value,
+      cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+      cuco::sentinel::empty_value<value_t>{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<size_t>(static_cast<double>(unique_minor_keys.size()) / load_factor),
                  static_cast<size_t>(unique_minor_keys.size()) + 1),
-        invalid_vertex_id<vertex_t>::value,
-        invalid_vertex_id<vertex_t>::value,
+        cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+        cuco::sentinel::empty_value<value_t>{0},
         stream_adapter,
         handle.get_stream());
 

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<size_t>(
                static_cast<double>(thrust::distance(map_key_first, map_key_last)) / load_factor),
              static_cast<size_t>(thrust::distance(map_key_first, map_key_last)) + 1),
-    invalid_vertex_id<vertex_t>::value,
-    invalid_vertex_id<vertex_t>::value,
+    cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+    cuco::sentinel::empty_value<value_t>{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<size_t>(static_cast<double>(unique_keys.size()) / load_factor),
              unique_keys.size() + 1),
-    invalid_vertex_id<vertex_t>::value,
-    invalid_vertex_id<vertex_t>::value,
+    cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+    cuco::sentinel::empty_value<value_t>{0},
     stream_adapter,
     stream_view);
   {
@@ -233,8 +233,8 @@ collect_values_for_unique_keys(
     std::max(static_cast<size_t>(
                static_cast<double>(thrust::distance(map_key_first, map_key_last)) / load_factor),
              static_cast<size_t>(thrust::distance(map_key_first, map_key_last)) + 1),
-    invalid_vertex_id<vertex_t>::value,
-    invalid_vertex_id<vertex_t>::value,
+    cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+    cuco::sentinel::empty_value<value_t>{0},
     stream_adapter,
     stream_view);
   {

cpp/src/structure/relabel_impl.cuh

@@ -113,8 +113,8 @@ void relabel(raft::handle_t const& handle,
                     std::max(static_cast<size_t>(
                                static_cast<double>(rx_label_pair_old_labels.size()) / load_factor),
                              rx_label_pair_old_labels.size() + 1),
-                    invalid_vertex_id<vertex_t>::value,
-                    invalid_vertex_id<vertex_t>::value,
+                    cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+                    cuco::sentinel::empty_value<vertex_t>{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<size_t>(static_cast<double>(unique_old_labels.size()) / load_factor),
                    unique_old_labels.size() + 1),
-          invalid_vertex_id<vertex_t>::value,
-          invalid_vertex_id<vertex_t>::value,
+          cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+          cuco::sentinel::empty_value<vertex_t>{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<size_t>(static_cast<double>(num_label_pairs) / load_factor),
                  static_cast<size_t>(num_label_pairs) + 1),
-        invalid_vertex_id<vertex_t>::value,
-        invalid_vertex_id<vertex_t>::value,
+        cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+        cuco::sentinel::empty_value<vertex_t>{0},
         stream_adapter,
         handle.get_stream());
 

cpp/src/structure/renumber_edgelist_impl.cuh

@@ -747,8 +747,8 @@ renumber_edgelist(
                      static_cast<double>(partition.local_edge_partition_major_range_size(i)) /
                      load_factor),
                    static_cast<size_t>(partition.local_edge_partition_major_range_size(i)) + 1),
-          invalid_vertex_id<vertex_t>::value,
-          invalid_vertex_id<vertex_t>::value,
+          cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+          cuco::sentinel::empty_value<vertex_t>{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<size_t>(static_cast<double>(segment_size) / load_factor),
                               static_cast<size_t>(segment_size) + 1),
-                     invalid_vertex_id<vertex_t>::value,
-                     invalid_vertex_id<vertex_t>::value,
+                     cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+                     cuco::sentinel::empty_value<vertex_t>{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<size_t>(
                               static_cast<double>(renumber_map_minor_labels.size()) / load_factor),
                             renumber_map_minor_labels.size() + 1),
-                   invalid_vertex_id<vertex_t>::value,
-                   invalid_vertex_id<vertex_t>::value,
+                   cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+                   cuco::sentinel::empty_value<vertex_t>{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<size_t>(static_cast<double>(renumber_map_labels.size()) / load_factor),
                renumber_map_labels.size() + 1),
-      invalid_vertex_id<vertex_t>::value,
-      invalid_vertex_id<vertex_t>::value,
+      cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+      cuco::sentinel::empty_value<vertex_t>{0},
       stream_adapter,
       handle.get_stream()};
   auto pair_first = thrust::make_zip_iterator(

cpp/src/structure/renumber_utils_impl.cuh

@@ -199,8 +199,8 @@ void unrenumber_local_int_edges(
                      std::max(static_cast<size_t>(
                                 static_cast<double>(edge_partition_major_range_size) / load_factor),
                               static_cast<size_t>(edge_partition_major_range_size) + 1),
-                     invalid_vertex_id<vertex_t>::value,
-                     invalid_vertex_id<vertex_t>::value,
+                     cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+                     cuco::sentinel::empty_value<vertex_t>{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<size_t>(static_cast<double>(segment_size) / load_factor),
                               static_cast<size_t>(segment_size) + 1),
-                     invalid_vertex_id<vertex_t>::value,
-                     invalid_vertex_id<vertex_t>::value,
+                     cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+                     cuco::sentinel::empty_value<vertex_t>{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<size_t>(
                               static_cast<double>(renumber_map_minor_labels.size()) / load_factor),
                             renumber_map_minor_labels.size() + 1),
-                   invalid_vertex_id<vertex_t>::value,
-                   invalid_vertex_id<vertex_t>::value,
+                   cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+                   cuco::sentinel::empty_value<vertex_t>{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<vertex_t, vertex_t, cuda::thread_scope_device, decltype(stream_adapter)>>(
     size_t{0},
-    invalid_vertex_id<vertex_t>::value,
-    invalid_vertex_id<vertex_t>::value,
+    cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+    cuco::sentinel::empty_value<vertex_t>{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<size_t>(static_cast<double>(sorted_unique_ext_vertices.size()) / load_factor),
         sorted_unique_ext_vertices.size() + 1),
-      invalid_vertex_id<vertex_t>::value,
-      invalid_vertex_id<vertex_t>::value,
+      cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+      cuco::sentinel::empty_value<vertex_t>{0},
       stream_adapter,
       handle.get_stream());
 
@@ -438,8 +438,8 @@ void renumber_ext_vertices(raft::handle_t const& handle,
       std::max(static_cast<size_t>(
                  static_cast<double>(local_int_vertex_last - local_int_vertex_first) / load_factor),
                static_cast<size_t>(local_int_vertex_last - local_int_vertex_first) + 1),
-      invalid_vertex_id<vertex_t>::value,
-      invalid_vertex_id<vertex_t>::value,
+      cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+      cuco::sentinel::empty_value<vertex_t>{0},
       stream_adapter,
       handle.get_stream());
 
@@ -504,8 +504,8 @@ void renumber_local_ext_vertices(raft::handle_t const& handle,
     std::max(static_cast<size_t>(
                static_cast<double>(local_int_vertex_last - local_int_vertex_first) / load_factor),
              static_cast<size_t>(local_int_vertex_last - local_int_vertex_first) + 1),
-    invalid_vertex_id<vertex_t>::value,
-    invalid_vertex_id<vertex_t>::value,
+    cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+    cuco::sentinel::empty_value<vertex_t>{0},
     stream_adapter,
     handle.get_stream());
 
@@ -673,8 +673,8 @@ void unrenumber_int_vertices(raft::handle_t const& handle,
                    std::max(static_cast<size_t>(
                               static_cast<double>(sorted_unique_int_vertices.size()) / load_factor),
                             sorted_unique_int_vertices.size() + 1),
-                   invalid_vertex_id<vertex_t>::value,
-                   invalid_vertex_id<vertex_t>::value,
+                   cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
+                   cuco::sentinel::empty_value<vertex_t>{0},
                    stream_adapter,
                    handle.get_stream()};