Skip to content

Commit 20b07bf

Browse files
PointKernelPointKernel
authored
Update cuCollections for thrust upgrade (#10983)
Depends on NVIDIA/cuCollections#165 and NVIDIA/cuCollections#171 This is a preparation step to finally address #10841. It fetches the latest version of `cuco` that [adds missing thrust headers ](NVIDIA/cuCollections#161) and [improves `static_map::retrieve_all`](NVIDIA/cuCollections#169). Authors: - Yunsong Wang (https://github.com/PointKernel) Approvers: - Bradley Dice (https://github.com/bdice) - Vyas Ramasubramani (https://github.com/vyasr) URL: #10983
1 parent ec93336 commit 20b07bf

File tree

9 files changed

+34
-27
lines changed

9 files changed

+34
-27
lines changed

‎cpp/cmake/thirdparty/get_cucollections.cmake‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@ function(find_and_configure_cucollections)
2222
GLOBAL_TARGETS cuco::cuco
2323
BUILD_EXPORT_SET cudf-exports
2424
CPM_ARGS GITHUB_REPOSITORY NVIDIA/cuCollections
25-
GIT_TAG ebaba1ae378a5272116414b6d7ae5847e5cf5715
25+
GIT_TAG 55029034c3f82bca36148c9be29941b37492394d
2626
EXCLUDE_FROM_ALL ${BUILD_SHARED_LIBS}
2727
OPTIONS "BUILD_TESTS OFF" "BUILD_BENCHMARKS OFF" "BUILD_EXAMPLES OFF"
2828
)

‎cpp/src/io/parquet/chunk_dict.cu‎

Lines changed: 14 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -125,8 +125,11 @@ __global__ void __launch_bounds__(block_size)
125125
column_device_view const& data_col = *col->leaf_column;
126126

127127
// Make a view of the hash map
128-
auto hash_map_mutable = map_type::device_mutable_view(
129-
chunk->dict_map_slots, chunk->dict_map_size, KEY_SENTINEL, VALUE_SENTINEL);
128+
auto hash_map_mutable =
129+
map_type::device_mutable_view(chunk->dict_map_slots,
130+
chunk->dict_map_size,
131+
cuco::sentinel::empty_key{KEY_SENTINEL},
132+
cuco::sentinel::empty_value{VALUE_SENTINEL});
130133

131134
__shared__ size_type total_num_dict_entries;
132135
size_type val_idx = s_start_value_idx + t;
@@ -184,9 +187,11 @@ __global__ void __launch_bounds__(block_size)
184187
auto& chunk = chunks[blockIdx.x];
185188
if (not chunk.use_dictionary) { return; }
186189

187-
auto t = threadIdx.x;
188-
auto map =
189-
map_type::device_view(chunk.dict_map_slots, chunk.dict_map_size, KEY_SENTINEL, VALUE_SENTINEL);
190+
auto t = threadIdx.x;
191+
auto map = map_type::device_view(chunk.dict_map_slots,
192+
chunk.dict_map_size,
193+
cuco::sentinel::empty_key{KEY_SENTINEL},
194+
cuco::sentinel::empty_value{VALUE_SENTINEL});
190195

191196
__shared__ cuda::atomic<size_type, cuda::thread_scope_block> counter;
192197
using cuda::std::memory_order_relaxed;
@@ -233,8 +238,10 @@ __global__ void __launch_bounds__(block_size)
233238

234239
column_device_view const& data_col = *col->leaf_column;
235240

236-
auto map = map_type::device_view(
237-
chunk->dict_map_slots, chunk->dict_map_size, KEY_SENTINEL, VALUE_SENTINEL);
241+
auto map = map_type::device_view(chunk->dict_map_slots,
242+
chunk->dict_map_size,
243+
cuco::sentinel::empty_key{KEY_SENTINEL},
244+
cuco::sentinel::empty_value{VALUE_SENTINEL});
238245

