Update cuCollections for thrust upgrade

cpp/cmake/thirdparty/get_cucollections.cmake

@@ -22,7 +22,7 @@ function(find_and_configure_cucollections)
     GLOBAL_TARGETS cuco::cuco
     BUILD_EXPORT_SET cudf-exports
     CPM_ARGS GITHUB_REPOSITORY NVIDIA/cuCollections
-    GIT_TAG ebaba1ae378a5272116414b6d7ae5847e5cf5715
+    GIT_TAG 55029034c3f82bca36148c9be29941b37492394d
     EXCLUDE_FROM_ALL ${BUILD_SHARED_LIBS}
     OPTIONS "BUILD_TESTS OFF" "BUILD_BENCHMARKS OFF" "BUILD_EXAMPLES OFF"
   )

cpp/src/io/parquet/chunk_dict.cu

@@ -125,8 +125,11 @@ __global__ void __launch_bounds__(block_size)
   column_device_view const& data_col = *col->leaf_column;
 
   // Make a view of the hash map
-  auto hash_map_mutable = map_type::device_mutable_view(
-    chunk->dict_map_slots, chunk->dict_map_size, KEY_SENTINEL, VALUE_SENTINEL);
+  auto hash_map_mutable =
+    map_type::device_mutable_view(chunk->dict_map_slots,
+                                  chunk->dict_map_size,
+                                  cuco::sentinel::empty_key{KEY_SENTINEL},
+                                  cuco::sentinel::empty_value{VALUE_SENTINEL});
 
   __shared__ size_type total_num_dict_entries;
   size_type val_idx = s_start_value_idx + t;
@@ -184,9 +187,11 @@ __global__ void __launch_bounds__(block_size)
   auto& chunk = chunks[blockIdx.x];
   if (not chunk.use_dictionary) { return; }
 
-  auto t = threadIdx.x;
-  auto map =
-    map_type::device_view(chunk.dict_map_slots, chunk.dict_map_size, KEY_SENTINEL, VALUE_SENTINEL);
+  auto t   = threadIdx.x;
+  auto map = map_type::device_view(chunk.dict_map_slots,
+                                   chunk.dict_map_size,
+                                   cuco::sentinel::empty_key{KEY_SENTINEL},
+                                   cuco::sentinel::empty_value{VALUE_SENTINEL});
 
   __shared__ cuda::atomic<size_type, cuda::thread_scope_block> counter;
   using cuda::std::memory_order_relaxed;
@@ -233,8 +238,10 @@ __global__ void __launch_bounds__(block_size)
 
   column_device_view const& data_col = *col->leaf_column;
 