239246
auto val_idx = s_start_value_idx + t;
240247
while (val_idx < end_value_idx) {

‎cpp/src/join/hash_join.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -286,8 +286,8 @@ hash_join<Hasher>::hash_join(cudf::table_view const& build,
286286
_composite_bitmask{cudf::detail::bitmask_and(build, stream).first},
287287
_nulls_equal{compare_nulls},
288288
_hash_table{compute_hash_table_size(build.num_rows()),
289-
std::numeric_limits<hash_value_type>::max(),
290-
cudf::detail::JoinNoneValue,
289+
cuco::sentinel::empty_key{std::numeric_limits<hash_value_type>::max()},
290+
cuco::sentinel::empty_value{cudf::detail::JoinNoneValue},
291291
stream.value(),
292292
detail::hash_table_allocator_type{default_allocator<char>{}, stream}}
293293
{

‎cpp/src/join/mixed_join.cu‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -127,8 +127,8 @@ mixed_join(
127127
// Don't use multimap_type because we want a CG size of 1.
128128
mixed_multimap_type hash_table{
129129
compute_hash_table_size(build.num_rows()),
130-
std::numeric_limits<hash_value_type>::max(),
131-
cudf::detail::JoinNoneValue,
130+
cuco::sentinel::empty_key{std::numeric_limits<hash_value_type>::max()},
131+
cuco::sentinel::empty_value{cudf::detail::JoinNoneValue},
132132
stream.value(),
133133
detail::hash_table_allocator_type{default_allocator<char>{}, stream}};
134134

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

‎cpp/src/join/mixed_join_semi.cu‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -159,8 +159,8 @@ std::unique_ptr<rmm::device_uvector<size_type>> mixed_join_semi(
159159
cudf::nullate::DYNAMIC{has_nulls}, *probe_view, *build_view, compare_nulls};
160160

161161
semi_map_type hash_table{compute_hash_table_size(build.num_rows()),
162-
std::numeric_limits<hash_value_type>::max(),
163-
cudf::detail::JoinNoneValue,
162+
cuco::sentinel::empty_key{std::numeric_limits<hash_value_type>::max()},
163+
cuco::sentinel::empty_value{cudf::detail::JoinNoneValue},
164164
detail::hash_table_allocator_type{default_allocator<char>{}, stream},
165165
stream.value()};
166166

@@ -397,8 +397,8 @@ compute_mixed_join_output_size_semi(table_view const& left_equality,
397397
cudf::nullate::DYNAMIC{has_nulls}, *probe_view, *build_view, compare_nulls};
398398

399399
semi_map_type hash_table{compute_hash_table_size(build.num_rows()),
400-
std::numeric_limits<hash_value_type>::max(),
401-
cudf::detail::JoinNoneValue,
400+
cuco::sentinel::empty_key{std::numeric_limits<hash_value_type>::max()},
401+
cuco::sentinel::empty_value{cudf::detail::JoinNoneValue},
402402
detail::hash_table_allocator_type{default_allocator<char>{}, stream},
403403
stream.value()};
404404

‎cpp/src/join/semi_join.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -91,8 +91,8 @@ std::unique_ptr<rmm::device_uvector<cudf::size_type>> left_semi_anti_join(
9191

9292
// Create hash table.
9393
semi_map_type hash_table{compute_hash_table_size(right_num_rows),
94-
std::numeric_limits<hash_value_type>::max(),
95-
cudf::detail::JoinNoneValue,
94+
cuco::sentinel::empty_key{std::numeric_limits<hash_value_type>::max()},
95+
cuco::sentinel::empty_value{cudf::detail::JoinNoneValue},
9696
hash_table_allocator_type{default_allocator<char>{}, stream},
9797
stream.value()};
9898

‎cpp/src/stream_compaction/distinct.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -64,8 +64,8 @@ std::unique_ptr<table> distinct(table_view const& input,
6464
auto const num_rows{keys_view.num_rows()};
6565

6666
hash_map_type key_map{compute_hash_table_size(num_rows),
67-
COMPACTION_EMPTY_KEY_SENTINEL,
68-
COMPACTION_EMPTY_VALUE_SENTINEL,
67+
cuco::sentinel::empty_key{COMPACTION_EMPTY_KEY_SENTINEL},
68+
cuco::sentinel::empty_value{COMPACTION_EMPTY_VALUE_SENTINEL},
6969
detail::hash_table_allocator_type{default_allocator<char>{}, stream},
7070
stream.value()};
7171

‎cpp/src/stream_compaction/distinct_count.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -130,8 +130,8 @@ cudf::size_type distinct_count(table_view const& keys,
130130
auto const has_null = nullate::DYNAMIC{cudf::has_nulls(keys)};
131131

132132
hash_map_type key_map{compute_hash_table_size(num_rows),
133-
COMPACTION_EMPTY_KEY_SENTINEL,
134-
COMPACTION_EMPTY_VALUE_SENTINEL,
133+
cuco::sentinel::empty_key{COMPACTION_EMPTY_KEY_SENTINEL},
134+
cuco::sentinel::empty_value{COMPACTION_EMPTY_VALUE_SENTINEL},
135135
detail::hash_table_allocator_type{default_allocator<char>{}, stream},
136136
stream.value()};
137137

‎cpp/src/text/subword/load_merges_file.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -106,9 +106,9 @@ std::unique_ptr<detail::merge_pairs_map_type> initialize_merge_pairs_map(
106106
// Ensure capacity is at least (size/0.7) as documented here:
107107
// https://github.com/NVIDIA/cuCollections/blob/6ec8b6dcdeceea07ab4456d32461a05c18864411/include/cuco/static_map.cuh#L179-L182
108108
auto merge_pairs_map = std::make_unique<merge_pairs_map_type>(
109-
static_cast<size_t>(input.size() * 2), // capacity is 2x;
110-
std::numeric_limits<cudf::hash_value_type>::max(), // empty key;
111-
-1, // empty value is not used
109+
static_cast<size_t>(input.size() * 2), // capacity is 2x;
110+
cuco::sentinel::empty_key{std::numeric_limits<cudf::hash_value_type>::max()},
111+
cuco::sentinel::empty_value{-1}, // empty value is not used
112112
hash_table_allocator_type{default_allocator<char>{}, stream},
113113
stream.value());
114114

0 commit comments

Comments
 (0)