-  auto map = map_type::device_view(
-    chunk->dict_map_slots, chunk->dict_map_size, KEY_SENTINEL, VALUE_SENTINEL);
+  auto map = map_type::device_view(chunk->dict_map_slots,
+                                   chunk->dict_map_size,
+                                   cuco::sentinel::empty_key{KEY_SENTINEL},
+                                   cuco::sentinel::empty_value{VALUE_SENTINEL});
 
   auto val_idx = s_start_value_idx + t;
   while (val_idx < end_value_idx) {

cpp/src/join/hash_join.cu

@@ -286,8 +286,8 @@ hash_join<Hasher>::hash_join(cudf::table_view const& build,
     _composite_bitmask{cudf::detail::bitmask_and(build, stream).first},
     _nulls_equal{compare_nulls},
     _hash_table{compute_hash_table_size(build.num_rows()),
-                std::numeric_limits<hash_value_type>::max(),
-                cudf::detail::JoinNoneValue,
+                cuco::sentinel::empty_key{std::numeric_limits<hash_value_type>::max()},
+                cuco::sentinel::empty_value{cudf::detail::JoinNoneValue},
                 stream.value(),
                 detail::hash_table_allocator_type{default_allocator<char>{}, stream}}
 {

cpp/src/join/mixed_join.cu

@@ -127,8 +127,8 @@ mixed_join(
   // Don't use multimap_type because we want a CG size of 1.
   mixed_multimap_type hash_table{
     compute_hash_table_size(build.num_rows()),
-    std::numeric_limits<hash_value_type>::max(),
-    cudf::detail::JoinNoneValue,
+    cuco::sentinel::empty_key{std::numeric_limits<hash_value_type>::max()},
+    cuco::sentinel::empty_value{cudf::detail::JoinNoneValue},
     stream.value(),
     detail::hash_table_allocator_type{default_allocator<char>{}, stream}};
 
@@ -375,8 +375,8 @@ compute_mixed_join_output_size(table_view const& left_equality,
   // Don't use multimap_type because we want a CG size of 1.
   mixed_multimap_type hash_table{
     compute_hash_table_size(build.num_rows()),
-    std::numeric_limits<hash_value_type>::max(),
-    cudf::detail::JoinNoneValue,
+    cuco::sentinel::empty_key{std::numeric_limits<hash_value_type>::max()},
+    cuco::sentinel::empty_value{cudf::detail::JoinNoneValue},
     stream.value(),
     detail::hash_table_allocator_type{default_allocator<char>{}, stream}};
 

cpp/src/join/mixed_join_semi.cu

@@ -159,8 +159,8 @@ std::unique_ptr<rmm::device_uvector<size_type>> mixed_join_semi(
     cudf::nullate::DYNAMIC{has_nulls}, *probe_view, *build_view, compare_nulls};
 
   semi_map_type hash_table{compute_hash_table_size(build.num_rows()),
-                           std::numeric_limits<hash_value_type>::max(),
-                           cudf::detail::JoinNoneValue,
+                           cuco::sentinel::empty_key{std::numeric_limits<hash_value_type>::max()},
+                           cuco::sentinel::empty_value{cudf::detail::JoinNoneValue},
                            detail::hash_table_allocator_type{default_allocator<char>{}, stream},
                            stream.value()};
 
@@ -397,8 +397,8 @@ compute_mixed_join_output_size_semi(table_view const& left_equality,
     cudf::nullate::DYNAMIC{has_nulls}, *probe_view, *build_view, compare_nulls};
 
   semi_map_type hash_table{compute_hash_table_size(build.num_rows()),
-                           std::numeric_limits<hash_value_type>::max(),
-                           cudf::detail::JoinNoneValue,
+                           cuco::sentinel::empty_key{std::numeric_limits<hash_value_type>::max()},
+                           cuco::sentinel::empty_value{cudf::detail::JoinNoneValue},
                            detail::hash_table_allocator_type{default_allocator<char>{}, stream},
                            stream.value()};
 

cpp/src/join/semi_join.cu

@@ -91,8 +91,8 @@ std::unique_ptr<rmm::device_uvector<cudf::size_type>> left_semi_anti_join(
 
   // Create hash table.
   semi_map_type hash_table{compute_hash_table_size(right_num_rows),
-                           std::numeric_limits<hash_value_type>::max(),
-                           cudf::detail::JoinNoneValue,
+                           cuco::sentinel::empty_key{std::numeric_limits<hash_value_type>::max()},
+                           cuco::sentinel::empty_value{cudf::detail::JoinNoneValue},
                            hash_table_allocator_type{default_allocator<char>{}, stream},
                            stream.value()};
 

cpp/src/stream_compaction/distinct.cu

@@ -64,8 +64,8 @@ std::unique_ptr<table> distinct(table_view const& input,
   auto const num_rows{keys_view.num_rows()};
 
   hash_map_type key_map{compute_hash_table_size(num_rows),
-                        COMPACTION_EMPTY_KEY_SENTINEL,
-                        COMPACTION_EMPTY_VALUE_SENTINEL,
+                        cuco::sentinel::empty_key{COMPACTION_EMPTY_KEY_SENTINEL},
+                        cuco::sentinel::empty_value{COMPACTION_EMPTY_VALUE_SENTINEL},
                         detail::hash_table_allocator_type{default_allocator<char>{}, stream},
                         stream.value()};
 

cpp/src/stream_compaction/distinct_count.cu

@@ -130,8 +130,8 @@ cudf::size_type distinct_count(table_view const& keys,
   auto const has_null = nullate::DYNAMIC{cudf::has_nulls(keys)};
 
   hash_map_type key_map{compute_hash_table_size(num_rows),
-                        COMPACTION_EMPTY_KEY_SENTINEL,
-                        COMPACTION_EMPTY_VALUE_SENTINEL,
+                        cuco::sentinel::empty_key{COMPACTION_EMPTY_KEY_SENTINEL},
+                        cuco::sentinel::empty_value{COMPACTION_EMPTY_VALUE_SENTINEL},
                         detail::hash_table_allocator_type{default_allocator<char>{}, stream},
                         stream.value()};
 

cpp/src/text/subword/load_merges_file.cu

@@ -106,9 +106,9 @@ std::unique_ptr<detail::merge_pairs_map_type> initialize_merge_pairs_map(
   // Ensure capacity is at least (size/0.7) as documented here:
   // https://github.com/NVIDIA/cuCollections/blob/6ec8b6dcdeceea07ab4456d32461a05c18864411/include/cuco/static_map.cuh#L179-L182
   auto merge_pairs_map = std::make_unique<merge_pairs_map_type>(
-    static_cast<size_t>(input.size() * 2),              // capacity is 2x;
-    std::numeric_limits<cudf::hash_value_type>::max(),  // empty key;
-    -1,                                                 // empty value is not used
+    static_cast<size_t>(input.size() * 2),  // capacity is 2x;
+    cuco::sentinel::empty_key{std::numeric_limits<cudf::hash_value_type>::max()},
+    cuco::sentinel::empty_value{-1},  // empty value is not used
     hash_table_allocator_type{default_allocator<char>{}, stream},
     stream.value());